歡迎關注，轉載引用請注明 http://blog.csdn.net/leonwei/article/details/8909897
這裡将更深入的說明一些opencl api的功能
1. 建立buffer
涉及到記憶體與顯存的操作總是複雜麻煩的，這個函數也一樣。。。
<code>cl_memclcreatebuffer (</code>
cl_context context,
cl_mem_flags flags,
size_t size,
void *host_ptr,
cl_int *errcode_ret<code>)</code>
函數将建立（或配置設定）一片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間同步