当前位置：极客笔记 > OpenCL 教程 > OpenCL 同步及事件机制 > OpenCL 1.2中的原子操作

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）。

Camera课程

Camera API2 教程

API2 教程
Camera 性能教程

Perf 教程
Camera Native Framework 教程

Native FW 教程
Camera QCOM ISP 教程

ISP 教程
Camera QCOM Sensor 教程

Sensor 教程

Python教程

Python 教程

Python 教程
NumPy 教程

NumPy 教程
Django 教程

Django 教程

Java教程

Java 教程

Java 教程

Web教程

JavaScript 教程

JavaScript 教程
jQuery 教程

jQuery 教程
CSS 教程

CSS 教程
PHP 教程

PHP 教程
Laravel 教程

Laravel 教程
TypeScript 教程

TypeScript 教程
WordPress 教程

WordPress 教程

数据库教程

SQL 教程

SQL 教程
SQLite 教程

SQLite 教程
MySQL 教程

MySQL 教程
PL/SQL 教程

PL/SQL 教程
MongoDB 教程

MongoDB 教程

图形图像教程

OpenCV 教程

OpenCV 教程

办公软件教程

Excel 教程

Excel 教程

Linux教程

Linux命令全集

Linux命令全集
Linux内核API

Linux内核API
LVGL 教程

LVGL 教程

计算机教程

操作系统 教程

操作系统 教程
计算机网络 教程

计算机网络 教程
C语言 教程

C语言 教程
C++ 教程

C++ 教程
Swift 教程

Swift 教程

大数据教程

Hadoop 教程

Hadoop 教程
Spark 教程

Spark 教程
Scala 教程

Scala 教程

开发工具教程

Git 教程

Git 教程
Jenkins 教程

Jenkins 教程
ChatGPT 教程

ChatGPT 教程
IntelliJ 教程

IntelliJ 教程

OpenCL 同步及事件机制