OpenCL

GPU  OpenCL 

一、 CUDA vs OpenCL

1. 簡介

OpenCL: Open Computing Language，開放計算語言。
OpenCL和CUDA是兩種異構計算（此異構平台可由CPU，GPU或其他類型的處理器組成。）的編程模型。

1. CUDA只支持NVIDIA自家的GPU。
2. OpenCL最早是由Apple提出，后來交給了Khronos這個開放標准組織。OpenCL 1.0 在2008年底正式由Khronos發布，比CUDA晚了整整一年。

2012年移動圖形處理器市場份額，imagenation失去蘋果后一落千丈，已被別的公司收購：

enter description here

2. 操作步驟

CUDA C加速步驟：

1. 在device (也就是GPU) 上申請內存
2. 將host (也就是CPU) 上的數據拷貝到device
3. 執行CUDA kernel function
4. 將device上的計算結果傳回host
5. 釋放device上的內存

OpenCL操作步驟：

1. 檢測申請計算資源
   • 檢測platform, clGetPlatformIDs
   • 檢測platform對應的device, clGetDeviceInfo
   • 建立context, clCreateContextFromType
   • 建立command queue, clCreateCommandQueue
   • 在context內申請存儲空間, clCreateBuffer
2. 將host (也就是CPU) 上的數據拷貝到device, clCreateBuffer
3. OpenCL代碼編譯
   • 讀入OpenCL （kernel function） 源代碼，創立program 句柄, clCreateProgramWithSource
   • 編譯program, clBuildProgram
   • 創立一個 OpenCL kernel 句柄, clCreateKernel
   • 申明設置 kernel 的 參數, clSetKernelArg
   • 設置NDRange
4. 運行kernel , clEnqueueNDRangeKernel
5. 將device上的計算結果傳回host, clEnqueueReadBuffer
6. 釋放計算資源
   • 釋放kernel, clReleaseKernel
   • 釋放program, clReleaseProgram
   • 釋放device memory, clReleaseMemObject
   • 釋放command queue, clReleaseCommandQueue
   • 釋放context, clReleaseContext

procedure

整體架構如下：

enter description here

CUDA C語言與OpenCL的定位不同，或者說是使用人群不同。CUDA C是一種高級語言，那些對硬件了解不多的非專業人士也能輕松上手；而OpenCL則是針對硬件的應用程序開發接口，它能給程序員更多對硬件的控制權，相應的上手及開發會比較難一些。

OpenCL, CUDA C

3. 名詞比較

Block: 相當於opencl 中的work-group
Thread：相當於opencl 中的work-item
SP: 相當於opencl 中的PE
SM: 相當於opencl 中的CU
warp: 相當於opencl 中的wavefront(簡稱wave)，基本的調試單位

4. system tradeoff

各種硬件形態的開發效率與執行效率，而opencl在FPGA上作用就是綠色箭頭的方向，可以有效提高FPGA開發效率。

system

二、常用API

1. clEnqueueNDRangeKernel

clEnqueueNDRangeKernel

參數：

1. command_queue,
2. kernel,
3. work_dim，使用多少維的NDRange，可以設為1, 2, 3, ..., CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS。
4. global_work_offset(GWO), 每個維度的偏移，如果不設置默認為0
5. global_work_size(GWS)，每個維度的索引長度，GWS(1) * GWS(2) * ... * GWS(N) 應該大於等於需要處理的數據量
6. local_work_size(LWS), 每個維度work-group的大小，如果不設置，系統會自己選擇一個合適的大小
7. num_events_in_wait_list: 執行kernel前需要等待的event個數
8. event_wait_list: 需要等待的event列表
9. event: 當前這個命令會返回一個event，以供后面的命令進行同步
   返回：

函數返回執行狀態。如果成功， 返回CL_SUCCESS

2. clCreateBuffer

clCreateBuffer

1. context

