CUDAでGPGPUをやってみた（簡単なコードを書いてみた）
GPGPU（汎用GPU）を触ってみました。
使用した開発環境はもちろんNVIDIAのCUDA（クーダ）です。
というか、現時点でGPGPUをやろうと思ったらCUDAしか選択の余地はありません。今のGPU市場はNVIDIAの一強だからです。
簡単なコードを書いてみたので、そういうのも交えながら、ご紹介しようと思います。
難しいことはまだやっていません。簡単なことしかやっていませんのでご了承ください^_^;
目次
スポンサーリンク
GPGPU（汎用GPU）について
GPGPU（汎用GPU）とは？
GPGPU（汎用GPU）とは、General-purpose computing on graphics processing units の略です。
これまでGPUというのは画像処理にしか使われていませんでした。
GPUの演算資源を画像処理以外の用途で使おう、というのがGPGPUです。
GPUを演算に使うメリットとは？
GPUはCPUと比べるとスレッドの数が圧倒的に多いのです。
それゆえ、単純な並列計算ではGPUの方が高速で計算できるのです。それもCPUとは比べものにならない速さで計算できます。
ちなみに言っておきますが、CPUはGPUに比べると何もかもダメだというわけではありません。用途の違いです。GPUは単純な計算は得意ですが、複雑な計算は苦手です。そいういうのはCPUが計算します。
なぜGPGPUが注目されているのか？
やはりディープラーニングが流行っているからでしょう。(ディープラーニングについては、もはや説明不要だと思うのでしません。)
ディープラーニングは単純な計算の連続です。それはまさにGPUが得意とする処理なのです。
CUDA（クーダ）とは？
CUDA（クーダ）はNVIDIAが開発しているGPGPUのための統合開発環境です。
みんなが大好きな（？）C言語で開発することができます。（私は大好きですよC言語。。。）
今、GPGPUをやろうと思ったらCUDAしか選択の余地はありません。今のGPU市場はNVIDIAの一強状態だからです。
GPUはAMDも作っていますが、NVIDIAに遅れをとっていてGPGPU向けに開発環境を整える余裕はないのです。
なぜGPGPUをやろうと思ったのか
実は最近、ディープラーニングの勉強を始めました。（上でも述べた通り、これはGPUが大活躍する場面なのです。）
ディープラーニングの勉強をしたいだけならば、フレームワークを導入すれば簡単にできます。代表例はCaffeやTensorFlow、Chainerなどでしょう。
これら導入すればGPUとの連携も自動でやってくれますが・・・。
せっかくならGPGPUの勉強もしてしまおうと思ったわけです・・・。
理由はそれだけですw
開発環境について
開発環境は私の自作PCを使用します。
OSはWindows 10 PROです。
GPUはGTX970です。
＠GTX970
10世代が出たので１つ前の型ということになりますが、それでも最強クラスのGPUだと思っています。
CUDAのインストールはとても簡単だった
WindowsのCUDAのインストールはとても簡単でしたね。（MacやLinuxはどうなのか知りません^_^;）
Windowsの場合はCUDA Toolkitというのをダウンロードしてインストールするだけです。（事前にVisual Studioを入れておく必要があるかもしれない）
すると、Visual Studioを起動すると新しいプロジェクトからCUDAのプロジェクトを作成するようになっていました。
とっても簡単ですね。
コードを書いてみた
参考にした本
参考にしたのは「CUDA by Example 汎用GPUプログラミング入門」という本です。
この本はNVIDIAのエンジニアが書いた本です。
難易度としては中級者向けの本というところでしょうか？少なくとも一通り理解するためにはC言語はマスターしている必要があります。
ポインタ（メモリの確保）などはバンバン出てきます。
書いたコード
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <stdlib.h>
#define N 1000000
__global__ void add(int *a, int *b, int*c)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
while (tid < N)
{
c[tid] = a[tid] + b[tid];
tid += blockDim.x * gridDim.x;
}
}
int main(void)
{
int *a, *b, *c;
int *dev_a, *dev_b, *dev_c;
// CPU側でメモリを割り当てる
a = (int *)malloc(N * sizeof(int));
b = (int *)malloc(N * sizeof(int));
c = (int *)malloc(N * sizeof(int));
// GPU側でメモリを割り当てる
cudaMalloc((void**)&dev_a, N * sizeof(int));
cudaMalloc((void**)&dev_b, N * sizeof(int));
cudaMalloc((void**)&dev_c, N * sizeof(int));
// CPU側で配列aと配列bを設定する
for (int i = 0; i < N; i++)
{
a[i] = i;
b[i] = 2 * i;
}
// 配列aと配列bをGPUにコピーする
cudaMemcpy(dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice);
add << <128, 128 >> >(dev_a, dev_b, dev_c);
// 配列cをGPUからCPUにコピーする
cudaMemcpy(c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost);
// 要求した処理をGPUが行ったことを確認する
bool success = true;
for (int i = 0; i < N; i++)
{
if ((a[i] + b[i]) != c[i])
{
printf("Error: %d + %d != %d\n", a[i], b[i], c[i]);
success = false;
}
}
if (success)
{
printf("We did it !\n");
}
// GPU側で割り当てたメモリを開放する
cudaFree(dev_a);
cudaFree(dev_b);
cudaFree(dev_c);
// CPU側で割り当てたメモリを解放する
free(a);
free(b);
free(c);
int i;
scanf("%d", &i);
return 0;
}
書いたコードについて
任意の長さのベクトル和を求めるプログラムです。
マクロ宣言されているNの値がベクトルの長さということになります。
ベクトルの値は35行目のfor文で決まっています。ここを変えれば変更することも可能ですよ。
並列処理が行われていることについて
9行目の__global__で始まっている部分が実際に並列処理を行なっている関数です。
この関数はadd <<<128, 128 >>>(dev_a, dev_b, dev_c);という風に呼び出されています。
<<<128, 128>>>というのは、128ブロック128スレッドで処理するという意味です。
add関数内のint tid = threadIdx.x + blockIdx.x * blockDim.x;というのは、一意な番号を得るために行っている処理です。
そして、同関数内のtid += blockDim.x * gridDim.x;は、同じスレッドで次の演算を行う一意な番号を得るための処理です。
わかってしまえば簡単ですね(^ ^)
今後について
今後については全く未定です。本はまだ読み途中ですので、とりあえず全て読みきろうと思っています。
しかし、今後使う予定がありません。CUDAを勉強しても使う機会がないですからね(´･ω･`)
以上です！
ノシ
スポンサーリンク
関連記事