CUDA カーネルを実行すると何が起きるか（35 分読）

これはシンプルな CUDA プログラムです。2 つのベクトルを加算します。 __global__ void vadd(const float* a, const float* b, float* c, int n) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) c[i] = a[i] + b[i]; } int main() { int n = 1 << 20; // 100 万個の浮動小数点数（1,048,576） size_t bytes = n * sizeof(float); float *a = (float*)malloc(bytes), *b = (float*)malloc(bytes), *c = (float*)malloc(bytes); for (int i = 0; i >>(da, db, dc, n); // 4096 * 256 = n スレッド、浮動小数点数ごとに 1 つ cudaMemcpy(c, dc, bytes, cudaMemcpyDeviceToHost); // デバイスからホストへデータ転送 printf("c[0]=%f c[n-1]=%f\n", c[0], c[n-1]); } RTX 4090 向けにコンパイルして実行すると、正しく 1+1=2 が計算されることが確認できます。これを 100 万回行います。すべてをチェックしたわけではありませんが..。 $ nvcc -arch=sm_89 -o vadd vadd.cu && ./vadd c[0]=2.000000 c[n-1]=2.000000 CPU 命令が数千万回、デバイスファイルが数個、IOCTL が 900 回、そしてメモリマップドされたドアベルレジスタが 1 つ関与していることをお伝えしました。今回の記事では、このカーネルをコードからウォープまで、そして再び答えへと遡る過程を追跡します。 余談ですが、この記事はエージェントが生み出した「可視性の転換」の一例です。好奇心と（機械支援された）粘り強ささえあれば、コンピュータに関する事実で知ることができないことはほとんどありません。AI が私たちに何を教えることができるかという点における可視性の影響についての興味深い議論 [こちら](https://resobscura.substack.com/p/ai-legibility-archives-future-of-research)。 ## nvcc でプログラムをコンパイルする§ まず、この CUDA プログラムをデバイスが実際に読み取れる形に変換する方法から始めましょう。そのためにはコンパイラが必要です。実際には、多くのコンパイラが必要になります。 nvcc は、他の複数のコンパイラを実行してその出力を結合するドライバプログラムです。--keep オプションを指定すると、処理パイプライン全体がディスク上に残され、読み取れるようになります： $ nvcc --keep -arch=sm_89 -o vadd vadd.cu && ls ... vadd.ptx # PTX 形式のデバイスコード (cicc から) vadd.sm_89.cubin # SASS 形式のデバイスコード (ptxas から) vadd.fatbin # cubin と PTX をバンドルしたもの (fatbinary から) vadd.cudafe1.stub.c # ホスト起動用スタブとカーネル登録 vadd.o # 最終的なホストオブジェクト、fatbin が埋め込まれる ... ホストコードはホストコンパイラに渡されます。一方、デバイスコード（vadd）はより多くのステップを経ます：cicc という [LLVM](https://en.wikipedia.org/wiki/LLVM) ベースのコンパイラがこれを [PTX](https://developer.nvidia.com/blog/understanding-ptx-the-assembly-language-of-cuda-gpu-computing/) に変換し、その後 ptxas が PTX を [SASS](https://modal.com/gpu-glossary/device-software/streaming-assembler) に変換します。 PTX は仮想の [ISA（命令セットアーキテクチャ）](https://en.wikipedia.org/wiki/Instruction_set_architecture) です。無限に存在する型付きレジスタを持ち、ハードウェアが実際にいくつのレジスタを持っているかという概念はありません。以下は PTX における vadd の本体（一部省略）です： $ cat vadd.ptx ... mad.lo.s32 %r1, %r3, %r4, %r5; // レジスタ r1 に ctaid*ntid + tid を設定 setp.ge.s32 %p1, %r1, %r2; // i >= n の場合、述語 p1 を true に設定 @%p1 bra $L__BB0_2; // 範囲外の場合、終了ラベルへジャンプ cvta.to.global.u64 %rd4, %rd1; // 汎用ポインタ %rd1 をグローバルアドレスに変換し、%rd4 に格納 mul.wide.s32 %rd5, %r1, 4; // r1 に 4 を乗算し、その結果を %rd5 に格納 add.s64 %rd6, %rd4, %rd5; // %rd4 と %rd5 を加算し、結果を %rd6 に格納 ld.global.f32 %f2, [%rd6]; // a[i] を %f2 にロード ... add.f32 %f3, %f2, %f1; %f1 と %f2 を加算し、結果を %f3 に格納 st.global.f32 [%rd10], %f3; // c[i] = ... をグローバルメモリにストア 仮想レジスタは%rd1～%rd10、%f1～%f333のように表示されます。プレフィックスは型を示し、%r は 32 ビット整数、%rd は 64 ビット整数、%f は 32 ビット浮動小数点数、%p は 1 ビットの述語を表します。 PTX は予想以上に「手書き」的な記法です。例えば、%rd6 にアドレスを形成するには 3 つの PTX 命令が必要です。これは、PTX がデバイス非依存であるためです。 なぜ 3 つなのか？CUDA ポインタはデフォルトで「ジェネリック」であり、グローバルメモリ、共有メモリ、ローカルメモリのいずれも指し得ます。cvta.to.global はポインタがグローバルウィンドウ内にあることを宣言し、その後より安価な ld.global を使用可能にします。次に mul.wide.s32 がインデックス i に 4（sizeof(float)）を乗算してバイトオフセットに変換すると同時に、32 ビットから 64 ビットへワンステップで拡張します。最後に add.s64 でこれをベースポインタに加算します。 次に ptxas は、デバイス非依存である PTX を、アーキテクチャに依存する SASS に変換します。生成される SASS は以下のように異なります： $ cuobjdump -sass vadd /*0000*/ MOV R1, c[0x0][0x28] ; // set up the stack pointer (ABI; unused here) /*0010*/ S2R R6, SR_CTAID.X ; // R6 = blockIdx.x /*0020*/ S2R R3, SR_TID.X ; // R3 = threadIdx.x /*0030*/ IMAD R6, R6, c[0x0][0x0], R3 ; // i = ctaid*ntid + tid /*0040*/ ISETP.GE.AND P0, PT, R6, c[0x0][0x178], PT ;// P0 = (i >= n) /*0050*/ @P0 EXIT ; // if so, exit /*0060*/ MOV R7, 0x4 ; // load literal 4 (sizeof(float)) into R7 as multiplier /*0070*/ ULDC.64 UR4, c[0x0][0x118] ; // uniform load of a driver-provided system value /*0080*/ IMAD.WIDE R4, R6, R7, c[0x0][0x168] ; // &b[i] /*0090*/ IMAD.WIDE R2, R6, R7, c[0x0][0x160] ; // &a[i] /*00a0*/ LDG.E R4, [R4.64] ; // b[i] /*00b0*/ LDG.E R3, [R2.64] ; // a[i] /*00c0*/ IMAD.WIDE R6, R6, R7, c[0x0][0x170] ; // &c[i] /*00d0*/ FADD R9, R4, R3 ; // a[i] + b[i] /*00e0*/ STG.E [R6.64], R9 ; // c[i] = ... /*00f0*/ EXIT ; What the S2R lines are doing S2R は「特殊レジスタからレジスタへ」を意味し、ハードウェアがスレッドごとに維持する*特殊*レジスタ（ここでは SR_CTAID.X：ブロックのインデックスである blockIdx.x および SR_TID.X：ブロック内のレーンのインデックスである threadIdx.x）の内容を、通常のレジスタにコピーします。これにより、IMAD がその値に対して算術演算を実行できるようになります。 10 弱の仮想レジスタが 7 つの実態レジスタに圧縮されました。ncu のレポートによると、スレッドあたりの起動レジスタ数は 16 です。ディスアセンブリでは R9 までしか名前が示されていませんが、アロケータは ABI と整列のためにさらにいくつかを予約しています。 2 つの mul.wide に加算するシーケンスが融合され、単一の IMAD.WIDE になりました。cvta 変換はなくなり、アドレス指定に吸収されています。 c[0x0][…]オペランドは**定数バンク 0**にあり、これはドライバが管理する小領域にあります。これらはカーネルの引数であり、ポインタ a, b, c とサイズ n、および起動ジオメトリです。このバンクへの書き込みは、起動時に GPU に渡される QMD という構造体の役割です。これは起動自体がカードに到達した時点で詳しく説明します。 なぜ引数が定数バンク 0 に置かれているのか、そしてどこにあるのかについて：これらは定数メモリ内にあるのは、これが*ブロードキャスト*読み取りであるためです。グリッド内のすべてのスレッドが同一のポインタを必要とし、定数キャッシュは 32 ラーンすべてを一度にサービスできます。レイアウトは固定されており、0x160, 0x168, 0x170 がそれぞれポインタ a, b, c を表し、0x178 は n です。起動ジオメトリも隣接して 0x0 に配置されています（blockDim.x）。バンク 0 には ABI パラメータも保持されており、例えば c[0x0][0x28] はスタックベースであり、MOV R1, c[0x0][0x28] がエントリー時に読み込む値です。ホストのスタブが起動用に引数をパックする際にも、これらのオフセットを再度確認することになります。 この SASS を保持する「cubin」ファイルは、[ELF](https://en.wikipedia.org/wiki/Executable_and_Linkable_Format) ファイルです。これは Linux が通常の実行可能ファイルや共有ライブラリに使用するのと同じオブジェクトファイルコンテナです55.cuobjdump -elf を実行すると、シンボルテーブルが表示され、機械コードを保持する .text.vadd セクションに加え、CUDA 固有のセクションである .nv.callgraph.. などが見られます。fatbinary 実行可能ファイルは、cubin と PTX を単一の「fatbin」にバンドルしており、その結果に対して cuobjdump を実行すると、バイナリに埋め込まれた fatbin が*両方*を内包していることが明らかになります。 $ cuobjdump vadd ... Fatbin elf code: arch = sm_89 # 先ほど読み取った SASS Fatbin ptx code: arch = sm_89 compressed # 同梱される PTX SASS は実際にこの 4090 で実行されるものですが、PTX は将来互換性のフォールバックとして一緒に持ち運ばれます。その後、このバイナリを cubin がカバーしていないアーキテクチャの GPU に持っていくと、ドライバーがロード時に PTX を JIT コンパイルして新しい SASS を生成します。 最後に、その fatbin はホスト実行可能ファイルの中にネストされており、readelf -S で確認すると、独自のセクションとして占有していることがわかります。 $ readelf -S vadd ... [18] .nv_fatbin PROGBITS ... [19] __nv_module_id PROGBITS ... [29] .nvFatBinSegment PROGBITS ... ... nvcc が出力する vadd バイナリは、ホストコードを含む単一の実行可能ファイルであり、Ada SASS を完全含む ELF オブジェクトと、PTX のコピーを内包しています。PTX は冗長なプレーンテキストであるため、バイナリのサイズを小さく保つために nvcc はデフォルトでこれを圧縮します。ドライバーは、事前コンパイルされた SASS でカバーされていないアーキテクチャ上でこのバイナリが実行された場合にのみ、これを展開して JIT コンパイルを実行します。 ## ホストが GPU をトリガーする方法§ コンパイル済みの GPU マシンコードは、現在 ./vadd 実行ファイルの .nv_fatbin セクション内で静止状態となっています。ホスト上でプログラムを起動する際、私たちは PCIe バスを挟んで存在するホスト CPU と GPU の二つの世界を橋渡さなければなりません。 この橋を渡る方法を理解したホストバイナリを設定するために、フロントエンドコンパイラ（cudafe++）は main 関数が開始される前に実行される隠されたコンストラクタをコードに挿入します。その役割は、埋め込まれた fatbinary を CUDA ランタイムに登録し、ランタイムが後で使用するためのマッピングを記録することです。具体的には、ホスト側の関数ポインタ vadd と、fatbin 内の mangled な名前を持つコンパイル済みのデバイスカーネルとの関連付けを行います。 コンパイラが vadd >>(da, db, dc, n) という記述に出会うと、その高レベルな表現を生成されたホスト起動スタブに置き換えます。このスタブは、カーネル引数をホストメモリのバッファにパックします。ポインタ da, db, dc および整数 n は、バイトオフセット 0, 8, 16, 24 にそれぞれアラインされます。これらのオフセットは、先ほど SASS マシンコードが定数バンク 0 から読み込んでいた定数バンクオフセット 0x160, 0x168, 0x170, 0x178 です。 // from vadd.cudafe1.stub.c void __device_stub__Z4vaddPKfS0_Pfi(const float *__par0, const float *__par1, float *__par2, int __par3) { __cudaLaunchPrologue(4); __cudaSetupArgSimple(__par0, 0UL); // arg buffer offset 0 __cudaSetupArgSimple(__par1, 8UL); // offset 8 __cudaSetupArgSimple(__par2, 16UL); // offset 16 __cudaSetupArgSimple(__par3, 24UL); // offset 24 __cudaLaunch((char*)(void(*)(const float*, const float*, float*, int))vadd); } 引数がパッキングされると、スタブは __cudaLaunch を呼び出し、ホスト側のダミー vadd 関数のメモリアドレスを渡します。このホスト関数は CPU 上では単なる空のシェルに過ぎないため、そのホストメモリアドレスがルックアップキーとして機能します。ランタイムはこのアドレスを使用して登録テーブルを検索し、対応するデバイス側のシンボル名を見つけると、クローズドソースのユーザーモードドライバ（libcuda.so.1）77 の境界を越えて遷移します。このドライバのユーザーモード部分は CUDA ツールキットではなく、GPU のカーネルドライバに付属しています：strace で解決される libcuda.so.1 は、このマシン上のドライバリリースである libcuda.so.590.48.01 にリンクされます。これにより、そのカーネルの実行開始がトリガーされます。 ランタイムは、プログラム内の最初の GPU 呼び出し時にこのドライバを動的にオープンします。これは strace を使用して捕捉できます： $ strace -f -e trace=openat ./vadd ... openat(..., "/lib/x86_64-linux-gnu/libcuda.so.1", O_RDONLY|O_CLOEXEC) = 3 ... この最初の呼び出しが行われると、ドライバがデバイスと通信するために必要なすべてのインフラストラクチャを含む「コンテキスト」が作成されます。これには、CPU が GPU と通信するための *チャンネル* も含まれています。これについては次のセクションで詳しく説明します。 この段階では、コンパイルされたマシンコードはまだ GPU に到達していません。CUDA 12.2 以降、モジュール読み込みはデフォルトで遅延ロード（lazy loading）88 になっています。これは CUDA_MODULE_LOADING によって制御されており、CUDA 11.7 ではオプトインとして提供され、長年 EAGER がデフォルトでしたが、12.x シリーズではデフォルトが LAZY に切り替えられました（前方コストを前払いしたい場合は上書き可能です）。— ドライバは、特定のカーネルが実際に起動される最初の瞬間まで、そのカーネルの SASS cubin をカードのメモリにアップロードすることを先送りします。 libcuda の下にはカーネルモードドライバである nvidia.ko が存在し、libcuda はデバイスファイルに対して ioctl を呼び出すことでこれにアクセスします。cuLaunchKernel が実際に GPU に作業を割り当てる必要があるとき、それはそのカーネルモジュールとの対話になります。以下はその対話のメカニズムです。 ## GPU への転送§ GPU は CPU のように関数呼び出しを受け付けません。ジャンプするエントリポイントも存在せず、CPU から引数をスタックにプッシュすることもできません。GPU は PCIe バス上にあり、ホストメモリからドライバコマンドのストリームを読み取ります。この時点以降 cuLaunchKernel が行うことはすべて、完全に形成された起動コマンドをそのストリームに投入し、GPU に対してそれが完了したことを伝えることに役立ちます。 まず行わなければならないのは、GPU コードをデバイスにロードすることです。vadd を初めて実行する際、ドライバはカーネルのコードをコピーします：バッファを割り当てて SASS をコピーします。 コードが GPU に転送された後、CPU は GPU にそれを読み込んで実行を開始させる必要があります。これはホストとデバイスのメモリ間をまたぐ複雑な手順を通じて行われます。ホストと GPU の双方は互いのメモリ空間の領域をマップできますが、PCIe バスを経由したアクセスにはペナルティが発生します。カーネル起動を実現するため、両者は両方の空間に存在するさまざまな構造体に書き込みを行います。これらの構造体が *チャネル* を構成し、これは GPU の操作を実行する作業キューです。 ホスト RAM に存在する重要な構造体が 2 つあります。**プッシュバッファ**と**GPFIFO**で、これらが組み合わさって GPU が実行すべき作業リストを表しています。 **プッシュバッファ**は、ドライバーが GPU に対して *メソッド* と呼ばれるコマンドを書き込むメモリ領域です。メソッドとは、GPU のネイティブなコマンドエンコーディングにおけるレジスタアドレスと値のペアであり、この組み合わせによって GPU が実行すべきアクションが定義されます。 **GPFIFO**は、GPU と CPU が GPU がいまだに読み取る必要があるものと、すでに読み取ったものを調整するために使用するポインタのリングバッファです。GPFIFO の各エントリは 2 つの 32 ビットワードで構成され、プッシュバッファの一部（スパン）を記述します。99 この場合、base はホストメモリを指す GPU 仮想アドレスであり、(base, length) という形式をとります。 GPU は GPFIFO を継続的に走査して作業を見つけます。ドライバと GPU の間では、2 つのカーソルを維持する必要があります：GP_GET（GPU が消費した位置）と GP_PUT（ドライバが生成した位置）。両方のカーソルは USERD に存在し、これはチャネルごとの小さな構造体で、ここではデバイスメモリ上に配置されています。カーネルを実行するには、ドライバが関連するメソッドをプッシュバッファの範囲に埋め込み、GPFIFO エントリをその位置に指し示して GP_PUT を進めます。GPU がそのエントリを消費すると、GP_GET が進みます。 異なる部品がどこに存在するか。 CPU と GPU の PCIe を介したホスト RAM へのプッシュバッファへの書き込み、メソッドと QMD（Queue Meta Data）の GPFIFO リング、USERD — GP_GET / GP_PUT、ドアベル（MMIO）、HOST エンジン、DMA CPU · ホスト RAM → プッシュバッファ — メソッド + QMD の GPFIFO リング PCIe 経由での書き込み DMA → GPU USERD — GP_GET / GP_PUT、ドアベル（MMIO）→ HOST エンジン 私たちのカーネル起動は、一連のメソッドのバーストによってトリガーされます。まず SET_INLINE_QMD_ADDRESS_A/B が実行され（1010 行目）、libcuda がクローズドソースであるにもかかわらずこれがそのメソッドだと私が知っている理由は付録を参照してください。その後、LOAD_INLINE_QMD_DATA の一連の実行が続きます。 これらのメソッドは、「キューメタデータ」（QMD）と呼ばれるオブジェクトをプッシュバッファにストリーミングする役割を果たします。QMD は計算グリッドの起動記述子であり、グリッドとブロックの次元（.cu コードからの 4096 と 256）、スレッドあたりのレジスタ数、必要な共有メモリ、そして 2 つのアドレスを保持しています。その 1 つはプログラムの開始位置（GPU メモリに最初にロードされた SASS）であり、もう 1 つはカーネル引数を格納する定数バンクです。このバンクには、ホスト側のスタブがパッキングした引数が配置されます。ドライバがそれらをコピーし、QMD にそのバンクのアドレスを記録します。QMD は GPU に対して、SASS がどこにあるか、どのようにして SASS を並列プログラムに変換するか、そしてプログラムの完了をシグナルする場所を伝えます。 GPU が実行を開始するために必要なものはすべて整いました。問題は、GPU のホストエンジンがまだ動作していないことです。1111. GPU の制御論理のうち、ホストとインターフェースする部分。現代のカードではカーソルを監視していません1212. かつてはそうでした：古い GPU は USERD をスヌープしていたため、GP_PUT の書き込みだけで十分でした。Turing およびそれ以降のアーキテクチャではそのようではなく、そのためドライバが代わりにドアベルを鳴らします。したがって、GP_PUT の変更は、エンジンに確認するよう指示があるまでそのまま放置されます。 それはドアベルを通じて指示されます。GPU はそのレジスタの一部をプロセスにマッピングしており、その 1 つがドアベルです。ドライバはチャネルの作業提出トークンをそこに書き込みます。このトークンは、どのチャネルに新しい作業があるかを伝えます。 ドアベルが鳴ると、ホストエンジンは更新された GP_PUT を読み取り、新しい GPFIFO エントリをプッシュバッファのスパンへとたどり、DMA によってそこからメソッドを引き出します。計算メソッド（QMD を含む）に到達すると、その記述子を「計算作業配分器」に引き渡します。これについては後ほど詳しく説明します。 CPU の側から見ると、起動は完了しました。cuLaunchKernel はドアベルが鳴った瞬間に返却されます。この呼び出しは非同期であるため、制御はプログラムに戻り、GPU が動作している間も CPU は実行を続けます。カーネルの実行が終わったら、再びホスト側の処理を引き継ぎます。 いよいよ GPU がその役割を果たす時が来ました。 ## 命令ごとの実行§ CUDA カーネルを実行すると、GPU の各スレッドブロックがどのように動作するかを理解することが重要です。まず、カーネル呼び出しが発生すると、ホスト（CPU）はデバイス（GPU）にコマンドキューにタスクをプッシュします。その後、GPU のスケジューラが実行可能なブロックを選択し、SM（ストリーミングマルチプロセッサ）上でスレッドを起動します。 各スレッドは独立して命令を実行しますが、同じカーネル内のすべてのスレッドは同じコードパスを共有します。これは SIMT（Single Instruction, Multiple Threads）アーキテクチャの核心です。ただし、分岐条件によってスレッドが異なるパスをたどる場合、分岐回避（divergence）が発生し、パフォーマンスが低下する可能性があります。 メモリアクセスのパターンも重要です。グローバルメモリへのアクセスは遅延が高く、共役（coalesced）されたアクセスパターンで最適化することが推奨されます。ローカルメモリや共有メモリを適切に利用することで、データ転送のオーバーヘッドを削減できます。 また、カーネルの実行時間は、スレッドブロックの数、各ブロックのスレッド数、および各スレッドが実行する命令の数に依存します。最適な設定を見つけるには、プロファイリングツールを使用してボトルネックを特定し、パラメータを調整する必要があります。 最後に、エラーハンドリングも重要です。CUDA API やカーネル内で発生したエラーは、適切に検出・処理されなければ、予期せぬ動作やクラッシュの原因となります。エラーチェック関数を定期的に呼び出すことで、開発中の問題を早期に発見できます。 ホストエンジンが QMD を計算ワークディストリビュータに渡します。1313. 時折、ギガスレッドエンジンと呼ばれることもあります。このコンポーネントは GPU 全体に 1 つ存在します。VRAM 内には SASS 命令の単一の線形リストが存在し、計算ワークディストリビュータと QMD は、その線形のスレッド命令リストをすべてのストリーミングマルチプロセッサ (SMs) にわたる大規模並列プログラムに変換する方法をハードウェアに指示する最初のステップです。 スタックを下っていく私たちの旅において、現在の計算ワークディストリビュータは 256 スレッドからなる 4096 ブロックを記述する QMD を保持しています。私たちがターゲットとしているカードは、128 の SMs を持つ GeForce RTX 4090 チップです。1414. NVIDIA の AD102-300-A1 SKU では、製造歩留まりを最大化するために、フルダイ上の物理的な 144 の SMs のうち 16 が無効化されています。これは NVIDIA Ada GPU アーキテクチャのホワイトペーパーに詳述されています。ディストリビュータの任務は、すべての 128 を作業で飽和状態に保つことです。 コンパイルされたマシンコードは、グローバルメモリ内の単一の線形シーケンスとして存在します。各 SM は独自のローカル命令キャッシュ (I-cache) を持ち、GPU 上のすべてのアクティブなワープはそれぞれ独自のプライベートプログラムカウンタ (PC)1515. Volta 以降、モデルはさらに細分化されました — 各スレッドが独自のプログラムカウンタとコールスタック（独立スレッドスケジューリング）を持ち、ワープ内のスレッドが自由に分岐・再結合できるようになります。ただし、発行は依然としてワープ単位です — 各サイクルでスケジューラは 1 つのワープを選択し、現在共通の PC にあるレーンに発行します。SM 上のスケジューラはその後、その線形シーケンスから独立して命令を取得するため、異なるワープが同じ SASS コードを異なる速度で実行したり、異なる分岐パスを実行したりすることが可能になります。 VRAM 内の 1 つの命令ストリームが SM ごとにローカルにキャッシュされます。SM は最大 48 のワープ（グリッド）をレジデント状態に保ちますが、その 4 つのスケジューラは各サイクルで最大 1 命令ずつしか発行できません。ここではほぼすべてのワープが LDG.E ロード（オレンジ色）に待機しており、FADD（緑色）を発行しているのは 1 スロットのみです。 VRAM × 128 SMs の命令ストリーム (SASS) S2R R6, CTAID.X S2R R3, TID.X IMAD R6, R6, ntid, R3 ISETP.GE P0, R6, n@P0 EXIT IMAD.WIDE R4, … LDG.E R4, [R4] LDG.E R3, [R2] FADD R9, R4, R3 STG.E [R6], R9 EXIT SM (128 のうち 1 つ) I-cache: 48 レジデントワープ 4 スケジューラ · 1 命令/サイクル VRAM の命令ストリーム (SASS) S2R R6, CTAID.X S2R R3, TID.X IMAD R6, R6, ntid, R3 ISETP.GE P0, R6, n@P0 EXIT IMAD.WIDE R4, … LDG.E R4, [R4] LDG.E R3, [R2] FADD R9, R4, R3 STG.E [R6], R9 EXIT × 128 SMs SM (128 のうち 1 つ) I-cache: 48 レジデントワープ 4 スケジューラ · 1 命令/サイクル 私たちの SM のハードウェア制約が、同時に実行できるブロックの数を決定します。1616. cudaGetDeviceProperties はこの情報を提供します: +------------------------------------------------------------+ | AD102 SM リソースキャップ | +------------------------------------------------------------+ | Max Active Threads/SM | 1,536 スレッド (48 ワープ) | | Register File/SM | 65,536 32 ビットレジスタ (256 KB) | | Shared Memory/SM | 100 KB | +------------------------------------------------------------+ 私たちの起動構成は、256 スレッド（8 ワープ）からなるブロックを指定しており、ptxas はスレッドあたり 16 のレジスタを予約しました。 レジスタ容量：各ブロックには 256×16=4,096 のレジスタが必要です。 レジスタのみを考慮すると、SM は最大で 65,536/4,096=16 の定着ブロック（resident blocks）を収容できます。 - スレッド容量：ハードウェアは各 SM のアクティブスレッド数を 1,536 に制限しています。 この値をブロックサイズで割ると、1,536/256=6 の定着ブロックが得られます。 スレッド容量の方がより厳しいボトルネックであるため、各 SM が同時に保持できるのは最大で**6 ブロック（48 ワープ）**です。 ディス tributor はこれらの 6 つの定着ブロックを SM に割り当てます。各 SM は**4 つのプロセッシングブロック（サブパーティション）**に分割されています。各サブパーティションは独立した実行パイプラインです。 SM は、48 の定着ワープをこれら 4 つのサブパーティションに均等に分散します。したがって、SM が満杯の状態では、各ワープスケジューラが管理するアクティブワープ数は**12 ワープ**（48/4）となります。毎サイクル、ワープスケジューラは自身の 12 の候補を評価し、条件を満たす*eligible*なワープを 1 つ選択して、その実行スライスの 32 の物理レーンに対して次の命令をディスパッチします。 ## ワープが"eligible"（条件を満たす）であるとは何を意味するのか？§ GPU は命令の実行準備ができているタイミングを、CPU とは異なる方法で判断します。現代のアウト・オブ・オーダー CPU は、[Tomasulo のアルゴリズム](https://en.wikipedia.org/wiki/Tomasulo%27s_algorithm)のように実行時に動的に依存関係を発見し、[リオーダーバッファ](https://en.wikipedia.org/wiki/Re-order_buffer)や[レジスタ名付けロジック](https://en.wikipedia.org/wiki/Register_renaming)といった機構にシリコンを費やすことで、単一のスレッドから並列性を抽出します。一方、GPU はそのような機構を必要としません：多くのワープ（warp）を常駐させておき、それらがストールした際に切り替えることでレイテンシを隠蔽します。並列性が最優先される環境では、重厚な依存関係処理機構はシリコンの無駄遣いとなります。そのため、ハードウェアはタイミングを予測可能なすべての命令のスケジューリングをコンパイラに委ね、予測できないものについては軽量なハードウェアスコアボード（scoreboard）に頼ります。 128 ビットの SASS 命令には、ptxas1717 によって記述されたパケット化された制御コードペイロードが含まれています。最も明確な公開再構築例は、Citadel マイクロベンチマーク論文 ([Jia ら、「マイクロベンチマーキングによる NVIDIA Volta GPU アーキテクチャの解明」](https://arxiv.org/abs/1804.06826)) および Maxwell 向けの [maxas 制御コードノート](https://github.com/NervanaSystems/maxas/wiki/Control-Codes) です。これらのスケジューリング制御ビットはハードウェアのタイミングを直接規定しており、3 つの重要なディレクティブを含んでいます： - 静的ストールカウント：標準的な整数演算や浮動小数点計算のような固定レイテンシの命令の場合、コンパイラは ALU が結果を書き戻すタイミングを正確に把握しています。これは、このワープが次の命令を発行する前にどの程度待機すべきかをスケジューラに正確に伝えるサイクル数をエンコードしたものです。 - イールドヒント：このワープがスケジューリングの優先権を譲るべきかどうかをスケジューラに示す 1 ビットのフラグです。コンパイラがこのワープがボトルネックに直面しようとしていると判断した場合、このヒントを設定して、次のクロックサイクルで他のアクティブなワープを優先させるようにスケジューラに指示します。 - 依存関係バリアインデックス：コンパイル時に実行時間を予測できない可変レイテンシの演算、特にグローバルメモリの読み込み (LDG) や特殊関数 (MUFU) の場合、ハードウェアは各ワープに対して 6 つの物理的なスコアボードバリア（0 から 5 に番号付け）を提供します。 ディスアセンブルでこれらのビットが見られない理由 NVIDIA の標準ツールである nvdisasm を使用してバイナリをディスアセンブルすると、生の制御コードはデフォルトで隠蔽され、ツールはそれらを除去して標準的でクリーンな SASS マネモニクスを表示します。しかし、これらは命令のすぐ隣に直接保存されています。cuobjdump -sass を使用して生バイナリを検査し、16 進数の命令コメント（例：/* 0x... */）を注意深く確認すると、これらの制御ビットを格納したパッキングされた生の 16 進数ワードが見えます。 - A static stall count: For fixed-latency instructions—like standard integer or floating-point maths—the compiler knows exactly when the ALUs will write back. It encodes a precise cycle count telling the scheduler exactly how long to park this warp before issuing its very next instruction. - A yield hint: A single bit telling the scheduler whether this warp should yield its scheduling priority. If the compiler knows this warp is about to hit a bottleneck, it sets this hint to let the scheduler prioritize other active warps on the next clock cycle. - Dependency-barrier indices: For variable-latency operations whose duration cannot be predicted at compile time—most notably global memory loads (LDG) and special functions (MUFU)—the hardware provides six physical scoreboard barriers (numbered 0 to 5) per warp. Why you won't see these bits in the disassembly When you disassemble a binary using NVIDIA's standard nvdisasm tool, the raw control codes are hidden by default; the tool strips them away to show you standard, clean SASS mnemonics. However, they are stored directly alongside the instructions. If you inspect the raw binary using cuobjdump -sass and look closely at the hexadecimal instruction comments (e.g., /* 0x... */), you will see the packed, raw hex words that house these control bits. 彼らの正確なレイアウトに関する情報は、マイクロベンチマークコミュニティによる逆エンジニアリングの努力から得られています。ビットフィールドはMaxwell、Volta、Ampere、Ada Lovelaceの間でシフトし進化してきましたが、コアとなるアーキテクチャの構

CUDA カーネルを実行すると何が起きるか（35 分読）

背景や根拠まで確認しますか？

関連記事

調べる

選ぶ

サイト