Exprimiendo vLLM en AMD Strix Halo: 2,3 veces llama.cpp con carga

Por Mario CamposPublicado el
7 min de lecturaEN
Un pequeño cubo luminoso en lo alto de una colina bajo un cielo estrellado, del que sube un hilo de luz en espiral.

vLLM en la iGPU de Strix Halo, medido frente a llama.cpp con el mismo modelo de 35B: 2,1 veces más rápido por token con grafos, 2,3 veces el rendimiento con 16 peticiones a la vez, el arranque de 9 minutos a 97 segundos, y el único caso en el que llama.cpp sigue ganando.

Strix Halo, de AMD, junta una GPU integrada grande (gfx1151) con una memoria unificada que la GPU puede usar casi entera: ajustando el kernel de Linux, PyTorch ve 124 GB de memoria de GPU. Basta para servir un modelo de 35.000 millones de parámetros sin tarjeta dedicada.

El modelo es Qwen3.6-35B-A3B: una mezcla de expertos con 3.000 millones de parámetros activos por token y una arquitectura híbrida que combina atención normal con Gated DeltaNet, una atención lineal con estado recurrente. llama.cpp lo mueve bien. La pregunta era qué aporta vLLM en el mismo hardware, y qué hace falta para llegar ahí.

En resumen, con el mismo modelo y los mismos prompts:

  • 2,3 veces el rendimiento de llama.cpp con 16 peticiones a la vez (2,07 frente a 0,89 peticiones por segundo).
  • 2,1 veces más rápido por token generado con grafos en lugar del modo eager (39 frente a 81 ms).
  • El arranque, de 9 minutos a 97 segundos.
  • Respuestas idénticas para peticiones idénticas.
  • Y una excepción, para ser justos: con una sola petición larga, llama.cpp sigue siendo algo más rápido.

Lo conseguí con cuatro palancas, y ninguna toca el modelo ni sus kernels.

Palanca 1: un bucle de autotuning que duraba horas

En el primer arranque, vLLM se quedaba atascado en el calentamiento, con la GPU al 100 % y sin escribir nada en el log. py-spy dump (a través de podman unshare, porque dentro de un contenedor rootless no se permite ptrace) enseñó dónde estaba: _warmup_prefill_kernels → el autotuner de Triton → triton.testing.do_bench.

Por qué vLLM parecía colgado 1. Estimación: Mide 5 ejecuciones, → ≈ 0 ms en esta iGPU → 2. Repeticiones: rep ÷ estimación, → cientos de miles → 3. Cada repetición: Vacía un búfer de, 256 MiB y ejecuta, el kernel. × cientos de miles, por cada configuración. El arreglo: 3 de calentamiento + 10 medidas → autotuning en segundos. Por qué vLLM parecía colgado El do_bench de Triton con un kernel que termina casi al instante 1. Estimación Mide 5 ejecuciones → ≈ 0 ms en esta iGPU 2. Repeticiones rep ÷ estimación → cientos de miles 3. Cada repetición Vacía un búfer de 256 MiB y ejecuta el kernel × cientos de miles, por cada configuración El arreglo: 3 de calentamiento + 10 medidas → autotuning en segundos
Por qué vLLM parecía colgado 1. Estimación: Mide 5 ejecuciones, → ≈ 0 ms en esta iGPU → 2. Repeticiones: rep ÷ estimación, → cientos de miles → 3. Cada repetición: Vacía un búfer de, 256 MiB y ejecuta, el kernel. × cientos de miles, por cada configuración. El arreglo: 3 de calentamiento + 10 medidas → autotuning en segundos. Por qué vLLM parecía colgado El do_bench de Triton con un kernel que termina casi al instante 1. Estimación Mide 5 ejecuciones → ≈ 0 ms en esta iGPU 2. Repeticiones rep ÷ estimación → cientos de miles 3. Cada repetición Vacía un búfer de 256 MiB y ejecuta el kernel × cientos de miles, por cada configuración El arreglo: 3 de calentamiento + 10 medidas → autotuning en segundos

do_bench decide cuántas veces ejecutar un kernel midiendo cinco lanzamientos y dividiendo su presupuesto entre esa estimación, y antes de cada ejecución borra un búfer de 256 MiB. El calentamiento hace autotuning de kernels diminutos cuyo tiempo estimado en esta iGPU es prácticamente cero, así que la división da cientos de miles de repeticiones por cada configuración que se compara. Habría acabado, pero en horas.

