Obsidian-Main/04. Programming/OpenCL.md

OpenCL 可以使用 GPU 來實現異質運算。

OpenCL使用上的概念：

1. 選擇 platform
2. 由 platform 來選擇device
3. 由 device 來建立 context，透過 context 才能控制 device
4. 由 context 來建立 program，program 即是要執行在 GPU 上面的程式
5. 編譯 program
6. 建立 kernel 來執行 program
7. 建立 queue，用來跟 program 溝通
8. 建立 cl::Buffer，cl::Buffer 是 GPU 能使用的 memory，把 PC 端的資料 copy 進去
9. 呼叫 kernel.setArg(cl::Buffer) 來設定參數
10. 呼叫 queue.enqueueNDRangeKernel() 來執行 kerne l並設定工作組大小
11. 用 queue.enqueueReadBuffer() 來讀回處理好的資料

1. 選擇 platform

std::vector<cl::Platform> platforms;
cl::Platform::get(&platforms);
for (size_t i = 0; i < platforms.size(); ++i) {
    printf("Name: %s\n", platforms[i].getInfo<CL_PLATFORM_NAME>().c_str());
    printf("Vendor: %s\n", platforms[i].getInfo<CL_PLATFORM_VENDOR>().c_str());
    printf("Version: %s\n", platforms[i].getInfo<CL_PLATFORM_VERSION>().c_str());
    printf("Profile: %s\n", platforms[i].getInfo<CL_PLATFORM_PROFILE>().c_str());
    printf("Extensions: %s\n", platforms[i].getInfo<CL_PLATFORM_EXTENSIONS>().c_str());

    if (platformName.find("NVIDIA") != std::string::npos) {
        this->choosenPlatform = platforms[i];
        break;
    }
}

2. 選擇 device

std::vector<cl::Device> clDevices;

platform.getDevices(CL_DEVICE_TYPE_GPU, &clDevices); // 檢查裝置數目
if (!clDevices.empty()) {
    choosenDevice = clDevices[0];
    std::cout << "Devices size = " << clDevices.size() << std::endl;
    // Print device info
    std::cout << "********** DEVICE **********\n";
    for (auto& device : clDevices) {
        std::cout << "    Device name: " << device.getInfo<CL_DEVICE_NAME>() << std::endl;
        std::cout << "    Device vendor: " << device.getInfo<CL_DEVICE_VENDOR>() << std::endl;
        std::cout << "    Device version: " << device.getInfo<CL_DEVICE_VERSION>() << std::endl;
        std::cout << "    Device profile: " << device.getInfo<CL_DEVICE_PROFILE>() << std::endl;
        std::cout << "    Device extensions: " << device.getInfo<CL_DEVICE_EXTENSIONS>() << std::endl;
        std::cout << "\n";
    }
}

3. 建立 context

auto context = cl::Context({ this->choosenDevice });

4. 建立 program

auto program = cl::Program(this->context, textCode);

第1個參數就是前一步建立好的 context，第二個參數是你要在GPU上執行的程式碼。型別是 std::string。

這裡 textCode 的內容是：

static std::string textCode =
R"(
    __kernel void HistogramCalculator(
        __global const uchar* buffer, 
        const uint width, 
        const uint height, 
        __global uint* grayHistogram,
        __global uint* redHistogram,
        __global uint* greenHistogram,
        __global uint* blueHistogram)
    {
        uint id = get_global_id(0) + get_global_id(1) * width;

        uchar b = buffer[id * 3];
        uchar g = buffer[id * 3 + 1];
        uchar r = buffer[id * 3 + 2];
        uchar gray = (r + g + b) / 3;

        atomic_inc(&grayHistogram[gray]);
        atomic_inc(&redHistogram[r]);
        atomic_inc(&greenHistogram[g]);
        atomic_inc(&blueHistogram[b]);
    }
)";

5. 編譯 program

