rigayaの日記兼メモ帳

最近のCPUにはたいていGPUがついてきていて、


The Compute Architecture of Intel Processor Graphics Gen9より

まあ、大雑把にだいたい30～40%ぐらいGPUだったりする。このGPUを有効利用するには、OpenCLを使って、GPUで計算するのが良いとされるので、年末も近づき、暇になってきたので少し試してみることにした。

OpenCLは以前にもすこしだけやったことがあるが、すぐに忘れてしまうのでメモ書きを残すことにする。

---

まずはVisual Studio 2015 Communityをインストールし、さらにこのへんからダウンロードしたIntel OpenCL SDKをインストールする。

インストールすると、VC++のプロジェクトテンプレートの中にOpenCLのテンプレートが登場する。



Code Builder Project... を選ぶと、設定が出る。

ここでは、"Use explicit Local Work Group Size"にチェックを入れた。



テンプレートは配列加算のサンプルだが、その入出力の配列のメモリに関する設定。デフォルトは"Use Host Memory"になっているが、これはZero Copyを使って、CPU-GPU間転送を省略するものであり、OpenCLの基本とはいえない。そこで、"Allocate & Copy Memory"に変更した。



Finishとすると、OpenCLのサンプルコードがもう動かす事ができる形で自動生成される。



OpenCLといえば、まず動かすまでに長すぎる呪文詠唱で嫌になるという印象が強かったが、最近では便利になったものだ。

---

とはいえ、どういう呪文を詠唱すればよいのか確認していく。見通しをよくするため、自動生成されたサンプルコードからエラー処理の一切を省き、ある程度簡略化してある。

まず"platform"の取得と選択 FindOpenCLPlatform()

//今回はIntelのGPUを探したいので、
//deviceType = CL_DEVICE_TYPE_GPU
//preferredPlatform = "Intel"
cl_platform_id FindOpenCLPlatform(
  const char* preferredPlatform, cl_device_type deviceType) {
    //platformの数を取得
    cl_uint numPlatforms = 0;
    clGetPlatformIDs(0, NULL, &numPlatforms);

    //実際にplatformのリストを取得
    std::vector platforms(numPlatforms);
    clGetPlatformIDs(numPlatforms, &platforms[0], NULL);

    for (cl_uint i = 0; i < numPlatforms; i++) {
        //platformの名前を取得する
        char buf[4096];
        clGetPlatformInfo(platform, CL_PLATFORM_NAME,
            sizeof(buf), buf, NULL);
        
        //ここではIntelのplatformが見つかったら
        if (strstr(buf, preferredPlatform) != 0) {
            //そのplatformがdeviceType(ここではGPU)を持つか、確認する
            cl_int err = clGetDeviceIDs(
                platforms[i], deviceType, 0, NULL, &numDevices);
            if (err == CL_SUCCESS && numDevices > 0) {
                return platforms[i];
            }
        }
    }
    return NULL;
}

つぎに取得したplatformを使って、opencl "context"を作成し、その"device"を取得する。さらに、そのdeviceに命令を投入するためのキューを作成する。

    //先程の関数でplatformを作成
    cl_platform_id platformId
        = FindOpenCLPlatform("Intel", deviceType);
    
    //コンテキスト情報
    cl_context_properties contextProperties[] = 
        { CL_CONTEXT_PLATFORM, (cl_context_properties)platformId, 0 };
        
    //コンテキストを作成
    ocl->context = clCreateContextFromType(
        contextProperties, deviceType, NULL, NULL, &err);

    //コンテキストのdeviceを取得
    err = clGetContextInfo(ocl->context, CL_CONTEXT_DEVICES,
        sizeof(cl_device_id), &ocl->device, NULL);

    //プロファイリングを有効にしてキューを作成
    cl_command_queue_properties properties = CL_QUEUE_PROFILING_ENABLE;
    ocl->commandQueue = clCreateCommandQueue(
        ocl->context, ocl->device, properties, &err);

次に、kernel(GPUで実行する関数)をビルドして実行可能にする必要がある。通常、kernelはソースコードを文字列として、clCreateProgramWithSource()関数に渡してビルドする。

Intelのサンプルコードのように、ソースファイルからそのまま文字列を読み込んでもいいが、面倒なので、リソースとして実行ファイルに埋め込むことを考える。

OpenCLTest.rcを新規作成し、ここに

CLDATA    KERNEL_DATA DISCARDABLE "Template.cl"

と書いてプロジェクトに登録し、このリソースを対象としてビルドを行い、関数名を指定してkernelを取得する。リソースを対象とする場合、#includeなどは使えないという制約が出るが、そこは…気にしないことにする。

