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

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

我們之前介紹過變量的地址屬性。用__global修飾的變量存放在顯示存儲器中，特點(diǎn)是容量很大，但訪問速度很慢，并且所有工作項(xiàng)都能訪問；而用 __local修飾的變量存放在共享存儲器，其特點(diǎn)是速度比全局存儲要快很多，并且在同一工作組內(nèi)的工作項(xiàng)能夠?qū)ζ溥M(jìn)行訪問，而且每個工作組有自己獨(dú)立的共享存儲器；__private修飾或默認(rèn)狀態(tài)下定義的變量是私有的，即存放在寄存器中，其特點(diǎn)是訪問速度相當(dāng)快，基本上一次讀或?qū)憙H需要1個著色器周期，但它是工作項(xiàng)私有的，并且每個工作項(xiàng)只有若干個寄存器可以進(jìn)行訪問。

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

下面看求和的內(nèi)核代碼：

__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個工作項(xiàng)，共有8個工作組，這樣每個工作組就有512個工作項(xiàng)。這個算法很簡單，首先將每個工作組內(nèi)的工作項(xiàng)存放到共享數(shù)組中，等到一個工作組內(nèi)的所有工作項(xiàng)完成這個動作后，讓工作項(xiàng)0對共享存儲緩存中的數(shù)據(jù)進(jìn)行求和，完成后寫入到相應(yīng)的工作組索引的輸出緩存。

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

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

void barrier (cl_mem_fence_flags flags) 

這個內(nèi)建函數(shù)對應(yīng)于處理器的一條指令，其作用是同步一個工作組內(nèi)的所有工作項(xiàng)。我們現(xiàn)在把工作項(xiàng)看作為一個線程。當(dāng)其中一個線程執(zhí)行到barrier時，它會被處理器阻塞住，直到該工作組內(nèi)所有線程都執(zhí)行到這個barrier，然后這些線程才能繼續(xù)執(zhí)行下去。

這里有一個參數(shù)flags用于指示存儲器柵欄是局部的還是全局的，我們這里只需要局部的，因?yàn)檫@里不需要工作組之間的同步。

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

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

上述代碼是將每個工作組計(jì)算好的結(jié)果傳送給主機(jī)端。那么我們是否能讓GPU把這8個結(jié)果也一起解決掉呢？答案是肯定的。不過我們這里將會用到OpenCL1.0中的原子操作擴(kuò)展。這些基于int32位的原子操作在OpenCL1.1中將正式歸為語言核心，而不是擴(kuò)展。我們可以通過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) 

這個函數(shù)是先讀取p指針?biāo)傅刂返膬?nèi)容，然后將該內(nèi)容遞增1，最后寫回到這個地址中去，并且返回讀到的那個值（即更新以前的值）。整個操作都是不被打斷的，因此是一個原子操作。

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

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

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

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

下面要講一下關(guān)于Local Memory的一些高級話題。

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

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

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

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

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

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