# CUDA ## 定義 CUDA(Compute Unified Device Architecture)は、[[NVIDIA]] の GPU を汎用並列処理に使うためのプラットフォーム・プログラミングモデルである。ホスト(CPU)がデバイス(GPU)へメモリを割り当て、データを転送し、カーネル(kernel)と呼ばれる並列関数を多数のスレッドで起動し、結果をホストへ返すという実行モデルが基本となる。(Source: [[@2026__eunomia.dev__CUDA Events - eBPF-based CUDA API Tracing]]) ## CUDA のスレッド実行モデル CUDAはスレッド・スレッドブロック(cooperative thread array; CTA)・グリッドの3階層で並列作業を構造化する。スレッドはスレッドブロック(最大1,024スレッド)にまとめられ、スレッドブロックはグリッドを構成する。GPUの各SM(streaming multiprocessor)はワープ(32スレッド)単位でSIMT(Single Instruction, Multiple Threads)命令をロックステップ実行する。カーネルは`__global__`関数として記述し、`<<<blocksPerGrid, threadsPerBlock>>>`のchevron構文で起動する。ブロックサイズはワープサイズ32の倍数を選ぶのが基本原則で、Blackwell世代では256〜512スレッド/ブロックが占有率([[GPU占有率(Occupancy)]])とレジスタ・共有メモリ制約のバランス点として推奨される。(Source: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 6 GPU Architecture, CUDA Programming, and Maximizing Occupancy]] §Threads, Warps, Blocks, and Grids, §Choosing Threads-per-Block and Blocks-per-Grid Sizes) ## 主要な CUDA ランタイム API - `cudaMalloc` / `cudaFree` — デバイスメモリの割当と解放 - `cudaMemcpy` — ホスト・デバイス間・デバイス内のデータ転送 - `cudaLaunchKernel` — GPU カーネルの起動 - `cudaStreamCreate` / `cudaStreamSynchronize` — 非同期実行ストリームの管理と同期 - `cudaEventCreate` / `cudaEventRecord` / `cudaEventSynchronize` — イベントによるタイミング計測 - `cudaGetDevice` / `cudaSetDevice` — GPU デバイスの選択 - `cudaMallocAsync` / `cudaFreeAsync` — ストリーム順序保証つきの非同期メモリプール確保・解放。同期版の`cudaMalloc`/`cudaFree`はOSレベルの`mmap`/`ioctl`を伴い高コストなため、頻繁な細粒度確保・解放にはこちらが推奨される(Source: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 6 GPU Architecture, CUDA Programming, and Maximizing Occupancy]] §Asynchronous Memory Allocation and Memory Pools) - `cudaMallocManaged` / `cudaMemPrefetchAsync` / `cudaMemAdvise` — Unified Memory(CPU/GPU単一アドレス空間)の割当・プリフェッチ・配置ヒント ## 横断的知見 - **CUDA API は GPU 利用の「入口」として観測可能**: CUDA ランタイムライブラリ `libcudart.so` の関数呼び出しを eBPF uprobe で傍受することで、ソース改変なしに GPU 利用のマクロ視図を得られる。ただしこれは CPU 側のリクエスト単位であり、カーネル内部のスレッド動作までは届かない。(Source: [[@2026__eunomia.dev__CUDA Events - eBPF-based CUDA API Tracing]]) - **非同期実行がトレースとデバッグを難しくする**: CUDA API 呼び出しは多くの場合非同期であり、CPU が GPU への作業投入後も処理を継続する。これが従来の逐次デバッグツールでは追えない境界を生んでいる。(Source: [[@2026__eunomia.dev__CUDA Events - eBPF-based CUDA API Tracing]]) - **非同期性はAPIトレースの観測課題であると同時に、性能エンジニアリングの積極的な設計原則でもある**: eBPFベースのCUDA API トレース研究(§定義)は非同期実行を「追跡しにくい境界」として扱うのに対し、`AI Systems Performance Engineering`(ch.6)は同じ非同期性を`cudaMallocAsync`・`cudaMemcpyAsync`・非ブロッキングストリームという形で積極的に活用し、メモリ確保・転送・カーネル実行をオーバーラップさせることでスループットを最大化する。両者を合わせると、CUDAの非同期モデルは「観測を難しくする副作用」と「性能を引き出す主要機構」という二面性を持つことが分かる。(Source: [[@2026__eunomia.dev__CUDA Events - eBPF-based CUDA API Tracing]], [[@2025__OReilly__AI Systems Performance Engineering - Chapter 6 GPU Architecture, CUDA Programming, and Maximizing Occupancy]] §Asynchronous Memory Allocation and Memory Pools) - **エラー通知も非同期・遅延的である**: CUDAカーネルはスレッドごとの例外を持たず、不正なメモリアクセス等はグローバルなフォールトフラグをセットするのみで、ホストが次に同期API(`cudaGetLastError()`・`cudaDeviceSynchronize()`)を呼ぶまでエラーは表面化しない。これはCUDA API トレースがCPU側のAPI呼び出しを傍受するだけでは、実際の障害発生タイミングを正確に捉えられない可能性を示唆する。(Source: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 6 GPU Architecture, CUDA Programming, and Maximizing Occupancy]] §CUDA Programming Refresher) - **CUDAの非同期性・並行性モデルは、ストリーム間(ch.6)からカーネル内のwarp・ブロック単位(ch.10)へと階層的に深化する**。ch.6が`cudaMallocAsync`・非ブロッキングストリームによるカーネル外(inter-kernel)の重ね合わせを扱うのに対し、ch.10の CUDA Pipeline API(`producer_acquire`/`consumer_wait`等)・warp specialization([[Warp Specialization]])・thread block cluster([[Thread Block Cluster]])は同一カーネル内(intra-kernel)でのメモリ・計算の重ね合わせを扱う。両者は「非同期実行によるレイテンシ隠蔽」という同じ原則を、異なる粒度(カーネル間 対 カーネル内のwarp/ブロック間)で適用したものである。(Source: [[@2025__OReilly__AI Systems Performance Engineering - Chapter 6 GPU Architecture, CUDA Programming, and Maximizing Occupancy]], [[@2025__OReilly__AI Systems Performance Engineering - Chapter 10 Intra-Kernel Pipelining, Warp Specialization, and Cooperative Thread Block Clusters]]) ## 未解決の問い - CUDA API トレースで得られた enter/exit イベントを、分散トレーシングのスパンや ML フレームワークの演算子イベントと時系列で相関させるにはどのような識別子・タイムスタンプ同期が必要か。 - CPU 側 API トレースの情報だけで、GPU 内部のメモリアクセスパターンや分岐ダイバージェンスをどこまで推定できるか。 - 非同期カーネル起動のエラーがホスト側の同期タイミングまで遅延することを踏まえると、CUDA API トレース(eBPFベース)は`cudaGetLastError()`呼び出しのタイミングと実際のフォールト発生タイミングのズレをどう補正して障害箇所特定に使えるか。 ## 関連 - 概念: [[CUDA API トレース]] / [[GPU観測性]] / [[eBPF]] / [[動的計装]] / [[GPU占有率(Occupancy)]] / [[メモリ階層とキャッシュ]] / [[Rooflineモデル]] / [[Warp Specialization]] / [[Thread Block Cluster]] / [[Persistent Kernel]] - エンティティ: [[NVIDIA]] / [[eunomia-bpf]] - ソース: [[@2026__eunomia.dev__CUDA Events - eBPF-based CUDA API Tracing]] / [[@2025__OReilly__AI Systems Performance Engineering - Chapter 6 GPU Architecture, CUDA Programming, and Maximizing Occupancy]] / [[@2025__OReilly__AI Systems Performance Engineering - Chapter 10 Intra-Kernel Pipelining, Warp Specialization, and Cooperative Thread Block Clusters]]