GPUコンピューティングを始めよう

1. GPUとは？

GPUとは Graphics Processing Unit の略称です。

高速のVRAM（ビデオメモリ；グラフィックボード上のGPU専用メモリ）を持っています。

大量のデータを複数のプロセッサで同時かつ並列処理することができます。

2. GPUコンピューティングとは？

計算機での数値計算においてGPUを利用することをGPUコンピューティングといいます。

NVIDIA（エヌビディア）社が汎用計算に対応したGPUを開発し、CUDA（クーダ）を配信しGPUコンピューティングが普及しました。

3. CUDAとは？

CUDA（クーダ）とはCompute Unified Device Archtectureの略称です。

半導体メーカーNVIDIA（エヌビディア）社が提供するGPUコンピューティング向けの統合開発環境で、

プログラム言語はC言語ベースに拡張を加えたものを使用します。

4. CUDAに必要なもの（使用環境）

【GPU】

　CUDAに対応したNVIDIA製GPU

【OS】

　Windows 10（64bit版）

【コンパイラ】

　OSに対応したCコンパイラ（WindowsならVisual Studio）

　https://www.microsoft.com/ja-jp/dev/default.aspx

　Visual Studio community 2017を使用

　CUDA TOOLKITのダウンロード

　https://developer.nvidia.com/cuda-downloads

5. CUDAのインストール

①Visual Studio community 2017 をインストールします。

②CUDA TOOLKIT をインストールします。

6. Hello, World !

#include <stdio.h>

int main(int argc, char **argv) {

printf("Hello World!\n");

return 0;

}

コンパイルし実行すると、

**********************************************************************

** Visual Studio 2017 Developer Command Prompt v15.4.2

**********************************************************************

[vcvarsall.bat] Environment initialized for: 'x64'

C:\cuda>nvcc -o hello_world hello_world.cu

hello_world.cu

ライブラリ hello_world.lib とオブジェクト hello_world.exp を作成中

C:\cuda>hello_world

Hello World!

拡張子が「cu」であり、「nvcc」というCUDA用のコマンドを使用してコンパイルできることを確認します。

nvccコンパイラというのは、「ソースの中のCPUで動く部分とGPUで動く部分を分け、

CPU用コードをCコンパイラに渡し、GPU用コードをコンパイルする」というものです。

そのため、ソースがC言語のみで書かれていてもコンパイル可能、

その場合実際にコンパイルしているのはCコンパイラです。

7. Hello, World FROM GPU !

#include <stdio.h>

__global__ void helloFromGPU(void) {

printf("Hello World from GPU!\n");

}

int main(int argc, char **argv) {

helloFromGPU<<<1,10>>>();

return 0;

}

コンパイルし実行すると、

C:\cuda>nvcc -o hello_gpuN hello_gpuN.cu

hello_gpuN.cu

ライブラリ hello_gpuN.lib とオブジェクト hello_gpuN.exp を作成中

C:\cuda>hello_gpuN

Hello World from GPU!

先ほどのHello Worldと少し違う点がありますね。

・__global __から始まる helloFromGPU(void) という名前の空の関数

・<<<1,10>>>で修飾された関数の呼び出し

という、C言語では見慣れないコードが出てきました。

__global__（←アンダーバー２つ）は、

簡単に言えばコンパイラに「以下の関数をGPU側で実行します」と指定するコマンドです。

helloFromGPU()関数はGPU用コードを処理するコンパイラに渡され、

main()関数はCコンパイラに渡されます。

CUDA言語はC言語と似ており、同じようにプログラムを書くことができることが特徴です。

8. どれくらいスピードアップするの？

配列の加算処理を例にCPUとGPUの処理時間を比較してみる。

#include <stdio.h>

#include <helper_functions.h>

#include <helper_cuda.h>

#include <helper_timer.h>

#define CHECK(call) \

{ \

const cudaError_t error = call; \

if (error != cudaSuccess) \

{ \

printf("Error: %s:%d, ", __FILE__, __LINE__); \

printf("code:%d, reason: %s\n", error, \

cudaGetErrorString(error)); \

exit(1); \

} \

}

/*

この例では、GPUとホスト上の単純なベクトルの合計を示します。

sumMatrixOnGPU2Dは、GPU上のCUDAスレッド間でベクトル合計の作業を分割します。

2Dスレッドブロックと2Dグリッドが使用されます。

sumMatrixOnHostは、ホスト上のベクトル要素を順番に反復処理します。

*/

void initialData(float *ip, const int size){

int i;

for(i = 0; i < size; i++){

ip[i] = (float)(rand() & 0xFF) / 10.0f;

}

return;

}

void sumMatrixOnHost(float *A, float *B, float *C, const int nx, const int ny){

float *ia = A;

float *ib = B;

float *ic = C;

for (int iy = 0; iy < ny; iy++){

for (int ix = 0; ix < nx; ix++){

ic[ix] = ia[ix] + ib[ix];

}

ia += nx;

ib += nx;

ic += nx;

}

return;

}

