OpenCl 内核的屏障如何工作？

Question

内核代码：

#pragma OPENCL EXTENSION cl_khr_fp64: enable
#pragma OPENCL EXTENSION cl_amd_printf : enable

__kernel void calculate (__global double* in)
{
    int idx = get_global_id(0); // statement 1
    printf("started for %d workitem\n", idx); // statement 2
    in[idx] = idx + 100; // statement 3
    printf("value changed to %lf in %d workitem\n", in[idx], idx); // statement 4
    barrier(CLK_GLOBAL_MEM_FENCE); // statement 5
    printf("completed for %d workitem\n", idx); // statement 6
}

我正在使用 clEnqueueNDRangeKernel 调用内核，方法是传递一个具有 5 个元素且值初始化为 0.0 的双精度数据类型数组的参数

我使用 5 global_work_size 调用内核，因此我将在每个工作项上求解数组的每个元素。

但根据我对障碍的理论理解，为了同步工作组中的工作项，OpenCL 提供了类似的功能 具有屏障功能。这会强制一个工作项等到所有其他工作项 在群里到达壁垒。通过创建障碍，您可以确保每个工作项都达到相同的 点在其处理。当工作项目需要完成时，这是一个至关重要的问题 计算将在未来计算中使用的中间结果。

因此，我期待这样的输出：

started for 0 workitem
started for 1 workitem
value changed to 100.000000 in 0 workitem
value changed to 101.000000 in 1 workitem
started for 3 workitem
value changed to 103.000000 in 3 workitem
started for 2 workitem
value changed to 102.000000 in 2 workitem
started for 4 workitem
value changed to 104.000000 in 4 workitem

completed for 3 workitem
completed for 0 workitem
completed for 1 workitem
completed for 2 workitem
completed for 4 workitem

这些已完成的陈述，将在最后一起出现，因为障碍会限制其他工作项目直到达到那个点。

但是，我得到的结果，

started for 0 workitem
value changed to 100.000000 in 0 workitem
completed for 0 workitem
started for 4 workitem
value changed to 104.000000 in 4 workitem
completed for 4 workitem
started for 1 workitem
started for 2 workitem
started for 3 workitem
value changed to 101.000000 in 1 workitem
value changed to 103.000000 in 3 workitem
completed for 3 workitem
value changed to 102.000000 in 2 workitem
completed for 2 workitem
completed for 1 workitem

我在逻辑上遗漏了什么吗？那么，OpenCl Kernel 的屏障是如何工作的呢？

在内核中添加了更多检查，以便在 Barrier 之后交叉检查更新的值，而不是打印语句。

#pragma OPENCL EXTENSION cl_khr_fp64: enable
#pragma OPENCL EXTENSION cl_amd_printf : enable

__kernel void calculate (__global double* in)
{
    int idx = get_global_id(0);
    in[idx] = idx + 100;
    barrier(CLK_GLOBAL_MEM_FENCE);
    if (idx == 0) {
        in[0] = in[4];
        in[1] = in[3];
        in[2] = in[2];
        in[3] = in[1];
        in[4] = in[0];
    }
}

那么在数组之后应该是

after arr[0] = 104.000000
after arr[1] = 103.000000
after arr[2] = 102.000000
after arr[3] = 101.000000
after arr[4] = 100.000000

但是结果，我得到了：

after arr[0] = 0.000000
after arr[1] = 101.000000
after arr[2] = 102.000000
after arr[3] = 103.000000
after arr[4] = 104.000000

Answer 1

代码看起来非常好，我怀疑本地工作组的大小，如果你没有指定本地工作组大小，OpenCL 编译器会根据一些检查选择最好的（通常是 一个）。

检查您的 clEnqueueNDRangeKernel 调用 w.r.t 下面的调用

size_t global_item_size = 5; //Specifies no. of total work items
size_t local_item_size = 5; // Specifies no. of work items per local group
clEnqueueNDRangeKernel( command_queue, kernel, 1, NULL, &global_item_size,    &local_item_size, 0, NULL, NULL );

注意： 这个答案是假设您没有指定本地工作组的大小，或者没有按照您的要求正确设置。

关于工作组的更多信息：：

Barrier 将阻塞工作组中的所有线程，因为您没有指定工作组大小（其大小被视为一个），您将拥有 5 个工作组，每个工作组只有一个线程。

Answer 2

是的，您错过了添加printf() 会使所有结果订单无效的事实。

实际上，OpenCL 声明 printf() 的使用是实现定义的，In the case that printf is executed from multiple work-items concurrently, there is no guarantee of ordering with respect to written data. 简单的逻辑会告诉您队列将按顺序为每个 WI 刷新，因为这是序列化刷新的更简单方法在并行执行填充了许多缓冲区（每个 WI printf 一个）之后。

它们按照您期望的顺序执行，但是标准输出的输出刷新发生在内核已经完成之后，并且不遵循原始顺序。