尽管有障碍，但打开 CL 没有同步

Question

我刚开始通过 Python 的 PyOpenCL 接口使用 OpenCL。我试图创建一个非常简单的“循环”程序，其中每个内核中每个循环的结果取决于上一个循环周期中另一个内核的输出，但我遇到了同步问题：

__kernel void part1(__global float* a, __global float* c)
{
    unsigned int i = get_global_id(0);

    c[i] = 0;
    barrier(CLK_GLOBAL_MEM_FENCE);

    if (i < 9)
    {
        for(int t = 0; t < 2; t++){
            c[i] = c[i+1] + a[i];
            barrier(CLK_GLOBAL_MEM_FENCE);
       }
    }
}

宿主应用程序是

import pyopencl as cl
from numpy import *

ctx = cl.create_some_context()
queue = cl.CommandQueue(ctx)

#read in the OpenCL source file as a string
f = open('recurrent.cl', 'r')
fstr = "".join(f.readlines())

#create the program
program = cl.Program(ctx, fstr).build()

mf = cl.mem_flags

#initialize client side (CPU) arrays
a = array(range(10), dtype=float32)

#create OpenCL buffers
a_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=a)
dest_buf = cl.Buffer(ctx, mf.WRITE_ONLY, a.nbytes)

#execute program
program.part1(queue, a.shape, None, a_buf, dest_buf)
c = empty_like(a)
cl.enqueue_read_buffer(queue, dest_buf, c).wait()

print "a", a
print "c", c

结果是

a [ 0.  1.  2.  3.  4.  5.  6.  7.  8.  9.]
c [  0.   1.   5.   3.   4.  18.  13.   7.   8.   0.]

如您所见，一些结果值是正确的。例如。第三个位置 = 5 = 3 + 2 但例如第二个位置是 2 = 0 + 2。因此，尽管存在障碍，但总和超过了其他线程在不同时间点的结果。我认为屏障会确保所有线程都到达它并将其结果写入全局内存？

这可能是非常简单的事情，我将不胜感激任何提示和 cmets！

PS：我正在使用 Intel SDK 在 Sandy Bridge CPU 上运行它。

Answer 1

编辑： 用户blue2script 是对的，这是“所有本地线程都没有击中屏障”的问题。最重要的是，屏障无法在计算单元/工作组之间同步。

我的回答没有添加任何内容，也没有解决任何问题。所以不要在下面的内核函数中看到if。错了。

---

不完整

 __kernel void part1(__global float* a, __global float* c)
 {
      unsigned int i = get_global_id(0);

      c[i] = 0;
      barrier(CLK_GLOBAL_MEM_FENCE);

      if (i < 9)
      {
          for(int t = 0; t < 2; t++)
          {
              c[i] = c[i+1] + a[i];//c[i+1] is neighbour thread's variable
                                   //and there is no guarantee that
                                   //which one(ith or (i+1)st) computes first
                                   //so you need to get a copy of c[] first
              barrier(CLK_GLOBAL_MEM_FENCE);//thats why this line is not helping
          }
      }
 }

使用全局

 __kernel void part1(__global float* a, __global float* c,__global float* d)
 {
      unsigned int i = get_global_id(0);

      c[i] = 0;
      d[i]=c[i]; 
      barrier(CLK_GLOBAL_MEM_FENCE);

      if (i < 9)
      {
          for(int t = 0; t < 2; t++)
          {
              d[i] = c[i+1] + a[i];//it is guaranteed that no neighbour thread can
                                   //change this threads d[i] element before/after
                                   //execution
              barrier(CLK_GLOBAL_MEM_FENCE);
              c[i]=d[i];
              barrier(CLK_GLOBAL_MEM_FENCE);
          }
      }
      barrier(CLK_GLOBAL_MEM_FENCE);

 }

使用本地人（工作组大小为 256，总工作量是其倍数）：

 __kernel void part1(__global float* a, __global float* c)
 {
      unsigned int i = get_global_id(0);
      unsigned int Li=get_local_id(0);
      __local d[256];
      c[i] = 0;
      barrier(CLK_GLOBAL_MEM_FENCE);
      d[Li]=c[i]; 
      barrier(CLK_LOCAL_MEM_FENCE);

      if (i < 9)
      {
          for(int t = 0; t < 2; t++)
          {
              d[Li] = c[i+1] + a[i];//it is guaranteed that no neighbour thread can
                                   //change this threads d[i] element before/after
                                   //execution

             barrier(CLK_LOCAL_MEM_FENCE);
             c[i]=d[Li]; //guaranteed they dont interfere each other
             barrier(CLK_LOCAL_MEM_FENCE);
          }
      }

 }

工作组：

使用私有

 __kernel void part1(__global float* a, __global float* c)
 {
      unsigned int i = get_global_id(0);
      unsigned int Li=get_local_id(0);
      __private f1;
      c[i] = 0;

      if (i < 9)
      {
          for(int t = 0; t < 2; t++)
          {
              f1 = c[i+1] + a[i];

             barrier(CLK_GLOBAL_MEM_FENCE);
             c[i]=f1; //guaranteed they dont interfere each other
             barrier(CLK_GLOBAL_MEM_FENCE);
          }
      }

 }

Answer 2

我想我现在有了答案。 OpenCL 代码实际上完全没问题。然而，只有当所有线程都在一个工作组中时，障碍才会生效。情况并非如此，这很容易通过使用 get_local_id(0) 读出 local_id 来检查（如 Huseyin 所建议的）。在我的情况下，主机为每个线程创建了一个工作组 - 而不是将所有线程放在一个工作组中。在性能方面，这是有道理的，比较

Questions about global and local work size

然而，在我们的例子中，我们需要确保线程之间的数据是同步的，因此所有线程都应该在一个工作组中。为此我们需要改变程序1的执行，

program.part1(queue, a.shape, None, a_buf, dest_buf)

第二个参数指的是作业的 global_size（即创建的线程数），而第三个参数似乎指的是 local_size，即每个工作组的线程数。因此，这一行应该是

program.part1(queue, a.shape, a.shape, a_buf, dest_buf)

这将创建一个包含所有线程的工作组（但请注意一个工作组中允许的最大工作人员大小！）。现在，代码仍然不起作用。最后一个问题与 OpenCL 代码中的障碍有关：id = 10 的最后一个线程在循环中看不到障碍，因此所有线程都在等待最后一个遇到障碍（尽管我想知道为什么没有'不抛出异常？）。所以我们只需要减少线程总数（去掉最后一个），

program.part1(queue, (a.shape[0]-1,), (a.shape[0]-1,), a_buf, dest_buf)

这行得通！在此过程中吸取了一些教训...

再次感谢侯赛因！ blue2script