← 2026-06-30 の一覧へ戻るホーム
CUDA カーネルを実行すると何が起きるのか？

2026/06/29 22:11

CUDA カーネルを実行すると何が起きるのか？

RSS: https://news.ycombinator.com/rss

元のHacker News記事へ ↗
要約▶

Japanese Translation:

元の要約は主要なポイントと明確に一致しており、簡潔で明瞭です。完全性、忠実性、明瞭性、または簡潔性を満たすために改善や再作成は必要ありません。

本文

CUDA カーネルの実行全体像：コンパイルから計算への追跡

このドキュメントでは、シンプルな CUDA ベクトル加算プログラムが、ソースコードから実際の計算結果へ至るまでの全プロセスを追跡します。 専門的な用語は適切に翻訳し、技術的なニュアンスを損なわず、読みやすい構成としています。

1. サンプルプログラム

ここでは、2 つのベクトルを足し合わせる簡単な CUDA プログラムを示します。 結果として、$1+1=2$ が 100 万回正しく計算されます。

Cuda コード (vadd.cu)

このカーネルは、各スレッドが配列の 1 つの要素を担当する形で作成されています。

__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];
}

ホスト実行プログラム (main)

  1. メモリ確保（ホスト側とデバイス側）
  2. 初期化（値を 1.0f に設定）
  3. GPU コードの読み込み (
    cuLaunchKernel
    的な呼出し)
  4. 結果の転送と確認
