¿La forma correcta de informar los kernels OpenCL de muchos objetos de memoria?

Question

En mi programa OpenCL, voy a terminar con más de 60 memorias intermedias globales que cada kernel necesitará para poder acceder. ¿Cuál es la forma recomendada de permitir que cada núcleo conozca la ubicación de cada uno de estos almacenamientos intermedios?

Los búferes son estables durante toda la vida útil de la aplicación, es decir, asignaremos los búferes al inicio de la aplicación, llamaremos a múltiples kernels y luego desasignaremos los búferes al final de la aplicación. Sin embargo, su contenido puede cambiar a medida que los núcleos leen/escriben de ellos.

En CUDA, la forma en que lo hice fue crear más de 60 variables globales de ámbito de programa en mi código CUDA. Luego, en el host, escribiría la dirección de los almacenamientos intermedios del dispositivo asignados en estas variables globales. Entonces los kernels simplemente usarían estas variables globales para encontrar el buffer con el que necesitaban trabajar.

¿Cuál sería la mejor manera de hacer esto en OpenCL? Parece que las variables globales de CL son un poco diferentes a las de CUDA, pero no puedo encontrar una respuesta clara sobre si mi método CUDA funcionará, y si es así, cómo transferir los punteros del buffer a las variables globales. Si eso no funciona, ¿cuál es la mejor manera de lo contrario?

Answer 1

60 variables globales ¡seguro que es mucho! ¿Estás seguro de que no hay una forma de refactorizar tu algoritmo para usar trozos de datos más pequeños? Recuerde, cada núcleo debe ser una unidad mínima de trabajo, ¡no algo colosal!

Sin embargo, hay una posible solución. Suponiendo que sus 60 matrices son de tamaño conocido, puede almacenarlas todas en un gran búfer, y luego usar compensaciones para acceder a varias partes de ese gran conjunto. He aquí un ejemplo muy simple con tres matrices:

A is 100 elements 
B is 200 elements 
C is 100 elements 

big_array = A[0:100] B[0:200] C[0:100] 
offsets = [0, 100, 300] 

A continuación, sólo tiene que pasar big_array y las compensaciones a su núcleo, y se puede acceder a cada matriz. Por ejemplo:

A[50] = big_array[offsets[0] + 50] 
B[20] = big_array[offsets[1] + 20] 
C[0] = big_array[offsets[2] + 0] 

No estoy seguro de cómo esto afecta el almacenamiento en caché en el dispositivo, pero mi conjetura inicial es "no muy bien". Este tipo de acceso a arreglos también es un poco desagradable. No estoy seguro de si sería válido, pero podría comenzar cada uno de sus núcleos con algún código que extraiga cada desplazamiento y lo agregue a una copia del puntero original.

En el lado del host, para mantener sus matrices más accesibles, puede usar clCreateSubBuffer: http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clCreateSubBuffer.html que también le permitirá pasar referencias a matrices específicas sin la matriz de compensaciones.

No creo que esta solución sea mejor que pasar los 60 argumentos del kernel, pero dependiendo de la implementación de OpenCL, clSetKernelArgs, podría ser más rápido. Ciertamente reducirá la longitud de su lista de argumentos.

Answer 2

Tienes que hacer dos cosas. En primer lugar, cada núcleo que utiliza cada búfer de memoria global debería declarar un argumento para cada uno, algo como esto:

kernel void awesome_parallel_stuff(global float* buf1, ..., global float* buf60) 

de manera que cada búfer utilizado para ese kernel está en la lista. Y luego, en el lado del host , debe crear cada buffer y usar clSetKernelArg para adjuntar un búfer de memoria determinado a un argumento kernel dado antes de llamar al clEnqueueNDRangeKernel para iniciar la fiesta.

Tenga en cuenta que si los núcleos seguirán usando el mismo buffer con cada ejecución del kernel, solo necesita configurar los argumentos del kernel uno time. Un error común que veo, que puede sangrar el rendimiento del lado del host, es llamar repetidamente al clSetKernelArg en situaciones donde es completamente innecesario.