PyOpenCL indicizzazione di array 3D all'interno del codice del kernel

Question

Sto usando PyOpenCL per elaborare le immagini in Python e inviare un array 3D di numpy (height x width x 4) al kernel. Sto riscontrando problemi nell'indicizzazione dell'array 3D all'interno del codice del kernel. Per ora sono solo in grado di copiare l'intero array di input per l'output. Il codice attuale assomiglia a questo, in cui img è l'immagine con img.shape = (320, 512, 4):

__kernel void part1(__global float* img, __global float* results) 
{ 
    unsigned int x = get_global_id(0); 
    unsigned int y = get_global_id(1); 
    unsigned int z = get_global_id(2); 

    int index = x + 320*y + 320*512*z; 

    results[index] = img[index]; 
} 

Tuttavia, io non capisco come questo lavoro. Ad esempio, come faccio a indicizzare l'equivalente Python di img[1, 2, 3] all'interno di questo kernel? E inoltre, quale indice dovrebbe essere utilizzato in results per la memorizzazione di qualche elemento se voglio che sia nella posizione results[1, 2, 3] nell'array numpy quando ottengo i risultati di nuovo in Python?

Per eseguire questa Sto usando questo codice Python:

import pyopencl as cl 
import numpy as np 

class OpenCL: 
def __init__(self): 
    self.ctx = cl.create_some_context() 
    self.queue = cl.CommandQueue(self.ctx) 

def loadProgram(self, filename): 
    f = open(filename, 'r') 
    fstr = "".join(f.readlines()) 
    self.program = cl.Program(self.ctx, fstr).build() 

def opencl_energy(self, img): 
    mf = cl.mem_flags 

    self.img = img.astype(np.float32) 

    self.img_buf = cl.Buffer(self.ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=self.img) 
    self.dest_buf = cl.Buffer(self.ctx, mf.WRITE_ONLY, self.img.nbytes) 

    self.program.part1(self.queue, self.img.shape, None, self.img_buf, self.dest_buf) 
    c = np.empty_like(self.img) 
    cl.enqueue_read_buffer(self.queue, self.dest_buf, c).wait() 
    return c 

example = OpenCL() 
example.loadProgram("get_energy.cl") 
image = np.random.rand(320, 512, 4) 
image = image.astype(np.float32) 
results = example.opencl_energy(image) 
print("All items are equal:", (results==image).all()) 

Answer 1

Anche se questa non è la soluzione opitimal, ho linearizzato l'array in Python e mandata come 1D. Nel codice del kernel ho calcolato x, e z dall'indice lineare. Quando è tornato a Pyhon l'ho rimodellato alla forma originale.

Answer 2

Aggiornamento: L'OpenCL docs stato (a 3,5), che

"Memory objects are categorized into two types: buffer objects, and image objects. A buffer 
object stores a one-dimensional collection of elements whereas an image object is used to store a 
two- or three- dimensional texture, frame-buffer or image." 

così, un buffer è sempre lineare, o linearizzato come potete vedere dal mio esempio.

import pyopencl as cl 
import numpy as np 


h_a = np.arange(27).reshape((3,3,3)).astype(np.float32) 

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

mf = cl.mem_flags 
d_a = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=h_a) 

prg = cl.Program(ctx, """ 
__kernel void p(__global const float *d_a) { 
    printf("Array element is %f ",d_a[10]); 
} 
""").build() 

prg.p(queue, (1,), None, d_a) 

Mi dà

"Array element is 10" 

come output. Quindi, il buffer in realtà è l'array linearizzato. Tuttavia, l'approccio ingenuo [x, y, z] noto da numpy non funziona in questo modo. L'utilizzo di un'immagine 2 o 3-D invece di un buffer dovrebbe comunque funzionare.

Answer 3

Ho riscontrato lo stesso problema. https://lists.tiker.net/pipermail/pyopencl/2009-October/000134.html è un semplice esempio di come utilizzare gli array 3d con PyOpenCL che ha funzionato per me. Cito qui il codice di riferimento futuro:

import pyopencl as cl 
import numpy 
import numpy.linalg as la 

sizeX=4 
sizeY=2 
sizeZ=5 
a = numpy.random.rand(sizeX,sizeY,sizeZ).astype(numpy.float32) 

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

mf = cl.mem_flags 
a_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=a) 
dest_buf = cl.Buffer(ctx, mf.WRITE_ONLY, a.nbytes) 

prg = cl.Program(ctx, """ 
    __kernel void sum(__global const float *a, __global float *b) 
    { 
     int x = get_global_id(0); 
     int y = get_global_id(1); 
     int z = get_global_id(2); 

     int idx = z * %d * %d + y * %d + x; 

     b[idx] = a[idx] * x + 3 * y + 5 * z; 
    } 
    """ % (sizeY, sizeX, sizeX)).build() 

prg.sum(queue, a.shape, a_buf, dest_buf) 
cl.enqueue_read_buffer(queue, dest_buf, a).wait() 
print a