✍️ 필사 모드: CUDA GPUプログラミングモデル Deep Dive — SIMT、メモリ階層、Tensor Core、カーネル最適化完全ガイド (2025)

TL;DR

• CUDAはNVIDIAのGPUプログラミングプラットフォーム。2007年以来、AI/HPCの事実上の標準。ChatGPTやStable Diffusionなど主要AIワークロードはCUDA上で動く。
• SIMT (Single Instruction, Multiple Threads): プログラマはスカラーコードを書き、ハードウェアが32 threads = 1 warpにまとめて並列実行。SIMDの効率とスカラープログラミングの生産性を両立。
• スレッド階層: Grid, Block, Warp, Thread。各Blockは同じStreaming Multiprocessor (SM)上で実行、Warp単位スケジューリング。
• メモリ階層: Register (最速), Shared Memory (L1サイズ), L2 Cache, Global Memory (遅いが大容量)。各レベルで数十〜数百倍差。
• Memory Coalescing: Warpの32スレッドが連続アドレスにアクセスすれば1 transaction。ランダムアクセスは32 transactions。32倍の帯域差 — CUDA最適化の第1位。
• Warp Divergence: Warp内スレッドが異なるパスを取ると順次実行され並列性を失う。
• Tensor Core (Volta, 2017): 行列積を単一命令で。FP16基準125 TFLOPS。AI加速の鍵。
• 現代の流れ: Triton、cutlass、FlashAttentionといった高レベル抽象化が手動チューニングのCUDAを置き換えつつある。
• 代替: ROCm/HIP (AMD), SYCL/oneAPI (Intel), Metal (Apple)。エコシステム格差でCUDAが依然支配。

1. なぜGPUがAIを支配するのか

1.1 CPU vs GPU

CPU: 「少数の賢いコアが複雑な仕事を」。GPU: 「多数の単純コアが同じ仕事を並列で」。

1.2 AIワークロードの特性

Transformer推論1回:

MatMul x 数千回
Element-wise 演算 (GELU、LayerNorm)
Attention (QK^T, softmax, @V)

いずれも大規模な行列演算。数十億のMAC (multiply-accumulate) 演算。CPUは100-1000 GFLOPS、GPUは100,000+ TFLOPS (H100)。1000倍の差。

1.3 NVIDIAの独占

3つの要因:

1. CUDAエコシステム: 2007年以来蓄積されたライブラリ、ドキュメント、コミュニティ。
2. ハードウェア投資: Tensor Core, HBM, NVLink, NVSwitch。
3. Jensen Huangの長期ベット: 2012年のAlexNet以降AIに全振り。

AMD (MI300), Intel (Gaudi), Google (TPU), Apple (M-series), 中国 (Huawei Ascend) が競争中だが、2025年時点でCUDA互換性と成熟度でNVIDIAが独占。

2. GPUハードウェアアーキテクチャ

2.1 Streaming Multiprocessor (SM)

GPUの基本単位はSM。各SMは独立した実行ユニット。

H100:   132 SM
A100:   108 SM
RTX 4090: 128 SM

SM内部:

+---------------------------------+
|           SM                    |
|  +---------+  +---------+       |
|  | Warp    |  | Warp    | x 4   |
|  | Sched   |  | Sched   |       |
|  +---------+  +---------+       |
|                                 |
|  CUDA Cores (INT32, FP32, FP64) |
|  Tensor Cores                   |
|  Special Function Units         |
|  Load/Store Units               |
|                                 |
|  Register File (数万)           |
|  Shared Memory / L1 (数百KB)    |
+---------------------------------+

2.2 Warp

Warp = 32スレッドがlockstepで実行されるグループ。NVIDIA GPUの根本単位。A100 SMは最大64 warps (= 2048 threads) をアクティブに保持。Warp schedulerが毎cycle実行可能なwarpを選び、レイテンシを隠蔽。

2.3 SIMT — プログラミングモデル

