OpenCL 工作组内同步|极客笔记

OpenCL提供了一个工作组内所有工作项同步的接口。这对于目前的GPGPU而言往往就对应于一条membar指令。而这个功能就是基于当前大部分计算加速设备能在硬件层面上支持工作组内对所有工作项同步的特性。

计算加速设备对于支持工作组内同步机制是十分有必要的。因为计算加速设备的一个工作组内的每一个工作项是被并发发射来执行指令的。由于每个工作项所访问的存储空间可能不同，所做的计算方式也有可能不同，因此先被发射执行的工作项未必就会比最后发射执行的工作项先完成执行。所以，工作组中各个工作项对同一条指令的执行未必是按次序完成的，程序员可以认为完成次序是随机的。这种并发执行特性使得计算加速设备能更大地发挥大规模并行机制，对于若干执行得慢的工作项可以先不管它们，调度其他工作项去执行，从而可以以大规模并行计算来掩盖访存操作或对某些复杂计算，如计算超越函数（transcendental function）的延迟。

然而，在不少情况下，我们需要当前工作组内所有工作项都执行到某个点，然后才能继续执行下去。这个点被称为同步点。同步点的产生往往是因为在这个同步点之前，该工作组中的工作项需要一起配合来完成对某一个结果的计算。由于每个工作项的完成次序是不确定的，而唯一我们能确定的是，只有这些工作项全都完成同步点之前的所有指令之后，计算结果方能确定。因此，我们设立这个同步点来对该工作组内的所有工作项进行同步。OpenCL 2.0起，这个内建函数已经改名为work_group_barrier，而之前的版本，函数名为barrier。栅栏先被放在同步点上，对于当前工作组，当任一工作项运行到该同步点时均会被拦截住，不能往下执行，而只有当所有工作项都到达了该同步点，该工作组中的所有工作项才能通过栅栏继续往下执行。这里要注意的是，work_group_barrier无法对工作组中指定的某些工作项起作用，只能同时对该工作组中所有的工作项都起作用。另外，虽然OpenCL 2.0使用了work_group_barrier这一新的函数名，而为了兼容以前的版本，barrier依然可以使用。
work_group_barrier的函数原型定义如下：

void work_group_barrier(cl_mem_fence_flags flags,
                            memory_scope scope)

这里，f lags是用于指示在工作组内对工作项进行同步的同时，再对哪个存储空间做访存次序的同步。如果不想对存储空间做任何存储次序上的同步可以传0。这里再简单地介绍一下存储器访问次序。正如之前对工作组内所有工作项同步的设计原理一样，为了充分利用计算加速设备大规模并行计算，访存设计使用了弱次序访存机制。所谓的弱次序访存机制在CPU上主要是指，对于当前线程对某个非可Cache的存储空间用弱次序特性来写一组数据之后，另一个线程对这个存储空间用弱次序特性的读，那么数据可能是第一个线程写操作之前的。也就是说，使用弱次序的访存操作不同线程之间是不可见的，而对于同一个线程而言是可见的。即如果一个线程用弱次序特性对一个存储空间写一组数据，然后该线程再从这个存储空间用弱次序特性的读操作，那么读到的将一定是上次所写完的数据。因此，弱次序操作往往针对不同线程，因而对应到计算设备上就是不同的工作项。所以，对于同一个工作项对同一个存储空间先做写操作，然后再做读操作，读到的数据肯定是之前被写更新的；而如果是一个工作项对一个存储空间用了写操作，而另一个工作项再去对它读，那么后一个工作项未必就能读到前一个工作项写更新的数据，此时需要用存储次序同步操作来使得某一个工作项的写对其他所有工作项都可见。

这里，f lags是一个枚举类型变量，它具体可以取的枚举值以及含义如下定义，并且可由以下枚举值通过“按位或”（|）组合在一起使用。

