- html - 出于某种原因,IE8 对我的 Sass 文件中继承的 html5 CSS 不友好?
- JMeter 在响应断言中使用 span 标签的问题
- html - 在 :hover and :active? 上具有不同效果的 CSS 动画
- html - 相对于居中的 html 内容固定的 CSS 重复背景?
:)当我尝试管理内核资源时,我决定研究一下 PTX,但有一些事情我不明白。这是我编写的一个非常简单的内核:
__global__
void foo(float* out, float* in, uint32_t n)
{
uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
uint32_t one = 5;
out[idx] = in[idx]+one;
}
然后我使用:nvcc --ptxas-options=-v -keep main.cu
编译它,并在控制台上得到了以下输出:
ptxas info : 0 bytes gmem
ptxas info : Compiling entry function '_Z3fooPfS_j' for 'sm_10'
ptxas info : Used 2 registers, 36 bytes smem
生成的 ptx 如下:
.entry _Z3fooPfS_j (
.param .u64 __cudaparm__Z3fooPfS_j_out,
.param .u64 __cudaparm__Z3fooPfS_j_in,
.param .u32 __cudaparm__Z3fooPfS_j_n)
{
.reg .u16 %rh<4>;
.reg .u32 %r<5>;
.reg .u64 %rd<8>;
.reg .f32 %f<5>;
.loc 15 17 0
$LDWbegin__Z3fooPfS_j:
.loc 15 21 0
mov.u16 %rh1, %ctaid.x;
mov.u16 %rh2, %ntid.x;
mul.wide.u16 %r1, %rh1, %rh2;
cvt.u32.u16 %r2, %tid.x;
add.u32 %r3, %r2, %r1;
cvt.u64.u32 %rd1, %r3;
mul.wide.u32 %rd2, %r3, 4;
ld.param.u64 %rd3, [__cudaparm__Z3fooPfS_j_in];
add.u64 %rd4, %rd3, %rd2;
ld.global.f32 %f1, [%rd4+0];
mov.f32 %f2, 0f40a00000; // 5
add.f32 %f3, %f1, %f2;
ld.param.u64 %rd5, [__cudaparm__Z3fooPfS_j_out];
add.u64 %rd6, %rd5, %rd2;
st.global.f32 [%rd6+0], %f3;
.loc 15 22 0
exit;
$LDWend__Z3fooPfS_j:
} // _Z3fooPfS_j
现在有些事情我不明白:
used 2 registers
呢?u16
。这是CUDA规范中定义的吗?或者不同版本的 CUDA 驱动程序之间可能会有所不同?mul.wide.u16 %r1, %rh1, %rh2;
吗? %r1
是u32
,为什么使用wide
而不是u32
?%r
部分,但不理解 h
,(null),d
部分。是根据数据类型长度来选择的吗?即:h
代表 16 位,null 代表 32 位,d
代表 64 位?out[idx] = in[idx];
替换内核的最后两行,那么当我编译程序时,它会说使用了 3 个寄存器!现在怎么可能使用更多的寄存器?请忽略我的测试内核不检查数组索引是否越界的事实。
非常感谢。
最佳答案
PTX 是一种中间语言,旨在跨多个 GPU 架构进行移植。它由编译器组件 PTXAS 编译为最终的机器代码,也称为 SASS,适用于特定的体系结构。 nvcc 选项 -Xptxas -v
使 PTXAS 报告有关生成的机器代码的各种统计信息,包括机器代码中使用的物理寄存器的数量。您可以通过使用 cuobjdump --dump-sass
反汇编来检查机器代码。 .
因此,PTX 代码中使用的寄存器数量没有任何意义,因为这些是虚拟寄存器。 CUDA 编译器以所谓的 SSA 形式生成 PTX 代码(静态单一赋值,请参阅 http://en.wikipedia.org/wiki/Static_single_assignment_form)。这基本上意味着写入的每个新结果都会分配一个新寄存器。
指令mul.wide
PTX 规范中对此进行了描述,您可以在此处找到其当前版本 (3.1):http://docs.nvidia.com/cuda/parallel-thread-execution/index.html 。在您的示例代码中,后缀 .u16
意味着它将两个无符号 16 位量相乘并返回无符号 32 位结果,即它计算源操作数的完整双角乘积。
PTX 中的虚拟寄存器是有类型的,但它们的名称可以自由选择,与类型无关。 CUDA 编译器似乎遵循某些约定(据我所知)没有记录,因为它们是内部实现工件。看看一堆PTX代码,很明显当前生成的寄存器名称编码类型信息,这样做可能是为了便于调试: p<num>
用于谓词,r<num>
对于 32 位整数,rd<num>
对于 64 位整数,f<num>
对于 32 位 float ,和 fd<num>
对于 64 位 double 。您可以通过查看.reg
轻松地亲自看到这一点。 PTX 代码中创建这些虚拟寄存器的指令。
关于cuda - CUDA PTX 代码和寄存器内存的混淆,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/16975727/
我无法理解如何使用一些旧的 VGA 代码在这个示例中设置序列 Controller 寄存器: mov dx,SC_INDEX mov ax,0604h out dx,ax
我希望对 zmm 0-31 寄存器集的四字元素执行整数算术运算并保留这些运算产生的进位位。看来这只有在通用寄存器集中处理数据时才有可能。 因此,我想将信息从 zmm 0-31 寄存器之一复制到通用寄存
ARM 64中包含多种寄存器,下面介绍一些常见的寄存器。 1 通用寄存器 ARM 64包含31个64bit寄存器,记为X0~X30。 每一个通用寄存器,它的低32bit都可以被访问,记为W0~W
1.寄存器 组合逻辑存在一个最大的缺点就是存在竞争与冒险,系统会产生不定态;使用时序逻辑电路就会极大的改善这种情况 寄存器具有存储功能,一般是由D触发器构成,由时钟脉冲控制,每个D触发器能够
使用 $gp 是否存在危险?注册以存储值?我想我的问题是 $gp 的真正功能是什么?它是否以某种方式在幕后调用,以便如果我使用它,事情可能会变得非常非常错误? 最佳答案 那么,$gp register
我遇到了这段代码的问题,我无法弄清楚问题出在哪里。所以当我运行这段代码时:if $row["count"] > 0 else块运行和 $_SESSION["error"]设置。 当$row["coun
所以我正在做二进制炸弹的变体。这就是阶段 0x0000000000401205 : sub $0x8,%rsp 0x0000000000401209 : cmp $0x3,
我在一个名为 (EmployeeDetailKey - varchar(10)) 的字段中获得了一个值,其中包含顺序值,例如 00001, 00002, 00003.... 它位于 Employeed
我有一个要求,应该为每个调用的线程分配一个计数器变量。但我没有得到预期的结果,实际上计数器在线程中重复。我创建了一个虚拟表和一个过程来将计数器值插入到表中。无论如何,是否可以更改代码以便线程获得递增的
预期输出:需要打印第4季度的wage_amt +--------------+--------------+--------------+--------------+ | wages_amt_q1
如何匹配模式 abc_[someArbitaryStringHere]_xyz? 为了澄清,我希望正则表达式能够匹配以下性质的字符串: abc_xyz、abc_asdfsdf_xyz、abc_32rw
从下拉列表(自定义)中选择一个值而不是常规下拉列表,它有很多下拉值 我可以用代码选择第一个值 find('.selected', :text=>arg1,exact: false).click 但无法
这个问题不太可能帮助任何 future 的访问者;它只与一个小的地理区域、一个特定的时间点或一个非常狭窄的情况有关,这些情况并不普遍适用于互联网的全局受众。为了帮助使这个问题更广泛地适用,visit
我有 .csv 文件中的数据,它包含 2 列 x 轴和 y 轴。从 .csv 文件读取轴,然后使用拉伸(stretch)指数函数拟合数据,但显示错误。 这里我给出示例数据以方便理解。 我的函数是f(x
我正在尝试使用以下汇编代码将磁盘扇区加载到内存中,但正如我在终端中使用一些 int 0x10 时发现的那样,它不起作用的原因是它陷入了无限循环。我以为循环会自动为我递减 cx 寄存器。下面是代码,其中
我正在尝试编写一个脚本,该脚本将在 vim 中打开一个文件并将其中的特定行复制到 vim 的寄存器之一中。当脚本再次运行时,它会决定再次打开文件,然后将 vim 寄存器中的值粘贴回。实际上,脚本应该在
我目前正在尝试弄清楚如何将指针寄存器 SI 指向的内存中的第一个字节添加到 AX 寄存器的当前内容中。 因此,如果 SI 包含某个地址,并且该地址在内存中的值是:00 和 01,我希望将 00 添加到
我试图将两个 16 位数字与以下 NASM 代码相乘: mov ax, [input1] mov bx, [input2] mul bx 前面代码的结果存储在 DX:AX 我试图使用来自单独库“pri
我正在尝试修改 rip 寄存器(只是为了好玩)。 buffer 应该是内存地址,所以不知道为什么会得到Error: operand type mismatch for 'movq' #include
我可以告诉gcc-style inline assembly把我的__m512i变量到特定 zmm注册,如 zmm31 ? 最佳答案 就像在根本没有特定寄存器约束的目标(如 ARM)上一样,使用 lo
我是一名优秀的程序员,十分优秀!