SIMT (Single Instruction, Multiple Threads): プログラマは1スレッドのようにコードを書き、ハードウェアが32スレッドをまとめて同じ命令を異なるデータに発行。

SIMDとの違い:

2.4 GPU全体構造

Host (CPU + DRAM)
     | PCIe / NVLink
GPU
+-- L2 Cache (数十MB、全体共有)
+-- Global Memory (HBM、数十〜数百GB)
+-- SM x N
    +-- L1 / Shared Memory
    +-- Register File
    +-- Cores

Global Memory (HBM): GPU全体共有、40-192 GB、約500 cycle遅延、ただし帯域は巨大 (H100で3 TB/s)。L2 Cache: 全SM共有、数十MB。Shared Memory: SM内、プログラマが明示管理。Register: 最速、スレッドごとにローカル。

2.5 階層別アクセス時間

A100概算:

Register:       1 cycle
Shared Memory:  ~20 cycles
L1 cache:       ~30 cycles
L2 cache:       ~200 cycles
HBM:            ~500 cycles

ホストDRAM (PCIe): ~10,000+ cycles。データlocalityが性能の全て。

3. CUDAプログラミングモデル

3.1 Hello World

#include <stdio.h>

__global__ void hello_kernel() {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    printf("Hello from thread %d\n", tid);
}

int main() {
    hello_kernel<<<2, 32>>>();  // 2 blocks x 32 threads = 64 threads
    cudaDeviceSynchronize();
    return 0;
}

• __global__: GPUで実行されるカーネル関数。
• <<<grid, block>>>: 実行構成。
• threadIdx.x, blockIdx.x, blockDim.x: インデックス/サイズintrinsic。

3.2 スレッド階層

GridはBlockを含み、BlockはThreadを含む。同じBlock内のスレッドは同じSMで実行しShared Memoryを共有。2D/3Dも可:

dim3 grid(16, 16);
dim3 block(32, 32);
kernel<<<grid, block>>>();
// 合計 16x16x32x32 = 262,144 threads

3.3 グローバルインデックス

__global__ void add_kernel(float *a, float *b, float *c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) c[idx] = a[idx] + b[idx];
}

3.4 メモリ管理

float *d_a;
cudaMalloc(&d_a, n * sizeof(float));
cudaMemcpy(d_a, h_a, n * sizeof(float), cudaMemcpyHostToDevice);
kernel<<<grid, block>>>(d_a);
cudaMemcpy(h_a, d_a, n * sizeof(float), cudaMemcpyDeviceToHost);
cudaFree(d_a);

PCIeコピーは高コスト。最小化が原則。

3.5 Unified Memory

float *data;
cudaMallocManaged(&data, n * sizeof(float));

ランタイムがページフォルトを検知し自動でhost/device間転送。便利だが手動管理より若干遅い。

4. メモリ階層詳細

4.1 Global Memory

最大 (40-192 GB)、GPU全体共有、遅い (数百cycle)。

4.2 Shared Memory

SM内 (L1と同じ物理)。非常に高速。Block内スレッドで共有。サイズ制限 (~48-228 KB per SM)。

__global__ void kernel() {
    __shared__ float shared_data[256];
    shared_data[threadIdx.x] = threadIdx.x;
    __syncthreads();
    // 全スレッドがshared_dataを読める
}

複数スレッドが同じデータを繰り返しアクセスする場合に使う。

4.3 Register

最速 (1 cycle)、スレッドローカル、コンパイラが自動割当。スレッドあたりレジスタを使いすぎるとregister spill でGlobal Memoryに溢れ、遅くなる。

4.4 Constant Memory

__constant__ float coefficients[256];

64 KB、read-only、warp全スレッドが同じ値を読むbroadcastで高速。

4.5 Texture Memory

元々グラフィックス用。2D空間局所性最適化と補間ハードウェアを持つ。現代では __ldg() で部分代替。

