PyOpenCL 减少算法错误

Question

最近我一直在尝试使用 PyOpenCl 学习 gpu 编程，但尽管我尽了最大的努力，但我仍然无法运行下面代码中显示的缩减算法。相反，代码返回

RuntimeError：clEnqueueReadBuffer 失败：OUT_OF_RESOURCES

我对这个错误的理解是，它表示内存分配不足或内核中的索引超出范围。对于较小的全局尺寸（即较小的(N,A,t)），代码将成功运行，所以我怀疑是前者。但是，对于(1,1,t) 的工作组大小，我将np.dtype(np.float32).itemsize*t 字节分配给本地内存，我认为这应该足够了。有谁知道为什么我会收到这个错误？如果有帮助，我将在 NVIDIA GeForce GTX 960 上运行内核。

import numpy as np
import pyopencl as cl

np.random.seed(5)

N=2500*56
A=6
t=64

plat = cl.get_platforms()
devices = plat[0].get_devices()
ctx = cl.Context([devices[0]])
queue = cl.CommandQueue(ctx)

actions=np.random.randint(0,2,(N,A,t)).flatten(order='F')
tau=np.arange(1,np.add(t,1))
d=np.random.rand(N).astype(np.float32)
baseAct=np.empty((N,A)).astype(np.float32).flatten(order='F')

mf = cl.mem_flags
actions_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, 
hostbuf=actions)
tau_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=tau)
d_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=d)
loc_buf = cl.LocalMemory(np.dtype(np.float32).itemsize*t)
baseAct_buf = cl.Buffer(ctx, mf.WRITE_ONLY, baseAct.nbytes)

prg = cl.Program(ctx, """
    __kernel void calc_baseAct(__global const int *actions,
    __global const int *tau,
    __global const float *d,
    __local float *loc,
    __global float *baseAct,
    int N,
    int A,
    int t)
    {
      int xg = get_global_id(0);
      int yg = get_global_id(1);
      int zg = get_global_id(2);
      int xl = get_local_id(0);
      int yl = get_local_id(1);
      int zl = get_local_id(2);
      int xw = get_group_id(0);
      int yw = get_group_id(1);
      int zw = get_group_id(2);

      loc[xl+N*yl+N*A*zl] = actions[xg+N*yg+N*A*zg]*pow(tau[zg],-d[xg]);
      barrier(CLK_LOCAL_MEM_FENCE);


      for(uint s = t/2; s > 0; s >>= 1) {
        if(zl < s) {
          loc[xl+N*yl+N*A*zl] += loc[xl+N*yl+N*A*(zl+s)];
        }
        barrier(CLK_LOCAL_MEM_FENCE);
      }    
      if(zl == 0) baseAct[xw+N*yw+N*A*zw] = loc[xl+N*yl+N*A*zl];

    }
    """).build()

prg.calc_baseAct(queue, (N,A,t), (1,1,t), actions_buf, tau_buf, d_buf, 
loc_buf, baseAct_buf, np.int32(N), np.int32(A), np.int32(t))
cl.enqueue_copy(queue, baseAct, baseAct_buf)

baseAct=baseAct.reshape((N,A), order='F')

Answer 1

显然，loc 的访问权限越界，每个工作组分配有 64 个元素，并使用索引xl+N*yl+N*A*zl 进行访问，其中zl 在[0,63] 乘以N=2500*56 和A=6 的范围内。