Lo arregla sustituir do_bench desde un sitecustomize.py dentro de la imagen, que Python carga antes que cualquier otra cosa:

# do_bench de Triton calcula las repeticiones como rep / estimación. Con kernels
# diminutos la estimación sale ~0 ms en gfx1151 y se ejecuta cientos de miles de
# veces. Con 3 ejecuciones de calentamiento y 10 medidas basta para ordenarlos.
import triton.testing as tt

def do_bench(fn, warmup=25, rep=100, grad_to_none=None, quantiles=None, return_mode="mean", **_):
    import torch
    for _ in range(3):
        fn()
    torch.cuda.synchronize()
    times = []
    for _ in range(10):
        start, end = torch.cuda.Event(enable_timing=True), torch.cuda.Event(enable_timing=True)
        start.record(); fn(); end.record(); torch.cuda.synchronize()
        times.append(start.elapsed_time(end))
    return tt._summarize_statistics(times, quantiles, return_mode)

tt.do_bench = do_bench

Diez medidas sobran para ordenar las configuraciones del autotuning. El mismo autotuning tarda ahora segundos (uno de los kernels, merge_16x16_to_64x64_inverse_kernel, en 4,07 s).

Palanca 2: dejar de rehacer trabajo en cada arranque

Con el autotuning arreglado, el primer arranque tardaba 9 minutos y los siguientes entre 5 y 9. Había dos costes que se pagaban en cada arranque y se tiraban a la basura.

  • aiter compilaba su módulo principal cada vez. La librería de kernels de AMD compila module_aiter_core con hipcc la primera vez que se importa: unos 46 segundos, dentro del contenedor, que se perdían al pararlo. Ahora se compila una vez, al construir la imagen. La pega: aiter detecta la arquitectura de la GPU preguntando al dispositivo e ignora GPU_ARCHS, así que la imagen hay que construirla con la GPU a la vista (podman build --device /dev/kfd --device /dev/dri). 22 segundos, una sola vez.
  • Triton guardaba en caché los kernels, pero no los veredictos. Los kernels compilados estaban en caché, pero cada medida del autotuning se repetía en cada arranque. TRITON_CACHE_AUTOTUNING=1 guarda qué configuración ganó. Las cachés están fuera del contenedor: 239 MB de Triton y 642 MB de torch.compile.

Tiempo de arranque de vLLM Barras horizontales con escala lineal desde cero: antes, 540 s (9 min); 5–9 min, 9 min la primera vez. Después, 97 s. Desglose ampliado de los 97 s: pesos 33 s; compilación + autotuning + perfilado ~15 s; captura de grafos ~15 s; resto ~34 s. Tiempo de arranque de vLLM Qwen3.6-35B-A3B AWQ 4-bit, desde que se lanza hasta que está listo Escala: 0–540 s Antes: en cada arranque 9 min 5–9 min, 9 min la primera vez Después: aiter precompilado + autotuning de Triton en caché 97 s Desglose de los 97 s (ampliado) 1 2 3 4 1. pesos 33 s 2. compilación + autotuning + perfilado ~15 s 3. captura de grafos ~15 s 4. resto ~34 s Medido en un AMD Strix Halo (gfx1151)
Tiempo de arranque de vLLM Barras horizontales con escala lineal desde cero: antes, 540 s (9 min); 5–9 min, 9 min la primera vez. Después, 97 s. Desglose ampliado de los 97 s: pesos 33 s; compilación + autotuning + perfilado ~15 s; captura de grafos ~15 s; resto ~34 s. Tiempo de arranque de vLLM Qwen3.6-35B-A3B AWQ 4-bit, desde que se lanza hasta que está listo Escala: 0–540 s Antes: en cada arranque 9 min 5–9 min, 9 min la primera vez Después: aiter precompilado + autotuning de Triton en caché 97 s Desglose de los 97 s pesos 33 s compilación + autotuning + perfilado ~15 s captura de grafos ~15 s resto ~34 s Medido en un AMD Strix Halo (gfx1151)

Listo a los 97 segundos: pesos en 33 s, compilación, autotuning y perfilado en unos 15 s, captura de grafos en unos 15 s. La imagen final pesa 14,7 GB, frente a los 61,8 GB de la oficial de AMD.

