diff --git a/.obsidian/workspace.json b/.obsidian/workspace.json index efecd2b..dae9d5e 100644 --- a/.obsidian/workspace.json +++ b/.obsidian/workspace.json @@ -58,9 +58,21 @@ "source": true } } + }, + { + "id": "de9326d812f04a36", + "type": "leaf", + "state": { + "type": "markdown", + "state": { + "file": "04. Programming/OpenCL.md", + "mode": "source", + "source": true + } + } } ], - "currentTab": 3 + "currentTab": 4 } ], "direction": "vertical" @@ -118,7 +130,7 @@ "state": { "type": "backlink", "state": { - "file": "03. 專注Study/C++/chrono.md", + "file": "04. Programming/OpenCL.md", "collapseAll": false, "extraContext": false, "sortOrder": "alphabetical", @@ -143,7 +155,7 @@ "state": { "type": "outline", "state": { - "file": "03. 專注Study/C++/chrono.md" + "file": "04. Programming/OpenCL.md" } } }, @@ -194,10 +206,11 @@ "periodic-notes:Open today": false } }, - "active": "bedfb1c3a1e0e733", + "active": "de9326d812f04a36", "lastOpenFiles": [ - "03. 專注Study/C++/C++20.md", "03. 專注Study/C++/chrono.md", + "04. Programming/OpenCL.md", + "03. 專注Study/C++/C++20.md", "03. 專注Study/C++/C++17.md", "04. Programming/QT/timer.md", "00. Inbox/01. TODO.md", @@ -225,7 +238,6 @@ "02. 工作/01. Logitech/QA Sustaining Automation.md", "02. 工作/01. Logitech/Bolide.md", "02. 工作/01. Logitech/AutoStation.md", - "02. 工作/01. Logitech/AE Team.md", "attachments/android_mediacodec_life_cycle.png", "attachments/android_mediacodec_flow.png", "attachments/Pasted image 20230308105856.png", diff --git a/04. Programming/OpenCL.md b/04. Programming/OpenCL.md new file mode 100644 index 0000000..d77df9d --- /dev/null +++ b/04. Programming/OpenCL.md @@ -0,0 +1,166 @@ +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 platforms; +cl::Platform::get(&platforms); +for (size_t i = 0; i < platforms.size(); ++i) { + printf("Name: %s\n", platforms[i].getInfo().c_str()); + printf("Vendor: %s\n", platforms[i].getInfo().c_str()); + printf("Version: %s\n", platforms[i].getInfo().c_str()); + printf("Profile: %s\n", platforms[i].getInfo().c_str()); + printf("Extensions: %s\n", platforms[i].getInfo().c_str()); + + if (platformName.find("NVIDIA") != std::string::npos) { + this->choosenPlatform = platforms[i]; + break; + } +} +``` + +## 2. 選擇 device +```cpp +std::vector 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() << std::endl; + std::cout << " Device vendor: " << device.getInfo() << std::endl; + std::cout << " Device version: " << device.getInfo() << std::endl; + std::cout << " Device profile: " << device.getInfo() << std::endl; + std::cout << " Device extensions: " << device.getInfo() << 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(this->choosenDevice).c_str()); + break; +} +``` +OpenCL 是執行時才編譯，如果編譯錯誤則用 `program.getBuildInfo(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. 讀回資料 \ No newline at end of file