//リソースからOpenCLのカーネルファイルを取得
HMODULE hmdl = GetModuleHandle(NULL);
HRSRC hrsrc = nullptr;
HGLOBAL hdata = nullptr;
const char *source = nullptr;
size_t src_size = 0;
if ( nullptr == (hrsrc = FindResource(hmdl, "CLDATA", "KERNEL_DATA"))
  || nullptr == (hdata = LoadResource(hmdl, hrsrc))
  || nullptr == (source = (const char *)LockResource(hdata))
  || 0       == (src_size = SizeofResource(hmdl, hrsrc))) {
	return CL_INVALID_VALUE;
}

//ソースから"program"を作成
ocl->program = clCreateProgramWithSource(ocl->context,
    1, (const char**)&source, &src_size, &err);

//ビルドを実行
clBuildProgram(ocl->program, 1, &ocl->device, "", NULL, NULL);

//ビルドした関数のうち、"Add"関数をkernelとして取得
ocl->kernel = clCreateKernel(ocl->program, "Add", &err);

さて、配列加算 C = A + B を行うが、そのサイズを width * height とする。この時、それぞれの配列について、CPU側の配列とは別に、GPU側にバッファを作成してやる必要がある。

ここで、自動生成されたsampleプログラムは、clEnqueueMapBufferを使った効率的な方法を用いていたが、ここではより基本的な、明示的にバッファへの転送を行うclEnqueueWriteBuffer / clEnqueueReadBuffer を用いた方法を見ていく。

int CreateBufferArguments(
    ocl_args_d_t *ocl,
    cl_int* inputA,  //CPU側の入力配列へのポインタ 
    cl_int* inputB,  //CPU側の入力配列へのポインタ 
    cl_int* outputC, //CPU側の出力配列へのポインタ
    cl_uint width, cl_uint height //配列のサイズ
) {
    cl_int err = CL_SUCCESS;
    ocl->srcA = clCreateBuffer(ocl->context,
        CL_MEM_READ_ONLY, //GPUからはこのメモリに対し、読み込みのみ行う
        sizeof(cl_uint) * width * height, //配列のサイズ
        nullptr,
        &err);  //エラー情報を受け取る

    ocl->srcB = clCreateBuffer(ocl->context,
        CL_MEM_READ_ONLY, //GPUからはこのメモリに対し、読み込みのみ行う
        sizeof(cl_uint) * width * height, //配列のサイズ
        nullptr,
        &err);  //エラー情報を受け取る

    ocl->dstMem = clCreateBuffer(ocl->context,
        CL_MEM_ALLOC_HOST_PTR, ////GPUからはこのメモリに対し、書き込みのみ行う
        sizeof(cl_uint) * width * height, //配列のサイズ
        nullptr,   //CPUから転送するものはない
        &err);

    //inputAから作成したバッファに転送
    err = clEnqueueWriteBuffer(ocl->commandQueue, //
        ocl->srcA, //転送先
        CL_FALSE,  //転送が終了するまで待機するか -> しない
        0,         //オフセット
        sizeof(cl_uint) * arrayWidth * arrayHeight, //転送サイズ
        inputA,   //転送元
        0,         //この関数が待機すべきeventの数
        nullptr,   //この関数が待機すべき関数のリストへのポインタ
        nullptr);  //この関数の返すevent

    //inputBから作成したバッファに転送
    err = clEnqueueWriteBuffer(ocl->commandQueue, //
        ocl->srcB, //転送先
        CL_FALSE,  //転送が終了するまで待機するか -> しない
        0,         //オフセット
        sizeof(cl_uint) * arrayWidth * arrayHeight, //転送サイズ
        inputB,    //転送元
        0,         //この関数が待機すべきeventの数
        nullptr,   //この関数が待機すべき関数のリストへのポインタ
        nullptr);  //この関数の返すevent

    return err;
}

ここからいよいよ関数を実行…の前に、関数の引数を一つずつ設定する必要がある。引数がたくさんあると、まあ、簡単に悲惨なコードが出来上がるわけで、正直、このあたりがOpenCLの、CUDAと比べてなかなかに面倒なところだ。

引数を積み終わった後、clEnqueueNDRangeKernelにより、指定のカーネルをGPUで実行するよう指示する。

//関数引数を設定
clSetKernelArg(ocl->kernel, 0, sizeof(cl_mem), (void *)&ocl->srcA);
clSetKernelArg(ocl->kernel, 1, sizeof(cl_mem), (void *)&ocl->srcB);
clSetKernelArg(ocl->kernel, 2, sizeof(cl_mem), (void *)&ocl->dstMem);

//work sizeの指定
//ここでは1要素に対して1 work item
//またグループあたり1 work item (実は効率的でない)
//width * heightの2次元でwork itemを作成
size_t globalWorkSize[2] = { width, height };
size_t localWorkSize[2] = { 1, 1 };

