极客笔记介绍完了原子操作之后，本节我们再来介绍一下OpenCL从全局存储区域到局部存储区域以及从局部存储区域到全局存储区域的异步拷贝。这些操作在OpenCL 1.2中就已经被引入了。我们之前提到了OpenCL主机端的API提供了事件等待的函数接口，而对于大部分clEnqueue系API而言都会带有event_wait_list参数，用于指定等待哪些事件完成才能执行，以及event参数用来跟踪当前命令本身。在主机端，事件对象类型为cl_event。而OpenCL内核程序中也含有事件对象类型，被定义为event_t。事件对象可以用来跟踪异步拷贝的完成情况。
下面先介绍异步工作组基本拷贝功能的函数原型：
event_t async_work_group_copy (__local gentype ＊dst,
const __global gentype ＊src, size_t
num_gentypes,
event_t event)
event_t async_work_group_copy (__global gentype ＊dst,
const __local gentype ＊src, size_t
num_gentypes,
event_t event)
上一个函数是将全局存储器的数据拷贝到局部存储器中，下一个则是将局部存储器的数据拷贝到全局存储器中。这里，gentype是一个泛型类型，能支持OpenCL中所有标量及向量形式的基本类型。参数num_gentypes指明了需要拷贝多少个gentype的元素，因此一共需要拷贝的字节数为num_gentypes * sizeof（num_gentypes）。最后一个参数event是用于指定之前进行异步拷贝的事件对象。前后相继的异步拷贝操作可以共享一个事件对象。event参数被指定为前某一个事件对象时，该函数将直接返回该event事件对象；否则这个参数传0即可，函数返回对应该操作的一个事件对象。
这里需要注意的是，当我们调用异步拷贝函数时需要当前工作组的所有工作项参与进行操作，否则结果是未定义的。这意味着如果我们用以下代码来调用异步操作，结果是不确定的：
if(get_local_id(0) == 0)
async_work_group_copy(pDst, tmpBuffer, 64, 0);
然后，我们再介绍一下事件等待的内建函数，其函数原型如下：
void wait_group_events (int num_events, event_t ＊event_list)
这个函数非常简单，第一个参数num_events指定了第二个参数event_list中包含了多少个事件对象。第二个参数event_list就是指向事件对象的数组。这个函数与主机端的事件等待类似，只有当事件列表中的事件全都处于完成状态之后程序才能继续往下执行。当然，这个函数与async_work_group_copy一样，必须要当前工作组中所有工作项参与进来操作。另外，我们要注意的是，OpenCL标准中没有规定这个函数被调用之后，所有工作项都会处于同一个调用点，因此这个函数不能被当作一个barrier来使用。不过，这个函数可以保证当前工作项的拷贝工作能够切实完成。
下面，我们来举一个简单的例子来看一下基本功能的异步拷贝操作是如何具体使用的。我们将上述主机端代码从ret = clEnqueueNDRangeKernel（command_queue, kernel, 1, NULL,……{evt1,evt2},NULL）；这一行一直到FINISH标签，替换为如下代码：
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(contentLength);
//这里使用阻塞的方式读取数据
clEnqueueReadBuffer(command_queue, dstMemObj, CL_TRUE, 0,
contentLength,
pDeviceBuffer, 0, NULL, NULL);
//做数据校验
bool isOK = true;
for(int i = 0; i 〈 contentLength / sizeof(int); i++)
{
}
if(pHostBuffer[i] ＊ 2 != pDeviceBuffer[i])
{
isOK = false;
break;
}
puts(isOK ? "OK" : "NG");
主机端代码比较简单，不再赘述。我们马上看一下对应的OpenCL内核程序代码：
__kernel void kernel_test(__global int ＊pDst,
{
__global int ＊pSrc1, __global int ＊pSrc2)
local int tmpBuffer[GROUP_NUMBER_OF_WORKITEMS];
const size_t address_offset = get_group_id(0)
＊ GROUP_NUMBER_OF_WORKITEMS;
//这里使用了异步拷贝操作,
//将全局存储器的数据拷贝到相应的当前工作组的局部存储器中
//这里,函数原型中的gentype就是int类型
//一共拷贝GROUP_NUMBER_OF_WORKITEMS个元素
event_t event = async_work_group_copy(tmpBuffer,
&pSrc1[address_offset],
GROUP_NUMBER_OF_WORKITEMS, 0);
//这里抽空可以做其他事情,比如确定当前工作组中当前工作项的id等
const int index = get_local_id(0);
//等待拷贝结束
wait_group_events(1, &event);
//我们将对应元素乘以2,再写回局部存储器
tmpBuffer[index] ＊= 2;
//我们确保所有工作项都在调用后续的异步操作之前能够都完成对数据的修改操作
barrier(CLK_LOCAL_MEM_FENCE);
//将当前局部存储器中的数据再拷贝回对应的全局存储器中
event = async_work_group_copy(&pDst[address_offset], tmpBuffer,
GROUP_NUMBER_OF_WORKITEMS, 0);
//等待拷贝结束
wait_group_events(1, &event);
}
这部分代码非常简单。尽管我们可以直接用计算-修改的方式去做，不过作为一个简单的demo而言能作为使大家更容易理解的函数使用方式。这里再强调一下，async_work_group_copy作用于整个工作组，而不是某一工作项。因此在调用这个函数时，我们要清楚与其相关的地址偏移以及拷贝的数据大小需要参考的是工作组的区域范围而不是当前工作项的区域范围。
下面我们再介绍更具灵活性的，带有跨度的异步拷贝操作。下面先给出函数原型：
event_t async_work_group_strided_copy(__local gentype ＊dst,
const __global gentype ＊src,
size_t num_gentypes, size_t
src_stride,
event_t event)
event_t async_work_group_strided_copy(__global gentype ＊dst,
const __local gentype ＊src,
size_t num_gentypes, size_t
dst_stride,
event_t event)
这两个异步拷贝函数原型与上面两个基本差不多，不过这里分别多了一个参数。从全局存储器拷贝到局部存储器的函数中新增了src_stride参数，表示从全局存储器取数据时，一个元素与其后面的一个元素之间跨多少元素。例如，如果gentype是f loat类型，并且src_stride是4，那么第一个元素取的是src[0]，而第二个元素则取的是src[4]，第三个则是src[8]，而写入到局部存储器中则是按次序前后相继写进去的。也就是第一个元素写到dst[0]，第二个元素写到dst[1]，第三个写到dst[2]，对于从局部存储器异步拷贝到全局存储器的函数中，则增加了dst_stride参数，这个参数指明了从局部存储器所取出的元素写入到全局存储器中前一个与后一个相隔多少个元素。这意味着数据从局部存储器中取出时是逐个相继取出的，而写入到全局存储器时是以dst_stride跨度去写的。例如，取出的第一个数据是src[0]，第二个数据是src[1]，第三个数据是src[2]，如果dst_stride是4，那么写入的第一个数据是dst[0]，第二个数据是dst[4]，第三个数据则是dst[8]。
下面我们根据上面异步拷贝操作的例子进行稍许修改来给出带有跨度的异步拷贝操作的使用实例。下面，我们将上述代码中最后的校验部分替换为如下代码：
//做数据校验
bool isOK = true;
for(int i = 0; i 〈 contentLength / sizeof(int); i++)
{
//跨4个int元素进行比较校验
if((i & 3) == 0 && pHostBuffer[i] ＊ 2 != pDeviceBuffer[i])
}
{
isOK = false;
break;
}
puts(isOK ? "OK" : "NG");
由于我们这里写回到全局存储空间已经是间隔了4个int元素，所以我们在进行比较校验时也必须间隔4个int元素进行比较，否则结果肯定会失败。下面给出kernel程序代码：
__kernel void kernel_test(__global int ＊pDst,
{
__global int ＊pSrc1, __global int ＊pSrc2)
local int tmpBuffer[GROUP_NUMBER_OF_WORKITEMS];
const size_t address_offset = get_group_id(0) ＊
GROUP_NUMBER_OF_WORKITEMS;
//这里使用了异步拷贝操作,
//将全局存储器的数据拷贝到相应的当前工作组的局部存储器中
//这里,函数原型中的gentype就是int类型
//一共拷贝GROUP_NUMBER_OF_WORKITEMS / 4个元素
//从全局存储器获取的两个元素之间跨4个元素（即4个int大小）
event_t event = async_work_group_strided_copy(tmpBuffer,
&pSrc1[address_offset],
GROUP_NUMBER_OF_WORKITEMS / 4,
4, 0);
//这里抽空可以做其他事情,比如确定当前工作组中当前工作项的id等
const int index = get_local_id(0);
//等待拷贝结束
wait_group_events(1, &event);
//我们将对应元素乘以2,再写回局部存储器
tmpBuffer[index] ＊= 2;
//我们确保所有工作项都在调用后续的异步操作之前能够都完成对数据的修改操作
barrier(CLK_LOCAL_MEM_FENCE);
//将当前局部存储器中的数据再拷贝回对应的全局存储器中,并且相邻两个元素之间跨4个元素
event = async_work_group_strided_copy(&pDst[address_offset],
tmpBuffer,
GROUP_NUMBER_OF_WORKITEMS / 4,
4, 0);
//等待拷贝结束
wait_group_events(1, &event);
}