技術情報2024.09.13

CUDAによる処理の高速化実装編

この記事では、既存のCPUプログラムをGPUで高速化する方法を説明します。具体的には、行列の和を計算する部分をカーネル関数に変更し、メモリ確保やデータ転送、カーネル関数の実装と呼び出し、エラーチェック、チューニング方法について順を追って解説します。

「CUDA による処理の高速化基礎編」では用語の説明など基礎的な事項についてまとめたので、今回は具体的な実装について説明します。

書籍やネット上の解説記事では、プログラムを新規で書くことが普通だと思います。しかし、業務ではお客様から渡されたプログラムを修正する形で高速化することが多いので、今回はCPUで動いているプログラムがある前提で、それをどう高速化していくか、ということを順を追って説明します。

既存処理

まずは既存処理を見てみます。

void main() {
    int Width = 1024;
    int Height = 1024;
    int **Input1 = new int*[Height];
    int **Input2 = new int*[Height];
    int **Output = new int*[Height];

    for (int i = 0; i < Height; i++) {
        Input1[i] = (int*)malloc(Width * sizeof(int));
        Input2[i] = (int*)malloc(Width * sizeof(int));
        Output[i] = (int*)malloc(Width * sizeof(int));
    }

    // 入力データの初期化
    for (int i = 0; i < Height; i++) {
        for (int j = 0; j < Width; j++) {
            Input1[i][j] = 1; // すべての要素を1に設定
            Input2[i][j] = 2; // すべての要素を2に設定
        }
    }

    // 和を計算
    for (int i = 0; i < Height; i++) {
        for (int j = 0; j < Width; j++) {
            Output[i][j] = Input1[i][j] + Input2[i][j];
        }
    }

    // メモリ解放
    delete[] Input1; delete[] Input2; delete[] Output;

    return;
}

GPU化方針

まず、どの部分をカーネル関数にしてGPUに処理させるかを考えます。GPU を使って処理を高速化するには、「forループを外して並列に処理させる」ことが基本的な方針となります。今回は行列の和を計算している部分のみをカーネル関数に変更することにします。GPU化にあたって障害になりそうな部分を考えます。
今回は入出力データがダブルポインタになっているので、GPU化するためにシングルポインタにする必要があります。

まずは上記の対応を行い、CPUで動作するようなコードを作成します。面倒ですが、この手順を踏むことでデバッグする際に原因の切り分けができます。

void main(){
    int Width = 1024;
    int Height = 1024;
    int ArraySize = Width * Height;
    int *Input1, *Input2, *Output;

    Input1 = (int*)malloc(ArraySize * sizeof(int));
    Input2 = (int*)malloc(ArraySize * sizeof(int));
    Output = (int*)malloc(ArraySize * sizeof(int));

    // 入力データの初期化
    for (int i = 0; i < ArraySize; i++) {
        Input1[i] = 1; // すべての要素を1に設定
        Input2[i] = 2; // すべての要素を2に設定
    }

    // 和を計算
    for (int i = 0; i < ArraySize; i++) {
        Output[i] = Input1[i] + Input2[i];
    }

    // メモリの解放
    free(Input1); free(Input2); free(Output);
}

GPU化

ではGPU化していきます。ソースの全体はこの記事の最後にあります。

メモリ確保

cudaMalloc() を使用してデバイスメモリの確保を行います。ホストメモリと区別がつくようにデバイスメモリの変数には _d をつけておきます。

int *Input1_d, *Input2_d, *Output_d;
// デバイスメモリの確保
cudaMalloc(&Input1_d, ArraySize * sizeof(int));
cudaMalloc(&Input2_d, ArraySize * sizeof(int));
cudaMalloc(&Output_d, ArraySize * sizeof(int));

入力データコピー

ホストで作成した入力データをデバイスに cudaMemcpy() を使ってコピーします。memcpy() 同様、第2引数のデータを第1引数にコピーします。

