OpenCL 1.2中的原子操作

下面介绍下OpenCL 1.2中的原子操作。

原子加法

函数原型如下:

int atomic_add (volatile __global int *p, int val)
unsigned int atomic_add (volatile __global unsigned int *p,
                unsigned int val)
int atomic_add (volatile __local int *p, int val)
unsigned int atomic_add (volatile __local unsigned int *p,
                      unsigned int val)

原子加法的操作数与返回类型是int或unsigned int,而可操作的存储空间可以是全局存储空间也可以是局部存储空间。此操作过程为:将参数p所指的地址内容取出,然后与参数val相加,最后将相加后的结果再写回p所指的地址中,然后返回原来修改前的p所指地址的内容。整个操作是原子的,即不可被打断的。

原子减法

函数原型如下:

int atomic_sub (volatile __global int *p, int val)
unsigned int atomic_sub (volatile __global unsigned int *p,
                              unsigned int val)
int atomic_sub (volatile __local int *p, int val)
unsigned int atomic_sub (volatile __local unsigned int *p,
                              unsigned int val)

原子减法的操作数与返回类型是int或unsigned int,而可操作的存储空间可以是全局存储空间也可以是局部存储空间。此操作过程为:将参数p所指的地址内容取出,然后与参数val相减(即*p-val),最后将相减后的结果再写回p所指的地址中,然后返回原来修改前的p所指地址的内容。整个操作是原子的,即不可被打断的。

原子数据交换

函数原型如下:

int atomic_xchg (volatile __global int *p, int val)
unsigned int atomic_xchg (volatile __global unsigned int *p,
                              unsigned int val)
float atomic_xchg (volatile __global float *p, float val)
int atomic_xchg (volatile __local int *p, int val)
unsigned int atomic_xchg (volatile __local unsigned int *p,
                              unsigned int val)
float atomic_xchg (volatile __local float *p, float val)

原子交换函数的操作数与返回类型可以是int、unsigned int或f loat。其可操作的存储空间可以是全局存储空间,也可以是局部存储空间。此操作的过程很简单,就是将val值存放进参数p所指向的存储地址,然后将p所指向的存储地址中更新之前的值作为返回结果。整个过程是不可被打断的。原子交换操作常用于对某个全局变量进行初始化或是作为互斥体(mutex)使用。

原子递增

函数原型如下:

int atomic_inc (volatile __global int *p)
unsigned int atomic_inc (volatile __global unsigned int *p)
int atomic_inc (volatile __local int *p)
unsigned int atomic_inc (volatile __local unsigned int *p)

原子递增的操作数与返回类型是int或unsigned int,而可操作的存储空间可以是全局存储空间也可以是局部存储空间。此操作过程为:将参数p所指的地址内容取出,然后与1相加(即*p + 1),最后将相加后的结果再写回p所指的地址中,然后返回原来修改前的p所指地址的内容。整个操作是原子的,即不可被打断的。

原子递减

函数原型如下:

int atomic_dec (volatile __global int *p)
unsigned int atomic_dec (volatile __global unsigned int *p)
int atomic_dec (volatile __local int *p)
unsigned int atomic_dec (volatile __local unsigned int *p)

原子递减的操作数与返回类型是int或unsigned int,而可操作的存储空间可以是全局存储空间也可以是局部存储空间。此操作过程为:将参数p所指的地址内容取出,然后与1相减(即*p - 1),最后将相减后的结果再写回p所指的地址中,然后返回原来修改前的p所指地址的内容。整个操作是原子的,即不可被打断的。

原子比较与交换

函数原型如下:

int atomic_cmpxchg (volatile __global int *p, int cmp, int val)
unsigned int atomic_cmpxchg (volatile __global unsigned int *p,
                                  unsigned int cmp, unsigned int val)
int atomic_cmpxchg (volatile __local int *p, int cmp, int val)
unsigned int atomic_cmpxchg (volatile __local unsigned int *p,
                                  unsigned int cmp, unsigned int val)

原子比较与交换的操作数与返回类型可以是int或unsigned int,而可操作的存储空间可以是全局存储空间也可以是局部存储空间。此操作与原子交换操作比较类似,而有所不同的是,此操作是先用参数cmp的值与p所指的存储地址的内容进行比较,如果相等,则将val参数的值写到p所指的存储地址中去,否则p所指的存储地址的内容还是原来的值。结果返回p所指的存储地址中更新之前的值。这个原子操作与典型的比较与交换操作相比要鸡肋很多。因为它的返回值无法体现出值的交换是否真的发生,以至于适用范围大大缩减。不过在OpenCL 2.0中,新增的atomic_compare_exchange_strong函数的返回值就是布尔类型,能体现出比较结果。所以,我们将在OpenCL 2.0的原子操作中详细介绍原子比较与交换的操作。

