CUDA C サンプルプログラム

C++

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の中では以下のことが起きています。

1. GPU内部のスケジューラが、196個のブロックを各コア（SM）に割り当てる。

2. 各コアで256個ずつのスレッドが爆速で立ち上がる。

3. 各スレッドが計算式を使って「俺は3456番目の担当だ！」と計算する。

4. 一斉にメモリを読み書きする。

この「スレッドIDの計算」さえマスターすれば、どんなに複雑な並列計算も作れるようになります。

 

#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;
}