:)
当我尝试管理内核资源时，我决定研究一下 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
现在有些事情我不明白：
- 根据ptx汇编，使用了4+5+8+5=22个寄存器。那为什么它说
used 2 registers编译期间？
- 查看程序集，我意识到 threadId、blockId 等的数据类型是
u16。这是CUDA规范中定义的吗？或者不同版本的 CUDA 驱动程序之间可能会有所不同？
- 有人可以向我解释一下这一行：
mul.wide.u16 %r1, %rh1, %rh2;? %r1 is u32, why wide代替u32用来？
- 寄存器的名称是如何选择的？在我的花瓶里我明白
%r部分但我不明白h,(null),d部分。是根据数据类型长度来选择的吗？ IE：h16 位为空，32 位为空，d对于64位？
- 如果我用这个替换内核的最后两行
out[idx] = in[idx];，然后当我编译程序时它说使用了3个寄存器！现在怎么可能使用更多的寄存器呢？
请忽略我的测试内核不检查数组索引是否越界的事实。
非常感谢。
