omp 全局内存栅栏/屏障答案

【问题标题】：omp global memory fence / barrieromp 全局内存栅栏/屏障
【发布时间】：2022-09-23 23:53:41
【问题描述】：

在 GPU 上进行目标卸载的 OpenMP 是否包括类似于 OpenCL 的全局内存围栏/全局屏障？

barrier(CLK_GLOBAL_MEM_FENCE);

我尝试在团队构造中使用

#pragma omp target teams
{
    // Some initialization...

    #pragma omp distribute parallel for
    for (size_t i = 0; i < N; i += 1)
    {
        // Some work...
    }

    #pragma omp barrier

    #pragma omp distribute parallel for
    for (size_t i = 0; i < N; i += 1)
    {
        // Some other work depending on pervious loop
    }
}

然而，屏障似乎只在团队内部起作用，相当于：

barrier(CLK_LOCAL_MEM_FENCE);

我想避免将内核分成两个，以避免将团队本地数据发送到全局内存以再次加载它。

编辑：我已经能够使用全局原子计数器和团队的忙碌等待来强制执行所需的行为。然而，这似乎不是一个好的解决方案，我仍然想知道是否有更好的方法来使用正确的 OpenMP

标签： c++ gpu openmp

【解决方案1】：

barrier 构造仅同步当前团队中的线程。由teams 构造启动的来自不同线程组的线程之间的同步不可用。 OpenMP 的执行模型不保证这样的线程甚至会并发执行，因此使用atomic 构造在线程之间进行同步通常不会起作用：

初始线程是否并发执行团队地区是未指定，以及依赖于它们的并发执行的程序同步的目的可能会死锁。

请注意，OpenCL barrier 调用仅提供工作组内的同步，即使使用 CLK_GLOBAL_MEM_FENCE 参数也是如此。有关CLK_GLOBAL_MEM_FENCE 与CLK_LOCAL_MEM_FENCE 语义的更多信息，请参阅Barriers in OpenCL。

【讨论】：