Apple's OpenCL——变量地址限定符

这一讲我们将介绍更多的概念。 我们的示例程序是通过OpenCL来实现一个正方形的颜色渐变着色。 这里，我们将牵涉到变量存储属性，另外还引入了向量数据，向量数据是如何操作的，向量数据与标量数据是如何交叉操作的。
我先在首页贴上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);

附件: OpenCL_shading.zip (36 K) 下载次数:73

 

我们首先来谈谈向量类型。 上述代码中，我们引入了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个周期，呵呵。