Qué pasa internamente cuando ejecutas un kernel CUDA

 (fergusfinn.com)
2 puntos por GN⁺ 3 시간 전 | 1 comentarios | Compartir por WhatsApp
  • Incluso un programa CUDA simple de suma de vectores pasa por el pipeline de compilación, llamadas al driver, colas de comandos de la GPU, planificación de warps, jerarquía de memoria y semáforos de finalización antes de obtener el resultado 2.000000
  • nvcc separa el código host y el código device: genera PTX con cicc, SASS con ptxas, y empaqueta el cubin y el PTX en un fatbin que inserta dentro del ejecutable de Linux
  • La sintaxis de launch vadd<<<4096, 256>>> se transforma en un host launch stub, y los argumentos da, db, dc, n se entregan al driver pasando por el runtime de CUDA y libcuda.so.1
  • La ejecución en la GPU comienza con QMD, pushbuffer, GPFIFO, GP_PUT y una escritura MMIO al doorbell; los 128 SM de la RTX 4090 ejecutan la configuración de 4096 bloques y 256 threads por bloque a nivel de warp
  • Debido a su baja intensidad aritmética, que requiere transferir 12 bytes por cada suma float, este kernel queda limitado por el ancho de banda de memoria: en Nsight Compute marca 10.78 μs, 79.65% del pico de DRAM y 5.17% de warp issue

Kernel de ejemplo y alcance de la observación

  • El programa de ejemplo usa el kernel CUDA vadd, que suma dos arreglos de float y guarda el resultado en un tercer arreglo
    • Con n = 1 << 20, procesa 1,048,576 valores float
    • La configuración de launch es vadd<<<4096, 256>>>(da, db, dc, n) y usa 4096 * 256 = n threads
  • Al compilarlo para una RTX 4090 con nvcc -arch=sm_89 y ejecutarlo, imprime c[0]=2.000000 c[n-1]=2.000000
  • Incluso para obtener esa sola línea de resultado intervienen decenas de millones de instrucciones de CPU, un device file, unos 900 ioctl y un registro doorbell mapeado en memoria

Cómo nvcc crea el ejecutable

  • Con nvcc --keep se pueden inspeccionar directamente los artefactos del pipeline de compilación
    • vadd.ptx: PTX del código device generado por cicc
    • vadd.sm_89.cubin: SASS del código device generado por ptxas
    • vadd.fatbin: fatbin que empaqueta el cubin y el PTX
    • vadd.cudafe1.stub.c: host launch stub y código de registro del kernel
    • vadd.o: objeto host final que incluye el fatbin
  • El código host lo procesa el compilador host, mientras que el kernel device vadd pasa por las etapas cicc y ptxas
  • PTX es una ISA virtual que usa registros virtuales infinitos con tipo, y no refleja directamente la cantidad real de registros de hardware
    • El PTX del ejemplo incluye el cálculo blockIdx.x * blockDim.x + threadIdx.x, comprobación de límites, global load, suma float y global store
    • Los punteros CUDA son por defecto generic pointer, por lo que se convierten a global address con cvta.to.global antes de usar ld.global
    • mul.wide.s32 convierte el índice en un offset en unidades de 4 bytes, que es sizeof(float), y lo extiende de 32 a 64 bits
  • SASS son las instrucciones reales específicas de la arquitectura, y en la salida para RTX 4090 aparecen de forma más compacta que el PTX
    • S2R copia registros especiales como SR_CTAID.X y SR_TID.X a registros generales
    • La combinación de mul.wide y add de PTX se fusiona en IMAD.WIDE en SASS
    • La conversión cvta queda absorbida dentro del proceso de direccionamiento
  • Los operandos c[0x0][...] apuntan al constant bank 0 administrado por el driver
    • Los punteros a, b, c están en 0x160, 0x168, 0x170
    • n está en 0x178
    • La geometría del launch, como blockDim.x, y los valores ABI también están en el mismo banco
  • El cubin es un archivo ELF, el mismo formato de contenedor que un ejecutable de Linux
    • El fatbinary empaqueta juntos el cubin y el PTX
    • En esta RTX 4090 se ejecuta realmente el SASS, pero el PTX se incluye como fallback para que el driver pueda compilarlo con JIT en otras arquitecturas
    • Como el PTX es texto plano verboso, nvcc lo comprime por defecto

