TL;DR desde el principio, para que pueda dejar el punto: si alguna vez ha querido crear usted mismo un tiempo de ejecución de inferencia LLM (empaque sus propios pesos, sea dueño de cada barrera, capture sus propios gráficos CUDA), así es como se ve ese viaje en una NVIDIA H100 (Hopper, sm_90) para Qwen2.5-Coder-7B. Saldrá con opiniones sólidas sobre la especialización warp, TMA, __syncthreads() y el hecho de que un gráfico CUDA no es una ventaja, es una necesidad. Este artículo es un recorrido por un pequeño tiempo de ejecución de LLM llamado annotated-llm-runtime (aproximadamente dos docenas de archivos fuente CUDA/C++, cada ruta activa comentada por qué, no solo qué) y las tres historias de guerra (figurativas) que produjeron la mayoría de las anotaciones.
Github Repo (se utilizará mucho más adelante a lo largo de la publicación): https://github.com/AnubhabBanerjee/Annotated-LLM-Runtime
1. ¿Por qué crear su propio tiempo de ejecución de LLM?
La mayor parte de la inferencia de LLM de producción ya pasa por una ruta muy bien probada: GGUF → llama.cpp → hecho. Funciona, es rápido, tiene un ejército de contribuyentes y es la herramienta adecuada para el envío. Déjame ser claro ahora: nada en esta publicación intenta desafiar eso, y te recomiendo que sigas usándolo en producción, como lo hago yo.
La propiedad de la pila de decodificación es donde está el valor. Si alguna vez necesita cambiar el formato de cuantización, agregar un sampler personalizado, conectar una variante de atención novedosa o implementar una nueva arquitectura, realmente no podrá hacerlo desde dentro de una caja negra. Necesita poder abrir cada núcleo y saber qué se romperá y qué ya se rompió y se solucionó.
Creo que en algún momento ya nos hemos hecho o nos haremos la siguiente pregunta:
¿Puedo escribir y decodificar yo mismo, comprender cada lanzamiento, ser dueño de cada barrera y aun así obtener los tokens correctos?
La respuesta corta es sí. Si alguna vez ha sentido curiosidad por saber cómo se une un decodificador real (archivos de peso empaquetado, descuantificación INT4, KV paginado, especialización de deformación, gráficos CUDA, TMA), pero se ha sentido demasiado intimidado para leer las cincuenta mil líneas de GGML para descubrirlo, este repositorio (y esta publicación) es el mismo viaje comprimido en aproximadamente dos docenas de archivos fuente CUDA/C++ con comentarios densos. Cada ruta activa en el código conlleva el por qué. Cada barrera dice qué hilos llegan a ella y cuáles no. Cada núcleo dice para qué forma fue sintonizado y hacia qué forma retrocedió silenciosamente. Es la versión del código base que desearía que existiera cuando comencé.
Algunas de las anotaciones existen porque tuve que convencerme a mí mismo para arreglar el código. Otros existen porque tuve que convencerme de una mala idea que ya había implementado. Este es el segundo tipo de comentario, en la naturaleza:
// ejercicios de ruta prmt.b32 Tuberías de permutación de enteros de tolva: línea de base antes de la transmisión de escala fusionada. // El mapa de carriles ValueShuffle está congelado para coincidir con el empaquetador fuera de línea; el tiempo de ejecución no puede reinterpretar los nibbles.
Ese comentario es el epitafio de una tarde perdida.
2. Dónde se encuentra actualmente la pila
Antes de repasar los aspectos internos, el estado actual honesto del tiempo de ejecución. Esto se mide en la misma clase H100, en el protocolo benchmark_e2e bloqueado en la configuración (batch=1, pp512/tg128, –warmup 3 –prompt-tokens 512 –gen-tokens 128):
Si está reconstruyendo esta pila usted mismo, la columna llama.cpp es una implementación de referencia para el mismo protocolo, fijada en la confirmación b3040 con los indicadores de descarga estándar (-ngl 99 -c 8192 -b 512 -cb) y los pesos oficiales Qwen/Qwen2.5-Coder-7B-Instruct-GGUF. La misma clase de GPU, el mismo modelo, la misma decodificación codiciosa, diferente cantidad de ajustes detrás de ellos. Es el criterio para comprobar que tu propia estructura se comporta de manera razonable, no un rival al que perseguir tan pronto.
Para los lectores con sabor a CUDA: la palanca más importante dentro de esta pila era la decodificación de gráficos CUDA. La ansiosa decodificación por token en los mismos núcleos llegó a alrededor de 119 ms/token. Envolver la decodificación de estado estable en un gráfico capturado y reproducirlo redujo la cifra a ~17 ms/token. Los mismos núcleos, pero enviados de manera diferente. Toda la historia sobre eso en la Sección 6.
3. La pila en una página
Si cierras esta pestaña después de esta sección, aún habrás obtenido la forma. Cada token de entrada recorre el siguiente camino, una vez por token generado, en el H100:
Un poco de fundamento, porque los números importan para lo que viene después. Desde csrc/common/qwen25_model_dims.h:
constexpr estático int kHiddenSize = kNumAttentionHeads * kHeadDim; constexpr estático int kIntermediateSize = 18944; constexpr estático int kKvProjectionDim = kNumKvHeads * kHeadDim; // … constexpr estática int kNumHiddenLayers = 28; constexpr estático int kVocabSize = 152064;
28 capas de decodificador, tamaño oculto de 3584, ancho de MLP 18944 y vocabulario 152064. Las dimensiones estándar de Qwen2.5-Coder-7B-Instruct se capturan dentro de constexpr para que el compilador de C++ las posea y nunca tenga que analizar un YAML en tiempo de ejecución.
Para atención de consultas agrupadas (GQA), desde csrc/common/paged_kv_layout.h:
// GQA: 28 cabezales de consulta comparten cabezales de 4 KV: proporción de 7:1 en CTA de atención de decodificación. constexpr estático int kNumAttentionHeads = 28; constexpr estática int kNumKvHeads = 4; constexpr estática int kHeadDim = 128;
28 cabezales de consulta, 4 cabezales KV: una proporción de GQA de 7:1. Ese único número determina mucho: dice que la huella de KV por token es pequeña (4 cabezas × 128 × 2 bytes = 1 KB por token), razón por la cual la paginación de KV a 16 tokens por página en FP16 aterriza en 4 KiB limpios por porción de cabeza de KV y 16 KiB completos por página física. Cada uno de esos números se elige para alinearse con la línea de caché L2 de 128 bytes de Hopper, y el archivo le informará sobre ello en los comentarios si lo olvida:
static_assert(kKvBytesPerPage % 128 == 0, "la página KV debe alinearse con líneas L2 de 128 bytes");
Los pesos son INT4 simétricos por grupo, tamaño de grupo 128, una escala FP16 por grupo por fila de salida, escala = absmax / 7.0, punto cero fijado en 0. Eso cubre exactamente siete proyecciones por capa: q_proj, k_proj, v_proj, o_proj, gate_proj, up_proj, down_proj. Todo lo demás, incluidos embed_tokens, all rms_norm y lm_head, permanecen sin editar en FP16. Este es un diseño deliberado por elección. Cuando aparece un error numérico (y créanme, al menos uno lo hará), querrá que el espacio de búsqueda sea pequeño.
INT4 está empaquetado en uint32 palabras, ocho mordiscos por palabra, en un diseño específico “ValueShuffle” que hará que una sección entera sea culpa de este artículo cuando lleguemos al error n.° 3. Por ahora, solo sepa esto: todo el motor trata un archivo .nanoqwen como un blob binario compatible con mmap con un encabezado de 256 bytes, una cadena mágica y cargas útiles INT4 empaquetadas alineadas con límites L2 de 128 bytes. No hay YAML en el camino activo. Alguna vez.
4. El archivo de peso de cinco líneas
El empaquetador fuera de línea produce un archivo .nanoqwen. El tiempo de ejecución abre ese archivo, verifica ocho bytes y se niega a hablar con él si esos ocho bytes son incorrectos.
// Bytes mágicos de NANOQWEN: 0x4E414E4F5157454E ("NANOQWEN") afirmados antes de cualquier carga de peso de VRAM. // Primeros 8 bytes de cada archivo .nanoqwen: los archivos corruptos fallan rápidamente antes que cudaMemcpy grande. constexpr estático char sin firmar kNanoqwenMagicBytes[8]= { 0x4E, 0x41, 0x4E, 0x4F, 0x51, 0x57, 0x45, 0x4E, };
El archivo simplemente comienza con la palabra "NANOQWEN". Ese es todo el control de seguridad. Si un archivo no comienza exactamente con esas letras, el sistema se niega a cargar seis gigabytes de datos desconocidos en la GPU. No es un diseño glamoroso, pero ya me ha salvado de dos grandes dolores de cabeza (incluida una combinación de código desordenada).
El resto del archivo está organizado en una cuadrícula estricta y predecible. Dado que el espaciado está codificado, tanto el código Python como el C++ saben exactamente dónde buscar datos sin necesidad de leer instrucciones de configuración adicionales. Cada parte del proyecto (la herramienta de preparación de datos, las pruebas de Python y el motor C++) utiliza exactamente este mismo mapa. Si no lo hacen, simplemente se niegan a comunicarse.
El resto de la configuración (de config/nanoqwen.yaml) se refleja en csrc/common/nanoqwen_constants.h en el momento de la compilación. Tamaño del grupo, identificación de diseño de ValueShuffle, ancho de línea de caché L2, etiquetas de tipo de almacenamiento para tensores INT4 vs FP16: todo constexpr int estático. La ruta activa en tiempo de ejecución nunca ve un analizador YAML, que es la forma más sencilla de asegurarse de que la ruta activa no se asigne accidentalmente en algún lugar del montón.
Todo lo que pasa hasta este punto es sólo la historia de lo que sucede cuando un uint32 lleno de ocho enteros de 4 bits con signo golpea los tubos aritméticos de un H100. Y, lamentablemente, la historia es más interesante de lo que debería ser.
5. Error n.º 1: __syncthreads() que se comió el softmax en las colas de la página KV
Aquí es donde cae la abstracción y muestra el trabajo real.
El núcleo de atención paginada decodifica la atención en un caché KV almacenado en páginas de 16 tokens, por cabeza KV. En el H100, el plan está de moda y está especializado en warp: un productor warp emite copias masivas de TMA de Hopper (cp.async.bulk) para sorber páginas K y V de HBM a SMEM, y seis warps consumidores esperan, luego calculan la acumulación softmax + V ponderada en línea en registros FP32. Las etapas SMEM con triple buffer le dan al productor un lugar para escribir mientras los consumidores todavía están masticando el mosaico anterior.
Si nunca has visto esta forma dibujada, esta es la imagen que el núcleo intenta ser:
Ahora aquí está el error en una oración: una versión anterior de este kernel tenía un __syncthreads() dentro de la rama if (warp_id == 0), la deformación del productor, y en ningún otro lugar. Ese es un comportamiento indefinido en CUDA. __syncthreads() es una barrera que abarca todo el bloque. Si sólo llega una disformidad, las otras seis disformidades se la saltan por completo y avanzan alegremente. En este núcleo, "marchar hacia adelante" significaba: los warps del consumidor comenzaron a leer filas K y V de SMEM antes de que el TMA masivo del productor hubiera conseguido los bytes.
Las páginas completas de K/V se veían bien; de todos modos, la copia generalmente terminaba antes de que llegaran los consumidores. Pero en el momento en que el caché KV terminó en una página final parcial (digamos, longitud de secuencia 49 → el último bloque contiene un token, no dieciséis), el tiempo cambió, los consumidores leyeron cualquier basura obsoleta que aún estuviera en SMEM de etapa 0, y softmax produjo tokens erróneos, coherentes y seguros.
El kernel actual hace lo seguro. El trabajo de alcance warp está delimitado con __syncwarp() dentro de la rama del productor. La valla que abarca todo el bloque, a la que debe llegar cada warp, se coloca fuera de cualquier rama de warp-id, inmediatamente antes de que los consumidores toquen la loseta:
if (warp_id == 0) { // … emite TMA para el mosaico de la página K en smem_k_tile … __syncwarp(); // … emite TMA para el mosaico de la página V en smem_v_tile … __syncwarp(); } // Valla de todo el bloque: las deformaciones del consumidor deben ver los rellenos SMEM del productor antes de las lecturas de punto/softmax. __syncthreads();
Ese comentario no es decoración. Es el tejido cicatricial de una dolorosa sesión de depuración.
Las llamadas a __syncwarp() funcionan porque solo necesitan que los 32 hilos de una única deformación estén de acuerdo. La llamada a __syncthreads() funciona porque cada deformación en el bloque lo golpea. Pero si coloca __syncthreads() dentro de una verificación if (warp_id == 0), la GPU se bloquea porque las otras deformaciones nunca alcanzan esa barrera. Ese error me costó tres horas mirando la cola de un KV.
Para hacer las cosas más seguras, la sincronización entre el productor y los consumidores utiliza barreras (objetos de barrera con reconocimiento de transacciones de Hopper) en lugar de __syncthreads(). El productor le dice al operador exactamente cuántos bytes debe esperar. Los consumidores esperan en la misma barrera y solo se despiertan cuando el hardware confirma que la transferencia de datos ha finalizado por completo.
La lección principal: en un kernel especializado en warp, poner __syncthreads() dentro de cualquier rama warp-id es un error garantizado, no una optimización. El cálculo de softmax en las colas de KV es exactamente donde esta condición de carrera bloqueará su código primero.
Para un ingeniero de HPC de carrera, este error es tan estúpido que raya en lo ofensivo. Pero como arquitecto de telecomunicaciones en transición a GPU básicas, estoy acostumbrado a redes en las que los paquetes caídos finalmente descubren un apretón de manos y encuentran el camino a casa. Resulta que los hilos CUDA son mucho menos indulgentes: no retransmiten; simplemente se quedan ahí sentados en silencio, convirtiendo su costosa GPU Hopper en un calentador muy ineficiente.
6. Error n.º 2: 119 → 17 ms por token, debido a un truco de envío
Si el primer error se refería a la corrección, el segundo se refiere a por qué nunca se debe confiar en un número de latencia por token medido con lanzamientos ansiosos.
Aquí está la línea de tiempo original. Antes de implementar gráficos CUDA, calcular un solo token significaba realizar aproximadamente 10 lanzamientos de kernel por capa en 28 capas, además de las operaciones lm_head y argmax. Eso es más de 280 lanzamientos de kernel individuales por un solo token. Cada lanzamiento obliga a un viaje de ida y vuelta al controlador de la CPU. Esos retrasos de microsegundos se acumulan en milisegundos y, con miles de tokens, el tiempo de ejecución total se convierte en una absoluta vergüenza.
La ejecución de una decodificación ansiosa en esta configuración registró una dolorosa velocidad de 119 ms/token. Cuando miré el desglose de ITL (latencia entre tokens), apenas se registró el tiempo real de ejecución de la GPU. El programador permaneció inactivo el 90% del tiempo porque el hardware necesitaba instrucciones. El cuello de botella nunca fueron las matemáticas; fue la enorme sobrecarga de la CPU del host causada por cudaLaunchKernel.
La solución son los gráficos capturados. Una vez que su paso de decodificación sea completamente determinista (usando exactamente los mismos núcleos y formas sin ramificaciones a nivel de Python), puede capturar la secuencia completa en un solo cudaGraphExec_t. Simplemente le entrega un comando al controlador CUDA para reproducir toda la estructura, evitando por completo la sobrecarga de lanzamiento por kernel. La implementación de esto redujo mi decodificación en estado estable de ~119 ms/token a ~17 ms/token. Esto supone una aceleración 7 veces mayor simplemente al enviar exactamente los mismos núcleos de manera más eficiente.
Hay un problema crucial. El proceso de decodificación KV paginado tiene un punto de bifurcación oculto dependiendo de si el token actual cruza el límite de una página de memoria (posición % 16 == 0). Dado que no se pueden poner dos formas diferentes en un gráfico estático, el tiempo de ejecución resuelve esto capturando previamente ambas variantes del gráfico durante la configuración y seleccionando dinámicamente la correcta para cada paso.
constexpr int kPageBoundaryCapturePosition = 512; constexpr int kNoPageCapturePosition = 513;
Dos posiciones, una en el límite de una página y otra no, capturadas consecutivamente durante el calentamiento:
cudaError_t page_error = capture_one_graph( &state->decode_graphs.exec_with_page_boundary, state, kPageBoundaryCapturePosition, true, stream); cudaError_t no_page_error = capture_one_graph( &state->decode_graphs.exec_no_page_boundary, state, kNoPageCapturePosition, false, stream);
Y en el momento de la repetición, el seleccionador dice una sola línea:
const bool page_boundary = (position_offset % kPaggedKvTokensPerBlock) == 0; cudaGraphExec_t exec = límite_página? estado->decode_graphs.exec_with_page_boundary : estado->decode_graphs.exec_no_page_boundary;
Dos gráficos, elegidos por un módulo. Ese es todo el truco. cudaGraphLaunch(exec, stream) y listo. Cada escalar que necesitan los núcleos (desplazamiento de la posición actual, longitud de la secuencia de ejecución, identificación del token actual) es un puntero de dispositivo que se parchea con un pequeño cudaMemcpyAsync antes del lanzamiento, por lo que no es necesario volver a crear una instancia del gráfico en cada paso.
Hay una estructura de desglose de ITL por token en csrc/runtime/decode_itl_breakdown.h que rastrea once regiones (qkv_gemv, append_tail_kv, paged_attention, o_proj, gate_up_gemv, swiglu_down, etc.) además de una región kernel_gap que mide específicamente la sobrecarga del host en torno a cudaGraphLaunch. Este último existe porque la “brecha del núcleo” es exactamente lo que eliminan los gráficos CUDA, y si el número alguna vez vuelve a aumentar, alguien reintrodujo accidentalmente la decodificación ansiosa en alguna parte.
Lección aprendida: si un paso de decodificación contiene cientos de lanzamientos de kernel y cada uno de ellos toca el controlador, en realidad no estás comparando tus kernels. Estás comparando cudaLaunchKernel. Los gráficos CUDA no son un ajuste de rendimiento. Son la diferencia entre medir su GPU y medir su controlador.
7. Error n.º 3: prmt.b32 superó a __dp4a y las activaciones de INT8 se perdieron de todos modos
Cada una de las siete proyecciones cuantificadas son pesos INT4 × activaciones FP16. Lo cual suena simple hasta que te das cuenta de que las dos formas competitivas de hacer esto en Hopper parecen muy similares en papel y muy diferentes en un generador de perfiles.
La ruta A era __dp4a. Hopper todavía tiene acumulación rápida de producto punto INT8: empaquete cuatro operandos INT8 firmados en un uint32, alimente dos de ellos a __dp4a y obtendrá una actualización del acumulador INT32. El problema es que ambos lados de ese producto escalar deben ser INT8. Entonces, la Ruta A requirió un núcleo adicional por paso de decodificación para cuantificar las activaciones de FP16 a INT8 antes de cada GEMV, además de una escala de activación por grupo para deshacer esa cuantificación dentro del acumulador.
La ruta B mantuvo las activaciones en FP16 y utilizó un descomprimido INT4 con sabor a prmt.b32 (la instrucción PTX de permutación de bytes de Hopper) con un FMA de multiplicación y acumulación de FP32. El desembalaje en sí es breve:
__device__ __forceinline__ unsigned int prmt_b32( unsigned int source_low, unsigned int source_high, unsigned int selector) { unsigned int resultado = 0; asm volatile("prmt.b32 %0, %1, %2, %3;" : "=r"(resultado): "r"(source_low), "r"(source_high), "r"(selector)); resultado de devolución; }
Seguido de un ayudante de extensión de signo para un nibble firmado, porque INT4 en empaquetamiento simétrico vive en [-8, 7] y no encaja de forma nativa en un cambio uint32:
__device__ __forceinline__ int sign_extend_int4_nibble(unsigned int nibble_unsigned) { unsigned int nibble = nibble_unsigned & 0xFu; if (mordisco y 0x8u) { return static_cast(mordisco | 0xFFFFFFF0u); } devolver static_cast(mordisco); }
Ese es el contrato de decodificación completo para un peso INT4: enmascarar a cuatro bits, extender el signo si el bit 3 está configurado, multiplicar por la escala FP16 del grupo, usarlo como operando en el acumulador FP32 de la fila.
Sobre el papel, la Ruta A debería haber ganado en gate_proj / up_proj / down_proj: las formas gigantes de MLP GEMV donde el ancho de banda de activación de HBM es el muro. Las activaciones INT8 ocupan la mitad del ancho de banda de las activaciones FP16, así que es fácil ganar, ¿verdad?
No en Hopper. No para estas formas. La ruta de activación prmt.b32 FP16 – Ruta B – venció a las activaciones __dp4a INT8 – Ruta A – en las formas exactas de puerta/arriba/abajo en las que se suponía que iba a perder. Entonces, se probó, se comparó y luego se eliminó la ruta de activación de INT8. El código lleva la lápida: el asistente __dp4a todavía está presente en csrc/kernels/fused_gemv.cu como b1_accumulate_packed_word_dp4a, pero la ruta de lanzamiento de envío nunca lo activa.
Hay un detalle relacionado escondido en el diseño del embalaje. Los ocho nibbles INT4 dentro de un uint32 empaquetado no viven en un orden ingenuo [0..7]:
// Coincide con config/nanoqwen.yaml pack_layout [w0,w2,w4,w6,w1,w3,w5,w7]. // El orden de los carriles intercala los mordiscos pares/impares para que surjan patrones de bytes fáciles de usar en Hopper. cambiar (lane_in_octet) {caso 0: devolver 0; caso 1: retorno 16; caso 2: retorno 4; caso 3: devolver 20; caso 4: retorno 8; caso 5: retorno 24; caso 6: retorno 12; caso 7: retorno 28; predeterminado: devolver 0; }
Los carriles pares van a los 16 bits bajos, los carriles impares a los 16 bits altos. Eso no es estético. Se trata del empaquetador fuera de línea que dispone los mordiscos para que una futura variante de este núcleo pueda permutarlos en un diseño de archivo de registro que combine muy bien con la ruta de datos de permutación de bytes de Hopper. El diseño se llama ValueShuffle y se encuentra en el ID de diseño 1; Si el tiempo de ejecución alguna vez ve un archivo .nanoqwen con una identificación de diseño diferente, se niega a cargar, porque el mapa nibble dentro del dequant de C++ está compilado; no se pueden reinterpretar silenciosamente los pesos empaquetados y esperar lo mejor.
Lección aprendida: en Hopper, “INT8 tiene menos bytes que FP16, por lo tanto, las activaciones de INT8 ganan en GEMV vinculado a la memoria” no es un teorema. Depende de la forma exacta del mosaico, la ocupación exacta de SM y si el descomprimido INT4 que reemplaza es lo suficientemente barato como para que el núcleo cuantitativo de activación adicional se convierta en el nuevo cuello de botella. En este proyecto así fue.
8. El marcador honesto
En este punto, quizás te preguntes cómo puedo demostrar que este motor personalizado funciona. Las reglas se escriben directamente en el archivo config/nanoqwen.yaml y el tiempo de ejecución las aplica a través de una estricta escalera de validación:
Objetivo principal: lograr más del 80 % de utilización del ancho de banda de memoria (MBU) en las operaciones matemáticas principales (el INT4 GEMV fusionado). La GPU H100 SXM tiene un enorme ancho de banda de memoria máximo de 3,35 TB/s; MBU es el porcentaje real de esa velocidad que logramos alcanzar durante la ejecución. Objetivo secundario: hacer coincidir exactamente el punto de referencia de codificación (Pase HumanEval@1) del formato Q4_K_M de llama.cpp. La misma decodificación codiciosa, las mismas indicaciones, las mismas respuestas. Objetivo terciario: lograr una coincidencia de token del 99,5 % o mejor con una simulación estándar de PyTorch INT4.
Las pruebas se dividen en tres etapas que deben pasarse en orden: pruebas unitarias (que prueban operaciones matemáticas individuales), pruebas de capa (que prueban un bloque decodificador completo contra PyTorch) y pruebas de gráficos (que generan tokens para 100 mensajes que coincidan con la simulación de PyTorch). Lo más importante es que intencionalmente no intento forzar una coincidencia perfecta bit por bit con el diseño llama.cpp. Su formato Q4_K_M y mis pesos simétricos INT4 son fundamentalmente diferentes a nivel matemático, y pretender que se puedan alinear perfectamente es como un proyecto como este comienza a mentirse a sí mismo.
Cuando el archivo README dice "corrección: sí", significa que el motor pasa con éxito la prueba final de nivel gráfico en un H100 real. Los puntos de referencia de velocidad (TTFT, ITL y tokens) son los recibos que mostré anteriormente. Alcanzar ese objetivo principal del 80% de MBU es la persecución en curso.
Hay una regla más crucial para la honestidad. La configuración se niega explícitamente a bloquear los relojes de la GPU (lock_clocks: false). Bloquear los relojes en un H100 requiere sudo nvidia-smi (acceso raíz). Mi filosofía aquí es simple: si necesita privilegios de root para reproducir un punto de referencia, su punto de referencia no es reproducible. Las velocidades que ve se basan en el comportamiento predeterminado del reloj en silicio real y funcional, no en un experimento de laboratorio con reloj bloqueado artificialmente estable.
9. Conclusión: lo que realmente construí y aprendí
Escribí un motor de inferencia CUDA personalizado desde cero para Qwen2.5-Coder-7B en NVIDIA H100 que omite por completo los marcos estándar para usar un formato de archivo personalizado mapeado en memoria (.nanoqwen), maneja sus propias operaciones matemáticas INT4, administra un caché KV paginado y ejecuta toda la generación de tokens de estado estable usando solo dos gráficos CUDA capturados.
En este momento, el motor alcanza un tiempo hasta el primer token (TTFT) de ~128 ms y genera tokens a ~16,7 ms/token (~60 tokens/segundo). Para el contexto, el estándar de oro absoluto, llama.cpp, ejecuta exactamente la misma prueba a ~43 ms TTFT y ~4,95 ms/token. Mi tiempo de ejecución no intenta superar a llama.cpp hoy; esos números son un criterio para mostrar exactamente cuánta optimización extrema implica un motor de producción de clase mundial.
Construir esto me enseñó tres lecciones importantes sobre la IA básica:
Los errores de barrera son asesinos silenciosos: poner __syncthreads() dentro de una rama especializada en warp es un error fatal. Dañará silenciosamente sus matemáticas (específicamente el softmax en las colas de la página KV). Debes usar __syncwarp() para las ramas, o mejor aún, confiar en las barreras de Hopper para manejar el tiempo de forma segura. Los gráficos CUDA son obligatorios, no opcionales: ejecutar este motor con lanzamientos estándar por kernel resultó en unos terribles 119 ms/token. La captura exacta de las mismas operaciones en gráficos CUDA redujo la latencia a 17 ms/token. Se trata de una aceleración de 7 veces simplemente por eliminar la sobrecarga de envío de la CPU, sin cambios en las matemáticas subyacentes. Los formatos matemáticos importan: en la arquitectura Hopper, el uso de trucos específicos de desempaquetado de INT4 con activaciones de FP16 simplemente superó el intento de forzar activaciones de INT8. Construí la ruta INT8, la probé y finalmente la eliminé.
En última instancia, este proyecto está diseñado para ser leído, no sólo ejecutado. Es una guía paso a paso de lo que sucede exactamente entre su modelo de IA y el bus de memoria física de la GPU. Si eso se convierte en el comienzo de su propia pila de decodificación, será un resultado aún mejor.
Descargo de responsabilidad: las ilustraciones de este artículo se generaron utilizando IA (Claude Opus 4.8). Son ilustrativos, no fotográficos, y cualquier etiqueta visible dentro de las imágenes está estilizada en lugar de autorizada; consulte el cuerpo del artículo y el código mismo para obtener nombres precisos de funciones, valores métricos y detalles de arquitectura.