目前最火紅的GPU 語言為CUDA 和OpenCL，本文將只探討OpenCL 要如何使用？早期的GPU 是專注在圖形運算，如果要用GPU 作運算必須利用繪圖介面，例如OpenGL ，雖然不是不能 ...
Published
LinkedwithGitHub
Like1
Bookmark
Subscribe
Edit
#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能夠不同裝置上執行，但OpenCL設計之初主要是為了GPU，OpenCL的基本特性跟GPU有莫大關係。
首先我們會要先了解一點GPU在整個電腦架構中，它不是必需品，一些不需要圖形界面的電腦，CPU就足夠了，GPU作為一個獨立於電腦的物件，OpenCL必須考慮這點，因此在GPU實際執行程式之前，我們必須將資料送進GPU上，甚至還需要在GPU再次編譯程式。
而以上特性我們可以簡單歸納出幾點特性
1.OpenCL通用不同平台和設備，所以使用前要先告訴OpenCL我們使用甚麼平台和設備。
2.GPU設備是完全獨立的，在計算前必須先將資源從host轉到GPU裡。
##第一個OpenCL程式
在實際講解其他OpenCL特性之前，直接實做一個簡單的程式有助於理解。
####第零步下載並安裝OpenCL
在終端機依序輸入
$sudoaptinstallclinfo
$sudoaptinstallocl-icd-libopencl1
$sudoaptinstallopencl-headers
$sudoaptinstallocl-icd-opencl-dev
完成!
####第一步我們引入頭文件
```cpp=
#include
```
這裡可能會有人有疑問cl.hpp和以下有什麼不同
```cpp=
#include
```
因為OpenCL的是以CAPI的形式存在，cl.h定義OpenCLC的API，相比與用C++封裝的cl.hpp，OpenCLC的界面比較不方便使用，而且包含許多你可能用不到但又一定要寫的參數，因此這裡推薦使用cl.hpp。
####第二步獲取平台(platform)和裝置(device)
```cpp=
#include
#include
#include
intmain(intargc,char**argv){
std::vector<:platform>platforms;
cl::Platform::get(&platforms);//檢查平台數目
if(platforms.empty()){
std::cerr<Devices;
platform.getDevices(CL_DEVICE_TYPE_GPU,&Devices);//檢查裝置數目
if(Devices.empty()){
std::cerr<
#include
intmain(intargc,char**argv){
......
cl::Contextcontext({device});
......
}
```
在OpenCL裡的架構裡，context是整個執行的虛擬裝置，device是實際執行的裝置，基本上所有的執行指令都是對context操作的。
####第四步撰寫OpenCL的kernel
```cpp=
#include
staticstd::stringopencl_kernel=
R"(
__kernelvoidvecadd
(
__globalint*A,
__globalint*B,
__globalint*C
)
{
intid=get_global_id(0);
C[id]=A[id]+B[id];
}
)";
```
這段代碼代表A[:]和B[:]的vector依序相加到C[:]，其中get_global_id是獲取執行緒編號，假設有100個執行緒，那執行緒編號就會從0數到99。
####第五步編譯Kernel
```cpp=
#include
#include
#include
#include
intmain(intargc,char**argv){
......
cl::Programprogram(context,opencl_kernel);
if(program.build()!=CL_SUCCESS){//編譯程式
std::cerr<
#include
#include
intmain(intargc,char**argv){
......
cl::Bufferbuffer_A(context,CL_MEM_READ_WRITE,sizeof(int)*100);
cl::Bufferbuffer_B(context,CL_MEM_READ_WRITE,sizeof(int)*100);
cl::Bufferbuffer_C(context,CL_MEM_READ_WRITE,sizeof(int)*100);
vec_add.setArg(0,buffer_A);//連接bufferA到第一個參數位置
vec_add.setArg(1,buffer_B);//連接bufferB到第二個參數位置
vec_add.setArg(2,buffer_C);//連接bufferC到第三個參數位置
......
}
```
由於GPU和CPU的記憶體不共通，GPU只能使用自己的記憶體，cl::Buffer代表GPU的記憶體。
setArg連接kernel和buffer讓資料能進入function。
####第七步產生CommandQueue
```cpp=
#include
#include
#include
intmain(intargc,char**argv){
......
cl::CommandQueuequeue(context,device);
......
}
```
剛剛提到從我們所有操作都是針對context，但還是需要指定一實際的裝置執行指令，cl::CommandQueue的一個主要目的就是指定執行裝置。
####第八步運行
```cpp=
#include
#include
#include
intmain(intargc,char**argv){
......
std::vectorA(100,1);
std::vectorB(100,2);
std::vectorC(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),//設定執行緒數目
cl::NullRange);
queue.enqueueReadBuffer(buffer_C,CL_FALSE,0,
sizeof(int)*100,C.data());//讀出資料
queue.finish();//同步
return0;
}
```
通過cl::CommandQueue執行程式，需要注意的是cl::CommandQueue操作全部都是非同步指令，在使用buffer裡的參數前需要先同步。
####第九步編譯
編譯的參數如下
$g++main.cc-ofirst_opencl-std=c++11-lOpenCL
運行
$./first_opencl
####完整代碼
```cpp=
#include
#include
#include
#include
staticstd::stringopencl_kernel=
R"(
__kernelvoidvecadd
(
__globalint*A,
__globalint*B,
__globalint*C
)
{
intid=get_global_id(0);
C[id]=A[id]+B[id];
}
)";
intmain(intargc,char**argv){
std::vector<:platform>platforms;
cl::Platform::get(&platforms);
if(platforms.empty()){
std::cerr<Devices;
platform.getDevices(CL_DEVICE_TYPE_GPU,&Devices);
if(Devices.empty()){
std::cerr<()<<:endl cl::contextcontext cl::programprogram if std::cerr return-1 std::vector>A(100,1);
std::vectorB(100,2);
std::vectorC(100,0);
cl::Bufferbuffer_A(context,CL_MEM_READ_WRITE,sizeof(int)*100);
cl::Bufferbuffer_B(context,CL_MEM_READ_WRITE,sizeof(int)*100);
cl::Bufferbuffer_C(context,CL_MEM_READ_WRITE,sizeof(int)*100);
cl::CommandQueuequeue(context,device);
cl::Kernelvec_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(constauto&v:A){
std::cout<
