540
從零開始學習OpenCL開發（三）深入API
歡迎關注，轉載引用請注明 https://blog.csdn.net/leonwei/article/details/8909897
這裏將更深入的說明一些OpenCL API的功能
1. 創建buffer
涉及到內存與顯存的操作總是複雜麻煩的，這個函數也一樣。。。
cl_memclCreateBuffer ( |
cl_context context, |
| cl_mem_flags flags, | |
| size_t size, | |
| void *host_ptr, | |
cl_int *errcode_ret)
|
函數將創建（或分配）一片buffer，並返回。這裏創建的mem可以是globla也可以是local或private，具體要看kernal中怎樣聲明限定符。cl會根據執行情況自動管理global到更進一層如private的copy。這裏的buffer概念是用於kernal函數計算的（或者說是用於device訪問的，什麼是device？host是C++寫的那段控製程序，一定運行在CPU，device就是執行kernal計算的，運行在所有有計算能力的處理器上，有時你的CPU同時扮演host與device，有時用GPU做device），這裏模煳了host與device的內存，也就是說根據flag的不同，可以是在host上的，也可以是在device上的，反正隻有這裏分配的內存可以用於kernal函數的執行。
主要的參數在 flags,這些參數可以|
1 CL_MEM_READ_WRITE：在device上開辟一段kernal可讀可寫的內存，這是默認
2 CL_MEM_WRITE_ONLY：在device上開辟一段kernal隻可以寫的內存
3 CL_MEM_READ_ONLY：在device上開辟一段kernal隻可以讀的內存
4 CL_MEM_USE_HOST_PTR：直接使用host上一段已經分配的mem供device使用，注意：這裏雖然是用了host上已經存在的內存，但是這個內存的值不一定會和經過kernal函數計算後的實際的值，即使用clEnqueueReadBuffer函數拷貝回的內存和原本的內存是不一樣的，或者可以認為opencl雖然借用了這塊內存作為cl_mem，但是並不保證同步的，不過初始的值是一樣的，（可以使用mapmem等方式來同步）
5 CL_MEM_ALLOC_HOST_PTR：在host上新開辟一段內存供device使用
6 CL_MEM_COPY_HOST_PTR：在device上開辟一段內存供device使用，並賦值為host上一段已經存在的mem
7 CL_MEM_HOST_WRITE_ONLY:這塊內存是host隻可寫的
8 CL_MEM_HOST_READ_ONLY:這塊內存是host隻可讀的
9 CL_MEM_HOST_NO_ACCESS:這塊內存是host可讀可寫的
談談這些flag，這些flag看起來行為比較複雜和亂，因為Opencl是一個跨硬件平台的框架，所以要照顧到方方麵麵，更統一就要更抽象。
首先456的區別，他們都是跟host上內存有關，區別是，4是直接使用已有的，5是新開辟，6是在device上開內存，但是初值與host相同（45都是在host上開內存）
然後看看123 和789,123是針對kernal函數的訪問說的，而789是針對host的訪問說的，kernal函數是device的訪問，而除了kernal函數的訪問基本都是host的訪問（如enqueueRead/write這些操作）
通常使用host上的內存計算的效率是沒有使用device上的效率高的,而創建隻讀內存比創建可寫內存又更加高效（我們都知道GPU上分很多種內存區塊，最快的是constant區域，那裏通常用於創建隻讀device內存）
通常用各種方式開內存你的程序都work，但這裏就要考驗不同情況下優化的功力了
size參數：要開的內存的大小
host_ptr參數：隻有在4.6兩種情況用到，其他都為NULL
當然這些內存都要使用clReleaseMemObject釋放
內存的call_back:
有些方式 ，如CL_MEM_USE_HOST_PTR,cl_mem使用的存儲空間實際就在host mem上，所以我們要小心處理這塊主存，比如你刪了它，但是cl_mem還在用呢，就會出現問題，而clReleaseMemObject並不一定會馬上刪除這個Cl_mem，它隻是一個引用計數的消減，這裏需要一個回調，告訴我們什麼時候這塊主存可以被放心的清理掉，就是clSetMemObjectDestructorCallback
CL的規範中特別說明最好不要在這個callback裏麵加入耗時的係統和cl API。
2.內存操作
1 從Cl_mem讀回host mem（就算Cl_mem是直接使用host mem實現的，想讀它的內容，還是要這樣讀回來，可以看做cl_mem是更高一層封裝）
clEnqueueReadBuffer
2 使用host_mem的值寫cl_mem
clEnqueueWriteBuffer
3 在Cl_mem和host mem之間做映射
clEnqueueMapBuffer
這個函數比較特殊，回顧上麵一節在創建buf時有一種方法CL_MEM_USE_HOST_PTR，是直接讓device使用host上已有的一塊的mem（p1）做buf，但是這個產生的CL_mem（p2）經過計算後值會改變,p2改變後通常p1不會被改變，因為雖然用的一塊物理空間，但是cl_mem是高層封裝，和host上的mem還是不一樣的，要想使p1同步到p2的最新值，就要調用這句map
MAP與CopyBack的性能對比
後來我想了想，這和使用clEnqueueReadBuffer從p2read到p1有什麼區別呢？map的方法按道理更快，因為p1p2畢竟一塊物理地址嗎，map是不是就做個轉換，而read則多一遍copy的操作。而且應該在CPU做device時map速度更快，但是事實是這樣的嗎？本著刨根問題的精神，我真的做了一下實驗，
我的實驗結果是這樣的，如果使用CPU做host，GPU做device，那麼CopyBack反而更快，但是如果使用CPU做host，CPU也做device，那麼MAP更快（不跨越硬件），而且總體上CPU+GPU的方式更快。
這個實驗結果徹底顛覆了我最初的一些想法，實驗數據說明1.不考慮硬件差異，MAP確實比CopyBack更快，跟我理解一樣，從CPU做device的兩組數據就可看出。2.至少在我的這個實驗中，主存與顯存間的數據copy比主存到主存自己的數據copy更快，所以在CPU+GPU的架構中，由於CopyBack方式采用的是主存顯存拷貝，而MAP值涉及主存上的操作，所以CopyBack更快。不過這裏我仍存在疑慮，我的分析很可能不對，或存在其他因素沒考慮，關於這點，要再繼續查查關於pinned memory和內存顯存傳遞數據的一些知識。
所以在這種異構計算領域，性能和你的硬件架構、性能、組合有著非常重要的關聯，所以最好的方法就是實際做實驗對比。
4 在Cl_mem直接做copy
clEnqueueCopyBuffer
這些函數都跟執行kernal一樣是投入到device的command queue裏的，但是他們又都帶有一個參數blocking_read，可以指定函數是否在執行完畢後返回。
3.Program
3.1.compile build link
有兩種從文本創建program的方式
- 直接build：clBuildProgram
- 先complie好，根據情況動態的link，即把上麵的過程拆分為兩個步驟
clCompileProgram clLinkProgram
但是1.2的方式不保險，這是CL1.2中加入的，而目前不是所有的platform都支持到1.2，NVIDIA好像就才1.1
opencl實際上會根據不同的硬件把通樣一份代碼編譯成不同的機器語言，如CPu匯編或GPU匯編
4.Kernal的執行
這裏是精華
1.設置kernal的參數
clSetKernelArg
2.執行kernal
clEnqueueNDRangeKernel
先給一段kernal代碼，方便下邊參數的解釋，另外這裏需要一些空間想象能力~
kernal代碼
__kernel void adder(__global const float* a, __global const float* b, __global float* result)
{
int idx = get_global_id(0);//得到當前單元格的0維度上的序號
result[idx] = a[idx] + b[idx];
}
參數說明：
command_queue ：執行那個device的命令序列
kernel：待執行的kernal obj
work_dim：我們知道CL的執行是放在一個個獨立的compute unit中進行的，你可以想像這些unit是排成一條線的，或是一個二維方陣，甚至是一個立體魔方，或著更高維，這裏參數就描述這個執行的維度，從1到CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS之間
global_work_size ：每個維度的unit的數量，這樣總共擁有的計算單元的數量將是global_work_size[0]*global_work_size[1]...
global_work_offset ：這裏就是規定上麵代碼裏每個維度上第一個get_global_id()0得到的id，默認為0，例如計算一個一維的長度為255work_size的工作，cl會自動虛擬出255個計算單元，每個單元各自計算0-254位置的數相加，而如果你把他設為3，那麼cl會從3開始算，也就是說3-254位置的unit會計算出結果，而0 -2這些unit根本不去參與計算。
local_work_size ：前麵介紹過CL的的unit是可以組合成組的（同組內可以互相通信）這個參數就決定了Cl的每個組的各維度的大小，NULL時CL會自動給你找個合適的，這裏貼下我試著用不同大小的group做數組相加的效率，
這裏其實看不太出什麼，直覺對這個應用實例是組越少越快，但是其中也不是嚴格的線性關係，無論在CPU還是GPU上這個關係都是近似的，所以在實際開發中，我們選擇什麼維度？選擇什麼樣的組大小？我的答案是：多做實驗吧，或者要偷懶的話就置0吧，交給CL為你做（實時上Cl中很多函數都有這個NULL的自適應選項。。）
關於維度、偏移、worksize這裏有個原版的圖，說明的更加形象
後麵幾個參數就跟同步有關了
event_wait_list和num_events_in_wait_list：說明這個command的執行要等這些event執行了之後
event：將返回這個command相關聯的event
很明顯，通過這幾個參數的event可以控製command之間的執行順序。
5.指令執行順序和同步
command的執行默認都是異步的，這才有利於並行度提高效率，在並行的問題中我們有時經常要做些同步的事情，或者等待某個異步的操作完成，這裏有兩種方法：
- 使用EnqueueRead/write這些操作可以指定他們為同步的（即執行完畢才在host上返回）
- 使用event來跟蹤，像clEnqueueNDRangeKernel這樣的操作都會關聯一個event
event：
- clEnqueue這樣的操作都會關聯返回一個event
- 用戶可以自己創建一個自定義的event clCreateUserEvent，要使用clReleaseEvent釋放
關於event的操作：
正是通過event同步不同的command：
- 設置event狀態：
設置用戶自定義event的狀態，clSetUserEventStatus 狀態隻可以被設定一次，隻可以為CL_COMPLETE或者一個負值，CL_COMPLETE代表這個event完成了，等待它的那些command得以執行，而負值表示引起錯誤，所有等待他的那些command都被取消執行。其實event的狀態還有CL_RUNNING CL_SUBMITTED CL_QUEUED，隻是不能在這裏設置。
- 等待event
clWaitForEvents；可以在host中等待某些event的結束，如clEnqueueNDRangeKernel這樣的異步操作，你可以等待他的event結束，就標誌著它執行完了
- 查詢event信息：clGetEventInfo clGetEventProfilingInfo
- 設置回調：clSetEventCallback
不同device上的event：
clEnqueueNDRangeKernel這樣的操作等待的隻能是處於相同queue裏麵的event（也就是同一個device上的），而同步不同queue上的event則隻能用顯示的方法，如clWaitForEvents等。
marker：
marker是這樣一個object，它可以看做是一個投入queue的空指令，專門用於同步，它可以向其他comman一樣設定需要等待的event，操作有clEnqueueMarkerWithWaitList
barrier：
barrier和marker十分類似，但是從名字上就可以看出最大的不同點是：marker在等待到它的依賴event之後會自動執行完畢，讓後續指令執行，而barrier會阻塞在這裏，直到他關聯的event被顯示的設置成完成狀態
marker和barrier的實現在1.1和1.2版本上存在著較大的不同
同步是CL的大問題，關於同步，原版overview上也有一個非常生動的圖，貼在這裏吧：
在同一個device上同步
在多個device間同步
最後更新：2017-04-03 18:51:59