蘋果開(kāi)發(fā)語(yǔ)言OpenCL 變量地址限定符詳解是本文要介紹的內(nèi)容，首先我們先來(lái)了解OpenCL，即：Open Computing Language，是由蘋果公司起草設(shè)計(jì)的用于大規(guī)模并行計(jì)算的計(jì)算編程語(yǔ)言。CocoaChina 版主 “zenny_chen” 今天為我們帶來(lái)新的一篇 OpenCL 教程：變量地址限定符。

我們的示例程序是通過(guò)OpenCL來(lái)實(shí)現(xiàn)一個(gè)正方形的顏色漸變著色。這里，我們將牽涉到變量存儲(chǔ)屬性，另外還引入了向量數(shù)據(jù)，向量數(shù)據(jù)是如何操作的，向量數(shù)據(jù)與標(biāo)量數(shù)據(jù)是如何交叉操作的。

請(qǐng)先下載完整的工程文件 OpenCL_shading.zip (36 K) ，下面是 OpenCL 內(nèi)核代碼。

// Render a square  
// left-top:    red(1, 0, 0)  
// left-bottom: green(0, 1, 0)  
// right-top:   blue(0, 0, 1)  
// right-bottom:black(0, 0, 0)  
 
__constant float4 left_top        = (float4)(1.0f, 0.0f, 0.0f, 0.0f);  
__constant float4 left_bottom     = (float4)(0.0f, 1.0f, 0.0f, 0.0f);  
__constant float4 right_top       = (float4)(0.0f, 0.0f, 1.0f, 0.0f);  
__constant float4 right_bottom    = (float4)(0.0f, 0.0f, 0.0f, 0.0f);  
__kernel void ColorShading(  
                    __global float4 output[256][256]  
                    )  
{  
    int dimX = get_global_id(0);  
    int dimY = get_global_id(1);  
    __local float4 deltaLeft = (left_top - left_bottom) / 255.0f;  
    __local float4 deltaRight = (right_top - right_bottom) / 255.0f;  
 
    float4 left = left_bottom + deltaLeft * (float)dimY;  
    float4 right = right_bottom + deltaRight * (float)dimY;  
    float4 delta = (right - left) / 255.0f;  
    float4 result = left + delta * (float)dimX;  
 
// clamp  
    if(result.x > 1.0f)  
        result.x = 1.0;  
    if(result.y > 1.0f)  
        result.y = 1.0f;  
    if(result.x < 0.0f)  
        result.x = 0.0f;  
    if(result.y < 0.0f)  
        result.y = 0.0f;  
 
    output[dimY][dimX] = result + (float4)(0.0f, 0.0f, 0.0f, 1.0f);  
} 

我們首先來(lái)談?wù)勏蛄款愋?/p>

上述代碼中，我們引入了float4類型。它是一個(gè)向量類型。向量類型的定義規(guī)則是在基本類型后加n，n可以是2，4，8，16。比如：uchar8，float2，int16，long4等等。

而對(duì)于向量類型各分量的訪問(wèn)，如果向量的分量個(gè)數(shù)在4個(gè)以內(nèi)，我們可以依次用x，y，z，w來(lái)表示。這種標(biāo)識(shí)法是與OpenGL Shader中的vertex shader對(duì)向量分量的訪問(wèn)形式一樣。

另外，我們還可以用數(shù)值索引來(lái)訪問(wèn)一個(gè)向量的各個(gè)分量。這個(gè)時(shí)候，我們可以將一個(gè)向量變量視為一個(gè)數(shù)組。如果向量的元素個(gè)數(shù)是16，那么第0到第9個(gè)元素分別用索引0到9表示；而第10到第15個(gè)元素，我們用a到f或A到F（即十六進(jìn)制）來(lái)表示。而當(dāng)我們用索引來(lái)表示的話，向量變量的.的后面必須跟一個(gè)字母 s。

下面舉些例子：

int4 a = int4(1, 2, 3, 4); 

那么a.x是1；a.y是2；a.z是3；a.w是4。同樣，a.s0是1；a.s1是2；a.s2是3；a.s3是4。

對(duì)于向量變量，我們還能非常靈活地對(duì)其各個(gè)分量進(jìn)行賦值。我們這里再引入一個(gè)swizzle的概念。swizzle是指可以對(duì)一個(gè)向量用另一個(gè)向量所對(duì)應(yīng)的任意元素進(jìn)行賦值。比如說(shuō)：

int4 a = int4(1, 2, 3, 4);  
int4 b = a.wzyx; 

