# OpenCL 學習筆記 （ㄧ）： 第一個 OpenCL 程式 ###### tags: `OpenCL` 近年來由於深度學習的成功，使得 GPU 計算變得異常火熱，透過 GPU 強大的平行運算能力，通常能夠提供 10-20 倍以上運算速度，增益效果非常明顯。目前最火紅的 GPU 語言為 CUDA 和 OpenCL，本文將只探討 OpenCL 要如何使用？早期的 GPU 是專注在圖形運算，如果要用 GPU 作運算必須利用圖形程式庫，例如 OpenGL ，雖然不是不能計算，但在效率和彈性上都大大受限，直到 2007 年 Nvidia 發布 CUDA ，它能以極高的彈性做任何運算，這不是專為單一任務設計的計算，因此被稱為通用計算。CUDA 作為最早的通用計算語言，它的影響力也是所有語言裡最大的，還影響到後來 OpenCL 的設計方法，它們兩個都能做通用計算，不一樣的是， OpenCL 在設計之初就盡可能兼容不同平台和硬體的架構，使相同代碼能在 CPU 和 GPU 甚至在其他特殊平台上運行，這也是 OpenCL 最大的優勢之一。 ## OpenCL 的基本描述 OpenCL 是一套以由 Khronos Group 開發維護的標準，所以它可以在各種平台上運行，包括 CPU、GPU 或其它支援的平台，雖然 OpenCL 能夠不同裝置上執行，但 OpenCL 設計之初主要是為了平行裝置，即是我們常說的 GPU ，因此 OpenCL 的模型架構跟 GPU 有莫大關係。首先我們會要先了解一點，在歷史上 GPU 這個概念是比較晚才出來的，GPU 的前身叫做圖形加速器，它是用來輔助 CPU 加速圖形計算的，所以它並不是 CPU 計算單元的一部分，這類計算單元我們稱為外部裝置。外部裝置有下列特性，等等我們在寫程式時會看到 * 外部設備會有不同平台（AMD 或 Nvidia），使用前要先告訴 OpenCL 我們使用甚麼平台和設備 * GPU 需要從主記憶體傳送資料到 GPU 本身的記憶體，不能直接使用主記憶體資料 * GPU 和 CPU 計算是非同步的，在使用 GPU 記憶體的資料之前，要先確定 GPU 已經完成計算 ## 第一個 OpenCL 程式 以下的程式實做以 Ubuntu 為主，其它平台的環境搭建請參考網路資料 #### 第零步 下載並安裝 OpenCL 在終端機依序輸入 $ sudo apt install clinfo $ sudo apt install ocl-icd-libopencl1 $ sudo apt install opencl-headers $ sudo apt install ocl-icd-opencl-dev 完成! #### 第一步 引入頭文件 ```cpp= #include <CL/cl.hpp> ``` 這裡可能會有人有疑問 cl.hpp 和以下有什麼不同 ```cpp= #include <CL/cl.h> ``` 因為 OpenCL 的是以 C API 的形式存在，```cl.h``` 定義 OpenCL C 的 API，相比與用 C++ 封裝的 ```cl.hpp``` ，OpenCL C 的界面使用起來比較麻煩（但優點是文檔比較齊全），因此這裡推薦使用 ```cl.hpp```。 如果你的 OpenCL 驅動大於 2.0 ，會出現警告 ```#pragma message: This version of the OpenCL Host API C++ bindings is deprecated, please use cl2.hpp instead.``` ，則請用 ```cl2.hpp``` 替換 ```cl.hpp``` 頭文件相除警告。 #### 第二步 獲取平台 (platform) 和裝置 (device) ```cpp= #include <CL/cl.hpp> #include <iostream> #include <vector> int main(int argc, char** argv) { std::vector<cl::Platform> platforms; cl::Platform::get(&platforms); // 檢查平台數目 if (platforms.empty()) { std::cerr << "No platforms!" << std::endl; return -1; } cl::Platform platform = platforms[0]; std::vector<cl::Device> Devices; platform.getDevices(CL_DEVICE_TYPE_GPU, &Devices); // 檢查裝置數目 if (Devices.empty()) { std::cerr << "No Devices!" << std::endl; return -1; } cl::Device device = Devices[0]; ...... } ``` OpenCL 能夠兼容不同平台和裝置，必須要先收集所有可用的硬體。一般情況下，平台就是你目前使用的電腦，裝置就是你電腦的 CPU 和 GPU，設備這邊設定 ```CL_DEVICE_TYPE_GPU``` ，限制使用 GPU 。 #### 第三步 產生 context ```cpp= #include <CL/cl.hpp> #include <vector> int main(int argc, char** argv) { ...... cl::Context context({device}); ...... } ``` 在 OpenCL 裡的架構裡，```context``` 是所有裝置的溝通橋樑，對 ```device``` 的操作都要通過 ```context```。 #### 第四步 撰寫 OpenCL 的 kernel ```cpp= #include <string> static std::string opencl_kernel = R"( __kernel void vecadd ( __global int *A, __global int *B, __global int *C ) { int id = get_global_id(0); C[id] = A[id] + B[id]; } )"; ``` OpenCL 是典型的 SIMD，同一個程式會被多的核心執行，這段代碼代表 A 陣列和 B 陣列的值相加到 C 陣列中，而不同核心要靠 ```get_global_id``` 決定拿取拿陣列那一個位置的， ```get_global_id``` 功能是獲取執行緒編號，假設我們設定 100 個執行緒，那執行緒編號就會從 0 數到 99，所以這段程式碼就是編號 0 數到 99 的核心同時把 A 陣列 B 陣列相加（注意這裡的核心數目是虛擬的，實際上一次只能同時執行 40 個左右的核心）。 #### 第五步 編譯 Kernel ```cpp= #include <CL/cl.hpp> #include <iostream> #include <vector> #include <string> int main(int argc, char** argv) { ...... // 通過 context 傳送未編譯的 kernel cl::Program program(context, opencl_kernel); // 對所有 context 內裝置編譯程式 if (program.build() != CL_SUCCESS) { std::cerr << "Fail to build" << std::endl; return -1; } cl::Kernel vec_add(program, "vecadd"); ...... } ``` OpenCL 屬於執行期間編譯， ```vecadd``` 是 ```kernel``` 的函數名稱，裡面的程式碼尚未經過編譯， ```cl::Program``` 會先將文字讀入 ```context```，```program.build()``` 會對整段文字編譯，最終用 ```cl::Kernel``` 加載編譯好的 kernel，其中 ```vecadd``` 為 kernel 的名稱。 #### 第六步 設定 buffer 並連接 ```cpp= #include <CL/cl.hpp> #include <vector> #include <string> int main(int argc, char** argv) { ...... // 所有 buffer 往 context 連接 cl::Buffer buffer_A(context, CL_MEM_READ_WRITE, sizeof(int) * 100); cl::Buffer buffer_B(context, CL_MEM_READ_WRITE, sizeof(int) * 100); cl::Buffer buffer_C(context, CL_MEM_READ_WRITE, sizeof(int) * 100); vec_add.setArg(0, buffer_A); // 連接 buffer A 到第一個參數位置 vec_add.setArg(1, buffer_B); // 連接 buffer B 到第二個參數位置 vec_add.setArg(2, buffer_C); // 連接 buffer C 到第三個參數位置 ...... } ``` 由於 GPU 和 CPU 的記憶體不共通，GPU 只能使用自己的記憶體，```cl::Buffer``` 代表 GPU 的記憶體。 ```setArg``` 連接 ```kernel``` 和 ```buffer``` 讓資料能進入 function。 #### 第七步 產生 CommandQueue ```cpp= #include <CL/cl.hpp> #include <vector> #include <string> int main(int argc, char** argv) { ...... cl::CommandQueue queue(context, device); ...... } ``` 剛剛提到從我們所有操作都是針對 ```context``` ，但還是需要指定一實際的裝置執行指令，```cl::CommandQueue``` 可以指定哪些 GPU 執行指令。 #### 第八步 運行 ```cpp= #include <CL/cl.hpp> #include <vector> #include <string> int main(int argc, char** argv) { ...... std::vector<int> A(100, 1); std::vector<int> B(100, 2); std::vector<int> C(100, 0); queue.enqueueWriteBuffer(buffer_A, CL_FALSE, 0, sizeof(int)*100, A.data()); //寫入資料 queue.enqueueWriteBuffer(buffer_B, CL_FALSE, 0, sizeof(int)*100, B.data()); //寫入資料 queue.enqueueNDRangeKernel(vec_add, cl::NullRange, cl::NDRange(100), // 設定 100 個執行緒 cl::NullRange); queue.enqueueReadBuffer(buffer_C, CL_FALSE, 0, sizeof(int)*100, C.data()); // 讀出資料 queue.finish(); // 同步 return 0; } ``` 通過 ```cl::CommandQueue``` 執行程式，需要注意的是 ```cl::CommandQueue``` 操作全部都是非同步指令，它會將命令先堆積在自己的 Queue 裡，等到當指令完成後才會從 Queue 中拿下一個指令，因此在使用 ```buffer``` 裡的參數前需要先用 ```finish()``` 同步，等待 Queue 裡的指令都被執行完畢，才能確定 ```buffer``` 已經被更新。 #### 第九步編譯 編譯的參數如下 $ g++ main.cc -o first_opencl -std=c++11 -lOpenCL 之後運行 $ ./first_opencl #### 完整代碼 ```cpp= #include <iostream> #include <vector> #include <string> #include <CL/cl.hpp> static std::string opencl_kernel = R"( __kernel void vecadd ( __global int *A, __global int *B, __global int *C ) { int id = get_global_id(0); C[id] = A[id] + B[id]; } )"; int main( int argc , char **argv ) { std::vector<cl::Platform> platforms; cl::Platform::get(&platforms); if (platforms.empty()) { std::cerr << "No platforms!" << std::endl; return -1; } cl::Platform platform = platforms[0]; std::vector<cl::Device> Devices; platform.getDevices(CL_DEVICE_TYPE_GPU, &Devices); if (Devices.empty()) { std::cerr << "No Devices!" << std::endl; return -1; } cl::Device device = Devices[0]; std::cout << "Device : " << device.getInfo<CL_DEVICE_NAME>() << std::endl; cl::Context context({device}); cl::Program program(context, opencl_kernel); if (program.build() != CL_SUCCESS) { std::cerr << "Fail to build" << std::endl; return -1; } std::vector<int> A(100, 1); std::vector<int> B(100, 2); std::vector<int> C(100, 0); cl::Buffer buffer_A(context, CL_MEM_READ_WRITE, sizeof(int)*100); cl::Buffer buffer_B(context, CL_MEM_READ_WRITE, sizeof(int)*100); cl::Buffer buffer_C(context, CL_MEM_READ_WRITE, sizeof(int)*100); cl::CommandQueue queue(context, device); cl::Kernel vec_add(program, "vecadd"); vec_add.setArg(0, buffer_A); vec_add.setArg(1, buffer_B); vec_add.setArg(2, buffer_C); queue.enqueueWriteBuffer(buffer_A, CL_FALSE, 0, sizeof(int)*100, A.data()); queue.enqueueWriteBuffer(buffer_B, CL_FALSE, 0, sizeof(int)*100, B.data()); queue.enqueueNDRangeKernel(vec_add, cl::NullRange, cl::NDRange(100), cl::NullRange); queue.enqueueReadBuffer(buffer_C, CL_FALSE, 0, sizeof(int)*100, C.data()); queue.finish(); for (const auto &v : A){ std::cout << v << " "; } std::cout << std::endl; for (const auto &v : B){ std::cout << v << " "; } std::cout << std::endl; for (const auto &v : C){ std::cout << v << " "; } std::cout << std::endl; return 0; } ``` ## Linking [下一篇](https://hackmd.io/@yrHb-fKBRoyrKDEKdPSDWg/BJCUHajju)