Palanca 3: grafos en lugar de eager, 2,1 veces más rápido por token

El consejo habitual para vLLM en iGPU es --enforce-eager. En modo eager cada token generado costaba 81 ms, cinco veces lo que necesita llama.cpp en el mismo hardware: Python lanza uno a uno los kernels de cada capa para cada token, y con solo 3.000 millones de parámetros activos esa sobrecarga es casi todo el trabajo.

Sin --enforce-eager, vLLM compila el modelo (44 s la primera vez, después en caché) y captura siete grafos por tramos, así que cada paso reproduce un grafo grabado en lugar de lanzar cada kernel por separado.

Tiempo por token generado Barras con una escala lineal común desde cero. Eager (--enforce-eager): 81 ms; Grafos (piecewise): 39 ms. 2,1× más rápido. Tiempo por token generado vLLM, mismo modelo; menos es mejor Escala: 0–90 ms Eager (--enforce-eager) 81 ms Grafos (piecewise) 39 ms 2,1× más rápido Medido en un AMD Strix Halo (gfx1151)
Tiempo por token generado Barras con una escala lineal común desde cero. Eager (--enforce-eager): 81 ms; Grafos (piecewise): 39 ms. 2,1× más rápido. Tiempo por token generado vLLM, mismo modelo; menos es mejor Escala: 0–90 ms Eager (--enforce-eager) 81 ms Grafos (piecewise) 39 ms 2,1× más rápido Medido en un AMD Strix Halo (gfx1151)

39 ms por token. Las respuestas en mi conjunto de evaluación no cambiaron, y una petición repetida seguía saliendo idéntica, bit a bit.

Palanca 4: lotes más grandes, 2,3 veces llama.cpp con carga

El rendimiento con concurrencia es donde vLLM se gana su sitio. llama.cpp atiende un número fijo de slots y, en Vulkan, no procesa juntos los prefill de varias secuencias. vLLM mete todas las peticiones en curso en cada pasada del modelo.

La última palanca fue --max-num-batched-tokens: hasta 8.192 tokens por paso en lugar de 2.048. Con 8 peticiones a la vez el rendimiento pasó de 0,98 a 1,18 peticiones por segundo; con 16, de 1,77 a 2,07. La latencia de una petición suelta no empeoró.

Rendimiento con 16 peticiones a la vez Barras con una escala lineal común desde cero. llama.cpp (Q4, 4 slots): 0,89 pet/s; vLLM, eager: 1,56 pet/s; vLLM, grafos: 1,77 pet/s; vLLM, grafos + lotes de 8.192 tokens: 2,07 pet/s. 2,3× llama.cpp. Rendimiento con 16 peticiones a la vez Peticiones por segundo; más es mejor Escala: 0–2,2 peticiones/s llama.cpp (Q4, 4 slots) 0,89 pet/s vLLM, eager 1,56 pet/s vLLM, grafos 1,77 pet/s vLLM, grafos + lotes de 8.192 tokens 2,07 pet/s 2,3× llama.cpp Medido en un AMD Strix Halo (gfx1151)
Rendimiento con 16 peticiones a la vez Barras con una escala lineal común desde cero. llama.cpp (Q4, 4 slots): 0,89 pet/s; vLLM, eager: 1,56 pet/s; vLLM, grafos: 1,77 pet/s; vLLM, grafos + lotes de 8.192 tokens: 2,07 pet/s. 2,3× llama.cpp. Rendimiento con 16 peticiones a la vez Peticiones por segundo; más es mejor Escala: 0–2,2 peticiones/s llama.cpp (Q4, 4 slots) 0,89 pet/s vLLM, eager 1,56 pet/s vLLM, grafos 1,77 pet/s vLLM, grafos + lotes de 8.192 tokens 2,07 pet/s 2,3× llama.cpp Medido en un AMD Strix Halo (gfx1151)

2,07 peticiones por segundo con 16 a la vez: 2,3 veces llama.cpp con el mismo modelo (llama.cpp con un GGUF Q4 y 4 slots). Cada palanca suma: vLLM en eager ya gana a llama.cpp con carga, los grafos añaden un 13 % y los lotes más grandes, otro 17 %.

Una petición cada vez: donde llama.cpp sigue ganando

