¿Cómo juntar datos de hilos en CUDA?

Question

Tengo una aplicación que resuelve un sistema de ecuaciones en CUDA, estoy seguro de que cada hilo puede encontrar hasta 4 soluciones, pero ¿cómo puedo copiar y luego volver al host?

Estoy pasando una gran variedad con suficiente espacio para todas las soluciones de threads store 4 (4 dobles para cada solución) y otra con la cantidad de soluciones por subproceso, sin embargo es una solución ingenua y es el cuello de botella actual de mi kernel

Realmente me gusta optimizar esto. El principal problema es concatenar una cantidad variable de soluciones por hilo en una única matriz.

Answer 1

La funcionalidad que está buscando se llama compactación de flujo.

Probablemente necesite proporcionar una matriz que contenga espacio para 4 soluciones por subproceso, porque intentar almacenar directamente los resultados en una forma compacta es probable que cree tantas dependencias entre los subprocesos que el rendimiento obtenido al poder copiar menos datos de regreso al host se pierden por un tiempo de ejecución de kernel más largo. La excepción a esto es si casi todos los hilos no encuentran soluciones. En ese caso, es posible que pueda utilizar una operación atómica para mantener un índice en una matriz. Entonces, para cada solución que se encuentre, la almacenaría en una matriz en un índice y luego usaría una operación atómica para aumentar el índice. Creo que sería seguro usar atomicAdd() para esto. Antes de almacenar un resultado, el hilo usaría atomicAdd() para aumentar el índice en uno. atomicAdd() devuelve el valor anterior, y el hilo puede almacenar el resultado utilizando el valor anterior como índice.

Sin embargo, dada una situación más común, donde hay un buen número de resultados, la mejor solución será realizar una operación de compactación como un paso separado. Una forma de hacerlo es con thrust::copy_if. Consulte this question para obtener más información.