OpenCL 1.2：mem_fence() 或 barrier() 或两者兼有

Posted 2023-03-17

【中文标题】OpenCL 1.2：mem_fence() 或 barrier() 或两者兼有

【英文标题】：OpenCL 1.2: mem_fence() or barrier() or both

【发布时间】：2017-11-10 01:08:35

【问题描述】：

我刚刚开始 openCL C 编程。工作组的所有工作项都会更新本地内存的唯一位置。稍后，一个工作项的私有变量会根据其他两个工作项更新的本地数据进行更新。像这样的：

__kernel MyKernel(__global int *in_ptr)
          
           /* Define a variable in private address space */
           int priv_data;
           /* Define two indices in private address space */
           int index1, index2;

           /* index1 and index2 are legitimate local work group indices */  
           index1 = SOME_CORRECT_VALUE;
           index2 = ANOTHER_CORRECT_VALUE;

           /* Define storage in local memory large enough to cater to all work items of this work group */
           __local int tempPtr[WORK_GROUP_SIZE];
           tempPtr[get_local_id(0)] = SOME_RANDOM_VALUE;

           /* Do not proceed until the update of tempPtr by this WI has completed */
           mem_fence(CLK_LOCAL_MEM_FENCE);

           /* Do not proceed until all WI of this WG have updated tempPtr */
           barrier(CLK_LOCAL_MEM_FENCE);

           /* Update private data */
           priv_data = tempPtr[index1] + tempPtr[index2];
       

虽然上面的 sn-p 是保守的，但屏障不会像它在内部做击剑那样完成这项工作吗？

【问题讨论】：

是的，barrier 已经做了围栏。但是，在某些情况下，您可以使用单个围栏。如果您不关心本地工作人员不同步，但您只想完成之前的内存写入/读取。在你的情况下，栅栏就足够了。 （除非该代码在循环中运行并且您没有在示例中添加额外的代码）。

谢谢。您能否将其作为常规答案提交以便被接受？

【参考方案1】：

是的，屏障已经做了围栏。

屏障将在该点同步执行。因此，必须执行所有先前的指令，因此此时内存是一致的。 栅栏只会确保在执行任何进一步的读/写之前完成所有读/写，但工作人员可能正在执行不同的指令。

在某些情况下，您可以使用单个围栏。如果您不关心本地工作人员不同步，并且您只想完成之前的内存写入/读取。在你的情况下，栅栏就足够了。 （除非该代码在循环中运行并且您没有在示例中添加额外的代码）。