void checkResult(float *hostRef, float *gpuRef, const int N){

double epsilon = 1.0E-8;

bool match = 1;

for (int i = 0; i < N; i++){

if (abs(hostRef[i] - gpuRef[i]) > epsilon){

match = 0;

printf("host %f gpu %f\n", hostRef[i], gpuRef[i]);

break;

}

if (match)

printf("配列が一致します。\n\n");

else

printf("配列が一致しません。\n\n");

}

// grid 2D block 2D

__global__ void sumMatrixOnGPU2D(float *MatA, float *MatB, float *MatC, int nx, int ny){

unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x;

unsigned int iy = threadIdx.y + blockIdx.y * blockDim.y;

unsigned int idx = iy * nx + ix;

if (ix < nx && iy < ny){

MatC[idx] = MatA[idx] + MatB[idx];

}

int main(int argc, char **argv)

{

printf("%s 開始...\n", argv[0]);

// デバイスを設定する

int dev = 0;

cudaDeviceProp deviceProp;

CHECK(cudaGetDeviceProperties(&deviceProp, dev));

printf("使用するデバイス %d: %s\n", dev, deviceProp.name);

CHECK(cudaSetDevice(dev));

int nx = 12192;

int ny = 12192;

int nxy = nx * ny;

int nBytes = nxy * sizeof(float);

printf("配列サイズ: nx %d ny %d\n", nx, ny);

// mallocホストメモリ

float *h_A;

float *h_B;

float *hostRef;

float *gpuRef;

h_A = (float *)malloc(nBytes);

h_B = (float *)malloc(nBytes);

hostRef = (float *)malloc(nBytes);

gpuRef = (float *)malloc(nBytes);

//********************************************************************

// 時間計測(初期化)

//********************************************************************

StopWatchInterface *timer=NULL;

sdkCreateTimer(&timer);

sdkResetTimer(&timer);

sdkStartTimer(&timer);

initialData(h_A, nxy);

initialData(h_B, nxy);

sdkStopTimer(&timer);

float time = sdkGetTimerValue(&timer);

sdkDeleteTimer(&timer);

printf("配列初期化処理 %f 秒\n", time/1000);

memset(hostRef, 0, nBytes);

memset(gpuRef, 0, nBytes);

//********************************************************************

// 時間計測(CPU)

//********************************************************************

// 開始時刻の取得

sdkCreateTimer(&timer);

sdkResetTimer(&timer);

sdkStartTimer(&timer);

// 配列加算処理(CPU)

sumMatrixOnHost(h_A, h_B, hostRef, nx, ny);

// 終了時刻の取得

sdkStopTimer(&timer);

time = sdkGetTimerValue(&timer);

sdkDeleteTimer(&timer);

printf("配列加算処理(CPU) %f 秒\n", time/1000);

// mallocデバイスのグローバルメモリ

float *d_MatA;

float *d_MatB;

float *d_MatC;

CHECK(cudaMalloc((void **)&d_MatA, nBytes));

CHECK(cudaMalloc((void **)&d_MatB, nBytes));

CHECK(cudaMalloc((void **)&d_MatC, nBytes));

// ホストからデバイスにデータを転送する

CHECK(cudaMemcpy(d_MatA, h_A, nBytes, cudaMemcpyHostToDevice));

CHECK(cudaMemcpy(d_MatB, h_B, nBytes, cudaMemcpyHostToDevice));

// ホスト側でカーネルを起動する

int dimx = 32;

int dimy = 32;

dim3 block(dimx, dimy);

dim3 grid((nx + block.x - 1) / block.x, (ny + block.y - 1) / block.y);

//********************************************************************

// 時間計測(GPU)

//********************************************************************

// 開始時刻の取得

sdkCreateTimer(&timer);

sdkResetTimer(&timer);

sdkStartTimer(&timer);

// 配列加算処理(GPU)

sumMatrixOnGPU2D<<<grid, block>>>(d_MatA, d_MatB, d_MatC, nx, ny);

CHECK(cudaDeviceSynchronize());

// 終了時刻の取得

sdkStopTimer(&timer);

time = sdkGetTimerValue(&timer);

sdkDeleteTimer(&timer);

printf("配列加算処理(GPU) %f 秒\n", time/1000);

// カーネルエラーをチェックする

CHECK(cudaGetLastError());

// カーネルの結果をホスト側にコピーする

CHECK(cudaMemcpy(gpuRef, d_MatC, nBytes, cudaMemcpyDeviceToHost));

// デバイスの結果を確認する

checkResult(hostRef, gpuRef, nxy);

// デバイスのグローバルメモリを解放する

CHECK(cudaFree(d_MatA));

CHECK(cudaFree(d_MatB));

CHECK(cudaFree(d_MatC));

// ホストのメモリを解放する

free(h_A);

free(h_B);

free(hostRef);

free(gpuRef);

// デバイスをリセットする

CHECK(cudaDeviceReset());

return (0);

}

コンパイルし実行すると

Hello World from GPU!

C:\cuda>nvcc -o sumMatrixOnCPU sumMatrixOnCPU.cu

sumMatrixOnCPU.cu

ライブラリ sumMatrixOnCPU.lib とオブジェクト sumMatrixOnCPU.exp を作成中

C:\cuda>sumMatrixOnCPU

sumMatrixOnCPU 開始...

使用するデバイス 0: GeForce GT 710

配列サイズ: nx 12192 ny 12192

配列初期化処理 5.338648 秒

配列加算処理(CPU) 0.636599 秒

配列加算処理(GPU) 0.123255 秒

配列が一致します。

まとめ

　GPUの方が処理時間が約１／５になりました。

　単純な数値計算はすごく早いです。

　※CPU:AMD Ryzen 3 1300X プロセッサー（4コア 3.5GHz）

　※GPU:GeForce GT 710