vault backup: 2025-02-10 17:27:57
This commit is contained in:
186
21.01. Programming/OpenCL.md
Normal file
186
21.01. Programming/OpenCL.md
Normal file
@@ -0,0 +1,186 @@
|
||||
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。
|
||||
Reference in New Issue
Block a user