Lección 15: Diagnosticar y corregir el acceso con stride
En la lección anterior vimos qué es el coalescing: cuando hilos consecutivos de un warp tocan direcciones consecutivas, una transacción los atiende a todos, y cuando saltan con un stride grande el warp consume muchas transacciones desperdiciadas. Ahora aprendemos la habilidad práctica — detectar un
Imagina 4 trabajadores recogiendo manzanas de filas de árboles. Al principio cada trabajador corre a una columna distinta en un extremo lejano del huerto — 4 carreras largas. Si en cambio cada trabajador toma la siguiente manzana en una sola fila apretada, todos trabajan uno al lado del otro en una sola pasada. No cambiamos el huerto — solo quién-recoge-qué.
- acceso con stride
- Un patrón donde los hilos consecutivos saltan con un stride grande (por ejemplo data[threadIdx.x * N]), así que sus direcciones están muy separadas y cada una cae en una transacción de memoria distinta.
- coalesced vs. no coalesced
- Coalesced: las direcciones de los hilos vecinos son contiguas, una transacción para todo el warp. No coalesced: las direcciones están dispersas, el warp consume muchas transacciones, la mayoría desperdiciadas.
- corrección por remapeo del índice
- La corrección para el acceso disperso: hacer que los hilos vecinos toquen direcciones vecinas — darle al hilo t el elemento base + t en lugar de base + t*N, todo en memoria global.
- orden row-major
- Un diseño donde los elementos de la misma fila son vecinos en memoria. Recorrer por fila hace que los hilos vecinos toquen direcciones vecinas (coalesced); recorrer por columna salta según width (disperso).
