极客笔记OpenCL 2.0的原子函数集是C11（ISO/IEC 9899:2011）中原子函数操作的子集，所以对原子操作的大部分概念可以在C11的国际标准中找到。不过为了帮助读者更好地理解这些概念，在本节中，我们将给大家详细解析这些术语和相关概念，尽管对于很多细节而言，应用开发者可能并不会十分关心。但是倘若您是OpenCL的实现者，那么这些细节就可能对您而言十分重要了。这里还会再简单提一下。
我们先介绍一下，OpenCL 2.0官方手册中对原子操作函数的描述规则。我们就拿一个比较复杂的atomic_compare_exchange_strong_explicit函数原型为例来做介绍：
bool atomic_compare_exchange_strong_explicit(volatile A ＊object,
C ＊expected, C desired, memory_order success,
memory_order failure, memory_scope scope)
在OpenCL 2.0 C语言的所有原子函数声明中，A类型表示一个原子对象类型。在OpenCL 2.0中，如果要将一个变量作为原子对象，需要使用原子数据类型。OpenCL 2.0所支持的原子数据类型如下：
- atomic_int：原子的32位带符号整型；
- atomic_uint：原子的32位无符号整型；
- atomic_long：原子的64位带符号整型（此数据类型需要OpenCl扩展cl_khr_int64_base_atomics和cl_khr_int64_extended_atomics的支持）；
- atomic_ulong：原子的64位无符号整型（此数据类型需要OpenCl扩展cl_khr_int64_base_atomics和cl_khr_int64_extended_atomics的支持）；
- atomic_f loat：原子的32位单精度浮点型；
- atomic_double：原子的64位双精度浮点型（此数据类型需要OpenCL设备支持64位双精度浮点，并且需要OpenCl扩展cl_khr_int64_base_atomics和cl_khr_int64_extended_atomics的支持）；
- atomic_intptr_t：原子的int*长度类型；
- atomic_uintptr_t：原子的uint*长度类型；
- atomic_size_t：原子的size_t类型；
- atomic_ptrdiff_t：原子的void*长度类型。
对于atomic_intptr_t、atomic_uintptr_t、atomic_size_t和atomic_ptrdiff_t，如果计算设备的地址是64位的，那么只有当OpenCL设备支持了cl_khr_int64_base_atomics和cl_khr_int64_extended_atomics扩展才能支持这些类型。
而类型C表示相应的非原子类型，也就是一般的int、uint、f loat等。
memory_order是一个枚举类型，用于指明常规的（非原子的）存储器细致的同步操作。存储器次序往往表征了不同线程（在OpenCL中即为不同的工作项）之间对同一存储对象的访存操作的可见性。关于存储器次序的详细描述请见5.6节。这里我们为了避免读者前后翻阅，所以再简单罗列一下。memory_order有以下几个枚举值（它们都是相互排斥的，不能用按位或来联用）：
1）memory_order_relaxed：松弛的存储器次序。这个存储器次序意味着对存储器次序不做任何限制。如果用这个存储器次序进行原子操作（如使用atomic_fetch_add），原子对象也能安全地递增。但是该原子操作无法保证相对其他访存操作与它之间的次序。memory_order_relaxed也是次序最弱的。
2）memory_order_acquire：具有获得语义的一次同步操作（如栅栏或原子操作）从和它进行同步的释放操作那里“获得”副作用：如果一次获得操作跟一次释放操作进行同步，那么执行单元的获得操作将会看到在那个释放操作之前的所有副作用（并且也有可能看到后续的副作用）。而在使用memory_order_acquire存储器次序的当前执行单元中，其后面的访存操作不会被重新编排到此获得语义的访存操作之前执行。我们在编写OpenCL程序时可以使用一个“获得”语义以安全地观察到另一个执行单元对某些共享存储变量的修改情况。
3）memory_order_release：具有释放语义的一个同步操作（栅栏或原子操作）“释放”对与它进行同步的一个获得操作的副作用。在释放之前的所有副作用都被包含在这个释放中。在执行释放语义的当前执行单元中，执行释放语义的访存操作之前所有访存操作不会被重新安排到此释放语义访存操作之后执行。我们在编写OpenCL程序时可以使用一个“释放”语义以将当前执行单元对某些共享存储变量的存储操作暴露给其他执行单元，使得其他执行单元能安全地观察到这些共享存储变量已被修改。
4）memory_order_acq_rel：同时具备获得语义与释放语义的同步操作，具有memory_order_acquire与memory_order_release存储器次序的特性。该存储器次序一般用于读-修改-写操作。在使用memory_order_acq_rel存储器次序的当前执行单元上下文中，对某个原子对象使用memory_order_acq_rel存储器次序进行读-修改-写操作时，在加载该原子对象之后的所有访存操作都不能重新安排到该加载操作之前；而在修改完该原子对象进行存储操作时，所有在此存储操作之前的访存操作都不能重新安排到此存储操作之后。
5）memory_order_seq_cst：每个执行单元的加载和存储都能以程序次序被观察到。并且来自不同执行单元的加载和存储操作会以简单的交错形式被观察到。该存储器次序与memory_order_acq_rel语义上差不多，不过比memory_order_acq_rel存储器次序又多了一个单一总和次序，即在所有对同一个共享的原子对象使用memory_order_seq_cst存储器次序操作的执行单元中所观察到的修改都是以相同的次序完成的。memory_order_seq_cst被称为顺序一致性的（sequentially consistent）存储器次序。这也是最强的存储器次序。
存储器次序从广义上讲是针对在某个支持弱次序读写操作的处理器上，当一个线程对某一存储对象使用弱次序的写之后，在另一个线程上去读这个存储对象，所读到的值可能是在第一个线程写之前的值，即便保证了第二个线程是在第一个线程写操作之后再去读的。由于第一个线程对此存储对象的写对第二个对象而言不可见。因此，如果要确保第二个线程在稍后读的时候能读到第一个线程更新后的值，那么需要使用存储器栅栏操作或是第一个线程使用对其他线程可见的存储器次序的写操作。
另外，对于获得/释放存储次序操作在不同的处理器、编程语言上可能会有一些实现细节上的不同。例如，在ARMv8架构上，获得次序操作是与加载绑定使用的，这也被称为“load acquire”；而释放次序操作是与存储绑定使用的，这也被称为“tore release”。而Intel处理器目前只提供了MFENCE指令，同时对加载和存储做栅栏操作，相当于memory_order_seq_cst，而没有其他存储器次序的操作。然而在C11标准中，也明确指定了如下规则，这也是被OpenCL标准所部分采用的：
1）存储器写操作只能使用memory_order_release、memory_order_acq_rel或memory_order_seq_cst这些存储器次序来对指定的位置执行存储操作。
2）存储器加载操作只能使用memory_order_acquire、memory_order_acq_rel或memory_order_seq_cst这些存储器次序来对指定的位置执行加载操作。
下面介绍最后一个scope参数。memory_scope指定了存储器次序（memory_order）作用于哪个区域范围。memory_scope也是一个枚举类型。不过与memory_order不同，这个枚举变量的所有枚举值可以使用按位或（|）来联合使用。
1）memory_scope_work_item：当前原子操作的存储器次序作用于当前工作项，即只有当前工作项对此存储器次序可见。这个枚举值只有针对atomic_work_item_fence函数，并且其f lags参数被指定为CLK_IMAGE_MEM_FENCE时方可使用。在其他情况下，使用此枚举值是无效的。
2）memory_scope_work_group：当前原子操作的存储器次序作用于当前工作组，即只有当前工作组中的所有工作项对此存储器次序可见。
3）memory_scope_device：当前原子操作的存储器次序作用于当前计算设备，即只有当前计算设备中的所有工作项对此存储器次序可见。
4）memory_scope_all_svm_devices：当前原子操作的存储器次序作用于当前共享虚拟存储器上的所有设备，即当前共享虚拟存储空间上的所有计算设备，包括主机端对此存储器次序可见。
对原子对象的初始化
OpenCL 2.0提供了两种对原子对象初始化的方式，并且分别用于两种不同的情况。
1）对全局原子对象的初始化需要使用宏ATOMIC_VAR_INIT。例如：
global volatile atomic_int atom_obj = ATOMIC_VAR_INIT(10);
2）对定义在函数体内的临时原子对象进行初始化，需要使用atomic_init函数。这里要注意的是，atomic_init函数对原子对象的初始化并不是原子的，因此要用其他同步手段来对工作项进行同步，如果有需要的话。例如：
local volatile atomic_int local_atom_obj;
if (get_local_id(0) == 0)
atomic_init(&local_atom_obj, 10);
work_group_barrier(CLK_LOCAL_MEM_FENCE);
上述短小的示例代码中使用每个工作组的第一个工作项对当前工作组中的局部存储空间的local_atom_obj原子对象进行初始化为10。然后，使用work_group_barrier函数进行工作项的同步。
存储器栅栏操作
OpenCL 2.0中的存储器栅栏操作函数原型为：
void atomic_work_item_fence(cl_mem_fence_flags flags,
memory_order order, memory_scope scope)
这里，flags的取值可以是CLK_GLOBAL_MEM_FENCE、CLK_LOCAL_MEM_FENCE、CLK_IMAGE_MEM_FENCE，或是这些枚举值用按位或（|）相联合后的值。order可以用所有上述所描述的值，不过如果使用了memory_order_relaxed，那么这个函数调用将不会有任何效果。scope值可以使用上述所描述的所有枚举值。
在OpenCL 2.0之前的版本中，OpenCL C提供了mem_fence、read_mem_fence，以及write_mem_fence来分别对读写操作产生影响，写操作和读操作由f lags所指定的存储器访问次序同步。使用存储器栅栏操作可确保在栅栏摆放的这一点，之前所指定类型的所有访存操作都能在这点完成，并对所有工作项可见。而在OpenCL 2.0中，我们通过order参数来指定对哪种存储器次序语义做栅栏操作。如果是memory_order_relaxed，那么这个调用将毫无效果；如果是memory_order_acquire，那么将针对所有acquire语义的访存做栅栏操作，在一定程度上类似于read_mem_fence；如果是memory_order_release，那么将针对所有release语义的访存做栅栏操作，在一定程度上类似于write_mem_fence；如果是memory_order_acq_rel，那么将针对所有acquire+release语义，在一定程度上类似于mem_fence；如果是memory_order_seq_cst，那么将针对所有sequential consistent acquire+release语义，在一定程度上类似于mem_fence，而存储器次序将会严格按照串行顺序一致性的要求来执行。
这里还需要注意的是，如果我们使用了以read_write限定符修饰的一个图像对象（如image2d_t类型），那么必须调用atomic_work_item_fence（memory_scope_work_item）函数使得对该图像对象的写对于当前工作项而言可见，以至于后续再对该图像对象的读能确保读到的是更新后的值。这里通过采样器对图像对象的读写非常特殊。由于GPGPU通过专门的硬件纹理单元来对图像做采样，其中还包括了各种插值、裁减等算法，因此即便对于当前的工作项而言，对图像对象的读写的存储器次序可能也是不可见的。
原子加载与存储
原子加载和存储操作从字面上看似乎是对加载和存储做原子操作。但是“原子性”这个术语我们往往针对具有“读-修改-写”（RMW）特征的操作来说的，原子性确保整个操作是不可被打断的。那么在OpenCL 2.0中的原子加载和存储操作是怎么一回事呢？由于OpenCL 2.0新增了存储器次序这个特征，而我们往往可以通过原子加载操作来发动一次获得次序语义；而通过原子存储操作来发动一次释放次序语义。当然，原子加载和存储的指定对象必须是原子对象（即带有atomic_前缀类型的变量）。
原子存储的函数原型如下：
void atomic_store(volatile A ＊object, C desired)
void atomic_store_explicit(volatile A ＊object, C desired,
memory_order order)
void atomic_store_explicit(volatile A ＊object,
C desired,
memory_order order,
memory_scope scope)
这里，order参数只能使用memory_order_release、memory_order_seq_cst或memory_order_relaxed。这个函数是将参数desired的值存储到object所指向的原子对象的地址中去。
原子加载操作的函数原型如下：
C atomic_load(volatile A ＊object)
C atomic_load_explicit(volatile A ＊object, memory_order order)
C atomic_load_explicit(volatile A ＊object, memory_order order,
memory_scope scope)
这里，order参数只能是memory_order_acquire，memory_order_seq_cst或memory_order_relaxed。这个参数是将object所指向的原子对象地址的内容取出，然后作为返回值返回。
原子交换函数
OpenCL 2.0中的原子交换函数与OpenCL 1.2中的差不多。不过它增加了存储器次序以及工作项作用域范围的指定。函数原型如下：
C atomic_exchange(volatile A ＊object, C desired)
C atomic_exchange_explicit(volatile A ＊object, C desired,
memory_order order)
C atomic_exchange_explicit(volatile A ＊object, C desired,
memory_order order, memory_scope scope)
该函数的功能是将desired值存放到object所指向的原子对象中，然后将object所指向的原子对象修改之前的值作为返回值返回。
原子比较与交换
我们在OpenCL 1.2中的原子操作中介绍了OpenCL 1.2中关于原子比较与交换的函数，并且也谈到了在OpenCL 1.2中，此函数的实用性并不大。而到了OpenCL 2.0中，它的威力将被大大地发挥出来。我们在稍后介绍对数组元素做乘积的例子里也会看到我们将如何通过原子比较与交换来使得对整数的乘法操作变为一个原子操作。函数原型如下：
bool atomic_compare_exchange_strong(volatile A ＊object,
C ＊expected, C desired)
bool atomic_compare_exchange_strong_explicit(volatile A ＊object,
C ＊expected, C desired,
memory_order success,
memory_order failure)
bool atomic_compare_exchange_strong_explicit(volatile A ＊object,
C ＊expected, C desired,
memory_order success,
memory_order failure,
memory_scope scope)
这里，参数expected指向的是要同object所指向的原子对象进行比较的变量。这个函数的功能操作是：比较object所指向的原子对象的值与expected所指向的变量值是否相同。如果相同，那么将desired值存储到object所指向的原子对象中去，并返回true；否则，将object所指向的原子对象的值拷贝到expected所指向的变量中去，并返回false。整个操作都是原子的，即不可被打断的。
这里含有两个memory_order参数。success指明了比较相同时所采取的存储器次序；failure则指明了比较结果不相同时所采取的存储器次序。这里要注意的是，failure存储次序只能是memory_order_relaxed或memory_order_acquire。而且，failure的存储器次序强度不能超过success的强度。也就是说，如果success用了memory_order_relaxed，那么failure也必须只能使用memory_order_relaxed。
可以看到，这三个函数后面都有_strong后缀，那么有strong就会有weak。下面列出OpenCL 2.0中weak属性的原子比较与交换函数的原型：
bool atomic_compare_exchange_weak(volatile A ＊object,
C ＊expected, C desired)
bool atomic_compare_exchange_weak_explicit(volatile A ＊object,
C ＊expected, C desired,
memory_order success, memory_order failure)
bool atomic_compare_exchange_weak_explicit(volatile A ＊object,
C ＊expected, C desired,
memory_order success, memory_order failure,
memory_scope scope)
这里我们可以看到，weak属性的原子比较交换函数不管在返回值还是函数参数上都与strong的完全一样。那么，这里为何要引入weak的形式呢？因为有些处理器（如ARMv7架构的CPU）不支持原生的原子比较与交换操作，但是它们支持另一种更灵活的原子操作方式，被称为LL-SC（Load-Locked Store-Conditional）。这种原子操作是通过在存储器总线上设立一个监视器，在做上锁加载的时候监视器会观察到这个行为，然后将这部分存储器块进行上锁操作。当后续的一个带条件的存储对此存储器块进行操作之后，监视器进行解锁。因此，对于带条件存储操作而言，它是原子的。只要有一个带条件存储操作对锁定的存储器块操作过，那么该存储器块就变为非锁定状态（即正常状态）了，其后续的带条件存储操作都会失败，即不会将指定的数据写入该存储器块中去。如果对每个存储单元（如一个4字节的存储单元）进行观察，对于处理器的设计而言成本太大。因此，ARMv7核心的实现往往会对一小块存储空间进行监视，如一个128字节大小的存储块。那么这样就会引发一个问题。当我们定义了两个变量，比如：int objs[2]；objs是一个含有2个32位带符号的整型数组，它们的存储器空间都是毗邻的。这时候如果对它们同时使用LL-SC操作会发生什么呢？不管是objs[0]先被使用SC操作还是objs[1]先被SC操作，结论都一样。我们这里不妨假设objs[0]先做SC操作。那么objs[1]在使用SC的时候已经处于非锁定状态，尽管此时没有任何其他线程对objs[1]做任何修改，它的值仍然为交换之前的值，但是SC操作还是会失败。所以，这就导致了（*expected）与（*object）完全相同，但交换不成功的情况。采用weak形式的原子比较交换则允许这种情况发生。
当然，对于能够支持OpenCL 2.0的GPGPU而言，基本都能原生支持原子比较与交换操作。另外，在2015年夏季发布的ARMv8.1架构中也将引入原子比较与交换相关指令。这样不仅对于OpenCL 2.0，而且还能满足对C11标准的支持。所以，如果我们的OpenCL内核程序要求能支持的硬件范围更广，可使用weak形式；否则，用strong也行。一般来说，strong形式会更直接，对原子比较与交换的操作逻辑更能严格遵守，不过对于不支持它的实现而言就需要库自己做些额外的工作。
到目前为止，我们已经介绍了对一个数组中所有元素做乘积所要用到的原子操作函数。下面我们再继续介绍后续的一些原子操作之前就把这个程序例子呈现给大家。我们基于上面的代码，替换从ret=clEnqueueNDRangeKernel到FINISH之间的代码：
ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL,
(const size_t[]){
contentLength / sizeof(int) / 4},
(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 product = 1;
for(int i = 0; i 〈 contentLength / sizeof(int); i++)
product ＊= pHostBuffer[i];
if(product == ＊pDeviceBuffer)
puts("Result OK!");
else
puts("Result NG!");
这里，我们在内核程序中，只针对pSrc1的数据内容，即就对pSrc1所指向的int数组元素进行求积。另外，我们在分配工作项总数的时候，用的是contentLength/sizeof（int）/4，表示我们将在每个工作项中一次处理4个元素。这里也是给大家提供一个对某个数组的所有元素求和、求积等操作的优化思路。对于一个工作组，里面所有工作项所访问的存储空间最好是连续的。下面则列出相应的内核程序：
//这里将pDst声明为atomic_int＊,这样pDst[0]就能被当作为原子对象了
__kernel void kernel_test(__global atomic_int ＊pDst,
{
__global int ＊pSrc1, __global int ＊pSrc2)
local int tmpBuffer[GROUP_NUMBER_OF_WORKITEMS];
const int index = get_local_id(0);
//先对pDst[0]进行初始化为1,方便后续求积
if(get_global_id(0) == 0)
atomic_init(pDst, 1);
//由于后续是每个工作组的头一个元素对tmpBuffer元素做求积计算,
//因此本身需要使用栅栏操作
//再加上后续连续四次对全局存储空间的读,因此这里即便不加栅栏操作也问题不大
//work_group_barrier(CLK_GLOBAL_MEM_FENCE);
const int group_id = get_group_id(0);
//获得当前工作组对应的pSrc1的数组起始地址
//这里,我们一个工作组将计算4个工作组大小的数据,
//即每个工作项在垂直方向上访问4个元素
const int addr_index = group_id ＊ GROUP_NUMBER_OF_WORKITEMS ＊ 4
+ index;
int data0 = pSrc1[addr_index + GROUP_NUMBER_OF_WORKITEMS ＊ 0];
int data1 = pSrc1[addr_index + GROUP_NUMBER_OF_WORKITEMS ＊ 1];
int data2 = pSrc1[addr_index + GROUP_NUMBER_OF_WORKITEMS ＊ 2];
int data3 = pSrc1[addr_index + GROUP_NUMBER_OF_WORKITEMS ＊ 3];
data0 ＊= data1;
data2 ＊= data3;
data0 ＊= data2;
//先将每个工作组计算好的数据放进局部缓存
tmpBuffer[index] = data0;
//这里必须执行一次栅栏操作
work_group_barrier(CLK_LOCAL_MEM_FENCE);
if(index == 0)
{
int product = 1;
//对当前工作组中对应的所有向量元素求积
for(int i = 0; i 〈 GROUP_NUMBER_OF_WORKITEMS; i++)
product ＊= tmpBuffer[i];
//使用原子操作将结果进行求积到输出存储地址
//由于OpenCL 2.0没有原生的原子乘法,
//因此这里将借助原子比较与交换的操作
//先用原子加载获取当前pDst[0]的值,并使用acquire次序
int old_value = atomic_load_explicit(pDst,
memory_order_acquire,
memory_scope_device);
//求乘积
int new_value = old_value ＊ product;
//做原子比较与交换,并且设定存储器次序为release
bool ret = atomic_compare_exchange_strong_explicit(pDst,
&old_value, new_value,
memory_order_release,
memory_order_relaxed,
memory_scope_device);
while(!ret)
{
//由于做过一次原子比较与交换后,如果失败,
//pDst[0]的新的值会自动交给old_value
//所以,我们这里可以直接计算新的乘积
new_value = old_value ＊ product;
ret = atomic_compare_exchange_strong_explicit(pDst,
&old_value, new_value,
memory_order_acq_rel,
memory_order_relaxed,
memory_scope_device);
}
}
}
这里要注意的是，先有存储器次序获得操作时，后面一般会加一个存储器次序释放操作，当然前面已经提到过，存储器获得操作与释放操作不需要一定匹配在一起。如果单次操作想做存储器次序的可见操作，那么可以直接使用acq_rel存储器次序。
原子算术逻辑操作
OpenCL 2.0的原子算术逻辑操作与OpenCL 1.2的比起来也要丰富灵活很多，而且声明方式也更为统一。当然，两者的语义相差不多。其声明方式为：
C atomic_fetch_〈key>(volatile A ＊object, M operand)
C atomic_fetch_〈key>_explicit(volatile A ＊object, M operand,
memory_order order)
C atomic_fetch_〈key>_explicit(volatile A ＊object, M operand,
memory_order order, memory_scope
scope)
其中〈key〉表示各种具体的算术逻辑操作。可用的有：
- add：表示原子加法操作，函数名即为atomic_fetch_add。
- sub：表示原子减法操作，函数名即为atomic_fetch_sub。
- or：表示原子按位或操作，函数名即为atomic_fetch_or。
- xor：表示原子按位异或操作，函数名即为atomic_fetch_xor。
- and：表示原子按位与操作，函数名即为atomic_fetch_and。
- min：表示原子取最小值操作，函数名即为atomic_fetch_min。
- max：表示原子取最大值操作，函数名即为atomic_fetch_max。
这些原子操作都是将operand的值与object所指向的原子对象进行指定操作，然后将结果存入object所指的原子对象中，并且返回object所指的原子对象在修改之前的值。
原子标志类型与操作
原子标志类型提供了经典的测试与置一功能。这个功能常见于很多DSP中，用作为对互斥体（mutex）的操作。这里引入了OpenCL 2.0中新的基本类型——atomic_f lag。用这个类型所声明的变量也属于原子对象，并且与其他原子类型变量的初始化差不多，如果定义在全局，需要使用ATOMIC_FLAG_INIT；如果定义在函数体中作为局部变量，那么可以使用后面要介绍的atomic_f lag_clear函数。ATOMIC_FLAG_INIT的用法如下：
global atomic_flag mutex = ATOMIC_FLAG_INIT;
对于atomic_f lag变量有两种状态：一种是置一状态（set状态），可以被看作为相应于布尔值true；另一种就是清零状态（clear状态），可以看做为相应于布尔值false。在使用此类型的变量时，它一开始应该处于清零状态，所以要先对它进行初始化。下面介绍原子标志的测试与置一操作的函数原型：
bool atomic_flag_test_and_set(volatile atomic_flag ＊object)
bool atomic_flag_test_and_set_explicit(volatile atomic_flag
＊object,
memory_order order)
bool atomic_flag_test_and_set_explicit(volatile atomic_flag
＊object,
memory_order order,
memory_scope scope)
这个函数的操作逻辑如下：对object所指向的原子标志类型的对象置一，无论它之前是什么值，并且存储器次序order也会作用于object所指向的原子标志类型对象。这个函数返回object被置一之前的状态值。也就是说，如果object所指向的原子对象在此操作之前为清零状态，则返回false；否则返回true。所以，这个函数就可以被当作对互斥体的上锁操作。如果返回false，则意味着之前没有其他线程（工作项）对它进行上锁，那么上锁成功；否则，说明该互斥体已经处于置一状态，即已经被锁了，那么当前线程（工作项）必须等到它状态切换为清零状态之后才能访问后续的临界资源。
那么，有了上锁一定就有解锁操作。原子标志类型操作提供了清零操作，其函数原型如下：
void atomic_flag_clear(volatile atomic_flag ＊object)
void atomic_flag_clear_explicit(volatile atomic_flag ＊object,
memory_order order)
void atomic_flag_clear_explicit(volatile atomic_flag ＊object,
memory_order order,
memory_scope scope)
这个函数的功能很简单，就是将object所指向的原子标志类型对象清零。我们也可以使用这个函数对声明在函数内的局部原子标志变量进行初始化。这个函数就可以被当作对互斥体的解锁。这里，order参数只能是memory_order_release、memory_order_seq_cst或memory_order_relaxed。
下面我们还是利用上面的对数组所有元素的乘法求积的示例，把原本使用原子比较与交换的操作改为使用原子标志设置与清零操作。由于主机端代码无须修改，这里直接给出OpenCL内核程序代码：
global atomic_flag mutex = ATOMIC_FLAG_INIT;
__kernel void kernel_test(__global atomic_int ＊pDst,
__global int ＊pSrc1, __global int ＊pSrc2)
{
local int tmpBuffer[GROUP_NUMBER_OF_WORKITEMS];
const int index = get_local_id(0);
//先对pDst[0]进行初始化为1,方便后续求积
if(get_global_id(0) == 0)
pDst[0] = 1;
//由于后续是每个工作组的头一个元素对tmpBuffer元素做求积计算,
//因此本身需要使用栅栏操作。
//再加上后续连续四次对全局存储空间的读,因此这里即便不加栅栏操作也问题不大
//work_group_barrier(CLK_GLOBAL_MEM_FENCE);
const int group_id = get_group_id(0);
//获得当前工作组对应的pSrc1的数组起始地址
//这里,我们一个工作组将计算4个工作组大小的数据,即每个工作项在垂直方向上访问4个元素
const int addr_index = group_id ＊ GROUP_NUMBER_OF_WORKITEMS ＊ 4
+ index;
int data0 = pSrc1[addr_index + GROUP_NUMBER_OF_WORKITEMS ＊ 0];
int data1 = pSrc1[addr_index + GROUP_NUMBER_OF_WORKITEMS ＊ 1];
int data2 = pSrc1[addr_index + GROUP_NUMBER_OF_WORKITEMS ＊ 2];
int data3 = pSrc1[addr_index + GROUP_NUMBER_OF_WORKITEMS ＊ 3];
data0 ＊= data1;
data2 ＊= data3;
data0 ＊= data2;
//先将每个工作组计算好的数据放进局部缓存
tmpBuffer[index] = data0;
//这里必须执行一次栅栏操作
work_group_barrier(CLK_LOCAL_MEM_FENCE);
if(index == 0)
{
int product = 1;
//对当前工作组中对应的所有向量元素求和
for(int i = 0; i 〈 GROUP_NUMBER_OF_WORKITEMS; i++)
product ＊= tmpBuffer[i];
//使用原子操作将结果进行求积到输出存储地址
//这里借助原子标志测试与置一操作来对输出结果做乘积操作
bool isLocked;
do
{
//这里使用原子标签测试与置一对互斥体mutex进行上锁操作
//并且使用acquire存储器次序
isLocked = atomic_flag_test_and_set_explicit(&mutex,
memory_order_acquire,
memory_scope_device);
}
while(isLocked);
//先用原子加载获取当前pDst[0]的值,并使用acquire次序
int old_value = atomic_load_explicit(pDst,
memory_order_acquire,
memory_scope_device);
//求乘积
int new_value = old_value ＊ product;
//做原子存储操作,并且设定存储器次序为release
atomic_store_explicit(pDst, new_value, memory_order_release,
memory_scope_device);
//操作完之后别忘了对互斥体解锁,这里使用release存储器次序
atomic_flag_clear_explicit(&mutex, memory_order_release,
memory_scope_device);
}
}