OpenCL 地址空间修饰符

OpenCL C是在C语言语法上的扩展，其中一个就是对C语言中修饰符的扩展。OpenCL语言中的修饰符大致包含了地址限定符、函数限定符、存储类说明符（即static与extern，这两个与C语言中的作用类似，以下不再赘述），以及对象访问限定符。为了叙述方便，以下将限定符（qualif ier）与说明符（specif ier）统称为“修饰符”（modifer）。

在OpenCL存储器模型中，我们知道OpenCL设备有全局存储器、局部存储器、常量存储器和私有存储器。对于这四种存储器，对应的地址空间修饰符为：__global（或global）、__local（或local）、__constant（或constant）和__private（或private）。如果一个变量由一个地址空间修饰符限定，那么这个变量就在指定的地址空间中分配。

程序中的函数参数，或者函数中缺省地址修饰符的局部变量，它们的地址修饰符为private。所有函数参数一定在私有地址空间。在程序范围内的一个变量，或者程序内的一个static变量，它们在全局或常量地址空间。如果没有地址修饰符指定，默认为全局的。另外，在OpenCL 2.0之前，全局变量必须在常量地址空间，即必须用constant进行修饰。而从OpenCL 2.0开始，可以在程序范围直接用全局地址修饰符来定义一个全局变量。

通过几个例子来说明如何指定地址空间名：

//在private地址空间声明一个指针p,指针指向的是一个全局地址空间对象
global int ＊p;
//在private地址空间声明一个整数数组
float x[4]
//函数参数中:
//src:在私有地址空间声明了一个指向常量地址空间的一个指针
//v: 在私有地址空间声明
int my_func_a(constant float＊src, int4 v){
    float temp;                   //temp分配在私有地址空间
}

例子中的变量p是一个private变量，而其指向的数据保存在全局存储器中。数组变量x保存在private空间。函数my_func_a的参数变量和函数体内声明的局部变量都默认保存在private地址空间。

OpenCL 2.0增加了一个未命名通用（generic）地址空间。如果将一个指针声明为不指向任一命名地址空间，那么该指针指向通用地址空间。在访问一个指针所指向区域的内容前，该指针必须关联一个命名地址空间。如果函数参数的指针变量与指针类型的返回值没有声明地址空间，那么这些参数和返回值指向通用地址空间。除了常量地址空间，其他命名地址空间是通用地址空间的一个子集。

内核参数声明的指针类型必须指向global、local和constant三种类型之一。

全局地址空间

全局地址空间用来指示从全局存储器中分配的存储器对象（缓冲区或图像）。一个缓冲区对象在内核参数中声明为指针，可以指向标量、矢量和自定义结构体。内核可以读写缓冲区任何位置。存储器对象的大小在主机端调用API分配时确定。

下面给出几个例子：

global float4 ＊color;            //float4类型指针
typedef struct
{
    float a[4];
    int b[2];
} foot_t;
global foo_t ＊my_info;           //foo_t类型指针

一般在全局地址空间分配图像对象，但全局地址修饰符不能用于图像类型。对图像对象不能直接读/写，OpenCL提供了内建函数来支持对图像对象的读写。

程序范围内定义的变量以及函数内定义的静态变量可以声明在全局地址空间。在全局地址空间内的变量有整个程序的生命周期，它们的存储器空间在程序内一直保持。但是需要注意的是这些变量不能跨设备，不同设备间的全局地址空间变量在不同的存储空间，在程序范围内的全局地址空间中的变量或静态变量可以初始化。下面就给出几个例子：

global int foo;                          //ok
global uchar buf[512];                   //ok
global int baz = 12;                     //ok,允许初始化
constant int cst = 100;
constant int ＊global ptr = &baz;        //error ,因为baz在全局地址空间
constant int ＊global ptr = &cst;        //OK。ptr在全局地址空间,它指向常量地址空间
int＊constant constant_ptr = &foo;       //OK。constant_ptr在常量地址空间
global int ＊constant ptr = &baz;        //ok。在程序加载时初始化一个常量
global event_t ev;                       //error,非法的数据类型

局部地址空间

局部地址空间用来描述需要在局部存储器中分配的变量，这些变量被一个工作组中的所有工作项共享。局部地址空间的指针可以作为函数的参数和函数内声明的变量。内核函数中声明的变量可以在局部地址空间中分配，但是有一些限制：这些变量声明必须出现在内核函数作用域；这些变量不能被初始化。

下面给出几个合法和不合法的使用例子：

kernel void my_func(...)
{
    local float a;                       //ok。在局部地址空间内分配一个变量
    local float b[10];                   //ok。在局部地址空间分配一个10个float类型的数组
    local float c = 1;                   //error,不能在声明时初始化
    a = 1;                               //ok
    if(...)
    {
        local float c;                   //error。变量C的作用域不在内核函数作用域。
    }
}

