苹果开发语言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 多线程同步 附源码
