Nuestra publicación anterior siguió un kernel de suma de vectores desde nvcc hasta los warps. Esta vez, seguimos una instrucción SASS crítica (una carga global) a través del hardware de una RTX4090. El kernel CUDA investigado tiene dos líneas: __global__ void vadd(const float* a, const float* b, float* c, int n) {int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) c[i] = a[i] + b[i];}. El SASS compilado muestra: /*0080*/ IMAD.WIDE R4, R6, R7, c[0x0][0x168] ; // &b[i] /*00a0*/ LDG.E R4, [R4.64] ; // b[i]. Estas instrucciones cargan elementos del vector b desde la memoria global a los registros.
Una LDG.E solicita cuatro bytes en cada uno de los 32 carriles. Servirla requiere cuatro sectores de 32 bytes, una línea de caché, una traducción de direcciones, un cruce del crossbar, uno de los treinta y seis slices L2 y, cuando falla en todas partes, una activación y cuatro lecturas de columna en un chip DRAM. Nuestro warp vive en una de las cuatro subparticiones del SM junto con otros once warps residentes. La instrucción se emite a la unidad de carga/almacenamiento (LSU), que envía el código de operación, la máscara de carriles activos, las direcciones calculadas y el número de registro al coalescer.
El coalescer determina que se necesitan 4 solicitudes de sector contiguas para los 128 bytes que solicitó el warp. Al entrar en la caché L1, la solicitud se envía a la caché de direccionamiento virtual, asociativa por conjuntos de 4 vías. La búsqueda compara etiquetas en conjuntos de cuatro ranuras. Como es una carga por primera vez, falla y desciende más. Un acierto en L1 devuelve en unos 15,4 ns — 40 ciclos. La memoria virtual requiere traducción antes de salir del SM. La TLB del SM contiene dieciséis traducciones recientes. Los fallos cuestan unos 4,4 ns — once ciclos.
Después de la traducción, una solicitud por línea de 128 bytes avanza a través del crossbar hacia uno de los 36 slices L2 de 2 MiB, elegido por una función de su dirección física. Todos los slices pueden servir en paralelo, lo que da un ancho de banda agregado 36 veces el de un solo slice.
Fuente: Hacker News · Resumido por HeadlinesBriefing