CUDA入門 #1

この記事はNVIDIAのCUDAチュートリアル「An Even Easier Introduction to CUDA」学習時に生じた疑問点などをChatGPTのAdvanced Voice Modeで雑談した内容をもとにChatGPT 4oにブログ記事としてまとめてもらったものです

参考

本記事の内容をもとに行った勉強会の録画

https://www.youtube.com/watch?v=6paoohKS7oM

1. CUDAとは？

CUDA（Compute Unified Device Architecture）は、NVIDIAが開発した並列計算のためのプラットフォームです。通常のCPUで行う計算を、膨大な数のコアを持つGPU（Graphics Processing Unit）を使って処理することで、計算速度を劇的に向上させることができます。

GPUは、ゲームやグラフィックスの処理に使われるイメージが強いかもしれませんが、実は膨大なデータ処理や科学技術計算などにも非常に有効です。特に、画像処理、シミュレーション、機械学習などの分野では、GPUの並列処理能力が大いに活用されています。

ここで注目すべきは、GPUが大量のスレッド（軽量プロセス）を同時に実行できる点です。これにより、同時に多数のデータを処理することが可能となり、CPUでは得られない高速な演算が実現できます。

CUDAの基本構造

CUDAプログラムの基本的な流れは以下の通りです：

• データの準備: GPUに転送するためのデータを用意する。
• メモリの確保: GPU上にメモリを確保する。
• カーネル関数の定義: 並列実行される関数（カーネル）を定義する。
• データの転送: データをCPUからGPUに転送する。
• カーネルの実行: GPU上でカーネルを並列実行する。
• 結果の取得: GPUからCPUに結果を戻す。

次に、具体的なコード例を見ながらCUDAの基本的な使い方を理解していきましょう。

2. CUDAカーネルの基本構造

CUDAプログラムの肝となるのが、カーネル関数と呼ばれる、GPU上で並列実行される関数です。このカーネル関数は、特定の計算を多数のスレッドによって一斉に処理するものです。CUDAでは、グリッドとブロックという概念を用いて、スレッドの構成を指定します。

カーネル関数の定義と呼び出し
カーネル関数は、__global__修飾子を使って宣言されます。この修飾子を付けることで、関数がGPU上で実行されることが指定されます。以下に、非常にシンプルなカーネル関数の定義と、その呼び出し方を示します。

C++/CUDA サンプルコード

#include <iostream>
#include <cuda_runtime.h>

// カーネル関数の定義 (__global__ 修飾子を付ける)
__global__ void addKernel(int *data) {
    int idx = threadIdx.x;  // スレッドIDを取得
    data[idx] += 1;         // 各スレッドが1を加算
}

int main() {
    const int arraySize = 5;
    int hostData[arraySize] = {0, 1, 2, 3, 4};
    int *deviceData;

    // GPU上にメモリを確保 (cudaMalloc)
    cudaMalloc((void**)&deviceData, arraySize * sizeof(int));
    cudaMemcpy(deviceData, hostData, arraySize * sizeof(int), cudaMemcpyHostToDevice);

    // カーネルを起動 (<<<1, arraySize>>> はスレッド数を指定)
    addKernel<<<1, arraySize>>>(deviceData);

    // 結果をCPUにコピー (cudaMemcpy)
    cudaMemcpy(hostData, deviceData, arraySize * sizeof(int), cudaMemcpyDeviceToHost);

    // 結果を表示
    for (int i = 0; i < arraySize; i++) {
        std::cout << "Result[" << i << "]: " << hostData[i] << std::endl;
    }

    // GPUメモリを解放
    cudaFree(deviceData);

    return 0;
}

このコードでは、CUDAカーネルを用いて、各スレッドがhostData配列の要素に1を加算しています。threadIdx.xを使って、スレッドごとに異なる配列のインデックスを操作している点に注目してください。

Python (PyCUDA) サンプルコード

PythonでもCUDAを利用することが可能です。PyCUDAを使って同様の処理を行う例を示します。

import pycuda.autoinit
import pycuda.driver as cuda
import numpy as np
from pycuda.compiler import SourceModule

# CUDAカーネルコードを文字列として記述
mod = SourceModule("""
__global__ void addKernel(int *data) {
    int idx = threadIdx.x;
    data[idx] += 1;
}
""")

# カーネル関数の取得
addKernel = mod.get_function("addKernel")

# ホスト側のデータ
host_data = np.array([0, 1, 2, 3, 4], dtype=np.int32)

# デバイス側のメモリを確保し、データをコピー
device_data = cuda.mem_alloc(host_data.nbytes)
cuda.memcpy_htod(device_data, host_data)

# カーネルを起動
addKernel(device_data, block=(5, 1, 1))