5. Memory Coalescing — 最優先最適化

5.1 原理

Coalesced: 32スレッドが連続128バイトにアクセス → 1 memory transaction。

Uncoalesced: 32スレッドが散らばったアドレス → 32 transactions、各128バイト中4バイトしか使わない。

帯域差は32倍。GPU実性能の成否はここで決まる。

5.2 例

BAD:

__global__ void transpose_bad(float *in, float *out, int N) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    out[x * N + y] = in[y * N + x];  // stride N
}

GOOD (Shared Memoryを中間に):

__global__ void transpose_good(float *in, float *out, int N) {
    __shared__ float tile[32][33];  // +1でbank conflict回避
    int x = blockIdx.x * 32 + threadIdx.x;
    int y = blockIdx.y * 32 + threadIdx.y;
    tile[threadIdx.y][threadIdx.x] = in[y * N + x];
    __syncthreads();
    x = blockIdx.y * 32 + threadIdx.x;
    y = blockIdx.x * 32 + threadIdx.y;
    out[y * N + x] = tile[threadIdx.x][threadIdx.y];
}

20-30倍の性能差。

5.3 パターン

「Thread i が array[i] にアクセス」。これがCoalescingの基本。

6. Warp Divergence

Warp内の32スレッドが異なるパスを取ると?

if (threadIdx.x < 16) compute_a();
else compute_b();

ハードウェアは両方のパスを順次実行: まずthreads 0-15が compute_a()、16-31はidle。次に16-31が compute_b()、0-15がidle。並列性半減。

緩和:

1. Warp境界で分岐 (blockIdx.x % 2)。
2. 分岐の代わりに算術 (fabs, 三項演算子)。
3. 両パスを計算してselect。

Volta (2017) 以降のIndependent Thread Schedulingでより柔軟になったが、Divergence最小化が原則。

7. Shared Memory Bank Conflict

Shared Memoryは32 bankに分かれ、各bankは32-bit word単位。Warpの32スレッドが異なるbankにアクセス → 1 cycle。同じbankの異なるword → N-way conflict → N cycle。

__shared__ float data[32][32];
float y = data[0][threadIdx.x];  // 全スレッド → bank 0 → 32-way conflict

解決はパディング: __shared__ float data[32][33]; でcolumn accessが異なるbankに分散。

Broadcast (同じwordを複数スレッドが読む) はOK。

8. Occupancy

Occupancy = SMの最大warp数に対する現在アクティブwarp数の比。A100 SMは最大64 warps。

Warp switchingで遅延を隠蔽するため、高Occupancyは有利。制限要因:

1. スレッドあたりレジスタ数。
2. Blockあたり Shared Memory。
3. Blockあたりスレッド数。

__global__ __launch_bounds__(256, 4)
void kernel() { /* ... */ }

Occupancyが全てではない。Tensor Core利用コードは低Occupancyでも最高性能が出る。Nsight Computeで計測。

9. Tensor Core — AI加速の鍵

9.1 登場 (Volta, 2017)

V100 Tensor Coreは4x4 x 4x4行列積を1 cycleで実行: D = A x B + C。16 MACs per cycle (CUDA coreの1 MACに対して)。V100 FP16 = 125 TFLOPS (FP32の15 TFLOPSに対し)。

9.2 進化

• V100 (Volta, 2017): FP16 → FP32 matmul。
• A100 (Ampere, 2020): TF32, BF16, INT8, 2:4 sparsity。312 TFLOPS (BF16)。
• H100 (Hopper, 2022): FP8, Transformer Engine。2000 TFLOPS (FP8)。
• B100/B200 (Blackwell, 2024): FP4、20,000+ TFLOPS。

9.3 使い方

wmma APIで直接利用:

#include <mma.h>
using namespace nvcuda;