Una petición cada vez Barras agrupadas, escala lineal común de 0 a 2.000 ms. llama.cpp (Q4, 4 slots) primero; vLLM (AWQ, grafos, lote de 8.192) segundo. Petición corta: 876 ms / 629 ms; Media: 1.377 ms / 1.353 ms; Larga: 1.748 ms / 1.868 ms. Una petición cada vez Latencia en ms según el tamaño de la petición; menos es mejor llama.cpp (Q4, 4 slots) vLLM (AWQ, grafos, lote de 8.192) Escala: 0–2.000 ms Petición corta llama.cpp 876 ms vLLM 629 ms Media llama.cpp 1.377 ms vLLM 1.353 ms Larga llama.cpp 1.748 ms vLLM 1.868 ms Medido en un AMD Strix Halo (gfx1151)
Una petición cada vez Barras agrupadas, escala lineal común de 0 a 2.000 ms. llama.cpp (Q4, 4 slots) primero; vLLM (AWQ, grafos, lote de 8.192) segundo. Petición corta: 876 ms / 629 ms; Media: 1.377 ms / 1.353 ms; Larga: 1.748 ms / 1.868 ms. Una petición cada vez Latencia en ms según el tamaño de la petición; menos es mejor llama.cpp (Q4, 4 slots) vLLM (AWQ, grafos, lote de 8.192) Escala: 0–2.000 ms Petición corta llama.cpp 876 ms vLLM 629 ms Media llama.cpp 1.377 ms vLLM 1.353 ms Larga llama.cpp 1.748 ms vLLM 1.868 ms Medido en un AMD Strix Halo (gfx1151)

Con una sola petición, vLLM gana en los prompts cortos, empatan en los medios y llama.cpp sigue siendo algo más rápido en los largos. En solitario, la iGPU está limitada por el ancho de banda de memoria, y los kernels Q4 de llama.cpp están afinados justo para eso. Con concurrencia, agrupar peticiones gana de largo. Si tu tráfico es un usuario cada vez con prompts largos, llama.cpp sigue siendo una gran opción; si son muchos usuarios, vLLM.

Misma petición, misma respuesta

La cifra que no cabe en un gráfico de rendimiento: en vLLM, la misma petición repetida da la misma salida, bit a bit. En llama.cpp, con esta familia de modelos híbridos, el mismo prompt a temperatura 0 se movía hasta 0,44 en log-probabilidad entre ejecuciones. Para cualquier cosa que lea log-probabilidades, o que tenga que dar la misma respuesta dos veces, eso basta para decidir. Por qué varía llama.cpp da para un post propio.

Lo que no compensó

  • MTP (decodificación especulativa con la cabeza de borrador del propio modelo). Se aceptaba el 76 % de los tokens del borrador, con log-probabilidades completas, pero los aciertos de la caché de prefijo caían a cero y esta build no trae el kernel fusionado que verifica el borrador, así que ese paso se hace con Triton. Con 16 peticiones a la vez, el rendimiento bajaba a 0,80 por segundo. De momento, no.
  • Más reutilización de la caché de prefijo a base de ajustes. En un modelo híbrido, vLLM iguala el tamaño de página de la caché de atención al del estado recurrente, lo que fija los bloques de atención en 1.056 tokens. Los prompts de unos 1.200 tokens rara vez reutilizan más de un bloque: solo el 28 % de los tokens sale de la caché. --mamba-block-size no lo cambia.

Conclusiones

  • Antes de culpar al hardware, mira dónde está atascado. Un proceso al 100 % y sin logs no está colgado, está ocupado en algo, y py-spy dump te dice en qué con un solo comando.
  • Pon un suelo a los bucles de medición. Cualquier cosa que divida un presupuesto entre un tiempo medido explota cuando la medida es casi cero.
  • Guarda en caché los veredictos, no solo los artefactos. Los resultados del autotuning y la build de aiter eran decenas de segundos en cada arranque.
  • Prueba los grafos aunque lo que se recomiende sea eager. Aquí partieron por la mitad el tiempo por token.
  • Mide la concurrencia y la petición suelta por separado. Cada una tiene su ganador.
  • El tamaño de lote es una palanca de rendimiento. 8.192 tokens por paso dieron un 17 % más con carga, sin coste.

Con cuatro palancas y sin tocar el modelo, vLLM en la Strix Halo pasó de no terminar nunca el calentamiento a dar 2,3 veces el rendimiento de llama.cpp, ir el doble de rápido por token y estar listo en 97 segundos.