# 結果をホストにコピー
cuda.memcpy_dtoh(host_data, device_data)

# 結果を表示
print("Result:", host_data)

このPythonコードでは、PyCUDAを使ってC++と同様の処理を行っています。SourceModuleを使ってCUDAカーネルを定義し、get_functionで取得したカーネル関数をGPU上で実行しています。

3. メモリ管理とcudaMallocManagedの利用

GPUでの計算を行う際には、データをGPU側のメモリに転送する必要があります。CUDAではこれを効率的に行うために、いくつかのメモリ管理手法が提供されています。ここでは、基本的なメモリ管理関数であるcudaMallocと、より便利なcudaMallocManagedを紹介します。

cudaMallocによるメモリ確保
まず、cudaMallocを使ってデバイスメモリ（GPUメモリ）を確保する方法を確認しましょう。cudaMallocは、GPUに直接メモリを割り当て、ホスト（CPU）からGPUにデータを転送します。

C++/CUDA サンプルコード

#include <iostream>
#include <cuda_runtime.h>

// カーネル関数の定義
__global__ void multiplyByTwo(int *data) {
    int idx = threadIdx.x;
    data[idx] *= 2;  // 各スレッドが2倍にする
}

int main() {
    const int arraySize = 5;
    int hostData[arraySize] = {1, 2, 3, 4, 5};
    int *deviceData;

    // GPU上にメモリを確保
    cudaMalloc((void**)&deviceData, arraySize * sizeof(int));

    // データをホストからデバイスにコピー
    cudaMemcpy(deviceData, hostData, arraySize * sizeof(int), cudaMemcpyHostToDevice);

    // カーネルを実行
    multiplyByTwo<<<1, arraySize>>>(deviceData);

    // 結果をホストにコピー
    cudaMemcpy(hostData, deviceData, arraySize * sizeof(int), cudaMemcpyDeviceToHost);

    // 結果を表示
    for (int i = 0; i < arraySize; i++) {
        std::cout << "Result[" << i << "]: " << hostData[i] << std::endl;
    }

    // GPUメモリを解放
    cudaFree(deviceData);

    return 0;
}

このコードでは、cudaMallocを使ってGPU側にメモリを確保し、データをCPUからGPUへ転送しています。その後、カーネル関数を実行して各要素を2倍にし、結果をCPUに戻しています。

cudaMallocManagedの利便性
次に、cudaMallocManagedを紹介します。この関数は、**Unified Memory（統一メモリ）**を提供します。これにより、CPUとGPU間での明示的なデータ転送が不要になり、NVIDIAのメモリ管理システムが自動的にデータを必要な場所に移動します。

これにより、開発者はメモリの管理をシンプルにでき、特に複雑なアプリケーションの開発時に便利です。

Python (PyCUDA) サンプルコード

PyCUDAには、cudaMallocManagedに相当する関数はありませんが、メモリ管理は基本的に自動化されています。同じ処理をPyCUDAで実装した場合、以下のようになります。

import pycuda.autoinit
import pycuda.driver as cuda
import numpy as np
from pycuda.compiler import SourceModule

# CUDAカーネルコードを文字列として記述
mod = SourceModule("""
__global__ void addOne(int *data) {
    int idx = threadIdx.x;
    data[idx] += 1;
}
""")

# カーネル関数の取得
addOne = mod.get_function("addOne")

# ホスト側のデータ (統一メモリのように扱える)
data = np.array([0, 1, 2, 3, 4], dtype=np.int32)

# デバイス側のメモリを確保し、データをコピー
device_data = cuda.mem_alloc(data.nbytes)
cuda.memcpy_htod(device_data, data)

# カーネルを起動
addOne(device_data, block=(5, 1, 1))

# 結果をホストにコピー
cuda.memcpy_dtoh(data, device_data)

# 結果を表示
print("Result:", data)

PyCUDAでは、ホストデータを管理するnumpy配列を使って簡単にメモリ管理ができます。

4. スレッドとブロックの構造

CUDAプログラムの強力な点は、GPUの並列処理能力を最大限に活用できるスレッドとブロックの構造です。これにより、1つのカーネルで何百、何千ものスレッドを一斉に実行できます。

CUDAでは、スレッドは「ブロック」に分けられ、複数のブロックは「グリッド」を構成します。このようにスレッドは階層的に構成されており、各スレッドがthreadIdx、ブロックがblockIdx、グリッドがgridDimによって識別されます。

スレッドとブロックの定義
スレッド数やブロック数は、カーネル関数を呼び出すときに<<<グリッドサイズ, ブロックサイズ>>>の形で指定します。

例えば、次のように指定することで、1つのブロックに256スレッド、ブロック全体で1024スレッドが動作します。