__global__ void matmul_wmma(half *a, half *b, float *c) {
    wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::row_major> a_frag;
    wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::col_major> b_frag;
    wmma::fragment<wmma::accumulator, 16, 16, 16, float> c_frag;
    wmma::fill_fragment(c_frag, 0.0f);
    wmma::load_matrix_sync(a_frag, a, 16);
    wmma::load_matrix_sync(b_frag, b, 16);
    wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);
    wmma::store_matrix_sync(c, c_frag, 16, wmma::mem_row_major);
}

実際にはcuBLAS/cuDNN/PyTorch/Tritonを介して利用するのが一般的。

9.4 Mixed Precision

Tensor Coreは低精度入力 + 高精度累算。A, B: FP16、C, D: FP32。メモリ半減、帯域半減、Tensor Core速度、数値安全性を両立。PyTorchの torch.cuda.amp.autocast() (AMP) が自動化。

10. Streamと非同期実行

CUDA Stream は順序付き操作のキュー。異なるStreamは並列実行可能。

cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
kernel1<<<grid, block, 0, stream1>>>(...);
kernel2<<<grid, block, 0, stream2>>>(...);

コピーと計算をオーバーラップ。非同期コピーはpinned memoryが必要:

cudaMallocHost(&h_data, n * sizeof(float));
cudaMemcpyAsync(d_data, h_data, size, cudaMemcpyHostToDevice, stream);

Eventで Stream間同期。

11. CUDA Graph — Launch Overhead除去

毎回のカーネル呼び出しには数十μsのlaunch overheadがある。Transformer推論のように小カーネルが数百個あるとoverheadが支配的に。

cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
kernel1<<<..., stream>>>(...);
kernel2<<<..., stream>>>(...);
cudaStreamEndCapture(stream, &graph);
cudaGraphInstantiate(&graphExec, graph, NULL, NULL, 0);
cudaGraphLaunch(graphExec, stream);

数百のlaunchが単一graph launchに圧縮。Transformer推論token generationで2-3倍高速化。PyTorch 2.0+のCUDAGraphsが自動化。

12. マルチGPUとNCCL

LLM訓練は単一GPUメモリを超える。パラダイム: Data Parallel、Model Parallel、Pipeline Parallel。いずれもGPU間通信が必要。

• NVLink: H100で900 GB/s。
• NVSwitch: GPU間fully-connected switch、DGX H100に搭載。
• PCIe: ~64 GB/s。

NCCLはAllReduce, AllGather, Broadcast, Reduce, AllToAllを提供。NVLink/PCIe/InfiniBand/Ethernetを自動選択。PyTorch torch.distributed のデフォルトbackend。

P2P memcpy (NVLink経由):

cudaDeviceEnablePeerAccess(dev1, 0);
cudaMemcpyPeer(dst, dev1, src, dev0, size);

GPUDirect RDMAで他ノードのGPUと直接通信 (CPUを介さない)。

13. 高レベル抽象化

• cuBLAS / cuDNN / cuSPARSE / cuSOLVER: NVIDIAの標準数学ライブラリ。
• Thrust: STL風C++ GPUライブラリ。
• CUB: reduction/scan/sortの最適化ビルディングブロック。
• cutlass: GEMMビルディングブロック (FlashAttentionが利用)。
• Triton: OpenAIのPython風カーネル言語。

Triton例:

import triton
import triton.language as tl

@triton.jit
def add_kernel(x_ptr, y_ptr, output_ptr, n_elements, BLOCK_SIZE: tl.constexpr):
    pid = tl.program_id(axis=0)
    block_start = pid * BLOCK_SIZE
    offsets = block_start + tl.arange(0, BLOCK_SIZE)
    mask = offsets < n_elements
    x = tl.load(x_ptr + offsets, mask=mask)
    y = tl.load(y_ptr + offsets, mask=mask)
    tl.store(output_ptr + offsets, x + y, mask=mask)