2. flags參數共有9種：

   device權限，默認為可讀寫：
   CL_MEM_READ_WRITE: kernel可讀寫
   CL_MEM_WRITE_ONLY: kernel 只寫
   CL_MEM_READ_ONLY: kernel 只讀

   創建方式:
   CL_MEM_USE_HOST_PTR: device端會對host_ptr位置內存進行緩存，如果有多個命令同時使用操作這塊內存的行為是未定義的
   CL_MEM_ALLOC_HOST_PTR: 新開辟一段host端可以訪問的內存
   CL_MEM_COPY_HOST_PTR: 在devices新開辟一段內存供device使用，並將host上的一段內存內容copy到新內存上

   host權限，默認為可讀寫:
   CL_MEM_HOST_WRITE_ONLY：host 只寫
   CL_MEM_HOST_READ_ONLY： host只讀
   CL_MEM_HOST_NO_ACCESS: host沒有訪問權限

3. size是buffer的大小

4. host_ptr只有在CL_MEM_USE_HOST_PTR， CL_MEM_COPY_HOST_PTR時才有效。

一般對於kernel函數的輸入參數，使用CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR可以將host memory拷貝到device memory，表示device只讀，位置在device上並進行內存復制，host權限為可讀寫；
對於輸出參數，使用CL_MEM_WRITE_ONLY表示device只寫，位置在device上，host權限為可讀可寫。

如果進行host與device之間的內存傳遞，可以使用clEnqueueReadBuffer讀取device上的內存到host上, clEnqueueWriteBuffer可以將host上內存寫到device上。

3. clEnqueueWriteBuffer

clEnqueueWriteBuffer

1. command_queue,
2. buffer, 將內存寫到的位置
3. blocking_write, 是否阻塞
4. offset, 從buffer的多少偏移處開始寫
5. size, 寫入buffer大小
6. ptr, host端buffer地址
7. num_events_in_wait_list, 等待事件個數
8. event_wait_list, 等待事件列表
9. event, 返回的事件

4. clCreateImage

創建一個ImageBuffer:

clCreateImage

1. context
2. flags, 同clCreateBuffer里的flags
3. image_format, 圖像的屬性，包含兩個變量: image_channel_order, 指定通道數和形式，通常為RGBA；image_channel_data_type, 定義數據類型， CL_UNORM_INT8表示為unsigned規一化的INT8，CL_UNSIGNED_INT8
   表示 為非規一化的unsigned int8
4. image_desc, 定義圖像的維度大小，
5. host_ptr, 輸入圖像地址
6. errorce_ret, 返回狀態

5. clEnqueueWriteImage

clEnqueueWriteImage

1. command_queue
2. image, 目標圖像
3. block_writing, 是否阻塞，如果TRUE，則阻塞
4. origin, 圖像的偏移，通常為（0, 0, 0）
5. region, 圖像的區域，(width, height, depth)
6. input_row_pitch，每行字節數，可能有對齊；如果設為0,則程序根據每個像素的字節數 乘以 width 計算
7. input_slice_pitch，3D圖像的2D slice塊，如果是1D或2D圖像，這個值必須為0
8. ptr, host端輸入源圖像地址
9. num_events_in_wait_list, 需等待事件個數
10. evnet_wait_list, 需要等待的事件列表
11. event, 返回這個命令的事件，用於后續使用

Map buffer

將cl_mem映射到CPU可訪問的指針：

clEnqueueMapBuffer

1. command_queue
2. buffer, cl_mem映射的源地址
3. blocking_map, 是否阻塞
4. map_flags, CL_MAP_READ，映射的地址為只讀；CL_MAP_WRITE，向映射的地址寫東西；CL_MAP_WRITE_INVALIDATE_REGION, 向映射的地址為寫東西，host不會使用這段地址的內容，這時返回的地址處的內容不保證是最新的
5. offset, cl_mem的偏移
6. size, 映射的內存大小
7. num_events_in_wait_list, 等待事件個數
8. event_wait_list, 等待事件列表
9. event, 返回事件
10. errorcode_ret, 返回狀態

返回值是CPU可訪問的指針。

注意：

1. 當flag為CL_MAP_WRITE時，如果不使用unmap進行解映射，device端無法保證可以獲取到最新寫的值。
2. 如果不用unmap，那么device端無法釋放這部分內存

所以寫完內容后，要立馬解映射。

buffer

clEnqueueCopyBuffer: 從一個cl buffer拷貝到另一個cl buffer

event

cl_int clWaitForEvents(cl_uint num_events, const cl_event *event_list)
等待事件執行完成才返回，否則會阻塞