kernel<<<4, 256>>>(...);

この例では、グリッドに4つのブロックがあり、それぞれのブロック内で256スレッドが実行されます。

スレッドIDとブロックIDの利用
各スレッドには一意のIDが割り当てられており、これを使ってデータの処理を行います。例えば、次のようにblockIdxとthreadIdxを使って、スレッドごとに異なるデータを処理することができます。

C++/CUDA サンプルコード

#include <iostream>
#include <cuda_runtime.h>

// カーネル関数の定義
__global__ void processData(int *data) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x; // グローバルなスレッドIDを計算
    data[idx] += idx;  // スレッドIDを使ってデータを操作
}

int main() {
    const int arraySize = 1024;
    int hostData[arraySize];
    for (int i = 0; i < arraySize; i++) {
        hostData[i] = i;
    }
    int *deviceData;

    // GPU上にメモリを確保
    cudaMalloc((void**)&deviceData, arraySize * sizeof(int));

    // データをホストからデバイスにコピー
    cudaMemcpy(deviceData, hostData, arraySize * sizeof(int), cudaMemcpyHostToDevice);

    // 4ブロック x 256スレッドでカーネルを実行
    processData<<<4, 256>>>(deviceData);

    // 結果をホストにコピー
    cudaMemcpy(hostData, deviceData, arraySize * sizeof(int), cudaMemcpyDeviceToHost);

    // 結果を表示
    for (int i = 0; i < arraySize; i++) {
        std::cout << "Result[" << i << "]: " << hostData[i] << std::endl;
    }

    // GPUメモリを解放
    cudaFree(deviceData);

    return 0;
}

このコードでは、各スレッドが独自のID（blockIdx.x * blockDim.x + threadIdx.x）を使ってhostData配列の異なる要素を処理しています。

Python (PyCUDA) サンプルコード

PyCUDAでも同様にスレッドとブロックの階層構造を活用することができます。

import pycuda.autoinit
import pycuda.driver as cuda
import numpy as np
from pycuda.compiler import SourceModule

# CUDAカーネルコードを文字列として記述
mod = SourceModule("""
__global__ void processData(int *data) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    data[idx] += idx;
}
""")

# カーネル関数の取得
processData = mod.get_function("processData")

# ホスト側のデータ
data = np.arange(1024, dtype=np.int32)

# デバイス側のメモリを確保し、データをコピー
device_data = cuda.mem_alloc(data.nbytes)
cuda.memcpy_htod(device_data, data)

# 4ブロック x 256スレッドでカーネルを実行
processData(device_data, block=(256, 1, 1), grid=(4, 1))

# 結果をホストにコピー
cuda.memcpy_dtoh(data, device_data)

# 結果を表示
print("Result:", data)

PyCUDAでは、blockとgridを指定してスレッド数やブロック数をコントロールします。これにより、並列処理を効果的に行うことができます。

5. blockIdx、blockDim、threadIdxの理解

CUDAにおいて、各スレッドがどのデータを処理するかを決定するために、3つの重要な変数があります。それがblockIdx（ブロックID）、blockDim（ブロック内のスレッド数）、そしてthreadIdx（スレッドID）です。これらを組み合わせることで、スレッドごとに一意のインデックス（ID）を算出し、効率的にデータを処理することができます。

blockIdx、blockDim、threadIdxの関係

以下の式が、各スレッドのグローバルなインデックスを計算する基本式です：

グローバルインデックス = blockIdx.x * blockDim.x + threadIdx.x

• blockIdx.x: ブロックのインデックス（0から始まる）。
• blockDim.x: 1ブロック内のスレッド数。
• threadIdx.x: ブロック内のスレッドのインデックス（0から始まる）。
  この式を用いることで、GPU全体のどのスレッドがどのデータを処理するのかを決定できます。

具体例: 16要素の配列を並列に処理する
次に、16要素の配列を2つのブロック、各ブロックに8スレッドで処理する例を見てみましょう。このような設定では、各ブロックは8つのスレッドを持ち、全体で16スレッドが動作します。

C++/CUDA サンプルコード

#include <iostream>
#include <cuda_runtime.h>

// カーネル関数の定義
__global__ void calculateIndex(int *data) {
    int globalIdx = blockIdx.x * blockDim.x + threadIdx.x;  // グローバルスレッドIDを計算
    data[globalIdx] = globalIdx;  // グローバルIDをそのまま格納
}