cudaMemcpy(Input1_d, Input1, ArraySize * sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(Input2_d, Input2, ArraySize * sizeof(int), cudaMemcpyHostToDevice);

カーネル関数作成

行列の和を計算するカーネル関数を実装します。ホスト関数から呼び出すので、global で修飾し、戻り値は void である必要があります。
まず配列の添え字に使用するためのスレッドIDを決定します。
次に行列のサイズを超えたかの判定を行います。基本的に、デバイスメモリは alloc したエリアを超えてアクセスしていてもプログラムが落ちず、気づかずに他のエリアを壊していることがあります。
そのため、カーネル関数内ではインデックスのチェックを行うことをオススメします。あとはカーネル関数で実行したい処理を書きます。

__global__ void ArraySum(int *input1, int* input2, int *output, int width, int height) {

    int index_x = blockDim.x * blockIdx.x + threadIdx.x;
    int index_y = blockDim.y * blockIdx.y + threadIdx.y;

    int index = index_y * width + index_x;

    // 行列のサイズを超えたらreturn
    if (width * height <= index) return;

    output[index] = input1[index] + input2[index];
}

カーネル関数呼び出し

作成したカーネル関数を呼び出します。デバイスのスレッドを指定するときに、dim3型という3次元の変数を使って指定します。数学の座標と同じように、x成分、y成分、z成分の順です。
今回はzは指定しません。(z方向はスレッド数が少ないこともあり、使われないことが多いと思います。)スレッド数の割り当ては 32x4 としています。(THREAD_NUM_X が 32、THREAD_NUM_Y が 4)無駄なスレッドが生まれないように数を調整する必要があります。
カーネル関数の呼び出しは
関数名 <<< ブロック数, スレッド数 >>> (引数);
の形で呼び出します。

dim3 Block(std::ceil(Width / THREAD_NUM_X), std::ceil(Height / THREAD_NUM_Y));
dim3 Thread(THREAD_NUM_X, THREAD_NUM_Y);

// カーネルの実行
ArraySum << <Block, Thread >> > (Input1_d, Input2_d, Output_d, Width, Height);

エラーチェック

カーネル関数のエラーチェックを実装します。普通のプログラムだと戻り値を見てエラーチェックすることが多いと思いますが、カーネル関数の戻り値は void である必要があるため、戻り値で判定することができません。
そこで、 cudaDeviceSynchronize() の戻り値を使います。cudaDeviceSynchronize() はデバイスの同期を行う関数ですが、前のタスクのうち1つでも失敗した場合はエラーを返してくれます。
cudaDeviceSynchronize() でエラーコードを取得し、cudaGetErrorString() でエラーメッセージに変換します。cudaDeviceSynchronize() はデバイスの同期を行うため、使えば使うほど遅くなります。そのため、Debugモードでのみエラーチェックを行うなどしてパフォーマンスに悪影響を与えないように工夫する必要があります。

cudaError_t ret = cudaDeviceSynchronize();
if (ret != cudaSuccess) {
   printf("%s\n", cudaGetErrorString(ret));
}

計算結果のコピー

カーネル関数の結果をホストで参照したい場合は、ホストメモリにコピーします。デバイスへのコピーとは最後の引数が異なることに注意してください。

cudaMemcpy(Output, Output_d, ArraySize * sizeof(int), cudaMemcpyDeviceToHost);

メモリ解放

メモリの解放を忘れずに。

free(Input1); free(Input2); free(Output);
cudaFree(Input1_d); cudaFree(Input2_d); cudaFree(Output_d);

チューニング

GPUを使用して処理を高速化する場合、単にカーネル関数を書けば良いというものではなく、性能を最大限引き出すためのチューニングが必要になります。以下に、コーディングの際に気を付けることを挙げていきます。実装する際に頭に入れておくと、GPUを使ったのに遅い！ということが少なくなると思います。

cudaMalloc のオーバーヘッドは大きい
cudaMalloc はプログラムの途中で何度も呼び出さない方がパフォーマンスが安定します。メモリはどこか1か所で確保しておき、必要に応じて使うと良いです。ワーク用のメモリを複数取っておいて、それらを管理しながら使いまわすという方法も良いかもしれません。
カーネル関数内にif文、for文を入れると遅い
条件分岐によって異なる実行パスを取るスレッドが存在する場合、GPUはそれぞれのパスを順番に実行する必要があります。その結果、全てのスレッドが同じ命令を同時に実行できる場合に比べて、実行時間が長くなります。条件分岐したい場合は、条件によって実行するカーネル関数を変えるなどの工夫を行う必要があります。
メモリ転送は最小限にする
データをデバイスに置いたまますべての処理ができれば良いのですが、どうしてもホストに転送しなければいけない場面もあります。しかし、メモリ転送は同期的な処理なので多用すると遅くなります。(サイズにもよりますが、1ms～2ms かかると思っておくと良いと思います。)メモリ転送しなくて良いように処理を見直す、非同期でコピーできる cudaMemcpyAsync を使うなどの工夫が必要です。
スレッド数の割り当てを最適にする
プロファイラ等を使用しながら最適なスレッド数を設定してください。スレッド数に迷ったら32の倍数を設定しておくことをおすすめします。

まとめ

GPUを使えばどんな処理でも速くなるわけではなく、GPUの性能を最大限引き出せるようなコーディング、チューニングが大切です。

ソースの全体

#include <cuda_runtime.h>
#include "device_launch_parameters.h"
#include <iostream>
#include <stdio.h>
#include <Windows.h>

#define THREAD_NUM_X 32
#define THREAD_NUM_Y 4
#define USE_GPU
#define _TIME

__global__ void ArraySum(int *input1, int* input2, int *output, int width, int height) {

    int index_x = blockDim.x * blockIdx.x + threadIdx.x;
    int index_y = blockDim.y * blockIdx.y + threadIdx.y;

    int index = index_y * width + index_x;

    // 行列のサイズを超えたらreturn
    if (width * height <= index) return;

    output[index] = input1[index] + input2[index];
}

void main() {

#ifdef _TIME
    LARGE_INTEGER    frequency, timer_start, timer_end;
    double          theCompressTime;
#endif

    int Width = 1024;
    int Height = 1024;
    int ArraySize = Width * Height;        // 配列のサイズ
    int *Input1, *Input2, *Output;        // ホストメモリ
    int *Input1_d, *Input2_d, *Output_d;    // デバイスメモリ

    // ホストメモリの確保
    Input1 = (int*)malloc(ArraySize * sizeof(int));
    Input2 = (int*)malloc(ArraySize * sizeof(int));
    Output = (int*)malloc(ArraySize * sizeof(int));

    // デバイスメモリの確保
    cudaMalloc(&Input1_d, ArraySize * sizeof(int));
    cudaMalloc(&Input2_d, ArraySize * sizeof(int));
    cudaMalloc(&Output_d, ArraySize * sizeof(int));

    // 入力データの初期化
    for (int i = 0; i < ArraySize; i++) {
        Input1[i] = 1; // すべての要素を1に設定
        Input2[i] = 2; // すべての要素を2に設定
    }

    // デバイスにデータをコピー
    cudaMemcpy(Input1_d, Input1, ArraySize * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(Input2_d, Input2, ArraySize * sizeof(int), cudaMemcpyHostToDevice);

    dim3 Block(std::ceil(Width / THREAD_NUM_X), std::ceil(Height / THREAD_NUM_Y));
    dim3 Thread(THREAD_NUM_X, THREAD_NUM_Y);

#ifdef _TIME
    QueryPerformanceCounter(&timer_start);
#endif

    // カーネルの実行
    ArraySum << <Block, Thread >> > (Input1_d, Input2_d, Output_d, Width, Height);

#ifdef _TIME
    cudaDeviceSynchronize();
    QueryPerformanceCounter(&timer_end);
    QueryPerformanceFrequency(&frequency);
    theCompressTime = (double)(timer_end.QuadPart - timer_start.QuadPart) / (double)(frequency.QuadPart);
    printf("Time : %lf ms\n", theCompressTime * 1000);
#endif

#ifdef _DEBUG
    // エラーチェック
    cudaError_t ret = cudaDeviceSynchronize();
    if (ret != cudaSuccess) {
        printf("%s\n", cudaGetErrorString(ret));
    }
#endif

    // 結果をホストにコピー
    cudaMemcpy(Output, Output_d, ArraySize * sizeof(int), cudaMemcpyDeviceToHost);

    // メモリの解放
    free(Input1); free(Input2); free(Output);
    cudaFree(Input1_d); cudaFree(Input2_d); cudaFree(Output_d);

    return 0;
}