Ошибка алгоритма сокращения PyOpenCL - PullRequest
Новая
       23

Ошибка алгоритма сокращения PyOpenCL

0 голосов
/ 03 мая 2018

Недавно я пытался научиться программировать на GPU с помощью PyOpenCl, но, несмотря на все мои усилия, я не смог заставить работать алгоритм сокращения, показанный в приведенном ниже коде. Вместо этого код возвращает

RuntimeError: ошибка clEnqueueReadBuffer: OUT_OF_RESOURCES

Мое понимание этой ошибки состоит в том, что она указывает либо на недостаточное выделение памяти, либо на индексирование за пределами границ в ядре. Для небольших глобальных размеров (то есть малых (N,A,t)) код будет успешно выполнен, поэтому я подозреваю, что первый. Я выделяю np.dtype(np.float32).itemsize*t байт для локальной памяти, однако для размера рабочей группы (1,1,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')

1 Ответ

0 голосов
/ 03 мая 2018

Ясно, что вне привязанного доступа для loc, для которого выделено 64 элемента на рабочую группу, доступ осуществляется с индексом xl+N*yl+N*A*zl, где zl находится в диапазоне [0,63], умноженном на N=2500*56 и A=6.

...