GPUを使った並列プログラミング

はじめに

近年のGPUボードは拡張ディスプレイアダプターにとどまらず、GPUを搭載すると高画質の3Dグラフィックスや人工知能で画像認識ができるようになってきました。実際にどのような処理が行われることでこれらの高機能な処理を実現しているのかを考察します。

GPUの特徴として下記の２点があります。

• 浮動小数点演算を高速に行うことができる
• 複数の処理を並行して実行できる

GPUは特殊な計算処理を専用に行うコアを複数（1000コア以上）搭載しており、浮動小数点演算を高速に並行して演算することができます。そのため光の反射を計算する3Dグラフィックスや、画像の各要素に対応した複数の処理を並行に実行することで画像を判別することができるようになります。

また、GPUのコアも簡単な論理演算や判定処理を行うことができるため、小さいCPUが大量に存在すると考えることができます。

GPU演算の概念

GPUで並列プログラミングを行う上で、プログラムがどのように実行されるのかを現した図が下記のようになります。

CPUとGPUや直接命令のやり取りを行うことができず、CPUとGPUは全く独立して動作しています。

CPUとGPUの間ではGPUボード上のDRAMを介してデータをやり取りすることができ、CPU側から処理してほしいデータをGPUボードのDRAMに書き込み、GPUで処理された結果もDRAMを読み出すことで取得するようになっています。

CPU側で動作するプログラムを main.c とし、GPU側で動作するプログラムを kernel.cl とします。GPU側で動作するプログラムのことを OpenCL では「カーネル」と呼んでおり、カーネルはGPUの初期化時に動的にコンパイルされてGPUで実行可能な状態となります。

GPUの各コアにはインデックス（ｎ）が割り当てられており、プログラムを実行する際に自分が何番目のコアなのかを知ることができます。
GPUの各コアは自分のインデックスに応じたデータをDRAMから取り出し、処理した結果をDRAMに書き込みます。

自分がｎ番目のコアということが分かれば、それぞれのデータを独立して処理することができます。またコア同士は全く独立して動作することができるため、数千のコアが同時に処理することができ最大でコア数分の処理を並行して実行することができます。

プログラムの説明

CPU側の処理 main.c はC言語（実際にはC++）で書きます。これは通常のC言語のプログラムと同じで、コンソールへの文字列表示やファイル入出力、ネットワーク通信などができます。一方 GPU側の処理もC言語で書くのですが、C言語の高度なライブラリは使用することができず、printfすらできません。

それでは順番にプログラムの中心部分を読んでいきます。

main.cpp

    //Part 1 - handle command line options such as device selection, verbosity, etc.
    int platform_id = 0;
    int device_id = 0;

    //detect any potential exceptions
    try {
        //Part 2 - host operations
        //2.1 Select computing devices
        cl::Context context = GetContext(platform_id, device_id);

        //display the selected device
        std::cout << "Runinng on " << GetPlatformName(platform_id) << ", " << GetDeviceName(platform_id, device_id) << std::endl;

        //create a queue to which we will push commands for the device
        cl::CommandQueue queue(context);

        //2.2 Load & build the device code
        cl::Program::Sources sources;

        AddSources(sources, "kernel.cl");

        cl::Program program(context, sources);

        //build and debug the kernel code
        try {
            program.build();
        }
        catch (const cl::Error& err) {
            std::cout << "Build Status: " << program.getBuildInfo<CL_PROGRAM_BUILD_STATUS>(context.getInfo<CL_CONTEXT_DEVICES>()[0]) << std::endl;
            std::cout << "Build Options:\t" << program.getBuildInfo<CL_PROGRAM_BUILD_OPTIONS>(context.getInfo<CL_CONTEXT_DEVICES>()[0]) << std::endl;
            std::cout << "Build Log:\t " << program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(context.getInfo<CL_CONTEXT_DEVICES>()[0]) << std::endl;
            throw err;
        }

        //Part 3 - memory allocation
        //host - input
        std::vector<int> A = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 }; //C++11 allows this type of initialisation
        std::vector<int> B = { 0, 1, 2, 0, 1, 2, 0, 1, 2, 0 };

        size_t vector_elements = A.size();//number of elements
        size_t vector_size = A.size() * sizeof(int);//size in bytes

        //host - output
        std::vector<int> C(vector_elements);

        //device - buffers
        cl::Buffer buffer_A(context, CL_MEM_READ_WRITE, vector_size);
        cl::Buffer buffer_B(context, CL_MEM_READ_WRITE, vector_size);
        cl::Buffer buffer_C(context, CL_MEM_READ_WRITE, vector_size);

        //Part 4 - device operations

        //4.1 Copy arrays A and B to device memory
        queue.enqueueWriteBuffer(buffer_A, CL_TRUE, 0, vector_size, &A[0]);
        queue.enqueueWriteBuffer(buffer_B, CL_TRUE, 0, vector_size, &B[0]);

        //4.2 Setup and execute the kernel (i.e. device code)
        cl::Kernel kernel_add = cl::Kernel(program, "add");
        kernel_add.setArg(0, buffer_A);
        kernel_add.setArg(1, buffer_B);
        kernel_add.setArg(2, buffer_C);

        queue.enqueueNDRangeKernel(kernel_add, cl::NullRange, cl::NDRange(vector_elements), cl::NullRange);

        //4.3 Copy the result from device to host
        queue.enqueueReadBuffer(buffer_C, CL_TRUE, 0, vector_size, &C[0]);

        std::cout << "A = " << A << std::endl;
        std::cout << "B = " << B << std::endl;
        std::cout << "C = " << C << std::endl;
    }
    catch (cl::Error err) {
        std::cerr << "ERROR: " << err.what() << ", " << getErrorString(err.err()) << std::endl;
    }

ここでは行列の加算をGPUを使って並列に実行するというサンプルを取り上げます。
part1 ：ここではGPUボードが複数枚搭載されていた場合にどのGPUボードを使用するのかを定義しています。
part2 ：OpenCL のコンテキストをGPUボードのインデックスとGPU番号から取得します。
2.2 ：GPU で実行する OpenCL の kernel.cl をコンパイルします。
Part3 ：GPU で演算するためのデータを用意し、OpenCL の cl::Buffer でGPU上のメモリを確保します。
Part4(4.1) ：GPU で演算するためのデータを書き込みます。
4.2 ：kernel.cl の add 関数に引数を渡して実行します。
4.3 ：GPU で演算した結果を取得します。

kernel.cl

//a simple OpenCL kernel which adds two vectors A and B together into a third vector C
kernel void add(global const int* A, global const int* B, global int* C) {
    int id = get_global_id(0);
    C[id] = A[id] + B[id];
}

ちなみに kernel.cl は上記のような簡単な関数となっています。特殊な点としては関数定義の際に kernel 識別子を記載する必要があるという点です。

最後に

並列プログラミングで実際にGPUボードでプログラムを実行できた時は感動しました。何の変哲もないタダのコンソールの画面に数字が表示されているだけですが、GPUで並列に処理された結果だと思うと不思議なものがあります。これまでの開発環境の構築などWindows10を何回もインストールしたり、GPUボードのバージョンが合わなくてプログラムが実行できない原因が全く分からなかったりなどの苦難も感動の要因になっているような気もします。

実際にGPUでプログラムを実行することができるようになったので、次回以降は目的の処理をGPUで実行するためにはどうすればよいか、CPU側とGPU側での処理の切り分けなどを検討していきたいと思います。