OpenCL学习之深入（一）– 限定符
1.0 学习之始
走在成功的道路上，记录工作中的点点滴滴。
1.1 地址空间限定符
在opencl中的内核函数中，我会经常遇到__global、__local、__constant和__private。如：
__kernel void cl_add(__global int *a,__global int *b)
这些以前缀__的修饰符被称为地址空间限定符，声明变量是可以使用地址空间限定符来制定用来分配对象的内存区域。
- __global 用来指分配全局内存池的内存对象（缓存对象或图像图像），可以将缓存对象声明为一个指针，指向标量、矢量或用户自定义结构体。数组内存对象实际大小在宿主机代码中通过恰当的API调用分配此对象时确定。
- __local 可以用来描述分配局部内存的变量，这种变量由工作组中的所有工作项共享。此限定符可以用在声明为指针的函数参数上（包括 __kernel函数）,也可以用在 __kernel函数中声明的变量上。
- __constant 可以用来描述分配全局内存且内核中只读的变量。在内核执行时，这些只读变量可以被其所有全局工作项访问。此限定符可以用在声明为指针的函数参数上(包括 __kernel函数),也可以用在 __kernel函数中声明为指针的局部变量上，全局变量也可以。程序源码张所有声明带限定符 __constant的全局变量需要初始化。
- __private 函数(包括 __kernel函数)中的所有变量及传给此函数的所有参数都在 __private地址空间中声明。
1.2 图像访问限定符
内核参数中的图像对象可以声明为只读或只写的。内核不能只对一个图像对象即读又写。可以使用限定符 __read _only (或 read_only) 和 __write _only（或write _only）来表明图像对象是只读的还是只写的。默认的限定符是 __read _only
下面例子中
__kernel void foo ( __read_only image2d_t imageA, __ write_only image2d_t imageB)
{...}
imageA是只读的，imageB是只写的。
1.3 函数限定符
1.3.1 __kernel (或kernel)
限定符__kernel（或kernel）可以将函数声明为内核，此内核可以在OpenCL设备上由应用来执行。符合下列规则的函数才能使用此限定符：
* 只可以在此设备上执行。
* 可以被宿主机调用。
* 如果 __kernel函数被另一个内核函数调用，则仅仅是一个普通函数调用。
注意：
如果内核函数中声明了带有限定符__local 或local的变量，宿主机可以使用clEnqueueNDRangeKernel、clEnqueueTask等API来调用这些内核函数；而如果由其它内核函数调用它们，则其行为依赖于具体实现。
名字__kernel 和kernel被保留用作地址空间限定符，否则不能使用它们。
1.3.2 可选的特性（attribute ）限定符
如下面所说，用kernel声明内核函数时，可以加上 __attribute 来声明一些附加信息。
可选的 attribute ((vec_type_hint())) 可以给编译器一个暗示，表示__kernel的可计算的width,同时在编译器试图将代码自动矢量化时，它也是计算可利用的处理器带宽的基础。 vec_type_hint()应当是內建标量和矢量数据类型之一。而如果没有指定vec_type_hint()，默认值是int。
默认类型是 _ attribute_ ((vec_type_hint()))
例如,如果开发人员指定了一个float4的宽度，编译器假定通常使用4路浮点矢量进行计算，并且可能将工作项合并或分离到多个线程中以更好的匹配硬件性能。不要求实现一定会自动矢量化代码，但要支持这种暗示。即使没有提供暗示，编译器也可能自动矢量化代码。如果实现将N个工作项合并到一个线程中，那么它要负责正确处理这个情况：任一维度上的全局或局部工作项的数目模N不是0.
__kernel __attribute __((vec_type_hint(float4))) void foo( __global float4 *p) {....
__kernel __attribute __((vec_type_hint(double))) void foo( __global float4 *p) {....
__kernel void foo( __global float4 *p) {...
如果，声明_ kernel时带有 _ attribute __((vec_type_hint(float4)))（意味着 __kernel中大多数运算都明显使用float4矢量化），且内核在运行时使用IntelR高级矢量指令，实现了8个浮点宽的矢量单元。那么自动矢量化可能选择将两个工作项合并到一个线程中，第二个工作项在256位AVX寄存器的高128位中运行。
参考
- 倪庆亮 OpenCL规范
- Khronos OpenCL Working Group