AWSのGPUインスタンスでCUDAを動かす

はじめに

GPUプログラミングで事例を調べるとスパコンでの利用、物理シミュレーションでの利用などが出てくるため、敷居が高いイメージがあります。
しかし実際には、たくさんのスレッドを大量に走らせること自体はそれほど難しいことをする必要はなく、環境さえあれば意外と簡単に実行できます。
少し前までは、「環境さえあれば」の部分で敷居が高かったところもありますが、今はクラウドサービスで初期コストを抑えてGPUプログラミングを試すことができるようになってきました。
本記事では、CUDAを使ったいくつかのサンプルコードを紹介することで、基本的なCUDAの使い方をご紹介します。

ゴール

• GPUプログラミングは案外簡単だということを知ってもらう

AWS GPUインスタンス

今回はAWSのEC2の中でGPUが搭載されているP2インスタンスを使います。
P2インスタンスのGPUはNVIDIA Keplerアーキテクチャが採用されています。

CUDAのインストール

インストールは以前の記事を参考にしてください。

CUDAの書き方

CUDAは基本的にC/C++ベースでコードを記述することができます。
C/C++と違う点は、GPU側で動作させるためのコード(カーネル)が存在することです。

はじめは以下のような理解で十分かと思います。

• コンパイラは g++ではなくCUDA Toolkitに入っている nvccを使う
• ソースコードは .cpp、.cxxではなく .cuの拡張子を使う
• nvccがカーネルコード(GPU側)とホストコード(CPU側)を分けてコンパイルしてくれる
• カーネルコードには __global__を宣言する

サンプルを見た方が分かりやすいかと思いますのでさっそくCUDAでGPUを動かしましょう！

GPUを動かす

test.cu

#include <iostream>

__global__ void Kernel()
{
}

int main(void)
{
    Kernel<<<1, 1>>>();
    cudaDeviceSynchronize();

    std::cout << "done!" << std::endl;

    return 0;
}

これがGPUを動作させるための最も短いコードです。

コンパイル

$ nvcc -arch=sm_37 --std=c++14 test.cu

コンパイラはnvccを使います。-arch=sm_37という見慣れないコンパイルオプションが付いていますが、これはGPUのどのアーキテクチャを利用するのかという指定になります。P2インスタンスはKeplerアーキテクチャなので、sm_37を指定します。ちなみにP3インスタンスの場合、Voltaアーキテクチャなのでsm_70を指定することができます。
この指定によってGPUで使えるテクノロジーが変わってきます。

実行

$ ./a.out
done!

おめでとうございます！
無事GPUに処理を実行させることができました。

__global__ void Kernel()という部分でGPUで動作させるカーネルを定義しています。カーネルの呼び出しには Kernel<<<1, 1>>>()というちょっと特殊な記号をつけたものを使います。<<<1, 1>>>の部分で起動するスレッド数を指定します。

値を2つ設定していますが、<<<ブロック数, ブロック内で動作させるスレッド数>>>です。
ブロック数って何だ…？という気持ちになったかと思いますが、NVIDIAのGPUでは実行するスレッド群をブロック単位で処理しています。一つのブロックで扱えるスレッド数の上限は1024(Keplerの場合)となっており、かつプラクティスとして一つのブロックには256もしくは512スレッドを割り当てると最も高速に動作させられるため、例えば1024スレッド動作させようとすると 1024 / 256 = 4ブロック必要で、<<<4, 256>>>のように指定します。

cudaDeviceSynchronize()はカーネル実行の完了を待つCUDAのメソッドです。カーネルは非同期で実行されるため、このメソッドによってカーネル実行の完了を待つ必要があります。

Hello World

CUDAでHello Worldを出力してみましょう。
CUDAではカーネル側でprintf()実行が可能となっています。

test.cu

#include <iostream>

__global__ void Kernel()
{
    // スレッド番号の取得
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    printf("Hello World! thread:%d\n", i);
}

int main(void)
{
    Kernel<<<1, 10>>>();
    cudaDeviceSynchronize();

    std::cout << "done!" << std::endl;

    return 0;
}

せっかくなので10スレッドで実行してみましょう。
blockIdx.x * blockDim.x + threadIdx.xで実行されているスレッド番号を取得することができます。ブロック数、スレッド数ともに3次元に指定が可能なので3次元でブロック数、スレッド数を指定した場合はこの限りではないのですが、この例は1次元のスレッドなのでこれで大丈夫です。

コンパイル、実行

