166 lines
6.5 KiB
Markdown
166 lines
6.5 KiB
Markdown
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));
|
||
if (err != CL_SUCCESS) {
|
||
printf("OpenCL kernel error. err = %d\n", err);
|
||
}
|
||
```
|
||
|
||
## 11. 讀回資料 |