gpu-memory-model

GPU メモリモデル

目的

カーネルを最適化するエージェント向けに、GPU の実行およびメモリモデルを説明します。SIMT 実行、ワープ（32）とウェーブフロント（64）の分岐コスト、グローバルメモリのコアレス化ルール、共有メモリのバンク競合、L1/L2 キャッシュの動作、アトミック操作のメモリ順序、占有率とレイテンシ隠蔽のトレードオフを扱います。

使用するタイミング

• 理論上の帯域幅が高いにもかかわらず、カーネルがメモリバウンドになる理由を診断するとき
• 分岐によるワープ分岐を理解するとき
• タイル化アルゴリズムの共有メモリにおけるバンク競合を解消するとき
• 占有率とレジスター負荷のバランスを考慮してブロックサイズを選ぶとき
• NVIDIA（ワープ 32）と AMD（ウェーブフロント 64）の間でカーネルを移植するとき
• 並列リダクションにおけるアトミック競合を検討するとき

ワークフロー

1. SIMT 実行モデル

GPU hardware
├── Device
│   └── SM / CU (Streaming Multiprocessor / Compute Unit)
│       ├── Warp schedulers (NVIDIA) or Wavefront schedulers (AMD)
│       │   └── Warp/Wavefront (32 or 64 threads in lockstep)
│       ├── Register file (partitioned per thread)
│       ├── Shared memory / LDS (per SM)
│       └── L1 cache (often shared with shared memory)
└── L2 cache (device-wide) → DRAM/HBM

SIMT（Single Instruction, Multiple Threads）：1 つの命令ストリームがワープ／ウェーブフロントを駆動します。各スレッドは独自のレジスターとスレッド ID を持ちますが、同じ命令をロックステップで実行します。

2. ワープとウェーブフロントの比較

ベンダー	ユニットサイズ	名称
NVIDIA	32 スレッド	ワープ
AMD	64 スレッド	ウェーブフロント

影響：

• リダクションツリー：NVIDIA は 16→8→4→2→1、AMD は 32→16→8→4→2→1 の順に半減します
• ブロックサイズ：32（NVIDIA）または 64（AMD）の倍数を推奨します
• 占有率カウンターは、SM あたりのアクティブなワープ／ウェーブフロント数を報告します

3. ワープ分岐のコストモデル

ワープ内のスレッドが異なる分岐を取ると、ハードウェアは実行パスをシリアル化します。

// Divergent: half warp does A, half does B → 2x instruction issue
if (threadIdx.x % 2 == 0) {
    result = expensive_a(data[idx]);
} else {
    result = expensive_b(data[idx]);
}

// Non-divergent: all threads same path
result = expensive_a(data[idx]);

軽減策：

• 述語実行：両方を計算し、?: で選択します（均一性のために計算量を増やします）
• 異なるコードパスごとにカーネルを分ける
• ブロックレベルのデータのみで分岐する（ワープ内で均一）
• スレッドごとに結果がばらつく if (data[i] < threshold) の代わりに、ビンをループ処理する

分岐コスト ≈ 実行されたパスの合計（最大値ではありません）。

4. グローバルメモリのコアレス化

NVIDIA のコアレス化ルール（簡略版）：ワープ内のスレッドが連続した 4 バイトワードにアクセスすると、128 バイトのトランザクション 1 回で処理されます。

// Coalesced: consecutive threads → consecutive addresses
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float val = data[idx];

// Uncoalesced: stride access
float val = data[threadIdx.x * stride];  // stride > 1

// Partially coalesced: misaligned start
float val = data[base + threadIdx.x * 3];

AoS と SoA の影響：

// AoS — poor coalescing when reading one field
struct Particle { float x, y, z; };
float x = particles[i].x;  // threads read with stride 3

// SoA — coalesced
float x = pos_x[i];

5. 共有メモリのバンク競合

共有メモリは 32 個のバンク（4 バイトワード）に分かれています。同じバンク内の異なるアドレスへの同時アクセスはシリアル化されます。

__shared__ float tile[32][32];

// Bank conflict: all threads access tile[threadIdx.x][0]
// 32 threads, 32 banks, but column 0 → same bank per row offset
float val = tile[threadIdx.x][0];

// Fix: pad columns to break bank alignment
__shared__ float tile[32][33];  // +1 padding

検出方法：Nsight Compute の l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum、または NCU の共有メモリ読み込み競合メトリクスを確認します。

