极客笔记之前我们提到，对于某些CPU而言（如x86架构的CPU），存储器次序一般都会设计得比较严格。而对于很多GPGPU（如我们下面即将介绍的AMD Radeon R9 285，基于GCN 1.1架构），存储器次序设计得比较弱。两者实现可能都不会严格按照OpenCL 2.0标准中的存储器次序属性去做。例如，对于x86 CPU而言，memory_order_relaxed的存储器次序可能仍然会被设计成memory_order_seq_cst那么严格；而反观R9 285 GPU，即便使用memory_order_seq_cst，存储器次序可能仍然会跟memory_order_relaxed一样弱。而其实，这些计算设备对于存储器一致性的设计都是比较统一的，也就是说它们基本上无视memory_order属性以及memory_scope属性，而是直接根据当前操作的存储器类型（全局存储器还是本地存储器）来判定所使用的存储器次序以及存储器区域。因此，如果我们想用当前所写的OpenCL内核程序适配多种不同的计算设备，那么我们这里推荐各位读者直接使用类似于atomic_load、atomic_store等内建函数，而不是用它们的explicit版本（atomic_load_explicit、atomic_store_explicit），以忽略相关的存储器次序以及存储器区域的设定。
下面，为了让各位读者更清晰地了解这个事实，我们将基于AMD Radeon R9 285 GPGPU来展现这个实现特性。同时，我们也将会看到R9 285在存储器可见性上的特征。以下示例代码基于Windows 8.1操作系统，Visual Studio 2013开发环境，C语言。主机端的源文件为main.c（支持部分C99特性的C语言）。
以下是主机端代码，很简单：
#ifdef _APPLE_
#include 〈OpenCL/opencl.h>
#else
#include 〈CL/cl.h>
#endif
#include 〈stdio.h>
#include 〈stdlib.h>
#include 〈string.h>
int main(void)
{
cl_uint numPlatforms = 0;
cl_platform_id platform = NULL;
cl_context context = NULL;
cl_command_queue commandQueue = NULL;
cl_program program = NULL;
cl_mem outputMemObj = NULL;
cl_kernel kernel = NULL;
cl_int status = clGetPlatformIDs(0, NULL, &numPlatforms);
if (status != CL_SUCCESS)
{
puts("Error: Getting platforms!");
return 0;
}
/＊ 选择OpenCL平台 ＊/
if (numPlatforms > 0)
{
cl_platform_id ＊platforms = (cl_platform_id ＊)
malloc(numPlatforms ＊
sizeof(cl_platform_id));
status = clGetPlatformIDs(numPlatforms, platforms, NULL);
//在笔者的PC上platform[0]是Intel平台,
//这里使用platform[1]为AMD平台
platform = platforms[1];
free(platforms);
}
else
{
puts("Your system does not have any OpenCL platform!");
return 0;
}
/＊ 选择计算设备 ＊/
cl_uint numDevices = 0;
cl_device_id ＊devices;
status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, NULL,
&numDevices);
if (numDevices == 0) //no GPU available.
{
puts("No GPU device available.");
puts("Choose CPU as default device.");
status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_CPU, 0,
NULL, &numDevices);
devices = (cl_device_id ＊)malloc(numDevices ＊
sizeof(cl_device_id));
status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_CPU,
numDevices,
devices, NULL);
}
else
{
devices = (cl_device_id ＊)malloc(numDevices ＊
sizeof(cl_device_id));
status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU,
numDevices,
devices, NULL);
printf("The number of devices: %u\n", numDevices);
}
/＊ 创建上下文 ＊/
context = clCreateContext(NULL, 1, devices, NULL, NULL, NULL);
/＊ 创建命令队列 ＊/
commandQueue = clCreateCommandQueueWithProperties(context,
devices[0],
0, NULL);
/＊ 读内核代码文件 ＊/
//Read the kernel code to the buffer
FILE ＊fp = fopen("cl_kernel.cl", "rb");
if (fp == NULL)
{
puts("The kernel file not found!");
goto RELEASE_RESOURCES;
}
fseek(fp, 0, SEEK_END);
size_t kernelLength = ftell(fp);
fseek(fp, 0, SEEK_SET);
char ＊kernelCodeBuffer = (char ＊)malloc(kernelLength + 1);
fread(kernelCodeBuffer, 1, kernelLength, fp);
kernelCodeBuffer[kernelLength] = '\0';
fclose(fp);
const char ＊aSource = kernelCodeBuffer;
program = clCreateProgramWithSource(context, 1, &aSource,
&kernelLength, NULL);
/＊ 构建内核程序。注意这里使用了Open CL2.0编译选项,特指OpenCL内核代码使用
OpenCL 2.0 C语言 ＊/
status = clBuildProgram(program, 1, devices, "-cl-std=CL2.0",
NULL,NULL);
if (status != CL_SUCCESS)
{
size_t len;
char buffer[8 ＊ 1024];
printf("Error: Failed to build program executable!\n");
clGetProgramBuildInfo(program, devices[0],
CL_PROGRAM_BUILD_LOG,
sizeof(buffer), buffer, &len);
printf("%s\n", buffer);
goto RELEASE_RESOURCES;
}
int outputBuffer[2048];
memset(outputBuffer, 0, sizeof(outputBuffer));
/＊ 创建输出缓存对象 ＊/
outputMemObj = clCreateBuffer(context, CL_MEM_WRITE_ONLY,
sizeof(outputBuffer), NULL, NULL);
/＊ 创建内核程序 ＊/
kernel = clCreateKernel(program, "memory_order_test", NULL);
/＊ 设置内核程序的参数 ＊/
status = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void ＊)&outputMemObj);
/＊ 运行内核程序 ＊/
size_t maxWorkGroupSize = 0;
clGetDeviceInfo(devices[0], CL_DEVICE_MAX_WORK_GROUP_SIZE,
sizeof(maxWorkGroupSize), &maxWorkGroupSize,
NULL);
size_t global_work_size[1] = { 2048 };
status = clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL,
global_work_size,
(const size_t[]) {
maxWorkGroupSize
}, 0, NULL, NULL);
clFinish(commandQueue);
/＊ 读取缓存数据 ＊/
status = clEnqueueReadBuffer(commandQueue, outputMemObj,
CL_TRUE, 0,
sizeof(outputBuffer),
outputBuffer, 0, NULL, NULL);
printf("Local memory result: %d, global memory result: %d\n",
outputBuffer[0], outputBuffer[1]);
printf("Local wait counter is: %d\n", outputBuffer[2]);
RELEASE_RESOURCES:
/＊ 释放OpenCL资源 ＊/
status = clReleaseKernel(kernel);
status = clReleaseProgram(program);
status = clReleaseMemObject(outputMemObj);
status = clReleaseCommandQueue(commandQueue);
status = clReleaseContext(context);
free(devices);
puts("Program completed!");
getchar();
}
这里需要指出的是，如果我们要使用OpenCL 2.0 C编程语言来写OpenCL内核程序，必须在构建OpenCL内核程序时使用-cl-std=CL2.0编译选项。否则，默认可能会是OpenCL 1.2版本的。
我们这里就使用了一个输出缓存作为内核程序的参数，然后以此输出三个数据。第一个数据表示在另一个工作项中本地存储器数据是否可见；第二个数据表示在另一个工作项中全局存储器数据是否可见；第三个数据表示在另一个工作项中等待了多长时间才观察到了本地存储器的变化。
我们接下来看一下OpenCL内核程序代码：
global volatile atomic_int atomA = ATOMIC_VAR_INIT(0);
__kernel void memory_order__test(_global int ＊dst)
{
local volatile atomic_int flag;
int index = get_global_id(0);
atomic_init(&flag, 0);
work_group_barrier(CLK_LOCAL_MEM_FENCE);
if(index == 0)
{
atomic_store_explicit(&flag, 1, memory_order_seq_cst,
memory_scope_work_group);
atomic_store_explicit(&atomA, 1, memory_order_seq_cst,
memory_scope_device);
}
//atomic_work_item_fence(CLK_LOCAL_MEM_FENCE,
//memory_order_acq_rel, memory_scope_work_group);
if(index == 64)
{
int a, counter = 10000;
do
{
a = atomic_load_explicit(&flag, memory_order_seq_cst,
memory_scope_work_group);
counter--;
}
while(a == 0 && counter > 0);
dst[0] = a != 0 ? 1 : 0;
dst[2] = 10000 - counter;
counter = 10000;
do
{
a = atomic_load_explicit(&atomA, memory_order_seq_cst,
memory_scope_device);
counter--;
}
while(a == 0 && counter > 0);
dst[1] = a != 0 ? 1 : 0;
}
work_group_barrier(0);
}
以上代码也十分简单。我们这个内核程序就是要观察全局存储器中的原子对象与本地存储器中的原子对象对其他工作项的可见性。各位读者可以注意到，在此内核代码中，我们对所有原子对象的读写都使用了最最严格的memory_order_seq_cst存储器次序。通过这个程序也能看到，对于一个全局存储器上的原子对象，可以使用ATOMIC_VAR_INIT宏对其进行初始化。而在一个内核函数内的本地存储器原子对象只能通过atomic_init内建函数进行初始化。而且由于此内建函数对存储器次序没有任何一致性约束，因此我们必须显式使用work_group_barrier（CLK_LOCAL_MEM_FENCE）；来确保此原子对象的初始化对当前工作组内的所有工作项可见。当然，这里也可以使用atomic_work_item_fence来代替，效果上与work_group_barrier一样。
对于工作项0，我们分别对本地原子对象f lag与全局原子对象atomA，通过atomic_store_explicit来写1。对于工作项64，我们尝试使用atomic_load_explicit来读取本地原子对象f lag以及全局原子对象atomA。我们这里设置了超时循环。如果循环了10000次还没有观察到原子对象被修改，那么我们就断言在工作项64中无法观察到工作项0对该原子对象的修改操作。
最后一句work_group_barrier的同步必须添加，否则工作项0所做的副作用无法暴露给当前工作组中的其他工作项。我们通过执行这段程序后会发现，在工作项0中对本地存储器的原子对象修改以及对全局存储器原子对象的修改均能被工作项64观察到，并且循环4次即可被观察到。如果我们把最后一句work_group_barrier（0）给删除掉，那么我们会发现只有对全局存储器对象的修改可被观察到。这其实就跟我们之前所提到过的关于本地存储器的性质相关。由于本地存储器能提供极大的访存带宽，访问速度比全局存储器快得多，因此对于它的存储器次序，包括可见性的约束就远比全局存储器松弛。所以一旦我们最后不用栅栏操作把工作项0之前操作所产生的所有副作用暴露给其他工作项，那么对于其他工作项而言，工作项0对本地存储器的修改操作是无法被观察到的。而且，这里即便用atomic_work_item_fence也不行，只能通过work_group_barrier操作。因为atomic_work_item_fence操作只能对其后续的访存操作可见。因此从这里我们也能看到，memory_order_seq_cst对于R9 285的本地存储器次序而言不起作用。否则，我们即便不依赖于work_group_barrier操作也能使得工作项0中的本地存储器修改操作能在工作项64中可见。
下面，我们把if（index == 64）改成if（index == 32）再看看结果。我们此时会看到，同样的程序，仅仅就是把64改为32，工作项0对本地原子对象与全局原子对象的修改在工作项32中均无法被观察到。这个与AMD GCN GPGPU的架构设计有关。对于AMD Radeon HD Graphics的架构详细介绍请见后面的教程。对于R9 285，一个计算单元CU是以一个wavefront为单位（一个wavefront由id连续的64个工作项构成）进行执行调度的。而在一条wavefront内，CU是以一种固定模式将一条指令发射到64个工作项中执行。当遇到分支时，一条wavefront中的所有工作项都会执行该分支中的内容，满足分支条件的工作项将操作所获得的副作用做相应更新，而不满足分支条件的工作项将所有操作获得的副作用丢弃。因此在这种情况下，工作项32无法观察到处于同一条wavefront中的工作项0所产生的任何副作
用。除非，我们将被注释掉的一条atomic_work_item_fence调用语句启用。有了这一条存储器栅栏操作之后，那么工作项0所产生的任何访存操作都能被该工作组中所有工作项（无论其他工作项与工作项0是否处于同一个条wavefront）观察到。
由此我们可以看到，R9 285的存储器次序实现并非遵守以上存储器次序规则，而是使用自己固有的存储器访问特性。即便我们这里对所有访存操作都使用了最严格的memory_order_seq_cst也毫无作用。因此就目前而言，我们尚未到依赖存储器次序进行OpenCL内核程序设计开发的阶段，所以我们在此仍然建议各位读者使用默认的原子操作，即不含有explicit后缀的，使用默认的存储器次序与存储区域。我们将以上内核程序改写成以下形式，其表现也跟上面的一模一样：
global volatile atomic_int atomA = ATOMIC_VAR_INIT(0);
__kernel void memory_order_test(__global int ＊dst)
{
local volatile atomic_int flag;
int index = get_global_id(0);
atomic_init(&flag, 0);
work_group_barrier(CLK_LOCAL_MEM_FENCE);
if(index == 0)
{
atomic_store(&flag, 1);
atomic_store(&atomA, 1);
}
//atomic_work_item_fence(CLK_LOCAL_MEM_FENCE,
//memory_order_seq_cst, memory_scope_work_group);
if(index == 64)
{
int a, counter = 10000;
do
{
a = atomic_load(&flag);
counter--;
}
while(a == 0 && counter > 0);
dst[0] = a != 0 ? 1 : 0;
dst[2] = 10000 - counter;
counter = 10000;
do
{
a = atomic_load(&atomA);
counter--;
}
while(a == 0 && counter > 0);
dst[1] = a != 0 ? 1 : 0;
}
work_group_barrier(0);
}