meideru blog

家電メーカーで働いているmeideruのブログです。主に技術系・ガジェット系の話を書いています。

CUDAでGPGPUをやってみた(簡単なコードを書いてみた)

 

GPGPU(汎用GPU)を触ってみました。

使用した開発環境はもちろんNVIDIAのCUDA(クーダ)です。

というか、現時点でGPGPUをやろうと思ったらCUDAしか選択の余地はありません。今のGPU市場はNVIDIAの一強だからです。

CUDA(クーダ)

簡単なコードを書いてみたので、そういうのも交えながら、ご紹介しようと思います。

難しいことはまだやっていません。簡単なことしかやっていませんのでご了承ください^_^;

目次

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
グラフィックボード(GTX970)10世代が出たので1つ前の型ということになりますが、それでも最強クラスの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を勉強しても使う機会がないですからね(´・ω・`)

 

以上です!

ノシ

 - 技術系