//タスクをキューに積む
clEnqueueNDRangeKernel(
    ocl->commandQueue, //タスクを投入するキュー
    ocl->kernel,       //実行するカーネル
    2,                 //work sizeの次元
    NULL,              //NULLを指定すること
    globalWorkSize,    //全スレッド数
    localWorkSize,     //1グループのスレッド数
    0,                 //この関数が待機すべきeventの数
    NULL,              //この関数が待機すべき関数のリストへのポインタ
    NULL);             //この関数の返すevent

work size(work itemの数)で、並列処理する数を指定し、3次元まで設定できる。

work itemは階層構造になっており、通常は複数のwork itemがwork groupを形成し、work groupが集まってwork item全体を構成する。OpenCL 1.2までは各次元のglobalWorkSizeは各次元のlocalWorkSizeで割りきれる必要があることに注意する。

この時、work group内の個々のwork itemのIDがlocal idで、groupによらない一意に決まるIDをglobal id という。下記の図はwork itemが1次元の場合である。



実際には、work itemは3次元まで設定可能になっていて、サンプルコードは2次元の例になっている。

同じwork groupのwork itemは、GPUで実行される際に必ず同じ演算ブロック(Intelでいうsubslice、NVIDIAでいうSM)で実行され、__localメモリ(IntelでいうShared Local Memory、NVIDIAでいうsharedメモリ)を共有する。また、連続する8, 16, または32のwork itemは同時に実行されるという特徴があり、こうした特徴はメモリアクセス効率を考える際に重要となる。

GPU上で実行するkernelコードはこのようになっていて、単なる配列の加算である。

__kernel void Add(__global int* pA, __global int* pB, __global int* pC) {
    const int x     = get_global_id(0); //0次元目についてglobal idを取得
    const int y     = get_global_id(1); //1次元目についてglobal idを取得
    const int width = get_global_size(0);

    const int id = y * width + x;

    pC[id] = pA[id] + pB[id];
}

width * height 分の各work itemがkernelを実行することになるから、2次元のwork itemの指定により、CPUでいうと以下のような2重ループを並列化したことに等しい。

for (int y = 0; y < height; y++) { //<- 1次元目の並列化
    for (int x = 0; x < width; x++) { //<- 0次元目の並列化
        const int id = y * width + x;
        pC[id] = pA[id] + pB[id];
    }
}

さて、clEnqueueNDRangeKernelにより、指定のカーネルをGPUで実行するよう指示したが、これはタスクをキューに積んだだけなので、GPUが実際にこのkernelを実行し、完了するには時間がかかる。キューのすべてのタスクが終了するまで待機するには、キューを指定してclFinish()を呼ぶ。

clFinish(ocl->commandQueue);

その後、GPUの計算結果を回収する。出力情報のあるGPUバッファをclEnqueueMapBuffer()によりmappingし、mappingされたCPU側の配列へのポインタが得られる。内部でGPUメモリ→CPUメモリへの転送が行われる。

    //inputBから作成したバッフからoutputCに転送
    err = clEnqueueReadBuffer(ocl->commandQueue, //
        ocl->dstMem, //転送元
        CL_FALSE,    //転送が終了するまで待機するか -> しない
        0,         //オフセット
        sizeof(cl_uint) * arrayWidth * arrayHeight, //転送サイズ
        outputC,   //転送先
        0,         //この関数が待機すべきeventの数
        nullptr,   //この関数が待機すべき関数のリストへのポインタ
        nullptr);  //この関数の返すevent
    
    //clEnqueueMapBufferの終了を待機
    clFinish(ocl->commandQueue);

     //outputCに計算結果が入っていて、CPUから操作できる

---

というわけで、ここまででやっと配列の加算を行うことができた。

まず、deviceを開き、OpenCL環境を初期化するところまでをまとめると、こんな感じ。



そして、OpenCLバッファの確保から、出力の取り出しまでをまとめると以下のようになる。



実際にはエラー処理などもあるので、結構長いコードになってしまう。CUDAと比べると、とにかく呪文(手続き)が長い、というのが汎用性・標準化を重視したOpenCLのつらいところだと思う。

---

今回のコードはこのへん。

---

続き>>

---

OpenCLメモリスト
OpenCLメモ (その1) - かんたんな計算
OpenCLメモ (その2) - Intel GPUの構造
OpenCLメモ (その3) - work sizeの調整
OpenCLメモ (その4) - 転送(コピー)の排除 (USE_HOST_PTR)
OpenCLメモ (その5) - 転送(コピー)の排除 (SVM : OpenCL 2.0)
OpenCLメモ (その6) - imageの使用
OpenCLメモ (その7) [終] - reductionとshared local memory(SLM)の使用

---