OpenCL 原子函数

OpenCL 2.0中原子操作的用法与C11标准相同，与OpenCL 1.2中原子操作用法有差异，对于OpenCL 1.2支持的原子操作，通常以Tatomic_key命名，简单罗列如下表所示。

我们主要以OpenCL 2.0标准，对于OpenCL 2.0以下标准的原子操作本书就不展开阐述，需要的读者可以查看OpenCL 1.2标准文档。

在OpenCL 2.0中，完全把原子操作与其他操作分开。原子操作有自己的原子数据类型。常规的OpenCL操作（=、+、〉、〈及其他操作）对这些原子数据类型变量操作会导致内核代码编译失败。OpenCL 2.0中定义的原子数据类型有：atomic_int、atomic_uint、atomic_long、atomic_ulong、atomic_f loat、atomic_double、atomic_intptr_t、atomic_uintptr_t、atomic_size_t和atomic_ptrdiff_t。OpenCL支持的标量数据类型，在OpenCL 2.0中大部分都支持原子操作。不过需要注意的是，对于atomic_long、atomic_ulong、atomic_double、atomic_intptr_t、atomic_uintptr_t、atomic_size_t和atomic_ptrdiff_t， OpenCL设备是选择性支持的。

对于定义函数局部的原子变量，我们可以使用atomic_init（）来初始化它的值。对于全局原子变量，要使用宏ATOMIC_VAR_INIT（）来初始化它的值。如果要在常规操中使用原子变量的值，需要使用atomic_load（）函数，把原子变量的值拷贝到常规变量中再使用。如果要把常规变量中的值拷贝到原子变量中，需要使用atomic_store（）函数。OpenCL 2.0内建了对原子变量进行算术、逻辑、比较、交换等操作的内建函数。现在我们就来讲解下这些内建函数，如下表所示。

如下列举一些原子操作的例子：

kernel void DoubleTest(global int ＊a)
{
    local atomic_int guide;
    int id = get_global_id(0);
    a[id] = id;
    work_group_barrier(CLK_GLOBAL_MEM_FENCE);
    atomic_fetch_add((atomic_int ＊)&a[2], 3); //a[2]=386,即128＊3+2
    if(id == 0)
    {
        atomic_init(&guide, 50); //guide=50
    }
    work_group_barrier(CLK_LOCAL_MEM_FENCE);
    a[id] = atomic_load(&guide); //a[id]=50
    a[id] += 100;
    if(id == 0)
    {
        atomic_store(&guide, a[id]); //guide=150
        a[id] = atomic_exchange(&guide, 10) ;
        .//a[0]=150,guide=10;
    }
    work_group_barrier(CLK_LOCAL_MEM_FENCE);
    atomic_fetch_add(&guide, 22); //guide=2526,即(22＊128+10)
    atomic_fetch_add_explicit(&guide, 1, memory_order_relaxed,
                                    memory_scope_device); //guide=2954
    work_group_barrier(CLK_LOCAL_MEM_FENCE);
    if(id == 1)
    {
        a[1] = atomic_load(&guide); //a[1]=2954;
    }
}

上述例子展示了如何使用上表中的几个原子操作函数。细心的读者会发现，在表中对函数中的memory_order和memory_scope参数并未过多说明。

我们知道对于图像对象的访问修饰符，有read_write修饰符。对于read_write修饰符的图像，我们用如下函数：

void atomic_work_item_fence(cl_mem_fence_flags flags,
                                  memory_order order,
                                  memory_scope scope)

来确保一个工作项写入图像操作对这个工作项接下来的读图像操作是可见的。也就是设置一个栅栏，同步随后的读图像操作。例如：

kernel void ImageProcess(read_write image2d_t image,
                              int threshold)
{
}
int2 coord = (int2)(get_global_id(0), get_global_id(1));
float4 p00 = read_imagef(image, coord + (int2)(-1, -1));
float4 p10 = read_imagef(image, coord + (int2)(0, -1));
float4 p20 = read_imagef(image, coord + (int2)(1, -1));
float4 p01 = read_imagef(image, coord + (int2)(-1, 0));
float4 p21 = read_imagef(image, coord + (int2)(1, 0));
float4 p02 = read_imagef(image, coord + (int2)(-1, 1));
float4 p12 = read_imagef(image, coord + (int2)(0, 1));
float4 p22 = read_imagef(image, coord + (int2)(1, 1));
float3 gx = -p00.xyz + p20.xyz + 2 ＊ (p21.xyz - p01.xyz) -
   p02.xyz + p22.xyz;
float3 gy = -p00.xyz - p20.xyz + 2 ＊ (p12.xyz - p10.xyz) +
   p02.xyz + p22.xyz;
float3 g = native_sqrt(gx ＊ gx + gy ＊ gy);
write_imagef(image, coord, (float4)(g.x, g.y, g.z, 1.0f ));
atomic_work_item_fence(CLK_IMAGE_MEM_FENCE,
      memory_order_acq_rel,
      memory_scope_work_item);
float4 temp = read_imageui(image, imageSampler, coord);
temp.x = select(255, 0, (uint)(temp.x 〈 threshold));
write_imageui(image, coord, temp);