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

** Copyright (c) 2017 Microsoft Corporation

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

[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 from GPU!

Hello World from GPU!

Hello World from GPU!

Hello World from GPU!

Hello World from GPU!

Hello World from GPU!

Hello World from GPU!

Hello World from GPU!

 

Hello World from GPU!

 

 

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

 

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

 

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

 

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

 

__global__(←アンダーバー2つ)は、

 

簡単に言えばコンパイラに「以下の関数を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の方が処理時間が約1/5になりました。

 

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

 

 

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

 ※GPU:GeForce GT 710