# CUDA NVIDIAが開発したGPU向けプログラミング言語・実行環境(Compute Unified Device Architecture)。ホスト(システムプロセッサ)向けにC/C++、デバイス(GPU)向けにC/C++の方言を生成する。ベンダー非依存の類似言語としてOpenCLがある。NVIDIAは、マルチスレッディング・MIMD・SIMD・命令レベル並列性というGPUが持つあらゆる粒度の並列性の統一テーマとして**CUDA Thread**を採用し、このプログラミングモデルを**SIMT**(single instruction, multiple thread)と呼ぶ(Source: [[@2019__MorganKaufmann__Computer Architecture - A Quantitative Approach - Chapter 4 Data-Level Parallelism in Vector, SIMD, and GPU Architectures]] §4.4)。 ## プログラミングモデルの要点 - 関数修飾子 `__device__`/`__global__`(GPU側)と`__host__`(システムプロセッサ側)でコードの実行場所を区別する。 - `name<<<dimGrid, dimBlock>>>(...)` という拡張関数呼び出し構文でGrid(Thread Blockの集合)とThread Block(スレッドの集合)の次元を指定する。 - `blockIdx`(Thread Block識別子)・`threadIdx`(Thread Block内のスレッド識別子)・`blockDim`(Thread Blockあたりのスレッド数)がCUDA組み込みの識別子として提供される。 - Thread Blockは独立かつ任意の順序で実行可能であることが要求され、異なるThread Block間は直接通信できない(グローバルメモリのアトミック操作でのみ協調可能)。 - CUDA Threadは軽量だがPOSIX Threadとは異なり、任意のシステムコールを発行できない。 (Source: [[@2019__MorganKaufmann__Computer Architecture - A Quantitative Approach - Chapter 4 Data-Level Parallelism in Vector, SIMD, and GPU Architectures]] §4.4) ## ハードウェアとの関係 CUDAのGrid/Thread Blockという抽象は、ハードウェア側ではThread Block SchedulerとSIMD Thread Schedulerという2段のスケジューラによってmultithreaded SIMD Processor(NVIDIA公式名: Streaming Multiprocessor)へ動的に割り当てられる。CUDAプログラマは数千のCUDA Threadをプログラムしているつもりでも、実際には32スレッド単位のWarp(SIMD Thread)として多数のSIMDレーン上で実行される。効率的なCUDAコードを書くには、プログラミングモデルがMIMDのように見えても実際にはSIMD演算として振る舞うことを意識し、32スレッドを制御フロー上でまとめておくこと、隣接アドレスへのアクセスでAddress Coalescing Unitの恩恵を受けることが要求される(Source: 同章 §4.4)。 コンパイラのターゲットはNVIDIAの抽象命令セット[[PTX]]であり、PTXから実機命令への変換はロード時にソフトウェアで行われる。 ## 関連 - ソース: [[@2019__MorganKaufmann__Computer Architecture - A Quantitative Approach - Chapter 4 Data-Level Parallelism in Vector, SIMD, and GPU Architectures]] - エンティティ: [[NVIDIA]] / [[NVIDIA Pascal]] / [[PTX]] ## 出典 - John L. Hennessy, David A. Patterson, *Computer Architecture: A Quantitative Approach*, Sixth Edition, Morgan Kaufmann, 2019, Chapter 4, §4.4.