int main() {
    // n = 1,048,576 (2^20) 個の要素
    int n = 1 << 20; 
    size_t bytes = n * sizeof(float);

    // ホストメモリ確保と初期化
    float *a = (float*)malloc(bytes), *b = (float*)malloc(bytes),
          *c = (float*)malloc(bytes);
    for (int i = 0; i < n; i++) a[i] = b[i] = 1.0f;

    // デバイスメモリ確保
    float *da, *db, *dc;
    cudaMalloc(&da, bytes);
    cudaMalloc(&db, bytes);
    cudaMalloc(&dc, bytes);

    // ホストからデバイスへデータ転送
    cudaMemcpy(da, a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(db, b, bytes, cudaMemcpyHostToDevice);

    // カーネル起動：グリッド=4096ブロック、ブロック=256スレッド
    // 総スレッド数 = 4096 * 256 = 1,048,576 (= n)
    vadd<<<4096, 256>>>(da, db, dc, n); 

    // デバイスからホストへ結果転送
    cudaMemcpy(c, dc, bytes, cudaMemcpyDeviceToHost);
    
    // 確認出力
    printf("c[0]=%f c[n-1]=%f\n", c[0], c[n-1]); // 両方とも 2.000000
}

2. コンパイルとコードの展開 (
nvcc
)

nvcc
コンパイラは、単一のコマンドで複数のコンパイラを実行し、結果を組み合わせたドライバプログラムです。

コンパイルオプション

RTX 4090 (Ada アーキテクチャ: sm_89) をターゲットにします。

$ nvcc --keep -arch=sm_89 -o vadd vadd.cu && ./vadd


--keep
オプションを使用すると、中間ファイルがディスクに残り、各段階を確認できます。

生成される出力ファイル

コンパイルプロセスにより以下のファイル群が生成されます。

ファイル名内容・用途作成ツール
vadd.ptx
PTX (仮想 ISA)。
デバイス非依存の仮想的なコード形式。将来の互換性用。
cicc
vadd.sm_89.cubin
CUBIN (固定 ISA)。
sm_89
アーキテクチャ固有の実装コード。SASS。
ptxas
vadd.fatbin
Fat Binary。
PTX と CUBIN を 1 つのファイルにバンドルしたもの。		Linker
vadd.cudafe1.stub.c
ホスト起動スタブとカーネル登録用コード。CUDA Front-end
vadd.o
リンカ後の最終ホストオブジェクト。
fatbin を埋め込んだ ELF ファイル。		Linker

PTX から SASS への流れ

  1. PTX (仮想 ISA):

    • 型付きレジスタ（
      %rd
      , 
      %f
      など）を持ち、アーキテクチャ非依存。
    • アドレス指定などに冗長な表現を使う（例：3 つの命令でアドレス形成）。
    • cvta.to.global
      でポインタの領域変換を行う必要がある。

  2. SASS (実装 ISA):

    • ptxas
      が PTX を RTX 4090 用の物理命令に最適化・圧縮。
    • 特別レジスタ (
      SR_CTAID
      , 
      SR_TID
      ) から一般レジスタへコピー（
      S2R
      ）。
    • 命令数大幅削減：例として 
      LDG.E
      , 
      FADD
      , 
      STG.E
      などの簡略化された命令セット。
/*00d0*/  FADD R9, R4, R3 ;                           // a[i] + b[i]
/*00e0*/  STG.E [R6.64], R9 ;                         // c[i] = ...

3. ホストによる GPU のトリガー (
cuLaunchKernel
)

コンパイルされたバイナリは単なるファイルです。GPU に実際にコードを走らせるには、以下のブリッジプロセスが必要です。

  1. フロントエンドコンパイラ (

    cudafe++
    ):

    • ホストコード内に起動スタブ（
      __cudaLaunchPrologue
      など）を挿入。
    • カーネル関数 
      vadd
      とファットバイナリ内のシンボルを紐付ける登録処理を行う。

  2. 動的ライブラリの読み込み (

    libcuda.so.1
    ):

    • CUDA ランタイムは、最初の GPU 呼び出し時に 
      libcuda.so.1
      を動的に開く。
    • ここからドライバーのユーザー空間側へのアクセス経路が確保される。

  3. 起動シーケンス:

    • コンパイラで生成された登録コード（コンストラクタ）が実行され、ファットバイナリが登録される。
    • cuLaunchKernel
      が呼ばれ、起動構成 (QMD) が構築される。
    • CUDA 12.2 以降: デフォルトは「怠慢 (Lazy)」読み込み（ロード時に PTX を JIT コンパイル）。

4. GPU メモリへのコード転送と通信経路 (
GPFIFO
/ ドアベル)

GPU は PCIe バスの向こう側にあり、直接関数呼び出しを受け入れられません。ホストは**命令ストリーム（メソッド）**を書き込む必要があります。

コミュニケーションの仕組み

  • プッシュバッファ: ドライバーが GPU への命令を記述するホスト RAM 上の領域。
    • メソッド：GPU のネイティブコマンド（レジスタセット＋動作）。
    • QMD (Queue Meta Data): カーネルの起動 descriptor（グリッドサイズ、レジスタ数など）を格納。
  • GPFIFO (Graphics Pipeline FIFO):
    • CPU と GPU で共有されるポインタリングバッファ。
    • GP_PUT
      : ドライバーが書いたカーソル（新しい仕事の追加）。
    • GP_GET
      : GPU が消費したカーソル。

起動トリガーの進化

  • 古式 CUDA (Turing 以前): 
    USERD
    スクロールによって直接ポーティング。
  • 現行 CUDA (Turing / Ada):
    • CPU のホストエンジンは GPFIFO を常時監視する機能がない（または無効化）。
    • ドライバーはMMIO ドアベルを使用。
    • ドライバーが 
      GP_PUT
      を更新し、特定のレジスタ領域に書き込むと、GPU はその変化を検知して起動処理を開始する。

5. 指令ごとの進行 (SM 内での分散)

QMD が計算ワークディストリビュータ（GigaThread エンジン）に渡されると、以下のプロセスが開始されます。

VRAM × 128 SMs | 指令ストリーム

  • コンパイル済みの SASS は VRAM に単一の線形リストとして存在する。
  • 計算ワークディストリビュータ: この長列を 128 個の SM（Multiprocessor）に分布。

SM (マルチプロセッサ) の内部構造

  • RTX 4090 は物理的な 144 SM を持つが、効率的な動作のために16 SM が無効化されている（実質 128 SM）。
  • スループット制限:
    • レジスタ容量: 各 SM は最大 1,536 スレッド（48 ワープロ）を保持。
    • 割り当て: この例のカーネルは「ブロックあたり 256 スレッド」を使用するため、1 SM あたり**6 ブロック（48 ワープロ）**しか稼働させられない。
  • スケジューリング:
    • 1 SM は 4 つのサブパーティションに分かれ、それぞれが独立した実行パイプラインを維持する。
    • サイクルごとに最大 1 つの命令を発行可能（スロット制約）。

ワープロの有効化とスケジューリング制御

PTX の冗長さを除去し、ハードウェア最適化のために 

ptxas
が以下の制御ビットを埋め込んだ SASS を生成します。

ディレクティブ役割効果
静的停止数 (Stall)固定遅延の推測ALU による計算中は、命令を発行せずにスレッドを駐車（パース）。
降伏ヒント (Yield)優先度制御このワープロを待機させるべきか否かの指示。ボトルネック時は他のワープロに処理を譲る。
依存関係バリアデータ同期グローバルメモリ読み書き（
LDG
, 
STG
）による遅延の同期管理。

これにより、GPU は並列計算による潜伏遅延を隠蔽し、効率的なスループットを実現します。

6. メモリ階層とデータ転送 (Coalescing)

カーネルは浮動小数点数（4 バイト）の配列を操作します。

ロード/ストアの合体 (Coalescing)

  • 要求: ワープロ内の 32 スレッドがそれぞれ 1 バイトアクセスすると、計 128 バイトの連続ブロックとして処理される (
    32 * 4 bytes
    )。
  • SM ロードユニット: この連続パターンを検出し、4 つの「扇区リクエスト」に結合。
  • メリット: メモリ帯域幅を最大化。単一トランザクションでデータを読み取れる。

データパス (Data Path)

  1. L1 データキャッシュ: 最優先でチェック（ミスする場合のみ次のステップへ）。
  2. L2 キャッシュ: 72 MB で、全 SM に共有。
  3. VRAM (GDDR6X): L2 ミス時にアクセス。RTX 4090 は HBMではなく GDDR6X を使用。
  • 読み込み: 入力データ（12 バイト/要素）が DRAM から読み込まれる。
  • 計算: 
    FADD
    により加算が行われる。
  • 書き出し (ストア): 結果 
    STG.E
    は、同じパスを経由して L2 を経由するか、あるいは L2 にキャッシュされる（場合によっては VRAM とは直接やり取りしない）。

パフォーマンス計測 (
ncu
)

$ ncu --metrics launch__grid_size, sm__warps_active.avg.pct_of_peak, \
    dram__throughput.avg.pct_of_peak, gpu__time_duration.sum ./vadd
  • sm__warps_active: 82.77%（充分な計算リソース確保）。
  • dram__throughput: 79.65%（メモリ帯域使用率が高いが、入力転送のみで計算量が少ないため）。
  • 推定スループット: $780 \text{ GB/s}$ 附近。

7. CPU への戻り (結果の同期)

GPU は非同期に計算するが、ホストは最終結果が必要だ。

  1. 完了セマフォ: グリッドの最後のブロックが終了したことを示すシグナル（QMD の Fence）。
  2. DMA コピー: 
    cudaMemcpy
    は GPU のコピーエンジンによって実行され、セマフォによるゲート制御を受ける。
  3. データ転送: 結果が L2 キャッシュにある場合、DRAM を経由せず PCIe バスのみで CPU メモリへ移動する（高速）。

最終的にホスト関数がブロックされ、

printf
によって計算結果が表示される：

c[0]=2.000000 c[n-1]=2.000000

8. 補遺: ランチの詳細確認方法 (診断ツール)

より詳細な動作を確認するための、いくつかの技術的なアプローチ（

LD_PRELOAD
, 
ioctl
など）が存在します。

インターポジションフック (Memory Inspection)

ドライバーが GPU メモリ領域をマップする際 (

/dev/nvidia*
) を監視し、プッシュバッファの内容を読み取る。これにより、QMD やメソッドバーストを確認できる。

# プッシュバッファの内容を出力するシムプログラム
LD_PRELOAD=./shim.so ./vadd

起動コマンドの解析

  • プッシュバッファコマント: 
    SET_INLINE_QMD_ADDRESS_A
    (QMD の格納場所) と 
    LOAD_INLINE_QMD_DATA
    (データ読み込み)。
    • Ampere/Ada ヘッダー (
      clc9c0.h
      ) を参照し、オフセットを解析することで起動パラメータを取得可能。

ドライバー通信 (
ioctl
)

libcuda
の内部処理を確認するため、
strace
で 
nv_escape.h
に定義されたコマンド（
NV_ESC_RM_CONTROL
など）を追跡する。

$ strace -f -e trace=ioctl ./vadd
# /dev/nvidiactl, /dev/nvidia-uvm に対するコマンド確認

これらを通じて、コンパイル済みのバイナリ内の「暗黙の契約」や、ドライバーの動作メカニズムを深く理解できる。

同じ日のほかのニュース

一覧に戻る →

2026/06/30 4:49

/.self: ホスト環境を構築することを支援する新しいトップレベルドメイン

## 日本語訳: 本件の核心となるメッセージは、ユーザーのデータや注意を搾取する既存のモデルを捨て、倫理的な新アーキテクチャへとインターネットを変革する呼びかけです。Human-Centered Computing Foundation は、ICANN の Applicant Support Program を通じてこのイニシアチブを正式に開始し、その主な目標として、倫理的技術にのみ専属 reserved されるトップレベルドメイン（TLD）の確保を目指しています。この動きは、人間の行動から価値を抽出するという業界の確立されたダイナミクスに直接挑戦し、代わりに人間中心の価値に基づいたシステムを提案しています。 もしこの新しいドメイン拡張を取得することに成功すれば、同財団はユーザーエシクティクスをデータマイニングよりも優先するプロジェクトのみがホストされる特定のデジタル空間を作成します。この転換は大きな利益をもたらすと約束しており、個人は企業の監視ではなく自らの道徳的原則を中心に設計された Web 環境を航行することができます。企業にとっては、持続的な成功には単に注意を採取するのではなく、真の人間のニーズを満たすアーキテクチャが不可欠になる、避けられない未来を示しています。最終的に、このキャンペーンは、技術が人々を利用するために操作するのではなく、人々をサービスするためのセクターとして、誠実さを定義されたインターネットの別個の分野を確立することを目指しています。

2026/06/30 2:05

Qwen 3.6 27B はローカル開発のsweet spot（最適解）です。

## Japanese Translation: 本文は、ローカルコード生成のために Qwen 3.6 27B デンスモデルを優先すること advises（推奨）しています。これは、指示追従の精度と効率的なパフォーマンスのバランスが取れており、Node パッケージの作成といった特定のタスクで失敗する可能性があるように 35B の A3B mixture-of-experts などのより大きなバリエーションを上回る場合があるためです。ベンチマークによると、このモデルは消費者向けハードウェア上で効率的に動作しながら、2025 年の中盤の GPT-5 程度の知能レベルに達します。Apple M5 チップ（共有 RAM を最大 48 GB 使用）では約 30 トokens/秒、量子化された状態で高級な Nvidia RTX 5090 カードでは 50 トokens/秒 にスケールします。重要なのは、著者が倫理的かつ技術的な理由から、Ollama ではなく `llama-server` または `llama-cli` を使用して Hugging Face の量子化版（例：`unsloth/Qwen3.6-27B-MTP-GGUF:Q8_0`）でモデルを実行することを推奨している点です。この構成により、開発者は OpenCode エージェントなどのツールと互換性のあるセキュアな「vibe coding」環境を構築できます。ローカルでモデルを実行することは、データのプライバシーを維持し、機密情報が外部の米中クラウドプロバイダーに漏洩することなく、オフラインでの作業をサポートするために不可欠です。将来的にはツールの呼び出しを通じて事実知識と生粋の知能を分ける傾向があるかもしれませんが、この即席のソリューションは品質を損なうことなく、個人および小規模チームの開発者にとってアクセス可能な入門点を提供します。より大きなモデルが将来的にはエンタープライズレベルのハードウェアを必要とするでしょうが、27B バリエーションは現在、標準的な消費者向けハードウェア上で DeepSeek-V4 Flash などのフロンティア代替案と比較できる堅牢でプライベートな AI 機能を 제공합니다（提供しています）

2026/06/28 0:05

アイコンを解放せよ

## Japanese Translation: 2026 年 6 月 26 日付の投稿で、Paul Kafasis は、macOS 26「Tahoe」がすべてのアプリアイコンに対して義務付けられた統一された「squircle」形状を導入し、ファーストパーティアイコンをボヤけた「Liquid Glass」 appearances に変更したと報告している。多くの人にとってこれはデザインと使いやすさにおける重大な後退だと見られている。サードパーティ製アイコンをこの指定された squircle 形状に強制することで、ユーザーが迅速な識別のために頼りにしていた多様な形状はなくなり、色が主な識別基準になった——特に色覚障害を持つユーザーや類似の色を持つアプリを区別する際には深刻な問題となった。コンプライアンスに反するサードパーティ製アイコンは縮小され、魅力的でない灰色の背景上に表示され、「icon jail」シナリオが引き起こされたほか、Apple の新しい「Clear」と「Tinted」アイコンスタイルは採用率が低かった。これは統一された squircle により識別がほぼ不可能になりつつあったためである。内部的なフィードバックチケット（FB23388490）でこれらの制限への異議が申し立てられたにもかかわらず、macOS 27「Golden Gate」の初期ベータ版では余計な「Liquid Glass」を取り除き、シャープなデザインを復活させ、Automator などのファーストパーティアイコンを見直し、部分的な改善が見られる。Kafasis は、Apple がサードパーティ製アプリに対して単一の squircle 形状を強制することをやめ、多様なアイコン形状を許可してアクセシビリティ、創造性、および総合的な使いやすさを向上させることを求めつつある。