Cómo el código host prepara el launch

  • El frontend del compilador cudafe++ inserta un constructor oculto que se ejecuta antes de main
    • Ese constructor registra el fatbinary embebido en el runtime de CUDA
    • Conecta el puntero de función del lado host vadd con el nombre mangled del kernel device dentro del fatbin
  • La sintaxis vadd<<<4096, 256>>>(da, db, dc, n) se reemplaza por el host launch stub generado
    • da, db, dc, n se colocan alineados en el argument buffer de la memoria host, en los offsets 0, 8, 16, 24 respectivamente
    • Esos offsets corresponden a las posiciones 0x160, 0x168, 0x170, 0x178 que SASS lee desde constant bank 0
  • El stub llama a __cudaLaunch y le pasa la dirección de la función dummy vadd del lado host
    • Esa dirección no es una dirección de función para ejecutar en la CPU, sino que se usa como key para consultar la tabla de registro del runtime
    • El runtime encuentra el nombre del device symbol correspondiente y luego pasa al user-mode driver de código cerrado libcuda.so.1
  • En la primera llamada a la GPU, el runtime de CUDA abre dinámicamente libcuda.so.1 y crea un context
    • En strace se puede ver que se abre /lib/x86_64-linux-gnu/libcuda.so.1
    • El context incluye un channel mediante el cual la CPU se comunica con la GPU
  • Desde CUDA 12.2, la carga de módulos es lazy por defecto
    • La subida del cubin SASS se pospone hasta que se lanza por primera vez un kernel concreto
    • Se puede controlar con CUDA_MODULE_LOADING

La cola de comandos que entrega trabajo a la GPU

  • La GPU no recibe una llamada de función como la CPU ni hace jump a un entry point
    • Lee, a través del bus PCIe, un driver command stream ubicado en la memoria host
    • cuLaunchKernel coloca el launch command completo en ese stream y notifica a la GPU
  • En la primera ejecución, el driver copia el SASS del kernel a la memoria de la GPU
    • Asigna un code buffer y copia el SASS
  • El channel contiene dos estructuras clave ubicadas en la RAM host
    • pushbuffer: región de memoria donde el driver escribe los method, que son comandos para la GPU
    • GPFIFO: pointer ring buffer que apunta a spans del pushbuffer
  • Una entrada de GPFIFO se compone de dos words de 32 bits que representan (base, length) del span del pushbuffer
  • La GPU y el driver rastrean las posiciones de consumo y producción del trabajo con dos cursores
    • GP_GET: indica hasta dónde ha consumido la GPU
    • GP_PUT: indica hasta dónde ha producido el driver
    • Ambos están en una estructura por channel llamada USERD
  • Al lanzar un kernel, el driver escribe methods en un span del pushbuffer, hace que una entrada de GPFIFO apunte a ese span y luego avanza GP_PUT
  • En las GPU modernas, el host engine no vigila continuamente el cursor, por lo que se necesita un doorbell
    • La GPU mapea una pequeña ventana de registros en el process
    • El driver escribe el work-submit token del channel en el registro doorbell
    • Al recibir el doorbell, el host engine lee GP_PUT y trae por DMA la entrada GPFIFO y el span del pushbuffer

Información de ejecución contenida en el QMD

  • El launch comienza con un burst de methods SET_INLINE_QMD_ADDRESS_A/B y LOAD_INLINE_QMD_DATA
  • QMD(Queue Meta Data) es el launch descriptor del compute grid
    • Incluye los tamaños de grid y block: 4096, 256
    • Incluye la cantidad de registros por thread y los requisitos de shared memory
    • Incluye la dirección de inicio del programa y la dirección del constant bank que contiene los argumentos del kernel
    • También incluye la ubicación donde se notificará la finalización
  • Los argumentos empaquetados por el host stub son copiados por el driver al constant bank, y la dirección de ese banco se registra en el QMD
  • El QMD le indica a la GPU la ubicación del SASS, cómo está configurado el programa paralelo y dónde enviar la señal de finalización
  • cuLaunchKernel retorna en el momento en que se toca el doorbell
    • Como la llamada es asíncrona, la CPU puede seguir ejecutándose mientras la GPU trabaja