原子求最小值

函数原型如下:

int atomic_min (volatile __global int *p, int val)
unsigned int atomic_min (volatile __global unsigned int *p,
unsigned int val)
int atomic_min (volatile __local int *p, int val)
unsigned int atomic_min (volatile __local unsigned int *p,
unsigned int val)

原子求最小值的操作数与返回类型是int或unsigned int,而可操作的存储空间可以是全局存储空间也可以是局部存储空间。此操作过程为:将参数p所指的地址内容取出,然后与参数val比较获得最小值,最后将最小值结果再写回p所指的地址中,然后返回原来修改前的p所指地址的内容。整个操作是原子的,即不可被打断的。

原子求最大值

函数原型如下:

int atomic_max (volatile __global int *p, int val)
unsigned int atomic_max (volatile __global unsigned int *p,
                              unsigned int val)
int atomic_max (volatile __local int *p, int val)
unsigned int atomic_max (volatile __local unsigned int *p,
                              unsigned int val)

原子求最大值的操作数与返回类型是int或unsigned int,而可操作的存储空间可以是全局存储空间也可以是局部存储空间。此操作过程为:将参数p所指的地址内容取出,然后与参数val比较获得最大值,最后将最大值结果再写回p所指的地址中,然后返回原来修改前的p所指地址的内容。整个操作是原子的,即不可被打断的。

原子按位与

函数原型如下:

int atomic_and (volatile __global int *p, int val)
unsigned int atomic_and (volatile __global unsigned int *p,
                              unsigned int val)
int atomic_and (volatile __local int *p, int val)
unsigned int atomic_and (volatile __local unsigned int *p,
                              unsigned int val)

原子按位与的操作数与返回类型是int或unsigned int,而可操作的存储空间可以是全局存储空间也可以是局部存储空间。此操作过程为:将参数p所指的地址内容取出,然后与参数val做按位与操作(即*p & val),最后将结果再写回p所指的地址中,然后返回原来修改前的p所指地址的内容。整个操作是原子的,即不可被打断的。

原子按位或

函数原型如下:

int atomic_or (volatile __global int *p, int val)
unsigned int atomic_or (volatile _global unsigned int *p,
unsigned int val)
int atomic_or (volatile __local int *p, int val)
unsigned int atomic_or (volatile __local unsigned int *p,
                            unsigned int val)

原子按位或的操作数与返回类型是int或unsigned int,而可操作的存储空间可以是全局存储空间也可以是局部存储空间。此操作过程为:将参数p所指的地址内容取出,然后与参数val做按位或操作(即*p | val),最后将结果再写回p所指的地址中,然后返回原来修改前的p所指地址的内容。整个操作是原子的,即不可被打断的。

原子按位异或

函数原型如下:

int atomic_xor (volatile __global int *p, int val)
unsigned int atomic_xor (volatile __global unsigned int *p,
                              unsigned int val)
int atomic_xor (volatile __local int *p, int val)
unsigned int atomic_xor (volatile __local unsigned int *p,
                              unsigned int val)

原子按位异或的操作数与返回类型是int或unsigned int,而可操作的存储空间可以是全局存储空间也可以是局部存储空间。此操作过程为:将参数p所指的地址内容取出,然后与参数val做按位异或操作(即*p^val),最后将结果再写回p所指的地址中,然后返回原来修改前的p所指地址的内容。整个操作是原子的,即不可被打断的。
在OpenCL 1.2中已经定义了那么多种类丰富的原子操作。下面我们将通过一个简单的例子来介绍一下原子加法的使用。这个例子就是上述计算向量内积的更通用的版本。例如,要计算一对4096个元素向量的内积,我们仍然是先计算一个工作组中所有工作项对应两个向量数据元素的积,然后先保存到局部存储空间。最后每个工作组的第一个工作项计算所有当前工作组所有局部数据中数据元素的和,然后用原子加法把该结果与目的输出存储空间的值相加。我们将上述代码从program = clCreateProgramWithSource(context, 1,(const char**)&kernelSource,(const size_t*)&kernelLength, &ret);语句一直到FINISH标签之前,替换为以下代码:

program = clCreateProgramWithSource(context, 1,
(const char **)&kernelSource,
(const size_t *)&kernelLength,
&ret);
//获取最大工作组大小
size_t maxWorkGroupSize = 0;
clGetDeviceInfo(device_id, CL_DEVICE_MAX_WORK_GROUP_SIZE,
                  sizeof(maxWorkGroupSize), &maxWorkGroupSize, NULL);
/* 在编译选项中定义一个名为GROUP_NUMBER_OF_WORKITEMS的宏,
用于指定每个工作组一共有多少工作项*/
sprintf(kernelSource, "-D GROUP_NUMBER_OF_WORKITEMS=%zu",
        maxWorkGroupSize);
