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函数中如果有指针参数,可能需要声明多个,而引入通用空间修饰符后,只需要声明一个通用的版本即可。

赞(1)
未经允许不得转载:极客笔记 » OpenCL 地址空间修饰符

评论 抢沙发

  • 昵称 (必填)
  • 邮箱 (必填)
  • 网址
OpenCL 基本概念
OpenCL 是什么OpenCL 平台模型OpenCL 执行模型OpenCL 上下文简介OpenCL 命令队列简介OpenCL 在设备上执行内核OpenCL 存储器区域OpenCL 存储器对象OpenCL 共享虚拟存储器OpenCL 与OpenGL
OpenCL 基础教程
OpenCL 在Windows上搭建开发环境OpenCL 在Linux上搭建开发环境OpenCL 在OS X上搭建开发环境OpenCL 第一个程序OpenCL 平台OpenCL 设备OpenCL 创建上下文OpenCL 创建命令队列OpenCL 创建Program对象OpenCL 编译Program对象OpenCL 查询和管理Program对象OpenCL 创建内核对象OpenCL 设置内核参数OpenCL 查询和管理内核对象OpenCL 执行内核OpenCL 编写内核代码OpenCL 错误处理
OpenCL C特性
OpenCL 地址空间修饰符OpenCL 函数修饰符OpenCL 对象访问修饰符OpenCL 标量数据类型OpenCL 为什么要有矢量数据类型OpenCL 矢量初始化OpenCL 读取和修改矢量分量OpenCL 运算符OpenCL 维度和工作项OpenCL 工作组OpenCL 矢量数据拷贝OpenCL 异步拷贝和预取OpenCL 数学函数OpenCL 公共函数OpenCL 几何函数OpenCL 整数函数OpenCL 关系函数OpenCL 杂项矢量函数OpenCL 同步函数OpenCL 原子函数OpenCL 内建图像读函数OpenCL 内建无采样器图像读函数OpenCL 内建图像写函数OpenCL 内建图像查询函数OpenCL 工作组函数OpenCL 内建管道读/写函数OpenCL 内建工作组管道读/写函数OpenCL 内建管道查询函数OpenCL 设备队列OpenCL Blocks语法OpenCL 设备队列相关函数OpenCL 子内核存储器可见性OpenCL 设备队列的使用示例
OpenCL 存储器对象
OpenCL 存储器对象OpenCL 分配缓冲区对象OpenCL 创建子缓冲区对象OpenCL 图像对象和采样器对象OpenCL 图像对象OpenCL 图像格式描述符OpenCL 图像描述符OpenCL 图像对象查询OpenCL 采样器对象OpenCL 主机端采样器对象OpenCL 设备端采样器对象OpenCL 图像旋转示例OpenCL 管道OpenCL 创建管道对象OpenCL 管道对象查询OpenCL 主机与设备间数据传输OpenCL 图像对象主机与设备间数据拷贝OpenCL 缓冲区对象数据填充OpenCL 图像对象数据填充OpenCL 缓冲区对象间数据传输OpenCL 图像对象和缓冲区对象间数据拷贝OpenCL 缓冲区对象映射OpenCL 图像对象映射OpenCL 解映射OpenCL 共享虚拟存储器OpenCL SVM缓冲创建与释放OpenCL SVM缓冲映射与解映射OpenCL SVM缓冲填充与拷贝OpenCL SVM类型OpenCL SVM特性OpenCL 共享虚拟存储器示例OpenCL 存储器一致性模型OpenCL 存储器次序规则OpenCL 原子操作的存储器次序规则OpenCL 栅栏操作的存储器次序规则OpenCL 工作组函数的存储器次序规则OpenCL 主机端与设备端命令的存储器次序规则OpenCL 关于存储器次序在实际OpenCL计算设备中的实现
OpenCL 同步及事件机制
OpenCL 同步及事件机制OpenCL 主机端的OpenCL同步OpenCL OpenCL事件机制OpenCL 对OpenCL事件的标记和栅栏OpenCL 内核程序中的同步OpenCL 工作组内同步OpenCL 原子操作OpenCL 1.2中的原子操作OpenCL 2.0中的原子操作OpenCL 局部存储器与全局存储器间的异步拷贝OpenCL 工作组间同步
OpenCL 与OpenGL互操作
OpenCL 与OpenGL互操作OpenCL 从一个OpenGL上下文来创建OpenCL上下文OpenCL 使用OpenGL共享的缓存对象OpenCL 使用OpenGL纹理数据OpenCL 共享OpenGL渲染缓存OpenCL 从一个OpenCL存储器对象查询OpenGL对象信息OpenCL 访问共享对象的OpenCL与OpenGL之间的同步OpenCL AMD Cayman架构GPUOpenCL AMD GCN架构的GPUOpenCL NVIDIA CUDA兼容的GPUOpenCL NVIDIA GPU架构的执行模型OpenCL NVIDIA GPU的全局存储器OpenCL NVIDIA GPU的局部存储器OpenCL ARM Mali GPU硬件架构OpenCL ARM Mali GPU存储器层次OpenCL ARM Mali GPU OpenCL映射