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
```cpp
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
```cpp
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
```cpp
auto context = cl::Context({ this->choosenDevice });
```

## 4. 建立 program
```cpp
auto program = cl::Program(this->context, textCode);
```
第1個參數就是前一步建立好的 context，第二個參數是你要在GPU上執行的程式碼。型別是 `std::string`。

這裡 `textCode` 的內容是：
```cpp
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
```cpp
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 
```cpp
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
```cpp
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`：
```cpp
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 當作範例，其他三個行為都一樣：
```cpp
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。
```cpp
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 並設定工作組大小
```cpp
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. 讀回資料
```cpp
grayHistogram.resize(sizeof(uint32_t) * 256);
err = queue.enqueueReadBuffer(clGrayHistogram, CL_TRUE, 0, 256 * sizeof(uint32_t), grayHistogram.data());
if (err != CL_SUCCESS) {
    printf("OpenCL read clGrayHistogram error.\n");
}
```
第一行的 `grayHistogram` 是讀回 CPU 的記憶體。
[`enqueueReadBuffer`](https://registry.khronos.org/OpenCL/sdk/1.2/docs/man/xhtml/clEnqueueReadBuffer.html)的參數解釋：
第一個參數就是[[OpenCL#8. 建立 `cl::Buffer`]]所建立的其中一個 buffer。
第二個參數指定是否 **blocking_read**。
第三個參數是 `cl::Buffer` 的偏移量。
第四個參數要讀取的長度。
第五個參數是要寫入的 memory address。