Reducción final en varios bloques

Una reducción dentro de un bloque es sencilla, pero los bloques se ejecutan de forma independiente y no pueden sincronizarse entre sí durante la ejecución del kernel. Por eso, un solo lanzamiento no puede sumarlo todo.

Por tanto, cada bloque reduce su propio segmento a un número, una suma parcial, y la escribe en un pequeño array de salida indexado mediante blockIdx.

Con 1000 bloques, pasa de un millón de entradas a 1000 parciales. La parte difícil ya está hecha; solo queda combinar un array pequeño.

La forma más sencilla de terminar es realizar un segundo lanzamiento del mismo kernel sobre los parciales. Repita el proceso hasta que solo quede un valor.

Cada pasada reduce el tamaño del array según el tamaño del bloque. Unos pocos lanzamientos recursivos reducen millones de valores a una única suma final.

Como alternativa, el hilo 0 de cada bloque puede sumar su parcial directamente a un único total global mediante atomicAdd, evitando un segundo kernel.

Las operaciones atómicas son sencillas y solo requieren un lanzamiento, pero muchos bloques que compiten por la misma dirección pueden serializarse. Con pocos parciales, normalmente funcionan bien.

Un bucle grid-stride permite que cada hilo sume primero muchos elementos en un registro, de modo que se necesiten muchos menos bloques antes del paso final.

Realizar más trabajo por hilo al principio implica menos parciales y menos lanzamientos. Esto suele ser mejor que crear un hilo por elemento.

Si utiliza operaciones atómicas, recuerde poner a cero el total global antes del lanzamiento; de lo contrario, la suma comenzará con datos residuales de esa memoria.

Las entradas pequeñas aprovechan la sencillez de las operaciones atómicas; las entradas enormes favorecen un diseño de dos pasadas o grid-stride. Mida el rendimiento con sus datos para elegir.

Cada bloque produce una suma parcial y después usted las combina con un segundo lanzamiento, operaciones atómicas o grid-stride. Ahora puede reducir arrays de cualquier tamaño. 🏁