CLK_LOCAL_MEM_FENCE：使用此枚举值使得访存栅栏操作将确保之前对局部存储器的访问将对此工作组内的所有工作项可见。
CLK_GLOBAL_MEM_FENCE：使用此枚举值使得访存栅栏操作将确保之前对全局存储器的访问将对此工作组内的所有工作项可见。
CLK_IMAGE_MEM_FENCE：使用此枚举值使得访存栅栏操作将确保之前对图像存储器的访问将对此工作组内的所有工作项可见。该枚举值从OpenCL 2.0版本起可用。
从OpenCL 2.0起，work_group_barrier的函数还新增了scope这个参数，而之前版本的barrier是没有的。该参数指示了工作组内的所有工作项对由参数f lags所指定的存储空间的访问可作用于哪个范围。
memory_scope_work_group：表示访存可对当前工作组内的所有工作项可见。如果f lags使用了CLK_LOCAL_MEM_FENCE，那么scope只能用这个值。
memory_scope_device：表示访存可对当前设备的所有工作项可见。当f lags是CLK_GLOBAL_MEM_FENCE或CLK_IMAGE_MEM_FENCE时，scope可用这个值。
memory_scope_all_svm_devices：表示访存可对所有共享虚拟存储器的计算设备的工作项可见，也对主机端可见。只有当f lags为CLK_GLOBAL_MEM_FENCE时， scope才能用此值。

上面讲述了通过栅栏操作对工作组内工作项的同步，同时又提到了存储器次序的同步。下面我们将通过计算两个向量内积（也称为点积）的一个小示例来展示栅栏操作的基本用法。请大家在上述完整的代码中从program = clCreateProgramWithSource（context, 1,（const char**）&kernelSource,（const size_t*）&kernelLength, &ret）；这条语句开始一直到FINSH标签的所有代码替换为以下代码：

program = clCreateProgramWithSource(context, 1,
                                          (const char ＊＊)&kernelSource,
                                          (const size_t ＊)&kernelLength,
                                          &ret);
const int nWorkItems = 128;
/＊ 在编译选项中,定义一个名为TOTAL_NUMBER_OF_WORKITEMS的宏,
用于指定一共有多少工作项 ＊/
sprintf(kernelSource, "-D TOTAL_NUMBER_OF_WORKITEMS=%d",
                nWorkItems);
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_dot", &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[]){nWorkItems},
                                          (const size_t[]){nWorkItems}, 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, 4,
                                    pDeviceBuffer, 0, NULL, NULL);
int sum = 0;
for(int i = 0; i 〈 nWorkItems; i++)
        sum += pHostBuffer[i] ＊ pHostBuffer[i];
if(sum == ＊pDeviceBuffer)
        puts("Result is OK!");
else
        puts("Result not equal!");

这里，我们通过一个常量nWorkItems来控制计算内积的向量长度。当然，在这个例子中，向量长度不能超过一个工作组最大工作项的数量，否则得到的结果将会有问题。我们这里默认指定为128，这对大部分计算设备而言都能满足要求。而在计算内核端，我们在构造内核程序的时候使用了TOTAL_NUMBER_OF_WORKITEMS这个宏定义，使得我们只要在主机端代码修改了nWorkItems这个常量时，OpenCL内核程序也能获得相应修改。下面提供了OpenCL内核程序代码：

__kernel void kernel_dot(__global int ＊pDst, __global int ＊pSrc1,
__global int ＊pSrc2)
{
    __local int localBuffer[TOTAL_NUMBER_OF_WORKITEMS];
    int index = get_global_id(0);
    localBuffer[index] = pSrc1[index] ＊ pSrc2[index];
    barrier(CLK_LOCAL_MEM_FENCE);
    if(index == 0)
    {
         int sum = 0;
         for(int i = 0; i 〈 TOTAL_NUMBER_OF_WORKITEMS; i++)
             sum += localBuffer[i];
         pDst[0] = sum;
}
}

以上代码非常简单。首先，定义了我们要准备计算的向量长度大小的局部存储空间。然后将两个向量对应的元素进行相乘，把结果放入到对应的局部存储器中。然后，由第一个工作项做整合，把所有乘好的结果进行求和。这里，因为对于第一个工作项而言，其他工作项是否执行完乘法操作并把数据写入到自己对应的局部存储位置，是不可见的，所以需要使用栅栏操作做一次同步。只有当所有工作项都执行完乘法操作，并把结果写入对应的局部存储器才能做后续的求和操作。同时，这里对f lags变量设置为CLK_LOCAL_MEM_FENCE，要求局部存储器访问次序对后续的访存操作可见。