這表示用a的第3個(gè)元素賦給b的第0個(gè)元素；a的第2個(gè)元素賦給b的第1個(gè)元素；用a的第1個(gè)元素賦給b的第2個(gè)元素；用a的第0個(gè)元素賦給b的第3個(gè)元素。然后，我們還能這么做：

b.xz = a.s32; 

表示將a的第三個(gè)元素給b的第0個(gè)元素；a的第2個(gè)元素賦給b的第2個(gè)元素。是不是很靈活呢？呵呵。

向量變量之間的加、減、乘、除以及邏輯運(yùn)算都是針對(duì)向量所對(duì)應(yīng)的各個(gè)分量進(jìn)行的。比如說(shuō)上面的int4 a; int4 b;a *= b; 相當(dāng)于：a.x *= b.x; a.y *= b.y; a.z *= b.z; a.w *= b.w;

如果一個(gè)向量與一個(gè)標(biāo)量進(jìn)行計(jì)算，那么是將該標(biāo)量與向量的沒(méi)個(gè)分量做相同的運(yùn)算，比如：

int4 a;  
int i;  
a *= i; 

相當(dāng)于：a.x *= i; a.y *= i; a.z *= i; a.w *= i;

而i *= a;則是非法的。因此我們必須注意，在用向量與標(biāo)量做算術(shù)邏輯操作時(shí)，必須把向量放在操作符的左邊，而標(biāo)量要放在操作符的右邊。

上述代碼中：

// clamp  
    if(result.x > 1.0f)  
        result.x = 1.0;  
    if(result.y > 1.0f)  
        result.y = 1.0f;  
    if(result.x < 0.0f)  
        result.x = 0.0f;  
    if(result.y < 0.0f)  
        result.y = 0.0f;  
} 

這部分是分別對(duì)結(jié)果的r分量和g分量做飽和。那么這里，我們將引入一個(gè)OpenCL內(nèi)建函數(shù)來(lái)取代這些代碼。OpenCL內(nèi)建函數(shù)一般是GPU指 令集直接支持的，因此一個(gè)調(diào)用基本上只需要1條指令就能完成。所以我們?cè)趯慜penCL時(shí)可以盡量使用內(nèi)建函數(shù)。當(dāng)然，有些數(shù)學(xué)內(nèi)建函數(shù)為了效率會(huì)犧牲精 度，此時(shí)我們要自己判斷取舍。

我們接下去再介紹一個(gè)OpenCL的內(nèi)建函數(shù)——

gentype clamp (gentype x, gentype minval,  
gentype maxval) 

其語(yǔ)義是返回：fmin(fmax(x, minval), maxval)，即對(duì)一個(gè)向量的每個(gè)分量取minval和maxval范圍內(nèi)的值。如果超出下界，那么取minval；如果超出上界，那么取maxval。

那么我再對(duì)OpenCL程序做次更新：

// Render a square  
// left-top:    red(1, 0, 0)  
// left-bottom: green(0, 1, 0)  
// right-top:   blue(0, 0, 1)  
// right-bottom:black(0, 0, 0)  
   
__constant float4 left_top        = (float4)(1.0f, 0.0f, 0.0f, 0.0f);  
__constant float4 left_bottom     = (float4)(0.0f, 1.0f, 0.0f, 0.0f);  
__constant float4 right_top       = (float4)(0.0f, 0.0f, 1.0f, 0.0f);  
__constant float4 right_bottom    = (float4)(0.0f, 0.0f, 0.0f, 0.0f);  
   
__constant float4 minValue = (float4)(0.0f, 0.0f, 0.0f, 0.0f);  
__constant float4 maxValue = (float4)(1.0f, 1.0f, 1.0f, 0.0f);  
   
__kernel void ColorShading(  
                    __global float4 output[256][256]  
                    )  
{  
    int dimX = get_global_id(0);  
    int dimY = get_global_id(1);  
   
    __local float4 deltaLeft = (left_top - left_bottom) / 255.0f;  
    __local float4 deltaRight = (right_top - right_bottom) / 255.0f;  
   
    float4 left = left_bottom + deltaLeft * (float)dimY;  
    float4 right = right_bottom + deltaRight * (float)dimY;  
    float4 delta = (right - left) / 255.0f;  
    float4 result = left + delta * (float)dimX;  
   
    // clamp  
    result = clamp(result, minValue, maxValue);  
   
    output[dimY][dimX] = result + (float4)(0.0f, 0.0f, 0.0f, 1.0f);  
} 