6. L1/L2 キャッシュの動作

レベル	範囲	備考
L1	SM ごと	共有メモリと統合されていることが多く、分割を設定できます
L2	デバイス全体	キャッシュラインは通常 128 バイトです
テクスチャ／L1 読み取り専用	SM ごと	アクセスが均一な場合に使われる、キャッシュされた読み取り専用経路です

キャッシュに適したパターン：

• 空間局所性：連続したスレッドが連続したメモリにアクセスします
• 時間局所性：グローバルメモリから再取得する前に、共有メモリでデータを再利用します
• ランダムな散在アクセスを避ける：アトミック更新やポインター追跡はキャッシュの効果を損ないます

// Cache-friendly tile load
for (int t = 0; t < num_tiles; t++) {
    __shared__ float smem[TILE][TILE];
    smem[ty][tx] = global[row * N + t * TILE + tx];
    __syncthreads();
    // compute from smem — L1/L2 only hit on first load per tile
}

7. アトミック操作とメモリ順序

GPU のアトミック操作（atomicAdd、atomicCAS、atomicExch）は、同じアドレスを対象とするスレッド間で逐次一貫性を保証しますが、競合が激しいと実行がシリアル化されます。

// Bad: all threads atomic to one counter
atomicAdd(&global_sum, local_val);

// Better: per-block reduction, one atomic per block
__shared__ float block_sum;
// ... warp reduce to block_sum ...
if (threadIdx.x == 0)
    atomicAdd(&global_sum, block_sum);

HIP/CUDA のメモリフェンス：

__threadfence_block();  // visible to threads in same block
__threadfence();        // visible to all threads on device
__threadfence_system(); // visible to host (expensive)

8. 占有率とレイテンシ隠蔽

Occupancy tradeoff
├── High occupancy → more warps to hide memory latency
│   └── Costs: fewer registers/SM, less shared mem per block
└── Low occupancy + high ILP → enough independent instructions per warp
    └── Works for compute-bound kernels with deep pipelines

判断フローチャート：

Memory-bound kernel?
├── Yes → maximize active warps (occupancy), coalesce, tile with shared mem
└── No (compute-bound) → may lower occupancy if registers enable more ILP

# Measure achieved occupancy
ncu --metrics sm__warps_active.avg.pct_of_peak_sustained_active ./app

経験則：メモリバウンドのカーネルでは占有率 50% 以上が必要です。十分な命令レベル並列性があれば、コンピュートバウンドのカーネルは占有率 25% でも良好に動作する場合があります。

9. クイックリファレンス表

症状	考えられる原因	確認方法
DRAM スループットが低い	コアレス化されていないアクセス	NCU のメモリワークロード分析
共有メモリ読み込みのストール	バンク競合	配列にパディングを追加する、アクセスパターンを変更する
barrier によるストールが多い	__syncthreads の欠落、またはバリアでの分岐	synccheck
アトミック操作がボトルネック	競合するアトミック操作が多すぎる	階層的リダクション
占有率が低い	レジスター／共有メモリの負荷	launch__occupancy_limit_* メトリクス

よくある問題

症状	原因	対処法
SoA→AoS への変更後、10 倍遅くなる	ストライドアクセスによる合体アクセスが崩れている	ホットフィールドは SoA レイアウトに維持します
タイル化した行列乗算が素朴な実装より遅い	共有メモリのタイルでバンク競合が発生している	共有配列の列にパディングを追加します
コードは同じなのに NVIDIA と AMD で性能が異なる	Warp は 32、wavefront は 64	ブロックサイズとリダクションを再調整します
アトミック操作は正しいが遅い	グローバルメモリで競合している	まずブロック単位でリダクションします
L2 のヒット率は高いのに遅い	L2 の帯域幅が飽和している	移動する総バイト数を減らします

関連スキル

• skills/gpu/cuda — このモデルを使用した実践的なカーネルパターン
• skills/gpu/cuda-profiling — メモリアクセスの合体と占有率を検証するメトリクス
• skills/gpu/hip-rocm — AMD の wavefront 固有のチューニング
• skills/low-level-programming/cpu-cache-opt — CPU キャッシュの概念（類似）
• skills/low-level-programming/simd-intrinsics — CPU 側のベクトル幅
• skills/profilers/hardware-counters — キャッシュミスを汎用的に計測する方法