¿Cómo declarar la memoria local en OpenCL?

Question

estoy ejecutando el kernel OpenCL a continuación con un tamaño global de trabajo bidimensional de 1.000.000 x 100 y un tamaño de trabajo local de 1 x 100.

__kernel void myKernel(
     const int length, 
     const int height, 
     and a bunch of other parameters) { 

    //declare some local arrays to be shared by all 100 work item in this group 
    __local float LP [length]; 
    __local float LT [height]; 
    __local int bitErrors = 0; 
    __local bool failed = false; 

    //here come my actual computations which utilize the space in LP and LT 
} 

Sin embargo, esto se niega a compilar, ya que los parámetros y lengthheight no se conocen en tiempo de compilación. Pero no está claro para nada cómo hacer esto correctamente. ¿Debo usar punteros con memalloc? ¿Cómo manejar esto de manera que la memoria solo se asigna una vez para todo el grupo de trabajo y no una vez por elemento de trabajo?

Todo lo que necesito es 2 matrices de flotadores, 1 int y 1 booleano que se comparten entre todo el grupo de trabajo (por lo que los 100 elementos de trabajo). Pero no logro encontrar cualquier método que hace esto correctamente ...

Answer 1

Es relativamente simple, se puede pasar a las matrices locales como argumentos para el núcleo:

kernel void myKernel(const int length, const int height, local float* LP, 
        local float* LT, a bunch of other parameters) 

A continuación, establezca la kernelargument con un value de NULL y size igual al tamaño que desea asignar para el argumento (en byte). Por lo tanto, debe ser:

clSetKernelArg(kernel, 2, length * sizeof(cl_float), NULL); 
clSetKernelArg(kernel, 2, height* sizeof(cl_float), NULL); 

memoria local siempre es compartida por el grupo de trabajo (en contraposición a lo privado), así que creo que la bool y int debe estar bien, pero si no siempre se puede pasar a los que como argumentos también.

No realmente relacionado con su problema (y no necesariamente relevante, ya que no sé qué hardware planea ejecutar esto), pero al menos los gpus no son particularmente como tamaños de trabajo que no son múltiplos de una potencia particular de dos (creo que fue 32 para nvidia, 64 para amd), lo que significa que probablemente creará grupos de trabajo con 128 elementos, de los cuales los últimos 28 son básicamente desperdiciados. Por lo tanto, si está ejecutando opencl en gpu podría ayudar al rendimiento si usa directamente grupos de trabajo de tamaño 128 (y cambia el tamaño de trabajo global de forma apropiada)

Como nota al margen: Nunca entendí por qué todos usan la variante de subrayado para kernel, local and global, me parece mucho más feo

Answer 2

No tiene que asignar toda su memoria local fuera del kernel, especialmente cuando se trata de una variable simple en lugar de una matriz.

El motivo por el cual su código no se puede compilar es que OpenCL no admite la inicialización de la memoria local. Esto se especifica en el documento (https://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/local.html). Tampoco es factible en CUDA (Is there a way of setting default value for shared memory array?)

---

PD: La respuesta de Grizzly es lo suficientemente bueno y que sería mejor si me pueden enviar como un comentario, pero estoy limitado por la política de reputación. Lo siento.

Answer 3

También podría declarar matrices de la siguiente manera:

__local float LP[LENGTH]; 

y pasar la LONGITUD como definir en su compilación del núcleo.

clBuildProgram(program, 0, NULL, "-DLENGTH=128", NULL, NULL);