OpenCL 全局屏障工作项同步

Posted 2023-03-17

【中文标题】OpenCL 全局屏障工作项同步

【英文标题】：OpenCL global barrier work items synchronization

【发布时间】：2020-01-04 02:19:00

【问题描述】：

我正在测试一个带有四个工作项和一个工作组的 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] 行之前停止所有工作项吗？

【参考方案1】：

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

如果是nvidia gpu，可以使用inline ptx查询时钟周期，打印出来就知道是什么时间了。

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

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