if (program.build() != CL_SUCCESS) {
    printf("[ERROR] Fail to build program.\n");
    printf("  LOG: %s\n", this->program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(this->choosenDevice).c_str());
    break;
}

OpenCL 是執行時才編譯，如果編譯錯誤則用 program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(this->choosenDevice).c_str()) 取得錯誤訊息。

6. 建立 kernel

cl_int kernel_creation_result = 0;
this->kernel = cl::Kernel(program, kernelName.c_str(), &kernel_creation_result);
if (this->kernel() == NULL) {
    // Failed to create kernel object
    printf("Error: Failed to create kernel object! kernel_creation_result = %d\n", kernel_creation_result);
    break;
}

這邊將 program 與 kernel 連接起來，由 kernel 來執行。要注意的是第二個參數就是你要在 GPU 上執行的程式碼的「函式名稱」。不可以不一樣，否則這裡會報錯。 第三個參數用來接受錯誤代碼。

7. 建立 queue

auto queue = cl::CommandQueue(this->context, this->choosenDevice);

到了這邊，OpenCL 的執行單元已經建立完成，可以開始執行了。

8. 建立 cl::Buffer

因為是執行在 GPU 上，所以必須透過 cl::Buffer 來將資料送進 GPU。 參考OpenCL#4. 建立 program的 textCode ，可以知道我們要傳入7個參數，第一個是圖片的buffer，第二個與第三個分別是圖片的寬與高，第四、五、六、七參數則用來輸出算好的histogram。 寬與高是 uint32_t，可以直接傳入，但是buffer不行直接傳，所以要先建立 cl::Buffer。 以下建立要傳到 GPU 的 cl::buffer：

cl::Buffer bufferCl(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, imageWidth * imageHeight * byteDepth, imageBuffer.get());

第一個參數是OpenCL#3. 建立 context所建立的 context。 第二個參數是記憶體的屬性，因為不需要 GPU 寫回，所以這裡是 CL_MEM_READ_ONLY。 第三個參數是記憶體的長度。 第四個參數是記憶體的address。

再來建立要讀回的 cl::buffer，用 grayHistogram 當作範例，其他三個行為都一樣：

cl::Buffer clGrayHistogram(context, CL_MEM_WRITE_ONLY, 256 * sizeof(uint32_t));

第一個參數是OpenCL#3. 建立 context所建立的 context。 第二個參數是記憶體的屬性，這個 clGrayHistogram 因為需要 GPU 寫回，所以記憶體屬性是 CL_MEM_WRITE_ONLY。 第三個參數是記憶體的長度。 注意沒有第四個參數。

9. 設定參數

用OpenCL#6. 建立 kernel所建立的 kernel 來傳入剛剛建立好的 buffer。

this->kernel.setArg(0, bufferCl);
this->kernel.setArg(1, imageWidth);
this->kernel.setArg(2, imageHeight);
this->kernel.setArg(3, clGrayHistogram);
this->kernel.setArg(4, clRedHistogram);
this->kernel.setArg(5, clGreenHistogram);
this->kernel.setArg(6, clBlueHistogram);

注意這裡有7個參數，跟OpenCL#4. 建立 program的 textCode所定義的一樣。

10. 執行 kernel 並設定工作組大小

auto err = queue.enqueueNDRangeKernel(this->kernel, cl::NullRange, cl::NDRange(imageWidth, imageHeight), cl::NullRange);
if (err != CL_SUCCESS) {
    printf("OpenCL kernel error. err = %d\n", err);
}

第一個參數是OpenCL#6. 建立 kernel所建立的kernel。 第二個參數是偏移量，我們假設它在所有維度上都是 0。cl::NullRange 對象將滿足該 0 規範。 第三個參數是全局大小，它指定希望執行與內核對象K關聯的內核源代碼中指定的工作項 第四個參數是本地大小，它指定應將多少工作項分組到一個工作組中。

重點在於第三個參數，因為 GPU 可以平行運算，這裡指定說我們的資料需要幾個

11. 讀回資料