SM, warps y ocupación

  • El host engine entrega el QMD al compute work distributor
    • Este componente existe una sola vez en toda la GPU
    • Distribuye el stream lineal de instrucciones SASS entre los SM para ejecutarlo como un programa paralelo
  • La GPU objetivo, GeForce RTX 4090, usa 128 SM
    • El launch está compuesto por 4096 blocks y 256 threads por block
  • Cada SM tiene una local instruction cache, y los active warp mantienen un program counter
    • Desde Volta existe el modelo Independent Thread Scheduling, con program counter y call stack por thread
    • Aun así, el issue se realiza a nivel de warp
  • En el kernel de ejemplo, el resource limit determina la block residency
    • 256 threads = 8 warps por block
    • ptxas reserva 16 registros por thread
    • Por registros, serían posibles 16 blocks por SM
    • La capacidad de threads es de 1,536 active threads por SM, así que solo caben 1536 / 256 = 6 blocks
    • Por lo tanto, como máximo hay 6 blocks por SM, es decir, 48 warps en estado resident
  • Cada SM se divide en 4 processing blocks, o sub-partitions
    • Los 48 resident warp se distribuyen de manera uniforme entre las 4 sub-partitions
    • Cada warp scheduler administra 12 active warp cuando está lleno
    • En cada cycle elige un eligible warp y despacha la siguiente instrucción a 32 lanes

Condiciones para que un warp sea eligible

  • La GPU no extrae grandes cantidades de dependencias dinámicas de un solo thread como en la ejecución out-of-order de una CPU
    • Mantiene muchos resident warp y, cuando ocurre un stall, cambia a otro warp para ocultar la latencia
    • El compilador agenda los timings predecibles y el hardware scoreboard maneja las partes difíciles de predecir
  • Una instrucción SASS de 128 bits contiene el control-code payload escrito por ptxas
    • Para instrucciones de latencia fija incluye un static stall count
    • Un yield hint indica si se debe ceder la prioridad del scheduler
    • Las operaciones de latencia variable usan 6 barreras physical scoreboard por warp
  • En el tramo SASS del ejemplo, los dos LDG.E hacen set de la misma barrera scoreboard B2
    • FADD tiene B2 como wait-on
    • Hasta que las dos load vuelven y la barrera se limpia, ese warp queda en estado ineligible
    • Mientras tanto, el scheduler elige otros warp de la misma sub-partition
  • El tramo de FADD a STG.E se maneja como latencia fija
    • FADD tiene stall=5 y deja el warp estacionado unos cycles hasta que el resultado R9 esté listo
    • No hace falta una barrera separada
  • Este control payload queda oculto en la salida por defecto de nvdisasm
    • En el raw 128-bit encoding de cuobjdump -sass, está incluido en el segundo word de 64 bits
    • El layout no está documentado; se reconstruyó mediante microbenchmarking