cl_int clEnqueueWaitForEvents(cl_command_queue command_queue, cl_uint num_events, const cl_event *event_list)
和 clWaitForEvents 不同的是該命令執行后會立即返回，線程可以在不阻塞的情況下接着執行其它任務。而 clWaitForEvents 會進入阻塞狀態，直到事件列表 event_list 中對應的事件處於 CL_COMPLETE 狀態。

cl_int clFlush(cl_command_queue command_queue)
只保證command_queue中的command被commit到相應的device上，不保證當clFlush返回時這些command已經執行完。

cl_int clFinish(cl_command_queue command_queue)
clFinish直到之前的隊列命令都執行完才返回。clFinish is also a synchronization point.

cl_int clEnqueueBarrier(cl_command_queue command_queue)
屏障命令保證在后面的命令執行之前，它前面提交到命令隊列的命令已經執行完成。
和 clFinish 不同的是該命令會異步執行，在 clEnqueueBarrier 返回后，線程可以執行其它任務，例如分配內存、創建內核等。而 clFinish 會阻塞當前線程，直到命令隊列為空（所有的內核執行/數據對象操作已完成）。

cl_int clEnqueueMarker(cl_command_queue command_queue, cl_event *event)
將標記命令提交到命令隊列 command_queue 中。當標記命令執行后，在它之前提交到命令隊列的命令也執行完成。該函數返回一個事件對象 event，在它后面提交到命令隊列的命令可以等待該事件。例如，隨后的命令可以等待該事件以確保標記之前的命令已經執行完成。如果函數成功執行返回 CL_SUCCESS。

三、架構

1. Platform Model

1個host加上1個或多個device，1個device由多個compute unit組成，1個compute unit又由多個Processing Elemnet組成。

Platform Model

2. Execution Model

執行模型：

一個主機要使得內核運行在設備上，必須要有一個上下文來與設備進行交互。 一個上下文就是一個抽象的容器，管理在設備上的內存對象，跟蹤在設備上 創建的程序和內核。

主機程序使用命令隊列向設備提交命令，一個設備有一個命令隊列，且與上下文 相關。命令隊列對在設備上執行的命令進行調度。這些命令在主機程序和設備上 異步執行。執行時，命令間的關系有兩種模式：(1)順序執行，(2)亂序執行。

內核的執行和提交給一個隊列的內存命令會生成事件對象，可以用來控制命令的執行、協調宿主機和設備的運行。

有3種命令類型：
• Kernel-enqueue commands: Enqueue a kernel for execution on a device.（執行kernel函數）
• Memory commands: Transfer data between the host and device memory, between memory objects, or map and unmap memory objects from the host address space.（內存傳輸）
• Synchronization commands: Explicit synchronization points that define order constraints between commands.（同步點）

命令執行經歷6個狀態：

1. Queued: 將command放到CommandQueue
2. Submitted: 將command從CommandQueue提交到Device
3. Ready: 當所有運行條件滿足，放到Device的WorkPool里
4. Running: 命令開始執行
5. Ended: 命令執行結束
6. Complete: command以及其子command都結束執行，並設置相關的事件狀態為CL_COMPLETE

Execution Model

Mapping work-items onto an NDRange：

與CUDA里的grid, block, thread類似，OpenCL也有自己的work組織方式NDRange。NDRange是一個N維的索引空間(N為1, 2, 3...)，一個NDRange由三個長度為N的數組定義，與clEnqueueNDRangeKernel幾個參數對應：

1. global_work_size(GWS)，每個維度的索引長度，GWS(1) * GWS(2) * ... * GWS(N) 應該大於等於需要處理的數據量
2. global_work_offset(GWO), 每個維度的偏移，如果不設置默認為0
3. local_work_size(LWS), 每個維度work-group的大小，如果不設置，系統會自己選擇較好的結果

如下圖所示，整個索引空間的大小為，每個work-group大小為，全局偏移為。
對於一個work-item，有兩種方式可以索引：

1. 直接使用global id
2. 或者使用work-group進行相關計算，設當前group索引為，group里的local id分別為(s_x, s_y)，那么便有

NDRange index space

3. Memory Model

不同平台的內存模型不一樣，為了可移植性，OpenCL定義了一個抽象模型，程序的實現只需要關注抽象模型，而具體的向硬件的映射由驅動來完成。

Memory Model