$ nvcc -arch=sm_37 --std=c++14 test.cu
$ ./a.out
Hello World! thread:0
Hello World! thread:1
Hello World! thread:2
Hello World! thread:3
Hello World! thread:4
Hello World! thread:5
Hello World! thread:6
Hello World! thread:7
Hello World! thread:8
Hello World! thread:9
done!

10スレッド分のHello Worldが出力されました。
この例からも分かるように、実行スレッド数分標準出力に出力されるため、CUDAでprintデバッグする際には大量のスレッドから標準出力へ出力しないよう注意が必要です。CUDAには実行スレッドを指定できるgdbライクなデバッガcuda-gdbもありますので、デバッグ時にはこちらを使うのがオススメです。

大量のスレッドを動作させてみる

いよいよGPUプログラミングの醍醐味、大量のスレッド動作を行ってみます。

test.cu

#include <iostream>
#include <vector>
 
__global__
void Kernel(const int size, int *values)
{
    const int index = blockIdx.x * blockDim.x + threadIdx.x;
    if (index >= size) return;
 
    values[index] = index;
}
 
int main(void)
{
    const int kVectorNum = 1 << 24;
    const int kVectorSize = kVectorNum * sizeof(int);
    std::vector<int> host_values(kVectorNum);
 
    // デバイス側のポインタ情報を保持する変数
    int *device_values = nullptr;
    // デバイス側にメモリを確保
    cudaMalloc(&device_values, kVectorSize);
 
    // ホストからデバイスへメモリを転送
    cudaMemcpy(device_values, host_values.data(), kVectorSize, cudaMemcpyHostToDevice);

    // スレッド数とブロック数の指定 
    const int thread_num = 256;
    const dim3 block(thread_num);
    const dim3 grid((kVectorNum + block.x - 1) / block.x);
 
    // カーネル呼び出し
    Kernel<<<grid, block>>>(kVectorNum, device_values);
    // カーネル終了待ち
    cudaDeviceSynchronize();

    // デバイスからホストへメモリを転送
    cudaMemcpy(host_values.data(), device_values, kVectorSize, cudaMemcpyDeviceToHost);
 
    // デバイス側のメモリを開放
    cudaFree(device_values);
 
    std::cout << host_values.at(16000000) << std::endl;
    std::cout << "done!" << std::endl;
 
    return 0;
}

16,777,216個の要素を持つ配列を用意し、その配列にデバイス側で自身のスレッド番号を代入させています。

コンパイル、実行

$ nvcc -arch=sm_37 --std=c++14 test.cu
$ ./a.out
16000000
done!

1600万スレッド以上の大量のスレッド実行に成功しました！

デバイス(GPU)側からホスト(CPU)側のメモリを直接参照することはできないので、ホストに確保したメモリと同じサイズのメモリをデバイスでも確保する必要があります。それがcudaMallocです。デバイス側で確保したメモリはcudaFreeで忘れずに開放してください。

ホストとデバイス間のデータ通信にはcudaMemcpyを使います。

dim3という見慣れない変数の型がありますが、これがブロック数とスレッド数を3次元に指定するためのCUDA用の型です。1次元の場合は1つだけ値を渡すことができます。16,777,216スレッドなので、スレッド数が256の場合、ブロック数は65,536となります。

カーネルコードにif (index >= size) return;というサイズチェックのロジックがありますが、CUDAは実は32スレッドを一つの塊(Warp)として同時に実行するため、例えば35スレッド実行したつもりでも、実際には64スレッド動いていることになります。もし、配列が35しか確保されていなかった場合、0～34番目のスレッドは正しく確保された領域へアクセスしますが、35～63番目のスレッドは領域外アクセスとなってしまうため、サイズチェックを必ず入れ、領域外へのアクセスを避けましょう。

まとめ

CUDAの基本的な書き方からGPUプログラミングの醍醐味である大量のスレッド処理をCUDAで行うところまでを紹介しました。意外と簡単に大量のスレッドを動かせることが分かっていただけたかと思います。

GPUプログラミングの実行環境もクラウドサービスを利用することによって初期コストを低く用意できるため、GPUプログラミングを行うハードルは一昔前と比べて大幅に下がっています。

さらに、GPUの進化によって、学術利用だけではない、本当の意味での一般的なGPUプログラミングが可能になってきています。特に大量のデータを扱うようなケースではマッチするかと思いますので気軽にCUDAを試してみてください！

参考URL

Amazon EC2 P2インスタンス

参考書籍

CUDA C プロフェッショナル プログラミング