Lección 25: Scan paralelo — el kernel de duplicación
En la lección anterior vimos qué calcula una suma de prefijos y la idea de duplicación: en cada ronda cada elemento suma al vecino que está a una distancia d, y d se duplica (1, 2, 4...). Aquí lo construimos como un kernel real de CUDA. Trabajamos dentro de un solo bloque con un arreglo en memoria c
Cada hilo es dueño de una posición en el arreglo compartido. En cada paso, primero lee el valor que está d lugares a su izquierda y lo aparta, espera a que todos hayan leído (__syncthreads), y solo entonces lo suma a sí mismo, y de nuevo espera a que todos hayan escrito antes del siguiente paso. Las esperas son lo que mantiene el orden y evita la sobrescritura.
- Hillis-Steele
- El kernel de scan por duplicación en log2(n) pasos: en el paso d, cada hilo con tid>=d suma el valor de tid-d. d se duplica en cada paso.
- tile en memoria compartida
- Un arreglo __shared__ para todo el bloque. El scan carga la entrada en él, trabaja en el lugar a lo largo de todos los pasos, y finalmente escribe en out.
- barreras de lectura/escritura
- __syncthreads() después de la lectura (para que todos lean antes de que alguien escriba) y después de la escritura (para que las escrituras sean visibles en el siguiente paso).
- número de pasos log2(n)
- El bucle for(d=1; d<n; d*=2) se ejecuta log2(n) veces, porque d se duplica en cada paso hasta llegar a n.