ret = clBuildProgram(program, 1, &device_id, kernelSource, NULL,
                        NULL);
if (ret != CL_SUCCESS)
{
}
    size_t len;
    char buffer[8 * 1024];
    printf("Error: Failed to build program executable!\n");
    clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG,
                              sizeof(buffer), buffer, &len);
    printf("%s\n", buffer);
    goto FINISH;
//kernelSource后面不再使用,这里可以立即对它释放
free(kernelSource);
kernelSource = NULL;
//创建内核函数
kernel = clCreateKernel(program, "kernel_test", &ret);
if(kernel == NULL)
{
}
    puts("Kernel failed to create!");
    goto FINISH;
ret = clSetKernelArg(kernel, 0, sizeof(cl_mem),
                        (void *)&dstMemObj);
ret |= clSetKernelArg(kernel, 1, sizeof(cl_mem),
                          (void *)&src1MemObj);
ret |= clSetKernelArg(kernel, 2, sizeof(cl_mem),
                          (void *)&src2MemObj);
if(ret != CL_SUCCESS)
{
}
    puts("Set arguments error!");
    goto FINISH;
//这里指定将总共有nWorkItems个工作项
//然后,每个工作组也含有nWorkItems个工作项
//我们这里再复用evt1来跟踪内核程序1的执行状态
ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL,
                                  (const size_t[]){
                                        contentLength / sizeof(int)
                                  },
                                  (const size_t[])
                                  {
                                        maxWorkGroupSize
                                  }, 2,
                                  (const cl_event[])
                                  {
                                        evt1, evt2
                                  }, NULL);
if(ret != CL_SUCCESS)
{
    puts("kernel1 execution failed");
    goto FINISH;
}
//这里用clFinish做命令执行同步
clFinish(command_queue);
//准备做校验
pDeviceBuffer = malloc(sizeof(int));
//这里使用阻塞的方式读取数据
clEnqueueReadBuffer(command_queue, dstMemObj, CL_TRUE, 0,
                        1 * sizeof(int),
                        pDeviceBuffer, 0, NULL, NULL);
//做数据校验
int sum = 0;
for(int i = 0; i 〈 contentLength / sizeof(int); i++)
    sum += pHostBuffer[i] * pHostBuffer[i];
if(sum == *pDeviceBuffer)
    puts("Result OK!");
else
    puts("Result NG!");

在主机端代码,我们是将两个长度为16 × 1024 × 1024的int类型数据元素的向量准备做内积操作。我们在内核代码中定义了GROUP_NUMBER_OF_WORKITEMS宏来指定每个工作组包含多少工作项。尽管我们可以在内核程序中使用get_local_size内建函数来获得当前工作组包含多少工作项,但是由于OpenCL C中不能像C99那样定义可变长数组。因此,定义数组长度时必须用编译阶段即可确定的常量。所以我们这里使用宏来指定。下面给出内核程序:

__kernel void kernel_test(__global int *pDst,
                                __global int *pSrc1, __global int *pSrc2)
{
    local int tmpBuffer[GROUP_NUMBER_OF_WORKITEMS];
    int index = get_global_id(0);
    //让所有工作项的第一个对输出结果初始化为0,以便于后续累加计算
    if(index == 0)
    atomic_xchg(pDst, 0);
    //这里将索引变为当前工作组中的每个工作项的索引
    index = get_local_id(0);
    tmpBuffer[index] = pSrc1[index] * pSrc2[index];
    barrier(CLK_LOCAL_MEM_FENCE);
    if(index == 0)
    {
       int sum = 0;
       //对当前工作组中对应的所有向量元素求和
       for(int i = 0; i 〈 GROUP_NUMBER_OF_WORKITEMS; i++)
      sum += tmpBuffer[i];
      //将结果累加到输出存储地址
      atomic_add(pDst, sum);
      //pDst[0] += sum;
    }
}

对两个向量求内积的过程之前已经描述了。这里,对输出结果地址的内容求和再写回去必须使用原子操作,如果大家把“atomic_add(pDst, sum);”这条语句屏蔽掉,而换用下面的“pDst[0]+=sum;”,那么结果就很有可能是不正确的。
在这个demo中,我们不知不觉其实已经涉及了不同工作组之间的数据同步。在这个例子中,每个工作组的头一个工作项都会将本工作组计算好的求和结果与输出地址的内容相加,然后再写回去。由于OpenCL并没有提供任何对不同工作组之间工作项的栅栏同步操作,因此我们对工作组之间的同步往往也是借助于原子操作或是以原子操作作为技术前提的同步原语(synchronization primitive)。

赞(2)
未经允许不得转载:极客笔记 » OpenCL 1.2中的原子操作

评论 抢沙发

  • 昵称 (必填)
  • 邮箱 (必填)
  • 网址
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映射