***，我們討論一下本文章的主題——變量的地址空間限定符

OpenCL有四種地址空間限定符——全局的（__global或global），本地的（__local或local），常量的（__constant或constant），私有的（__private或private）。

全局地址空間用于引用從全局存儲(chǔ)空間池所分配的存儲(chǔ)器對(duì)象。該存儲(chǔ)器對(duì)象可以聲明為指向一個(gè)標(biāo)量的指針，指向一個(gè)向量的指針或用戶自定義的結(jié)構(gòu)的指針。這允許內(nèi)核程序讀或?qū)懺摼彺娴娜我馕恢?。這里要注意的是，__global（或global）所修飾的是指針變量所引用的地址。因此：

__global long4 g;    // Error  
__global image2d_t texture;    // OK. A 2D texture image  
 
void kernelMain(__global int *p    // OK  
                             )  
{  
    __global float4 a;    // Error  
} 

本地地址空間用于描述需要被分配在本地存儲(chǔ)空間的變量，并且能被一個(gè)工作組的所有工作項(xiàng)共享。該限定符可以被用于函數(shù)實(shí)參或在函數(shù)內(nèi)聲明的聲明的變量。而用于修飾函數(shù)實(shí)參時(shí)，變量必須是指針類型。

常量地址空間用于描述分配在全局存儲(chǔ)空間的變量，并且它們?cè)趦?nèi)核程序中是只讀的。這些全局只讀變量可以被所有工作組的所有工作項(xiàng)共享。
該限定符在可以用于修飾內(nèi)核函數(shù)的指針變量參數(shù)，或是在內(nèi)核函數(shù)中修飾指針變量，或是作為全局變量。在本例中，我們把__constant修飾全局變量。
這里還要注意的是，由于__constant變量不能被寫，因此，作為全局變量時(shí)，它必須聲明后立即用常量初始化。這里的常量是指在編譯時(shí)能計(jì)算出數(shù)值結(jié)果的表達(dá)式。

私有地址空間的范圍很廣。所有函數(shù)參數(shù)、函數(shù)內(nèi)定義的局部變量都是私有的。因此我們往往可以省略掉__private關(guān)鍵字。

這里要注意的是OpenCL支持const關(guān)鍵字。這個(gè)關(guān)鍵字只在編譯時(shí)進(jìn)行檢查，它所修飾的變量不能被修改，而對(duì)運(yùn)行時(shí)該變量分配在哪個(gè)存儲(chǔ)空間無(wú)關(guān)。

***對(duì)這些關(guān)鍵字與實(shí)際性能的影像做一下簡(jiǎn)單的介紹。當(dāng)前流行的GPU等HPC流處理器采用分層的存儲(chǔ)架構(gòu)。全局存儲(chǔ)空間非常大（就相當(dāng)于我們所說(shuō)的顯存，目前最少也有128MB，俺的Mac Mini就是這個(gè)大?。菐捄馨嘿F，因此數(shù)據(jù)傳輸也是最慢的。

而第二層是局部存儲(chǔ)空間，或稱為本地存儲(chǔ)空間。局部存儲(chǔ)空間只能被一個(gè)工作組的所有工作項(xiàng)共享，并且每個(gè)工作組都有自己獨(dú)立的一個(gè)局部存儲(chǔ)空間。而每個(gè)局部存儲(chǔ)空間比較小，一般在128KB左右吧。但是其數(shù)據(jù)傳輸性能要比全局存儲(chǔ)器要大很多。

私有存儲(chǔ)空間是每個(gè)工作項(xiàng)私有的。也就是說(shuō)每個(gè)工作項(xiàng)有自己獨(dú)立的私有存儲(chǔ)空間。這在GPU存儲(chǔ)架構(gòu)中實(shí)際上就是寄存器文件。寄存器文件比如一共 128KB，那么對(duì)所有工作項(xiàng)進(jìn)行劃分的話，那么每個(gè)工作項(xiàng)能分到的存儲(chǔ)空間就非常小。但是寄存器的訪問(wèn)是最快的，讀或?qū)懸淮沃恍枰?個(gè)周期。

小結(jié)：蘋果開(kāi)發(fā)語(yǔ)言OpenCL 變量地址限定符詳解的內(nèi)容介紹完了，希望本文對(duì)你有所幫助！

推薦一篇相關(guān)的文章：蘋果開(kāi)發(fā)語(yǔ)言O(shè)penCL 多線程同步 附源碼