主要分為host memory和device memory。而device memory 一共有4種內存：
private memory：是每個work-item各自私有
local memory: 在work-group里的work-item共享該內存
global memory: 所有memory可訪問
constant memory: 所有memory可訪問，只讀，host負責初始化

enter description here

4. Program Model

OpenCL支持數據並行，任務並行編程，同時支持兩種模式的混合。
分散收集（scatter-gather）：數據被分為子集，發送到不同的並行資源中，然后對結果進行組合，也就是數據並行；如兩個向量相加，對於每個數據的+操作應該都可以並行完成。
分而治之（divide-and-conquer）：問題被分為子問題，在並行資源中運行，也就是任務並行；比如多CPU系統，每個CPU執行不同的線程。還有一類流水線並行，也屬於任務並行。流水線並行，數據從一個任務傳送到另外一個任務中，同時前一個任務又處理新的數據，即同一時刻，每個任務都在同時運行。

enter description here

並行編程就要考慮到數據的同步與共享問題。

in-order vs out-of-order:
創建命令隊列時，如果沒有為命令隊列設置 CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE 屬性，提交到命令隊列的命令將按照 in-order 的方式執行。

OpenCL支持兩種同步：
同一工作組內(work-group)工作項(work-item)的同步(實現方式barrier)：
reduction的實現中，需要進行數據同步，所謂reduction就是使用多個數據生成一個數據，如tensorflow中的reduce_mean, reduce_sum等。在執行reduce之前，必須保證這些數據已經是有效的，執行過的，

命令隊列中處於同一個上下文中的命令的同步(使用clWaitForEvents，clEnqueueMarker, clEnqueueBarrier 或者執行kernel時加入等待事件列表)。

有2種方式同步：
鎖（Locks）：在一個資源被訪問的時候，禁止其他訪問；
柵欄（Barriers）：在一個運行點中進行等待，直到所有運行任務都完成；（典型的BSP編程模型就是這樣）

數據共享：
（1）shared memory
當任務要訪問同一個數據時，最簡單的方法就是共享存儲shared memory（很多不同層面與功能的系統都有用到這個方法），大部分多核系統都支持這一模型。shared memory可以用於任務間通信，可以用flag或者互斥鎖等方法進行數據保護，它的優缺點：
優點：易於實現，編程人員不用管理數據搬移；
缺點：多個任務訪問同一個存儲器，控制起來就會比較復雜，降低了互聯速度，擴展性也比較不好。
（2）message passing
數據同步的另外一種模型是消息傳遞模型，可以在同一器件中，或者多個數量的器件中進行並發任務通信，且只在需要同步時才啟動。
優點：理論上可以在任意多的設備中運行，擴展性好；
缺點：程序員需要顯示地控制通信，開發有一定的難度；發送和接受數據依賴於庫方法，因此可移植性差。

Experiment

1. 向量相加

guru_ge@dl:~/opencl/test$ ./cuda_vector_add
SUCCESS
copy input time: 15438.000000
CUDA time: 23.000000
copy output time: 17053.000000
CPU time: 16259.000000
result is right!
guru_ge@dl:~/opencl/test$ ./main
Device: GeForce GTX 1080 Ti
create input buffer time: 7
create output buffer time: 1
write buffer time: 4017
OpenCL time: 639
read buffer time: 30337
CPU time: 16197
result is right!

guru_ge@dl:~/opencl/test$ ./cuda_vector_add
SUCCESS
copy input time: 59825.000000
CUDA time: 36.000000
copy output time: 67750.000000
CPU time: 64550.000000
result is right!
guru_ge@dl:~/opencl/test$ ./main
Device: GeForce GTX 1080 Ti
create input buffer time: 7
create output buffer time: 1
write buffer time: 52640
OpenCL time: 1634
read buffer time: 80206
CPU time: 66502
result is right!
guru_ge@dl:~/opencl/test$

Reference

1. https://www.cnblogs.com/wangshide/archive/2012/01/07/2315830.html
2. https://www.cnblogs.com/hlwfirst/p/5003504.html
3. http://blog.csdn.net/leonwei/article/details/8909897
4. https://blog.csdn.net/babyfacer/article/details/6863572
5. https://blog.csdn.net/xbinworld/article/details/45949629
6. https://blog.csdn.net/Bob_Dong/article/details/70172165?locationNum=11&fps=1