CUDAよりはるかに簡潔、内部でMLIR経由でコンパイル。2024+はAI研究者の主流。FlashAttention (Tri Dao) はTensor Core + Shared Memory + fused kernelでメモリ効率の良いattentionを実現、GPT/LLaMA訓練に必須。

14. プロファイリングと最適化

• Nsight Systems: パイプライン全体のタイムラインプロファイル (nsys profile --stats=true ./my_app)。
• Nsight Compute: 単一カーネルの深い分析 (Occupancy, Memory throughput, Warp efficiency, Roofline) (ncu --set full ./my_app)。
• PyTorch Profiler: フレームワークレベルの視点。

Rooflineでカーネルがcompute-boundかmemory-boundか判定。Compute-bound → Tensor Core活用とOccupancy。Memory-bound → Coalescing、Shared Memory、reuse。

15. 競合

• AMD ROCm / HIP: hipify でCUDAコードをほぼそのまま移植可。MI300はH100に近い仕様。エコシステムで遅れる。
• Intel oneAPI / SYCL: Khronos標準、ベンダー横断。DPC++はportableだが性能は不均一。
• Apple Metal: PyTorch MPS backendでmacOS GPU訓練可。M-seriesのunified memoryに有利なワークロードあり。
• WebGPU: ブラウザでGPUアクセス。WebGPU + Wasmで小型LLMが動く。
• Google TPU: Google自社ASIC、JAX/XLA統合、Transformer最適化。Google Cloud専用。

16. 学習ロードマップ

1. 基本: CUDA C++ Programming Guide、簡単なカーネル (vector add, matmul)、nvcc、cudaMemcpy。
2. 最適化: Coalescing、Shared Memory、Bank conflict、Nsight Compute。
3. 高度: Warp-level primitive (__shfl_sync)、Tensor Core (wmma)、CUDA Graph、マルチGPU (NCCL)。
4. エコシステム: cuBLAS/cuDNN、Triton、PyTorch extension、FlashAttention/cutlass読解。

書籍: PMPP (Hwu/Kirk)、CUDA by Example、Professional CUDA C Programming。オンライン: NVIDIA Developer Blog、GPU MODE、PMPP講義。

17. Cheat Sheet

+------------------------------------------+
|          CUDA Cheat Sheet                |
+------------------------------------------+
| ハードウェア:                              |
|   GPU = N x SM                            |
|   SM = Cores + Tensor Cores + RegFile     |
|   H100: 132 SM                            |
|                                           |
| 実行モデル (SIMT):                         |
|   Grid -> Block -> Warp (32) -> Thread    |
|   スカラーコード、HWが32スレッドをまとめる   |
|                                           |
| メモリ階層:                                |
|   Register (最速)                          |
|   Shared / L1 (数十cycles)                 |
|   L2 (~200 cycles)                        |
|   HBM (~500 cycles)                       |
|                                           |
| 最優先最適化: Memory Coalescing            |
|   Thread i -> array[i]                    |
|                                           |
| 落とし穴:                                  |
|   Warp divergence, bank conflict,         |
|   register spill, uncoalesced access      |
|                                           |
| Tensor Core: Volta+, FP16/BF16/FP8/FP4    |
| Streams: copy/compute オーバーラップ       |
| CUDA Graph: launch overhead除去           |
| マルチGPU: NVLink + NCCL + P2P + RDMA     |
| ライブラリ: cuBLAS/cuDNN/Thrust/cutlass/  |
|            Triton/FlashAttention          |
| ツール: nvcc, Nsight Systems/Compute      |
| 代替: ROCm/SYCL/Metal/WebGPU/TPU          |
+------------------------------------------+

18. クイズ

関連記事:

• 「Transformer Architecture Deep Dive」 — GPUで実際に動くもの。
• 「Diffusion Models Deep Dive」 — 別のGPU-heavyワークロード。
• 「RDMA & NCCL」 — マルチGPU通信。
• 「LLVM Compiler Infrastructure」 — TritonがMLIRに compile される背景。