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


En mi programa OpenCL, voy a terminar con más de 60 memorias intermedias globales a las que cada kernel va a necesitar tener acceso. ¿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 mismos son estables a lo largo de la vida 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 del alcance del 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?


6
2018-06-16 11:20


origen


Respuestas:


60 variables globales seguro 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. Aquí hay 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]

Entonces, solo necesita pasar big_array y compensaciones a su kernel, y 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 afectaría el almacenamiento en caché de su dispositivo en particular, pero mi conjetura inicial es "no está 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ía 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 núcleo, 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.


1
2018-06-21 15:04



Tienes que hacer dos cosas. En primer lugar, cada kernel que utiliza cada memoria buffer global debe declarar un argumento para cada uno, algo como esto:

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

para que cada buffer utilizado para ese kernel aparezca en la lista. Y luego, en el anfitrión lado, necesitas crear cada buffer y usar clSetKernelArg para adjuntar un búfer de memoria dado a un argumento kernel dado antes de llamar clEnqueueNDRangeKernelpara comenzar 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 hora. Un error común que veo, que puede sangrar el rendimiento del lado del host, es llamar repetidamente clSetKernelArg en situaciones donde es completamente innecesario.


0
2018-06-16 21:23