OpenCL

OpenCL

GPU  OpenCL 

1、 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

2、經常使用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。

3、架構

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. http://www.javashuo.com/article/p-smellncn-dy.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