蘋果開發語言OpenCL 多線程同步 附源碼是本文要介紹的內容，首先我們先來了解一下OpenCL， 即：Open Computing Language，是由蘋果公司起草設計的用于大規模并行計算的計算編程語言。

今天我們將介紹OpenCL多線程同步技巧。我們下面的例子將是以一個簡單的求和算法來描述如何同步一個工作組內的線程以及工作組之間如何同步。

我們之前介紹過變量的地址屬性。用__global修飾的變量存放在顯示存儲器中，特點是容量很大，但訪問速度很慢，并且所有工作項都能訪問；而用 __local修飾的變量存放在共享存儲器，其特點是速度比全局存儲要快很多，并且在同一工作組內的工作項能夠對其進行訪問，而且每個工作組有自己獨立的共享存儲器；__private修飾或默認狀態下定義的變量是私有的，即存放在寄存器中，其特點是訪問速度相當快，基本上一次讀或寫僅需要1個著色器周期，但它是工作項私有的，并且每個工作項只有若干個寄存器可以進行訪問。

如果我們讓在一個工作組內的線程進行同步，那么我們可以借助共享存儲變量來幫我們達成這個目標；而如果是工作組之間的通信，則需要全局存儲變量。

下面看求和的內核代碼：

__kernel void solve_sum(  
                    __global int input[4096],  
                    __global int output[9]  
                    )  
{  
    __local int localBuffer[512];  
 
    size_t item_id = get_local_id(0);  
    size_t gid = get_global_id(0);  
 
    localBuffer[item_id] = input[gid];  
 
    barrier(CLK_LOCAL_MEM_FENCE);  
 
    if((item_id) == 0)  
    {  
        int s = 0;  
        for(int i = 0; i < 512; i++)  
            s += localBuffer;  
        output[get_group_id(0)] = s;  
        output[8] = get_num_groups(0);  
    }  
} 

在以上代碼中，一共有4096個工作項，共有8個工作組，這樣每個工作組就有512個工作項。這個算法很簡單，首先將每個工作組內的工作項存放到共享數組中，等到一個工作組內的所有工作項完成這個動作后，讓工作項0對共享存儲緩存中的數據進行求和，完成后寫入到相應的工作組索引的輸出緩存。

在上述代碼中，get_local_id獲得的是當前工作組中的當前工作項索引，在上述代碼環境中的范圍是0到511。因此，我們可以將localBuffer[item_id] = input[gid];這句改為：localBuffer[gid & 511] = input[gid];這兩條語句的語義完全等價。

這里要著重介紹的線程同步函數是：

void barrier (cl_mem_fence_flags flags) 

這個內建函數對應于處理器的一條指令，其作用是同步一個工作組內的所有工作項。我們現在把工作項看作為一個線程。當其中一個線程執行到barrier時，它會被處理器阻塞住，直到該工作組內所有線程都執行到這個barrier，然后這些線程才能繼續執行下去。

這里有一個參數flags用于指示存儲器柵欄是局部的還是全局的，我們這里只需要局部的，因為這里不需要工作組之間的同步。

我們把每個工作組計算出來的結果寫到輸出緩存中。由于輸出才8個32位數據，因此在CPU中再拿去計算也變成了小菜一碟。

下面附上整個工程的代碼 OpenCL_Basic.zip (17 K)

上述代碼是將每個工作組計算好的結果傳送給主機端。那么我們是否能讓GPU把這8個結果也一起解決掉呢？答案是肯定的。不過我們這里將會用到OpenCL1.0中的原子操作擴展。這些基于int32位的原子操作在OpenCL1.1中將正式歸為語言核心，而不是擴展。我們可以通過OpenCL查詢獲得

cl_khr_global_int32_base_atomics是否被支持。如果被支持，那么我們可以用下面的方法：

__kernel void solve_sum(  
                     __global int input[4096],  
                     __global int output[9]  
                     )  
 {  
     __local int localBuffer[512];       
     size_t item_id = get_local_id(0);  
     size_t gid = get_global_id(0);     
     localBuffer[item_id] = input[gid];      
     barrier(CLK_LOCAL_MEM_FENCE);  
     if(item_id == 0)  
     {  
         int s = 0;  
         for(int i = 0; i < 512; i++)  
             s += localBuffer[i];  
         output[get_group_id(0)] = s;      
         int index = atom_inc(&output[8]);  
         if(index == 7)  
         {  
             mem_fence(CLK_GLOBAL_MEM_FENCE);  
             s = 0;  
             for(index = 0; index < 8; index++)  
                 s += output[index];  
             output[8] = s;  
         }  
     }  
 } 

在上述代碼中，我們用了原子累積操作：

int atom_inc (__global int *p) 

這個函數是先讀取p指針所指地址的內容，然后將該內容遞增1，最后寫回到這個地址中去，并且返回讀到的那個值（即更新以前的值）。整個操作都是不被打斷的，因此是一個原子操作。

我們在上述代碼中，用一個索引來獲取返回值，如果索引為7，說明當前線程是最后一個寫結果的工作組中的第0個線程。于是，我們利用這個線程把8個結果累加，然后寫回到輸出緩存。

如果有兩個線程對同一地址同時執行atom_inc，那么GPU將會進行仲裁，它只允許其中一個執行這一操作，而等到這個操作完成之后，其它線程才能繼續，否則，其它要執行此操作的線程都將被處理器阻塞。

那么這里由于利用了輸出緩存作為全局存儲的計數器變量，因此它將不象第一份代碼那樣作為只寫參數，而是要設置為可讀可寫的參數，并且要把初始數據傳入給GPU設備端。

下面附上相應的工程和代碼 OpenCL_Basic.zip (17 K)

下面要講一下關于Local Memory的一些高級話題。

其實OpenCL中的local memory對應于CUDA中的shared memory。在訪問共享存儲器時，如果多個線程寫同一個共享存儲器段（memory bank），那么會導致段沖突（bank conflict）。

什么是共享存儲器段呢？一個共享存儲器段就是在共享存儲器中的一個32位字（當前主流的中低端GPU均是如此，高級點的則可能是64位或更大）。那么，如果一個工作組的共享存儲器空間是128KB的話，則共有128KB / 4B = 32 * 1024個段。

如果有兩個線程（即工作項）對同一個段進行寫操作，那么這些寫操作將由原來可以并行寫而變成串行化的寫，也就是說，總線控制器會對這些多個線程的寫進行串行 化，它會選擇其中一個線程先寫，完了之后再挑選下一個。那么這樣一來，多個線程的執行也就從原來的并行操作變成了串行操作，這樣會受到很大的性能懲罰。

因此，我們在設計算法時應該盡量保證每個線程只對自己相應的共享存儲器段進行寫操作，而避免有多個線程去寫同一個共享存儲器段。而像上面示例代碼中，由于讀寫的數據元素都是32位，正好是一個存儲器段的大小，并且一個工作組內的每個工作項都以自己id作為索引對共享存儲器進行寫，這樣每個工作項所寫的段都是相互獨立的，因此這里不會發生段沖突。

小結：蘋果開發語言 OpenCL 多線程同步 附源碼的內容介紹完了，希望本文對你有所幫助！

帖子地址 http://www.cocoachina.com/bbs/read.php?tid-37608.html，歡迎參與討論