int main() {
    const int arraySize = 16;
    int hostData[arraySize];
    int *deviceData;

    // GPU上にメモリを確保
    cudaMalloc((void**)&deviceData, arraySize * sizeof(int));

    // 2ブロック x 8スレッドでカーネルを実行
    calculateIndex<<<2, 8>>>(deviceData);

    // 結果をホストにコピー
    cudaMemcpy(hostData, deviceData, arraySize * sizeof(int), cudaMemcpyDeviceToHost);

    // 結果を表示
    for (int i = 0; i < arraySize; i++) {
        std::cout << "Result[" << i << "] = " << hostData[i] << std::endl;
    }

    // GPUメモリを解放
    cudaFree(deviceData);

    return 0;
}

実行結果

Result[0] = 0
Result[1] = 1
Result[2] = 2
...
Result[15] = 15

ここでは、blockIdx.xがブロックのIDを表し、threadIdx.xがブロック内でのスレッドのインデックスです。例えば、1つ目のブロック（blockIdx.x = 0）のスレッドはglobalIdxが0～7、2つ目のブロック（blockIdx.x = 1）のスレッドはglobalIdxが8～15を処理します。

Python (PyCUDA) サンプルコード

次に、同じ処理をPyCUDAで行う例です。

import pycuda.autoinit
import pycuda.driver as cuda
import numpy as np
from pycuda.compiler import SourceModule

# CUDAカーネルコードを文字列として記述
mod = SourceModule("""
__global__ void calculateIndex(int *data) {
    int globalIdx = blockIdx.x * blockDim.x + threadIdx.x;
    data[globalIdx] = globalIdx;
}
""")

# カーネル関数の取得
calculateIndex = mod.get_function("calculateIndex")

# ホスト側のデータ
data = np.zeros(16, dtype=np.int32)

# デバイス側のメモリを確保
device_data = cuda.mem_alloc(data.nbytes)

# カーネルを実行
calculateIndex(device_data, block=(8, 1, 1), grid=(2, 1))

# 結果をホストにコピー
cuda.memcpy_dtoh(data, device_data)

# 結果を表示
print("Result:", data)

実行結果

Result: [ 0  1  2  3  4  5  6  7  8  9 10 11 12 13 14 15]

説明
blockDim.x = 8は、1ブロックに8つのスレッドがあることを示します。
gridDim.x = 2は、2つのブロックがあることを示します。
threadIdx.xは各ブロック内で0から7の値を取り、blockIdx.xに基づいてスレッドごとのグローバルなインデックスを決定します。
このように、CUDAではスレッドごとに異なるデータを並列に処理するために、blockIdx、blockDim、threadIdxの関係を用いてデータのインデックスを計算します。

ここまでが第5節です。blockIdx、blockDim、threadIdxの関係を例を通して説明しました。内容をご確認ください。

6. cudaDeviceSynchronizeの重要性

GPUとCPUは通常、非同期的に動作します。つまり、GPUでカーネルが実行されている間に、CPUは次の処理を進めることができます。この非同期処理は効率的ですが、CPUがGPUでの計算が終わる前に結果を参照しようとすると、まだ計算が完了していない不正確なデータを参照してしまう可能性があります。

そこで使用するのが、cudaDeviceSynchronize()です。この関数を使うことで、GPU上で実行中のカーネルのすべての処理が終了するまで、CPUが待機するように指示できます。

cudaDeviceSynchronizeの使い方
以前の例（第5節のサンプルコード参照）では、各スレッドが配列の要素をインクリメントしていますが、cudaDeviceSynchronize()を使ってGPUでのすべての処理が完了するまで待機することで、正しい結果を取得できるようにしています。

cudaDeviceSynchronize()なしでの動作
一方で、cudaDeviceSynchronize()を使用しない場合、CPUがGPUの処理が完了する前に結果を読み取ることがあります。その場合、正しい結果が取得されない可能性があります。

実行結果（cudaDeviceSynchronize()なし）

• 結果はあくまで一例です（実際にはほぼランダムで変わります）

Result[0] = 2
Result[1] = 1
Result[2] = 0
Result[3] = 7
Result[4] = 11
...
Result[15] = 9

この結果から分かるように、cudaDeviceSynchronize()を使わなかった場合、正しい結果が得られないことがあります。これは、GPUのカーネルがすべての処理を完了する前に、CPUが結果を取得しようとしているためです。

正しい結果を得るために
cudaDeviceSynchronize()を使ってGPUの処理が完了するのを待つことで、すべてのスレッドが完了した後に結果を取得できるようになります。第5節のサンプルコードでは、cudaDeviceSynchronize()を使用することで、すべてのスレッドが正しく配列の値を更新するまで待機してから、結果を取得しています。

cudaDeviceSynchronize();  // この行があることで、GPU処理の完了を待つ

この関数を使うことで、データの整合性が保たれ、GPUでの並列処理の結果を正確に取得することができるのです。