CUDA（Compute Unified Device Architecture）は、NVIDIAが開発した並列コンピューティングプラットフォームおよびプログラミングモデルである。GPU（Graphics Processing Unit）を汎用計算に活用するための技術基盤を提供し、従来のCPUでは困難な大規模並列処理を効率的に実現する。CUDAはC、C++、Fortran、Pythonなどの言語をサポートし、科学技術計算、機械学習、画像処理など多岐にわたる分野で広く利用されている。
開発の背景
2000年代初頭、GPUは主にグラフィックス処理に特化していたが、その並列処理能力が汎用計算にも応用可能であることが認識され始めた。2006年、NVIDIAはGeForce 8800 GTXと共にCUDAを発表し、GPUを汎用計算に利用するための統一的なプログラミングモデルを提供した。これにより、従来のシェーダー言語を用いた煩雑な手法に代わり、C言語ベースの直感的な開発が可能となった。
主要なバージョンと機能追加
CUDAはバージョンアップを重ねるごとに機能を拡充してきた。CUDA 1.0（2007年）は基本機能を提供し、CUDA 2.0（2008年）では倍精度浮動小数点演算をサポート。CUDA 3.0（2010年）ではクラスター向けのGPU Direct機能が追加され、CUDA 5.0（2012年）では動的並列処理やC++11の一部対応が実現された。CUDA 7.0（2015年）ではテンソルコアを導入し、機械学習向けの効率的な行列演算を可能にした。以降もCUDA 11.x（2020年）ではAmpereアーキテクチャ対応、CUDA 12.x（2023年）ではHopperアーキテクチャ対応など、ハードウェアの進化に合わせて機能が拡張されている。
ハードウェアモデル
CUDAのハードウェアモデルは、複数の階層から構成される並列処理アーキテクチャに基づく。GPUは多数の演算ユニットを内蔵し、大規模なスレッド並列処理を実行する。
ストリーミングマルチプロセッサ
ストリーミングマルチプロセッサ（SM）はGPUの基本処理単位であり、複数のCUDAコアを内蔵する。各SMは独立した命令スケジューリングと共有メモリを備え、ブロック単位でスレッドグループを割り当てられる。SMの数と構成は世代ごとに異なり、例えばVoltaアーキテクチャでは1SMあたり64個のCUDAコアを搭載する。
CUDAコア
CUDAコアはSM内の最小演算ユニットであり、単精度浮動小数点演算や整数演算を実行する。各コアは独立した演算パイプラインを持ち、ワープ単位で命令を実行する。現代のCUDAコアはFMA（Fused Multiply-Add）やテンソルコア（バージョン依存）などの特殊演算もサポートする。
メモリ階層
CUDAのメモリモデルは、容量と速度のトレードオフを考慮した階層構造を採用する。プログラマは各メモリの特性を理解し、適切に使い分けることで性能を最大化できる。
グローバルメモリ
グローバルメモリはGPU全体で共有される大容量メモリであり、ホストとデバイス間のデータ転送にも使用される。遅延は大きいが、すべてのスレッドからアクセス可能。GDDR6やHBMなどの高速メモリ技術により、帯域幅は数百GB/sに達する。
共有メモリ
共有メモリはSM内に存在する高速なオンチップメモリであり、同一ブロック内のスレッド間で共有される。遅延はレジスタに次いで低く、手動キャッシュとして利用することでグローバルメモリアクセスを削減できる。容量はSMあたり数十KBから数百KBである。
レジスタとローカルメモリ
レジスタは各スレッドに割り当てられる最速の記憶領域であり、コンパイラにより自動管理される。レジスタ不足が生じた場合、データはローカルメモリ（グローバルメモリ上の一部）に退避されるが、レイテンシが増大するため注意が必要である。
テクスチャメモリ
テクスチャメモリは読み取り専用の特殊メモリであり、空間的局所性を活かしたキャッシュ最適化やハードウェア補間機能を提供する。画像処理や数値計算の特定パターンで有効であり、不規則な読み取りアクセスにも一定の性能を発揮する。
カーネル関数
カーネル関数はGPU上で並列実行される特別な関数であり、__global__修飾子で宣言される。ホスト側から呼び出され、指定されたスレッド構成で実行される。カーネル内では、スレッドIDやブロックIDなどの組み込み変数を用いて並列処理を記述する。
スレッド構成
グリッドとブロック
スレッドは階層的に編成される。グリッドはブロックの集合であり、各ブロックはスレッドの集合である。ブロック内のスレッドは共有メモリを通じて協調し、ブロック間の通信はグローバルメモリを介して行われる。グリッドとブロックの次元は1次元から3次元まで指定できる。
ワープとスケジューリング
ワープは32スレッドからなるスケジューリング単位であり、SM内でSIMT（Single Instruction, Multiple Thread）方式で並列実行される。同一ワープ内のスレッドは同一命令を実行するが、分岐によりダイバージェンスが発生すると性能が低下する。効率的なプログラムではワープ内で条件分岐を最小限に抑える設計が求められる。
同期と通信
ブロック内同期
ブロック内の全スレッドを同期させるには__syncthreads()を使用する。これにより共有メモリの読み書きの一貫性を保証できるが、過剰な同期はパフォーマンスを低下させるため、必要な箇所に限定して使用する。
アトミック操作
アトミック操作はグローバルメモリや共有メモリ上で競合を回避しながら読み書きを行う機能である。例えばatomicAddは、複数スレッドが同一アドレスに加算する場合に正しい結果を保証する。アトミック操作は負荷が高いため、可能な限りローカライズして使用する。
CUDA Toolkit
CUDA ToolkitはNVIDIAが提供する公式開発キットであり、コンパイラ、デバッガ、プロファイラ、ライブラリなどを含む。インストールにより、CUDA対応アプリケーションの開発と実行が可能となる。
nvccコンパイラ
nvccはCUDA C/C++コードをコンパイルするドライバコンパイラである。ホストコード（CPU）とデバイスコード（GPU）を分離し、デバイスコードをPTX（中間表現）またはSASS（機械語）に変換する。最適化オプションやアーキテクチャ指定を行うことで、ターゲットGPUに最適なバイナリを生成できる。
プロファイラ（Nsight）
NsightはCUDAアプリケーションのパフォーマンス解析ツールである。Nsight Systemsはシステム全体のトレースを提供し、Nsight Computeはカーネル単位の詳細な分析を行う。これにより、メモリアクセスパターン、ワープ占有率、命令スループットなどのボトルネックを特定し、最適化に役立てる。
ライブラリ
cuBLAS
cuBLASは線形代数演算の高速実装を提供するライブラリである。BLAS（Basic Linear Algebra Subprograms）のGPU版として、行列積やベクトル演算などをCUDA上で最適化して実行する。深層学習や科学計算で広く利用される。
cuFFT
cuFFTは離散フーリエ変換（FFT）の高速実装である。多次元の実数・複素数変換、バッチ処理などをサポートし、信号処理や画像処理における計算を効率化する。
cuDNN
cuDNNは深層学習に特化したライブラリであり、畳み込み演算、プーリング、活性化関数、正規化などを最適化する。多くの機械学習フレームワーク（TensorFlow、PyTorchなど）のバックエンドとして使用され、GPU上での学習・推論を高速化する。
科学技術計算
分子動力学
分子動力学シミュレーションは多数の粒子の相互作用を計算するため、GPUの並列性能が有効である。AMBER、GROMACS、NAMDなどの主要なソフトウェアはCUDAをサポートし、従来のCPU比で数十倍の高速化を実現している。
流体シミュレーション
流体力学の数値解法（例：格子ボルツマン法、有限体積法）は大規模メッシュを用いるため、CUDAによる並列化が効果的である。OpenFOAMや独自実装により、航空宇宙、気象予測、工業設計などで利用されている。
機械学習
深層学習の高速化
深層学習では大量の行列演算と勾配計算が必要であり、CUDAとcuDNNが中核的な役割を果たす。特に畳み込みニューラルネットワーク（CNN）やトランスフォーマーモデルの学習時間を大幅に短縮し、研究開発の効率を向上させた。
推論の最適化
推論フェーズでは、テンソルコアを用いた混合精度演算（FP16、INT8）により、電力効率とスループットを改善する。TensorRTなどの最適化ツールと組み合わせることで、リアルタイム推論が要求されるアプリケーション（自動運転、音声認識など）に適用される。
画像・動画処理
リアルタイムレンダリング
CUDAは従来のグラフィックスパイプラインに加え、物理シミュレーションやポストエフェクトなどのGPU処理を可能にする。NVIDIA GameWorksやOptiXなどのライブラリは、リアルタイムなレイトレーシングや物理演算を実現する。
画像認識
画像認識では畳み込み処理や特徴抽出に大量の並列演算が必要であり、CUDAが高速化に貢献する。OpenCVのCUDAモジュールや独自実装により、リアルタイムな物体検出やセマンティックセグメンテーションが可能となる。
OpenCL
OpenCL（Open Computing Language）はKhronos Groupが策定したオープンな並列プログラミング標準であり、GPUだけでなくCPUやFPGAなど多様なデバイスを対象とする。CUDAに比べて移植性は高いが、NVIDIAプラットフォームではCUDAの方が性能やツールの充実度で優位とされる。
ROCm
ROCm（Radeon Open Compute）はAMDが推進するオープンソースのGPUコンピューティングプラットフォームである。HIP（Heterogeneous Interface for Portability）を提供し、CUDAコードをAMD GPU向けに移植しやすくする。性能面ではCUDAとの差が縮まりつつあり、特にHPCや深層学習領域で競合関係にある。
OneAPI
OneAPIはIntelが提唱する統一プログラミングモデルであり、CPU、GPU、FPGA、AIアクセラレータを横断するアプリケーション開発を目指す。DPC++（Data Parallel C++）を用いてSYCLベースのコード記述を行う。Intel GPUの普及に伴い、科学技術計算分野での採用が進んでいるが、エコシステムの成熟度はCUDAに及ばない。