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:%02dn", 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.6fn",gid_x,gid_y,l2_norm);
barrier(CLK_GLOBAL_MEM_FENCE);
l2_norm+=1;
barrier(CLK_GLOBAL_MEM_FENCE);
printf("%d,%d after: %2.6fn",gid_x,gid_y,l2_norm);
printf("testing %d,%dn",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，您可以使用内联 ptx 查询时钟周期以打印它们并知道它发生的时间。

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

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