Gemini先生の教えメモ
このプログラムは、2つの配列（ベクトル）の各要素をGPUの数千ものスレッドで同時に足し合わせるもの。
|
1 2 3 4 5 6 7 8 |
// GPUで実行される「カーネル」関数 __global__ void vectorAdd(const float *A, const float *B, float *C, int numElements) { int i = blockDim.x * blockIdx.x + threadIdx.x; // スレッドIDの計算 if (i < numElements) { C[i] = A[i] + B[i]; } } |
1. __global__ 修飾子
これは**「カーネル (Kernel)」**であることを宣言するキーワードです。
-
意味: 「CPUから呼び出され、GPU上で実行される関数」であることを示します。
-
役割: コンパイラ（nvcc）はこのキーワードを見て、「この関数はx86（CPU）用ではなく、NVIDIAのアーキテクチャ（PTXコード）用にビルドせよ」と判断します。
-
制約: 戻り値は必ず
voidである必要があります。
2. カーネルの引数 (const float *A など)
引数はすべて「GPUのメモリ（ビデオメモリ/VRAM）」を指すポインタです。
-
CPU側のメモリを直接渡すことはできず、事前に
cudaMallocで確保し、cudaMemcpyでデータを転送しておいたアドレスを渡します。 -
numElementsは、配列のサイズを伝えて、範囲外アクセスを防ぐために必要です。
3. スレッドIDの計算（最重要ポイント）
C++
|
1 |
int i = blockDim.x * blockIdx.x + threadIdx.x; |
ここが並列計算の魔法の種明かしです。GPUは、この関数を数万個のスレッドで同時に実行します。各スレッドは「自分は誰か？」を知るために、組み込み変数（構造体）を参照します。
| 変数名 | 意味 | 例（今回のケース） |
blockIdx.x |
現在のスレッドが所属しているブロックの番号 | 0 〜 195 |
blockDim.x |
1つのブロックの中に含まれるスレッドの数 | 256 |
threadIdx.x |
ブロック内でのスレッドの通し番号 | 0 〜 255 |
4. if (i < numElements) によるガード
なぜこれが必要なのでしょうか？
GPUは「ブロック（今回なら256単位）」でスレッドを起動します。
-
計算したい要素数：50,000
-
起動したスレッド数：196ブロック × 256 ＝ 50,176スレッド
-
余り：176スレッド
もしこの if 文がないと、50,001番目以降の存在しないメモリ領域にアクセスしようとして、プログラムがクラッシュ（セグメンテーション違反）してしまいます。これを防ぐために「自分の番号 i が配列サイズ内にあるときだけ仕事をする」という条件分岐を入れています。
5. 計算処理：C[i] = A[i] + B[i];
ここまで来て、ようやく実際の計算です。
-
CPUの場合:
for(int i=0; i<N; i++)と書いて、1つずつ順番に足していきます。 -
GPUの場合: 50,000人の作業員（スレッド）が、一斉に自分の担当番号
iの箱を開け、中身を足して、結果をC[i]に書き込みます。
論理的には**「一瞬」**で全要素の足し算が終わります（実際にはメモリ帯域などの物理的制約がありますが、CPUのループより遥かに高速です）。
まとめ：GTX 1060の中での動き
あなたがプログラムを実行したとき、GTX 1060の中では以下のことが起きています。
-
GPU内部のスケジューラが、196個のブロックを各コア（SM）に割り当てる。
-
各コアで256個ずつのスレッドが爆速で立ち上がる。
-
各スレッドが計算式を使って「俺は3456番目の担当だ！」と計算する。
-
一斉にメモリを読み書きする。
この「スレッドIDの計算」さえマスターすれば、どんなに複雑な並列計算も作れるようになります。
|
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 |
#include <stdio.h> #include <cuda_runtime.h> // GPUで実行される「カーネル」関数 __global__ void vectorAdd(const float *A, const float *B, float *C, int numElements) { int i = blockDim.x * blockIdx.x + threadIdx.x; // スレッドIDの計算 if (i < numElements) { C[i] = A[i] + B[i]; } } int main() { int numElements = 50000; size_t size = numElements * sizeof(float); // ホスト（CPU）用メモリ確保 float *h_A = (float *)malloc(size); float *h_B = (float *)malloc(size); float *h_C = (float *)malloc(size); // データの初期化 for (int i = 0; i < numElements; ++i) { h_A[i] = 1.0f; h_B[i] = 2.0f; } // デバイス（GPU）用メモリ確保 float *d_A, *d_B, *d_C; cudaMalloc((void **)&d_A, size); cudaMalloc((void **)&d_B, size); cudaMalloc((void **)&d_C, size); // CPUからGPUへデータをコピー cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice); cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice); // カーネルの実行設定（256スレッド/ブロック） int threadsPerBlock = 256; int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock; printf("CUDA kernel launch with %d blocks of %d threads\n", blocksPerGrid, threadsPerBlock); vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, numElements); // 結果をGPUからCPUへコピー cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost); // 結果の検証 for (int i = 0; i < 10; ++i) { printf("Element %d: %.1f + %.1f = %.1f\n", i, h_A[i], h_B[i], h_C[i]); } // メモリ解放 cudaFree(d_A); cudaFree(d_B); cudaFree(d_C); free(h_A); free(h_B); free(h_C); return 0; } |