Accesos a memoria y medición de rendimiento

  • Cuando un warp ejecuta LDG.E, los 32 threads calculan cada uno su dirección
    • En el ejemplo se accede a arreglos float consecutivos, así que todo el warp solicita un bloque contiguo de 32 * 4 = 128 bytes
  • La load/store unit del SM realiza request coalescing
    • Combina 32 solicitudes de 4 bytes en 4 sector request de 32 bytes
    • Si los accesos no fueran consecutivos, podría leer más datos de los necesarios
  • La solicitud coalesced primero revisa la L1 Data Cache local del SM
    • Si hay miss, va a través del crossbar interconnect hacia un slice de la L2 Cache de 72 MB
    • Si también hay miss en L2, pasa por el memory controller y el memory bus hasta la VRAM GDDR6X
  • El store STG.E sigue en principio el mismo camino en sentido contrario
  • Las mediciones de Nsight Compute muestran que este kernel está limitado por memoria
    • launch__grid_size: 4,096
    • launch__block_size: 256
    • launch__registers_per_thread: 16
    • launch__waves_per_multiprocessor: 5.33
    • sm__warps_active.avg.pct_of_peak: 82.77%
    • smsp__issue_active.avg.pct_of_peak: 5.17%
    • dram__throughput.avg.pct_of_peak: 79.65%
    • gpu__time_duration.sum: 10.78 μs
  • El kernel tiene una intensidad aritmética muy baja
    • Ejecuta 1 suma float por dos loads de 4 bytes y un store de 4 bytes, 12 bytes transferidos en total
    • Desde el punto de vista de lecturas DRAM, lee 8.4 MB en 10.78 μs, unos 780 GB/s, alrededor de 4/5 del pico
    • La salida c de 4 MB cabe en la L2 de 72 MB, así que no se hace flush a DRAM hasta que una copia device-to-host la lee

Cómo el resultado vuelve a la CPU

  • Como el kernel launch retorna a la CPU en el momento en que toca el doorbell, la GPU debe notificar la finalización por separado
  • Cuando los 4096 blocks se han retirado, la GPU postea el completion semaphore contenido en el QMD
    • El fence field del QMD está en los words 23–24
  • En el default stream, cudaMemcpy(c, dc, ...) queda ubicado después del kernel
    • El copy engine de la GPU queda gated hasta que se levanta el semaphore
    • Como c todavía está dirty en la L2 de 72 MB, la lectura del copy engine se resuelve desde L2 sin un viaje de ida y vuelta a DRAM
    • Los datos cruzan PCIe hacia la memoria host
  • Cuando termina la copia, el copy engine postea su propio semaphore
    • Finaliza la espera de cudaMemcpy en el host
    • c vuelve a ser memoria host normal
    • printf lee c[0] y c[n-1] desde RAM y los imprime en stdout

Cómo mirar dentro del launch

  • Leer los open kernel modules no basta para comprobar directamente algunas operaciones, porque libcuda es de código cerrado
  • La escritura de methods no pasa por syscalls; se escribe directamente en un buffer write-combined ya mapeado, así que para ver el pushbuffer hay que leer la memoria
  • Con un shim LD_PRELOAD que envuelva mmap, se pueden registrar las regiones mapeadas desde /dev/nvidia*
    • Si el programa de prueba llama a la función dump del shim justo después del launch, se puede imprimir el pushbuffer mapeado
    • El dump busca el method burst correspondiente a SET_INLINE_QMD_ADDRESS_A
  • El header de un pushbuffer method contiene opcode, payload count, subchannel index y register offset como bit fields
    • 0x0318 es SET_INLINE_QMD_ADDRESS_A
    • 0x0320 + i * 4 es LOAD_INLINE_QMD_DATA(i)
    • En el dump se ve un increasing-method burst con count 66, que incluye 2 address words y 64 QMD words, es decir, un QMD inline de 256 bytes en total
    • El word 12 dentro del QMD es 0x1000 y el word 18 es 0x100, correspondientes a los 4096 y 256 del launch
  • La configuración del driver se realiza mediante ioctl
    • En un programa con un solo kernel, strace registra 948 ioctl
    • La mayoría son configuración inicial de una sola vez
    • Los file descriptors principales son /dev/nvidiactl y /dev/nvidia-uvm
    • El magic byte de los ioctl del NVIDIA resource manager es 0x46, es decir, 'F'
    • El command number 0x2A se interpreta como NV_ESC_RM_CONTROL, y 0x2B como NV_ESC_RM_ALLOC
  • En vadd.cudafe1.stub.c, generado con nvcc --keep, también se puede ver el código de registro de startup
    • Una función con __attribute__((__constructor__)) se ejecuta antes de main
    • A través de __cudaRegisterBinary y __cudaRegisterEntry, se conecta el host function pointer vadd con el device entry point _Z4vaddPKfS0_Pfi

Lecturas relacionadas

