CUDA プログラミング - 基礎知識

CUDA 実行プログラム（関数）の修飾子

__global__

• デバイス（GPU）で実行される関数
• ホスト CPU から呼び出される
• ダイナミックパラレリズムを使用する場合はデバイスから呼び出すことも可能

__device__

• デバイス（GPU）で実行される関数
• デバイス（GPU）から呼び出される

__host__

• ホスト CPU で実行される関数
• ホストからのみ呼び出しが可能

__shared__

• デバイス（GPU）で実行される関数

__managed__

• ユニファイドメモリのサポートに伴い追加された修飾子

メモリ構成

グローバルメモリ

• GPU で実行されるすべてのスレッドで共用される

コンスタントメモリ

• すべてのスレッドで共用される定数を格納する

テクスチャメモリ

• グラフィックスのテクスチャパターンを記憶する

ローカルメモリ

• スレッドごとに割り当てられる（プライベート）

シェアードメモリ

• GPU で実行されるスレッドブロック単位に割り当てられる
• 1つのスレッドブロックに含まれるスレッド間では共用
• 他のスレッドブロックとは独立していて、スレッドブロック間のアクセスはできない
• デバイスメモリは DRAM で作られているため大容量だが、アクセス時間は数百サイクル掛かるのに対し、シェアードメモリは高速の SRAM（Static RAM）で作られ GPU チップに内蔵されているので、数サイクルでアクセスできる高速メモリ
• 各スレッドブロック内で共有され、頻繁にアクセスされる変数はシェアードメモリに置く
• 容量が小さいため、必要なデータのみ配置する

メモリ領域の修飾子

• 変数や配列がどこに配置されているか、どのように使用されるのかを指定する
• ホストメモリ（CPU 側）、デバイスメモリ（グローバル・コンスタント・テクスチャ・ローカル）、シェアードメモリで大きく分離される
• 修飾子の指定がない場合は、グローバル領域に配置される
• __device__ や __constant__ と指定された領域や変数はデバイスメモリに配置される
• デバイスメモリのどの領域（グローバル・コンスタント・テクスチャ・ローカル）に配置するか詳しく指定することも可能
• 指定がない変数はグローバルメモリに配置され、アプリケーションの実行開始から終了まで留まる
• __constant__ は定数扱いのため値の変更は不可
• __shared__ はシェアードメモリに配置される

CUDA プログラムで使用される変数

• char
• short
• int
• long
• longlong
• float
• double
• uchar
• ushort
• uint
• ulong
• ulonglong

ベクトル型の変数

• 例えば、int は 4 バイト（32 ビット）の整数型の変数だが、int3 のように末尾に数字を付与すると、3 つの 4 バイトの整数をまとめたベクトル変数を表すことが可能
• ベクトル長の範囲は 1 ~ 4
• dim3 という特別な変数タイプも存在する
• dim3 は uint3 と同等だが、dim3 の中の指定のない要素の値は自動的に 1 になる
• dim3 は GPU で実行するカーネルのグリッドやスレッドブロックのサイズを指定する場合に使われる
• 各スレッドで自分の位置がわかるようにするための、組み込み変数（blockIdx, threadIdx）も存在する

デバイスメモリの確保 / 解放 とホストメモリとのデータ転送

• CPU 側から GPU 側に入力データを転送したり、GPU 側で計算した結果を CPU 側に転送する必要があるため、ホストプログラムでは大まかに以下の処理ステップを実行する必要がある

  • カーネルの入出力となるホストメモリ領域の確保
  • デバイス側に入力を受け取るデバイスメモリ領域の確保
  • デバイス側に出力や作業用のデバイスメモリ領域の確保
  • ホストメモリからデバイスメモリへの入力データのコピー
  • カーネルを起動して計算処理を実行
  • 計算結果をデバイスメモリからホストメモリにコピー
  • 使用済みのメモリを解放

• ホスト側のメモリの確保は melloc() 関数

• デバイス側のメモリの確保は cudaMalloc() 関数

• ホスト側のメモリからデバイス側のメモリにデータをコピーするのは cudaMemcpy() 関数

• カーネルを呼び出して計算処理を実行後、使用済みのメモリを解放するは cudaFree() 関数

• ホスト側のメモリの解放は free() 関数

• カーネルの呼び出し（起動）カーネル関数名<<<ブロック数, ブロックサイズ>>>(引数)

  add カーネルを1ブロック - 256スレッドで実行

  __global__ void add(int *a, int *b, int *c) {
      int index = threadIdx.x;
      c[index] = a[index] + b[index];
  }
   
  int main() {
      // カーネルの起動
      add<<<1, 256>>>(d_a, d_b, d_c);
  }

  // デバイスで実行するカーネル
  __global__ void VecAdd(float* A, float* B, float* C, int N)
  {
      int i = blockDim.x * blockIdx.x + threadIdx.x;
      if (i < N)
      C[i] = A[i] + B[i];
  }
   
  // ホストで実行するコード
  int main()
  {
      int N = ...;
      size_t size = N * sizeof(float);
   
      // 入力ベクトルh_Aとh_B をホストメモリに獲得
      float* h_A = (float*)malloc(size);
      float* h_B = (float*)malloc(size);
   
      // 入力ベクトルの初期化
      ...
   
      // 入出力のベクトルの領域をデバイスメモリに獲得
      float* d_A;
      cudaMalloc(&d_A, size);
      float* d_B;
      cudaMalloc(&d_B, size);
   
      float* d_C;
      cudaMalloc(&d_C, size);
   
      // ホストメモリからデバイスメモリに入力をコピー
      cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
      cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
   
      // カーネルの呼び出し
      int threadsPerBlock = 256;
      int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
      VecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
   
      // 結果をデバイスメモリからホストメモリにコピー
      // 加算結果h_Cをホストメモリにコピー
      cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
   
      // 使用を終わったデバイスメモリを解放
      cudaFree(d_A);
      cudaFree(d_B);
      cudaFree(d_C);
   
      // ホストメモリを解放
      ...
  }

GPUメモリの管理

CUDAの関数 cudaMalloc、cudaMemcpy、および cudaFree は、それぞれGPUメモリの管理に関連している

cudaError_t cudaMalloc(void **devPtr, size_t size);

cudaError_t cudaMemcpy(void *dst, const void *src, size_t count, cudaMemcpyKind kind);

cudaError_t cudaFree(void *devPtr);