c - 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 由 clfnish 刷新，因此它处于内核级同步。这就是为什么您不应该依赖输出文本的顺序，而应该依赖数据本身。

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

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

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