1 comentarios

 
GN⁺ 3 시간 전
Opiniones en Hacker News

  • Fue un artículo interesante, y también me pareció divertida la explicación de los semáforos del stream predeterminado.
    Me gusta que CUDA maneje implícitamente la sincronización de comandos y permita usar comandos paralelos de forma opcional mediante streams.
    Contrasta con Vulkan, que desde el principio le pasa al usuario toda la complejidad de la sincronización.

  • Del lado del hardware hay algo de documentación pública.
    No hace falta leer el código fuente del kernel para encontrar la documentación de métodos o el formato QMD.
    Consultar https://github.com/NVIDIA/open-gpu-doc/blob/master/classes/c...

  • Fue muy útil.
    En particular, la parte de doorbell y QMD fue la más útil porque muestra cómo la sintaxis de ejecución de CUDA se conecta con lo que realmente se envía a la GPU.
    La mayoría de las explicaciones se detienen cerca de kernels, bloques y warps, pero este artículo hace mucho más fácil seguir la ruta CPU→driver→GPU.

  • El código de control es un poco más complejo de lo que describe el artículo.
    En realidad, se parece más a una consulta de tabla que a un bit dentro de la palabra de control.

  • Hoy hay empresas cuyo trabajo principal es optimizar kernels para que corran más rápido.
    Me pregunto si algún día esas empresas serán desplazadas por alguna biblioteca open source que haga esto muy bien.
    Nvidia podría sacar algo así en cualquier momento.
    O tal vez les vaya mejor si los grandes proveedores las adquieren para usarlo como moat que aumente la velocidad de inferencia.

    • A corto plazo, las adquisiciones por talento parecen bastante probables.
      Sin embargo, viendo cómo avanzan los modelos en benchmarks relacionados como kernelbench, creo que inevitablemente terminarán apareciendo soluciones más generalizadas.
      El problema es que cada nueva generación de hardware suele traer restricciones o funciones que los modelos existentes nunca habían visto.
      Por ejemplo, tcgen05 de Blackwell alguna vez fue un caso fuera de distribución.
      Si los modelos empiezan a generalizar mejor, quizá no sea una barrera fatal, pero al menos por ahora sigue siendo un obstáculo.
      [1] https://kernelbench.com/
    • Ejecutar CUDA a gran escala consume una cantidad asquerosamente grande de horas de ingeniería lidiando con bugs de drivers y bibliotecas de Nvidia.
      No he visto a mucha gente con ganas de depender más de las bibliotecas de Nvidia.
    • Probablemente no.
      Porque los detalles de la carga de trabajo —los parámetros exactos, la representación de los datos en memoria, los rangos de valores, etc.— hacen que las estrategias de optimización difieran mucho.

  • Acabo de terminar una maestría en HPC y tomé clases de CUDA, MPI+CUDA y OpenCL; creo que habría sido mucho más útil haber leído algo así antes de esas clases.
    En especial, me gustó el contexto alrededor de la parte sobre qué significa que un warp sea ejecutable.

  • Para empezar, es un buen artículo que explora muy bien varios rincones.
    Dicho eso, si no pasas por la runtime API de CUDA, desaparece gran parte del vudú en espacio de usuario.
    Si usas la Driver API, tomas el código fuente del kernel como string y lo compilas con el compilador en tiempo de ejecución de NVIDIA, puedes ver mejor qué ocurre.
    No todo, pero sí una buena parte, se vuelve transparente.
    Aquí hay una versión más “primitiva”:
    https://github.com/NVIDIA/cuda-samples/tree/master/cpp/0_Int...
    Si quieres ver lo mismo en una API moderna de C++, mucho más legible y aun así completamente transparente, mira esto:
    https://github.com/eyalroz/cuda-api-wrappers/blob/master/exa...
    Es un programa de ejemplo de mi biblioteca header-only CUDA API wrappers.

    • La Driver API es buena porque permite tratar los kernels CUDA como shaders con hot reload.
      Es divertido poder desarrollar cambiando el código mientras se ejecuta.

  • ¿En bare metal?