OpenCL 全局屏障工作项同步

Question

我正在测试具有四个工作项和一个工作组的 opencl 内核。内核是：

__kernel void pgs(__global float l2_norm)
{
    int gid_x=get_global_id(0);
    int gid_y=get_global_id(1);
    if (gid_x==0 && gid_y==0) printf("[INFO] local_size_x:%02d, local_size_y:%02d, global_size_x:%02d, global_size_y:%02d, group_size_x:%02d, group_size_y:%02d\n", get_local_size(0), get_local_size(1), get_global_size(0), get_global_size(1), get_group_size(0), get_group_size(1));
    barrier(CLK_GLOBAL_MEM_FENCE);

    printf("%d,%d before: %2.6f\n",gid_x,gid_y,l2_norm);
    barrier(CLK_GLOBAL_MEM_FENCE);
    l2_norm+=1;
    barrier(CLK_GLOBAL_MEM_FENCE);
    printf("%d,%d after: %2.6f\n",gid_x,gid_y,l2_norm);

    printf("testing %d,%d\n",gid_x,gid_y);
}

输出为：

1,1 before: 0.000000
0,1 before: 0.000000
1,0 before: 0.000000
[INFO] local_size_x:01, local_size_y:01, global_size_x:02, global_size_y:02, group_size_x:01, group_size_y:01
1,1 after: 1.000000
0,1 after: 2.000000
1,0 after: 3.000000
testing 1,1
0,0 before: 3.000000
testing 0,1
testing 1,0
0,0 after: 4.000000
testing 0,0

我的问题是：为什么没有先打印以[INFO]开头的行？在工作项 0 打印 [INFO] 行之前，全局屏障不应该停止所有工作项吗？

Answer 1

关卡仅用于组内等待。 Printf 由 clfinish 刷新，因此它处于内核级同步。这就是为什么你不应该依赖于输出文本的顺序，而是依赖于数据本身。

如果是nvidia gpu，可以使用inline ptx查询时钟周期打印出来，知道发生的时间。

对于其他供应商，您可以拥有一个全局原子变量并在屏障之间递增它。原子变化不会跨越障碍。这样，第一个线程的增量将在屏障之后的其他线程之前发生。由于这只是数据，您仍然需要在 printf 之前在主机环境中重新排序。但这仅提示不同的同步区域。您仍然无法知道同一同步区域中的顺序。你只能知道在屏障之前或之后发生了一些事情。

也许制作自己的格式化程序更容易。创建一个不会溢出许多文本的长缓冲区。有一个全局原子游标计数器变量。在每个线程中，使用类似于 printf 的格式化程序函数，但让它自动递增其光标并用给定的格式化文本填充尾随区域。然后在主机环境中逐行写入整个字符串或您在其中使用的任何定界符来分隔输入。

OpenCL 全局屏障工作项同步

OpenCL global barrier work items synchronization

c

synchronization

opencl

barrier