:) カーネル リソースを管理しようとしていたときに、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は、代わりにが使用されるのはu32なぜですか?wideu32 - レジスターの名前はどのように選ばれますか? 私の花瓶では、その部分は理解できますが、(null) の部分
%rは理解できません。データ型の長さに基づいて選択されていますか? すなわち: 16 ビットの場合、32 ビットの場合は null、64 ビットの場合は?hdhd - カーネルの最後の 2 行をこれに置き換えると
out[idx] = in[idx];、プログラムをコンパイルすると、3 つのレジスタが使用されていると表示されます。現在、より多くのレジスタを使用するにはどうすればよいですか?
テスト カーネルが配列インデックスが範囲外かどうかをチェックしないという事実を無視してください。
どうもありがとうございました。