8

:) カーネル リソースを管理しようとしていたときに、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 つのレジスタが使用されていると表示されます。現在、より多くのレジスタを使用するにはどうすればよいですか?

テスト カーネルが配列インデックスが範囲外かどうかをチェックしないという事実を無視してください。

どうもありがとうございました。

4

1 に答える 1

12

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は、2 つの符号なし 16 ビット量を乗算し、符号なし 32 ビットの結果を返すことを意味します。つまり、ソース オペランドの全倍幅の積を計算します。

PTX の仮想レジスタには型が付けられますが、その名前は型に関係なく自由に選択できます。CUDA コンパイラは、内部実装のアーティファクトであるため、(私の知る限り) 文書化されていない特定の規則に従っているようです。一連の PTX コードを見ると、現在生成されているレジスタ名が型情報をエンコードしていることは明らかです。これは、デバッグを容易にするために行われます。p<num>述語、r<num>32 ビット整数、rd<num>64 ビット整数、f<num>32ビットに使用されます。 float、およびfd<num>64 ビット double の場合。.regこれらの仮想レジスタを作成する PTX コード内のディレクティブを見ると、これを簡単に確認できます。

于 2013-06-07T05:24:50.623 に答える