Pasar la estructura con los miembros del puntero al núcleo OpenCL usando PyOpenCL

Supongamos que tengo un kernel para calcular la sum de dos arreglos en cuanto a elementos. En lugar de pasar a, b y c como tres parámetros, los hago estructurar miembros de la siguiente manera:

typedef struct { __global uint *a; __global uint *b; __global uint *c; } SumParameters; __kernel void compute_sum(__global SumParameters *params) { uint id = get_global_id(0); params->c[id] = params->a[id] + params->b[id]; return; } 

Hay información sobre las estructuras si usted RTFM de PyOpenCL [1], y otros han abordado esta pregunta también [2] [3] [4]. Pero ninguno de los ejemplos de estructura OpenCL que he podido encontrar tiene punteros como miembros.

Específicamente, me preocupa si los espacios de direcciones del host / dispositivo coinciden, y si los tamaños del puntero del host / dispositivo coinciden. Alguien sabe la respuesta?

[1] http://documen.tician.de/pyopencl/howto.html#how-to-use-struct-types-with-pyopencl

[2] Alineación de Struct con PyOpenCL

[3] http://enja.org/2011/03/03/30/adventures-in-opencl-part-3-constant-memory-structs/

[4] http://acooke.org/cute/Somesimple0.html

No, no hay garantía de que los espacios de direcciones coincidan. Para los tipos básicos (float, int, …) tiene un requisito de alineación (sección 6.1.5 de la norma) y tiene que usar el nombre cl_type de la implementación de OpenCL (cuando se progtwig en C, pyopencl hace el trabajo bajo el capó I lo diría).

Para los indicadores es aún más simple debido a esta falta de coincidencia. El comienzo de la sección 6.9 del estándar v 1.2 (es la sección 6.8 para la versión 1.1) establece:

Los argumentos a las funciones del kernel declarados en un progtwig que son punteros se deben declarar con el calificador __global, __constant o __local.

Y en el punto p .:

Los argumentos de las funciones del núcleo que se declaran como una estructura o unión no permiten que los objetos OpenCL se pasen como elementos de la estructura o unión.

Tenga en cuenta también el punto d .:

Las matrices de longitud variable y las estructuras con matrices flexibles (o sin tamaño) no son compatibles.

Por lo tanto, no hay forma de hacer que el kernel se ejecute como se describe en su pregunta y es por eso que no ha podido encontrar algunos ejemplos de estructuras OpenCl que tengan punteros como miembros.
Todavía puedo proponer una solución que aproveche el hecho de que el kernel está comstackdo en JIT. Todavía requiere que usted empaque sus datos correctamente y que preste atención a la alineación y, finalmente, que el tamaño no cambie durante la ejecución del progtwig. Sinceramente, me gustaría que un kernel tomara 3 buffers como argumentos, pero de todos modos, ahí está.

La idea es usar la opción del preprocesador –D como en el siguiente ejemplo en python:

Núcleo:

 typedef struct { uint a[SIZE]; uint b[SIZE]; uint c[SIZE]; } SumParameters; kernel void foo(global SumParameters *params){ int idx = get_global_id(0); params->c[idx] = params->a[idx] + params->b[idx]; } 

Código de host:

 import numpy as np import pyopencl as cl def bar(): mf = cl.mem_flags ctx = cl.create_some_context() queue = cl.CommandQueue(self.ctx) prog_f = open('kernels.cl', 'r') #a = (1, 2, 3), b = (4, 5, 6) ary = np.array([(1, 2, 3), (4, 5, 6), (0, 0, 0)], dtype='uint32, uint32, uint32') cl_ary = cl.Buffer(ctx, mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=ary) #Here should compute the size, but hardcoded for the example size = 3 #The important part follows using -D option prog = cl.Program(ctx, prog_f.read()).build(options="-D SIZE={0}".format(size)) prog.foo(queue, (size,), None, cl_ary) result = np.zeros_like(ary) cl.enqueue_copy(queue, result, cl_ary).wait() print result 

Y el resultado:

 [(1L, 2L, 3L) (4L, 5L, 6L) (5L, 7L, 9L)] 

No conozco la respuesta a mi propia pregunta, pero hay tres soluciones que se me pueden ocurrir. Considero que la solución 3 es la mejor opción.

Solución 1: aquí solo tenemos 3 parámetros, por lo que podríamos simplemente hacer los parámetros del kernel a, b y c. Pero he leído que hay un límite en la cantidad de parámetros que puede pasar a un kernel, y personalmente me gusta refactorizar cualquier función que requiera más de 3-4 argumentos para usar structs (o, en Python, tuplas o argumentos de palabras clave) . Así que esta solución hace que el código sea más difícil de leer y no se puede escalar.

Solución 2: Volcar todo en una sola matriz gigante. Entonces el kernel se vería así:

 typedef struct { uint ai; uint bi; uint ci; } SumParameters; __kernel void compute_sum(__global SumParameters *params, uint *data) { uint id = get_global_id(0); data[params->ci + id] = data[params->ai + id] + data[params->bi + id]; return; } 

En otras palabras, en lugar de usar punteros, use las compensaciones en una sola matriz. Esto se parece mucho a los inicios de la implementación de mi propio modelo de memoria, y se siente como si estuviera reinventando una rueda que existe en algún lugar de PyOpenCL, OpenCL, o ambos.

Solución 3: Hacer los núcleos de la incubadora. Me gusta esto:

 __kernel void set_a(__global SumParameters *params, __global uint *a) { params->a = a; return; } 

y ídem para set_b, set_c. Luego ejecute estos kernels con el tamaño de trabajo 1 para configurar la estructura de datos. Aún necesita saber qué tan grande es un bloque para asignar a los parámetros, pero si es demasiado grande, no ocurrirá nada malo (excepto un poco de memoria desperdiciada), por lo que diría que supongamos que los punteros son de 64 bits.

El rendimiento de esta solución es probablemente horrible (me imagino que una llamada al kernel tiene una gran sobrecarga), pero afortunadamente eso no debería importar demasiado para mi aplicación (mi kernel se ejecutará por unos segundos a la vez, no es una cosa gráfica que tenga que ejecutar a 30-60 fps, así que imagino que el tiempo que requieren las llamadas adicionales al kernel para establecer los parámetros terminará siendo una pequeña fracción de mi carga de trabajo, sin importar qué tan alta sea la sobrecarga de llamadas por kernel).