局部地址空间中的变量如果作为指针参数传递，或者在一个内核函数中声明，那么仅在执行这个内核的工作组的生命周期内存在。

常量地址空间

常量地址空间用来描述全局存储器中分配的一些变量。这些变量在内核中作为只读变量访问。这些只读变量在内核执行时允许所有工作项访问。常量地址空间的指针可以作为函数参数和函数内声明的变量（这一点和CUDA有所区别）。

程序作用域中的变量和内核函数作用域中的变量可以在常量地址空间分配。这些变量都要求定义时直接初始化，而且用来初始化这些变量的值必须是编译时常量。对这些变量写操作会导致一个编译时错误。另外，程序中声明的所有字符串也会存储在常量地址空间。

下面给出几个合法和不合法的使用例子：

//ok,程序范围变量可以在常量地址空间分配
constant float A[] = {0, 1, 2, 3, 4, 5}
kernel void my_func(constant float ＊A, constant float ＊B)
{
    constant float4 ＊p = A;        //ok。p为私有变量,指向float4类型的常量指针
    constant float a;               //error,没有初始化
    constant float b = 2.0f;        //ok,在编译时初始化
    p[0] = (float4)(1.0f);          //error, 不能修改p的值
    char ＊c = "OpenCL";            //字符串"OpenCL"分配在常量地址空间
}

对于上述代码需要补充说明的是，contant f loat4*p=A；这一句。需要注意的是，OpenCL C语言的地址标识符与C语言的const、volatile等修饰符的用法一样，因此放在不同位置含义不同。这里，p是指向一个f loat4类型常量的指针，但它本身是分配在私有地址空间的（或直接作为一个寄存器）。例如：

//ok,程序范围变量可以在常量地址空间分配
constant float A[] = {0, 1, 2, 3, 4, 5}
kernel void my_func(constant float ＊A, constant float ＊B)
{
    constant float4 ＊p = A;        //ok。p为私有变量,指向float4类型的常量指针
    constant float a;               //error,没有初始化
    constant float b = 2.0f;        //ok,在编译时初始化
    p[0] = (float4)(1.0f);          //error, 不能修改p的值
    char ＊c = "OpenCL";            //字符串"OpenCL"分配在常量地址空间
}

e4在私有地址空间，而不是常量地址空间。

私有地址空间

内核函数中未使用地址空间修饰符的变量、非内核函数中声明的所有变量，以及函数参数，都在私有地址空间中。私有地址空间用来描述一个工作项的私有变量，这些私有变量不能在一个工作组中的不同工作项之间或者跨工作组访问。

通用地址空间

全局、局部和私有地址空间是通用地址空间的子集，指向这三个地址空间的指针可以隐式地类型转换为指向未命名的通用地址空间的指针。但是反过来，通用地址空间不能隐式地转换为这三个地址空间。可以把指向全局、局部和私有地址空间的指针转换为一个未命名通用地址空间，反之也可以。需要注意的是上述两者的区别，一个是隐式的类型转换，一个是强制类型转化。指向常量地址空间的指针既不能强制也不能隐式地将类型转换为通用地址空间。

下面给出几个合法的使用例子：

//foo函数参数a没有地址空间修饰符,它指向通用地址空间
void foo(int ＊a)
{
    ＊a = ＊a + 2;
}
kernel void k1(local int ＊a)
{
    foo(a);               //a指向局部地址空间,隐式地转化为通用地址空间
}
kernel void k2(global int ＊a)
{
    foo(a);               //a指向全局地址空间,隐式地转化为通用地址空间
}
kernel void bar(global int ＊g, local int ＊l)
{
    int ＊var;             //根据不同条件分支, var指向全局或局部地址空间
    if(is_even(get_global_id(0)))
        var = g;
    else
        var = l;
    ＊var = 42;
}
int ＊ptr;
global int g;
ptr = &g;                 //ok
local int l;
ptr = &l;                 //ok
private int p;
ptr = &p;                 //ok
constant int c;
ptr = &c;                 //error
global int ＊gp;
local int ＊lp;
private int ＊pp;
int ＊p;
p = gp;                   //ok
p = lp;                   //ok
p = pp;                   //ok
//通用地址空间不能隐式地转换全局、局部和私有地址空间
gp = p;                   //error
lp = p;                   //error
pp = p;                   //error

在OpenCL引入通用空间修饰符之前，一个OpenCL C函数中如果有指针参数，可能需要声明多个，而引入通用空间修饰符后，只需要声明一个通用的版本即可。