CUDAプログラミングの簡単な入門 この記事は、NVIDIAが提供する並列計算プラットフォームおよびプログラミングモデルであるCUDAについての初心者向け入門です。CUDA C++を使用することで、C++言語を用いて高性能アルゴリズムを開発し、GPU上で数千の並列スレッドで処理を加速できます。多くの開発者がこの方法でバンド폭幅や計算量の大きなアプリケーションを加速しており、特にDeep Learning革命の基礎となっているライブラリやフレームワークで活用されています。 基本的なプログラムの作成 まずは簡単なC++プログラムから始めましょう。このプログラムは、100万個の要素を持つ2つの配列を加算します。 ```cpp include void add(int a, int b, int* c, int n) { for (int i = 0; i < n; ++i) { c[i] = a[i] + b[i]; } } int main() { const int N = 1000000; int a = new int[N]; int b = new int[N]; int* c = new int[N]; for (int i = 0; i < N; ++i) { a[i] = i; b[i] = i; } add(a, b, c, N); bool error = false; for (int i = 0; i < N; ++i) { if (c[i] != 2 * i) { error = true; break; } } std::cout << "Error in sum: " << error << std::endl; delete[] a; delete[] b; delete[] c; return 0; } ``` このプログラムをadd.cppというファイル名で保存し、コンパイラ（Linuxではg++、WindowsではMSVCまたはWSLでのg++）でコンパイルして実行します。期待通り、出力は「Error in sum: 0」を表示して終了します。 GPU上での並列処理 次に、この計算をGPU上で並列に実行します。add関数をGPUで実行可能なカーネルに変更します。そのためには、関数に__global__修飾子を追加します。 cpp __global__ void add(int* a, int* b, int* c, int n) { for (int i = 0; i < n; ++i) { c[i] = a[i] + b[i]; } } これにより、add関数はCUDAカーネルとしてGPU上で実行されます。メモリの割り当てには、システム内のすべてのGPUとCPUからアクセス可能な単一のメモリ空間を提供するUnified Memoryを使うことができます。メモリをUnified Memoryで割り当てるには、cudaMallocManaged()を使用します。メモリを解放する際は、ポインタをcudaFree()に渡します。 コードの修正とコンパイル 配列のメモリ割り当てをnewからcudaMallocManaged()に変更し、配列の解放をdelete[]からcudaFree()に変更します。 ```cpp int a; cudaMallocManaged(&a, N * sizeof(int)); int b; cudaMallocManaged(&b, N * sizeof(int)); int* c; cudaMallocManaged(&c, N * sizeof(int)); // ...初期化など... cudaFree(a); cudaFree(b); cudaFree(c); ``` カーネルの起動 次に、カーネルをGPU上で実行するためのコードを追加します。CUDAのカーネル起動は、三連角括弧 <<<>>> を使用します。 cpp add<<<1, 1>>>(a, b, c, N); cudaDeviceSynchronize(); これにより、CPUスレッドがGPUスレッドの完了を待つことができます。コンパイルする際は、CUDA C++コンパイラnvccを使用します。 並列スレッドの利用 現在のコードは1つのスレッドのみで動作しますが、並列化するためにカーネルの実行構成 <<>> を変更します。例として、256スレッドのスレッドブロックを1つだけ利用してみましょう。 cpp add<<<1, 256>>>(a, b, c, N); それぞれのスレッドが配列の一部を処理するように、カーネルのコードも更新します。 ```cpp global void add(int a, int b, int* c, int n) { int index = threadIdx.x; int stride = blockDim.x; for (index; index < n; index += stride) { c[index] = a[index] + b[index]; } } ``` 複数のスレッドブロック CUDA GPUは、多くの並列プロセッサがSM（Streaming Multiprocessor）にグループ化されています。SMごとに複数の並列スレッドブロックを実行できます。たとえば、NVIDIA T4 GPUは40SMと2560CUDAコアを持ち、各SMは最大1024のアクティブスレッドをサポートします。この並列プロセッサを最大限に活用するために、複数のスレッドブロックでカーネルを起動します。 ```cpp int blockSize = 256; int numBlocks = (N + blockSize - 1) / blockSize; add<<>>(a, b, c, N); ``` さらに、 Unified Memory のページミグレーションによるボトルネックを解消するため、データを事前にGPUにプレフェッチします。 cpp cudaMemPrefetchAsync(a, N * sizeof(int), cudaCpuDeviceId); cudaMemPrefetchAsync(b, N * sizeof(int), cudaCpuDeviceId); cudaMemPrefetchAsync(c, N * sizeof(int), cudaCpuDeviceId); コンパイルと実行後、 profilierの出力を確認すると、カーネルの実行時間が大幅に短縮されていることがわかります。 | バージョン | 時間（ns） | 単一スレッド比の高速化 | バンディWITH | |------------|------------|------------------------|--------------| | 単一スレッド | 91,811,206 | 1x | 137 MB/s | | 単一ブロック（256スレッド） | 2,049,034 | 45x | 6 GB/s | | 複数ブロック | 47,520 | 1932x | 265 GB/s | データがメモリにロードされた後、単一ブロックから複数ブロックへの移行はGPU内のSM数（40）に比例して高速化します。また、CUDAは非常に高い帯域幅（265 GB/sはT4のピーク帯域幅320 GB/sの80%以上）を達成できます。 今後何をするか この記事があなたのCUDAに対する興味を引き、自身の計算でCUDA C++の活用に興味を持っていただければ幸いです。 CUDA Toolkitドキュメンテーションの閲覧や、プロファイラでのコード解析、printf()を使用したデバッグなどを試してみてください。 NVIDIA DLI（Deep Learning Institute）では、様々なレベルのCUDAプログラミングのコースを提供しています。 Pythonプログラマ向けのFundamentals of Accelerated Computing with CUDA Python や、より詳しい内容を学べる Advanced CUDA Courses もあります。 GPUプログラミングの世界へぜひ飛び込んでください。 業界からのコメント NVIDIAの技術者によると、「CUDAの初歩的な導入は重要ですが、実践的な経験を通じてその真の力を理解することが不可欠です。DLC提供のコースや豊富なチュートリアルは、CUDAを初めて触る人々にとって最適なスタート地点となります」とのことです。 NVIDIAは、HPC（高性能計算）やAI分野での並列計算の普及に尽力しており、CUDAを介した高性能なGPU利用は今後益々重要になるとされています。