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

Question

我刚刚开始 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];
       }

虽然上面的代码片段是保守的，但 barrier 不会像它在内部做 fencing 那样完成这项工作吗？

Answer 1

是，barrier 已经做 fencing 了。

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

在某些情况下，您可以使用单个栅栏。如果你不关心本地worker不同步，你只是想让之前的内存writes/read完成。在您的情况下，栅栏就足够了。（除非该代码是运行在一个循环中，并且您还没有在示例中放入额外的代码）。

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

OpenCL 1.2: mem_fence() or barrier() or both

memory-fences

opencl

memory-barriers

barrier