OpenCL 函数修饰符

函数修饰符用来修饰OpenCL内核函数及一般函数的一些特性，以帮助编译器确定一些信息。

kernel修饰符

__kernel（或kernel） 修饰符声明一个函数为一个内核函数，这个内核函数将会在OpenCL设备上执行。

下面的例子展示了内核函数限定符的用法：

kernel void add1_gpu(global float ＊a, global float ＊b,
                        global float ＊c)
{
    ……
}
//错误,内核不能有返回值
kernel int add2_gpu(global float ＊a, global float ＊b,
                        global float ＊c)
{
    ……
}

内核函数返回类型必须是void类型，且只能在设备上执行。主机端可以调用这个函数。同时，如果一个内核函数调用另一个内核函数，那么被调的内核函数作为一个普通的函数调用。需要注意的是，如果内核函数中声明了local修饰符的变量，则在其他内核函数中调用此内核函数会有什么结果，这取决于OpenCL实现。

kernel void func_add1(global float ＊a, global float ＊b)
{
    local float var[32];
    ……
}
kernel void func_add2(global float ＊a, global float ＊b)
{
    func_add1(a, b);
}

上述代码结果依赖于OpenCL实现，所以这些代码不能保证跨平台的可移植性，因此尽量避免这种用法。

内核可选属性修饰符

kernel修饰符可以与关键字__attribute__联用来声明以下关于内核的额外信息：

• __attribute__（（vec_type_hint（〈type〉）））可以给编译器提示__kernel的计算的宽度，即内核正在处理的数据类型的大小。〈type〉为内建的矢量类型，或构成其元素的标量类型。如果没有指定，默认为__attribute__（（vec_type_hint（int）））。

• __attribute__（（work_group_size_hint（x,y,z）））可以给编译器提示可能使用的工作组的大小。也就是clEnqueueNDRangeKernel中local_work_size参数指定的值。

• __attribute__（（reqd_work_group_size（x,y,z）））用来指定必须使用的工作组的大小，即clEnqueueNDRangeKernel中local_work_size参数指定的值。这为编译器提供一个可以根据工作组大小完成特定优化的机会。
• __attribute__（（nosvm））修饰符，nosvm为一个指针变量，用来告诉编译器指针不会指向共享虚拟存储器区域。这个属性修饰符只能用于OpenCL 2.0或更高版本。