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

我們的示例程序是通過OpenCL來實現一個正方形的顏色漸變著色。這里，我們將牽涉到變量存儲屬性，另外還引入了向量數據，向量數據是如何操作的，向量數據與標量數據是如何交叉操作的。

請先下載完整的工程文件 OpenCL_shading.zip (36 K) ，下面是 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);  
__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);  
} 

我們首先來談談向量類型

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

而對于向量類型各分量的訪問，如果向量的分量個數在4個以內，我們可以依次用x，y，z，w來表示。這種標識法是與OpenGL Shader中的vertex shader對向量分量的訪問形式一樣。

另外，我們還可以用數值索引來訪問一個向量的各個分量。這個時候，我們可以將一個向量變量視為一個數組。如果向量的元素個數是16，那么第0到第9個元素分別用索引0到9表示；而第10到第15個元素，我們用a到f或A到F（即十六進制）來表示。而當我們用索引來表示的話，向量變量的.的后面必須跟一個字母 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。

對于向量變量，我們還能非常靈活地對其各個分量進行賦值。我們這里再引入一個swizzle的概念。swizzle是指可以對一個向量用另一個向量所對應的任意元素進行賦值。比如說：

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

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

b.xz = a.s32; 

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

向量變量之間的加、減、乘、除以及邏輯運算都是針對向量所對應的各個分量進行的。比如說上面的int4 a; int4 b;a *= b; 相當于：a.x *= b.x; a.y *= b.y; a.z *= b.z; a.w *= b.w;

如果一個向量與一個標量進行計算，那么是將該標量與向量的沒個分量做相同的運算，比如：

int4 a;  
int i;  
a *= i; 

相當于：a.x *= i; a.y *= i; a.z *= i; a.w *= i;

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

上述代碼中：

// 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;  
} 

這部分是分別對結果的r分量和g分量做飽和。那么這里，我們將引入一個OpenCL內建函數來取代這些代碼。OpenCL內建函數一般是GPU指 令集直接支持的，因此一個調用基本上只需要1條指令就能完成。所以我們在寫OpenCL時可以盡量使用內建函數。當然，有些數學內建函數為了效率會犧牲精 度，此時我們要自己判斷取舍。

我們接下去再介紹一個OpenCL的內建函數——

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

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

那么我再對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）。

全局地址空間用于引用從全局存儲空間池所分配的存儲器對象。該存儲器對象可以聲明為指向一個標量的指針，指向一個向量的指針或用戶自定義的結構的指針。這允許內核程序讀或寫該緩存的任意位置。這里要注意的是，__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  
} 

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

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

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

這里要注意的是OpenCL支持const關鍵字。這個關鍵字只在編譯時進行檢查，它所修飾的變量不能被修改，而對運行時該變量分配在哪個存儲空間無關。

***對這些關鍵字與實際性能的影像做一下簡單的介紹。當前流行的GPU等HPC流處理器采用分層的存儲架構。全局存儲空間非常大（就相當于我們所說的顯存，目前最少也有128MB，俺的Mac Mini就是這個大小），但是帶寬很昂貴，因此數據傳輸也是最慢的。

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

私有存儲空間是每個工作項私有的。也就是說每個工作項有自己獨立的私有存儲空間。這在GPU存儲架構中實際上就是寄存器文件。寄存器文件比如一共 128KB，那么對所有工作項進行劃分的話，那么每個工作項能分到的存儲空間就非常小。但是寄存器的訪問是最快的，讀或寫一次只需要1個周期。

小結：蘋果開發語言OpenCL 變量地址限定符詳解的內容介紹完了，希望本文對你有所幫助！

推薦一篇相關的文章：蘋果開發語言OpenCL 多線程同步 附源碼