En el capítulo anterior establecimos el modelo mental por capas de vLLM v1, y sabemos que una solicitud parte del API Server, atraviesa EngineCore y finalmente llega al Worker para su ejecución. Pero, ¿cómo se transforma la cadena JSON del cuerpo de una solicitud HTTP en un objeto que el motor interno puede programar, rastrear e interrumpir? Esta es la pregunta que la clase Request debe responder.

El sistema de especificaciones de KV Cache: de KVCacheSpec al registro

Request resuelve el problema de «quién debe calcular», mientras queKVCacheSpecresuelve el problema de «dónde calcular». En el mundo de PagedAttention, el KV cache de cada capa del modelo necesita describirse con precisión: cuántos heads tiene, cuánto ocupa cada head, cuántos tokens puede almacenar un bloque, si necesita cuantización. Esta información se codifica enKVCacheSpecel sistema de herencia de

Modelo intuitivo: KVCacheSpec es el «plano de distribución» de la memoria de video

Sin este sistema de especificaciones, la asignación de KV cache solo podría depender de suposiciones codificadas de forma rígida, incapaz de soportar las diversas necesidades de modelos que van desde MHA estándar hasta MLA, desde atención completa hasta ventana deslizante, y desde FP16 hasta cuantización FP8.

Estructura de datos: el árbol de herencia de KVCacheSpec y sus campos clave

KVCacheSpeces la clase base de todas las especificaciones; es una@dataclass(frozen=True) 📎 vllm/v1/kv_cache_interface.py:150-152. frozen significa que el objeto de especificación es inmutable una vez creado; esto garantiza que múltiples componentes (el planificador, el Worker, el KV Cache Manager) vean la misma especificación y no se produzcan inconsistencias por modificaciones en algún lugar.

La clase base define tres propiedades abstractas que deben ser implementadas por las subclases:num_heads、tokens_per_state、state_content_size_bytes 📎 vllm/v1/kv_cache_interface.py:182-183. Estas tres propiedades determinan conjuntamentepage_size_bytes, es decir, el número de bytes que ocupa un block.

AttentionSpeces la subclase más central; introducenum_kv_heads、head_size、dtype、kv_quant_modey otros campos📎 vllm/v1/kv_cache_interface.py:485-498. Entre ellos, el diseño del campotokens_per_statees especialmente ingenioso: su valor predeterminado es 1, lo que significa que un state corresponde a un token; pero puede establecerse como un entero mayor que 1 (como en el MLA disperso de DeepSeek-V4, que comprime múltiples tokens en un state), o como una fracción menor que 1 (como en el block pooling de Whisper, que usaFraction(1, block_pool_size)para indicar que un token corresponde a múltiples states)📎 vllm/v1/kv_cache_interface.py:501-501。

FullAttentionSpecsobre la base deAttentionSpecañadesliding_windowyattention_chunk_size 📎 vllm/v1/kv_cache_interface.py:566-566. Nótese que su docstring explica una decisión de diseño importante: cuando el asignador híbrido está deshabilitado, las capas de atención de ventana deslizante se tratan como atención completa en el KV Cache Manager (se asignan blocks para todos los tokens), pero en tiempo de ejecución del modelo todavía se calculan según la ventana deslizante📎 vllm/v1/kv_cache_interface.py:540-545. Esta es unaasignación conservadora, cálculo precisoestrategia.

MLAAttentionSpeces la especificación clave de la serie de modelos DeepSeek. Establecehead_size_vpor defecto en 0📎 vllm/v1/kv_cache_interface.py:670, porque MLA solo almacena un latent vector y no tiene una V independiente.alignmentEl campo se utiliza para el relleno de alineación de página📎 vllm/v1/kv_cache_interface.py:646-652, lo cual es crucial para backends como FlashMLA que requieren una alineación específica.

MambaSpecen cambio, no sigue en absoluto la ruta de attention. Utilizashapesydtypestuplas para describir la forma del tensor de estado📎 vllm/v1/kv_cache_interface.py:1027-1028，state_content_size_byteses la suma de los tamaños de todos los tensores de estado📎 vllm/v1/kv_cache_interface.py:1048-1052. Elmax_memory_usage_bytesde Mamba tiene tres formas diferentes de cálculo segúnmamba_cache_mode📎 vllm/v1/kv_cache_interface.py:1073-1084, lo que refleja la complejidad de la gestión de estado de Mamba: no crece linealmente como attention, sino que tiene un tamaño de estado fijo.

Impulsado por escenarios: la conversión de especificaciones a diseño de memoria de video

Cuando el motor se inicia, necesita convertir elKVCacheSpecde todas las capas en un diseño real de memoria de video. Este proceso lo realizanKVCacheTensorycreate_kv_cache_views.

KVCacheTensordescribe la posición de un grupo de capas de la misma forma en la asignación de KV cache📎 vllm/v1/kv_cache_interface.py:1406-1427. Sus campos centrales sonlayer_strideyblock_stride: el primero es la distancia en bytes entre capas adyacentes, y el segundo es la distancia en bytes entre blocks adyacentes. El docstring explica en detalle dos modos de diseño: el diseño con capas en el exterior (layer-outermost) da a cada capa una región contigua, y el diseño con blocks en el exterior (block-outermost) hace que cada block contenga las páginas de todas las capas📎 vllm/v1/kv_cache_interface.py:1416-1416。

flowchart LR
    subgraph spec["KVCacheSpec 层"]
        fas["FullAttentionSpec<br/>num_kv_heads=32<br/>head_size=128<br/>block_size=16"]
    end
    subgraph tensor["KVCacheTensor 层"]
        kt["KVCacheTensor<br/>size=2GB<br/>layer_stride=page*num_blocks<br/>block_stride=page"]
    end
    subgraph view["torch.Tensor 视图"]
        v1["layer_0: [B, H, N, C]"]
        v2["layer_1: [B, H, N, C]"]
        v3["layer_N: [B, H, N, C]"]
    end
    fas -->|"compute_layer_kv_cache_shape_bytes()"| kt
    kt -->|"create_kv_cache_views()"| v1
    kt -->|"create_kv_cache_views()"| v2
    kt -->|"create_kv_cache_views()"| v3

create_kv_cache_viewsLa función es el núcleo de este proceso📎 vllm/v1/kv_cache_interface.py:353-417. Recibe un buffer plano de int8 y, mediantetorch.as_strided, crea una vista 4D para cada capa[B, H, N, C]. El parámetro clave esstrides, que se calcula a partir decompute_layout_strides📎 vllm/v1/kv_cache_interface.py:314-350. Esta función calcula los pasos en bytes de cada dimensión en orden inverso, comenzando desde la dimensión más interna, según el orden de dimensiones especificado porlayout.stride_order.

Aquí hay una verificación de límites que vale la pena señalar: cuando kernel_block_size es menor que spec.block_size (es decir, un manager block se divide en múltiples kernel blocks), el código verifica si block_stride es igual a dense_page_size📎 vllm/v1/kv_cache_interface.py:381-382. Si no son iguales, significa que hay padding en el diseño y no se puede dividir de manera uniforme; en ese caso se lanza un ValueError con una sugerencia de corrección explícita.

Reflexiones de diseño: el patrón de registro y la extensibilidad

KVCacheSpecRegistryes un diseño clave para la extensibilidad de vLLM📎 vllm/v1/kv_cache_spec_registry.py:39-40. Mantiene dos diccionarios globales:_REGISTRY_KVCACHESPEC_LISTalmacena la asignación de clases spec a metadatos,_REGISTRY_ROLE_MANAGERSalmacena la asignación de roles a managers📎 vllm/v1/kv_cache_spec_registry.py:35-36。

get_manager_classEl método muestra la lógica central de búsqueda del registro: recorre hacia arriba el MRO (orden de resolución de métodos) de la clase spec y encuentra la primera clase base registrada📎 vllm/v1/kv_cache_spec_registry.py:129-130. Esto significa que unCustomFullAttentionSpecpersonalizado, si no se registra por separado, heredará automáticamente el manager deFullAttentionSpec. Estabúsqueda basada en herenciahace que al añadir un nuevo tipo de spec solo sea necesario registrar la parte diferencial.

check_kv_cache_spec_registryEl método valida en el arranque que los specs de todas las capas estén registrados📎 vllm/v1/kv_cache_spec_registry.py:165-174. Nótese que usaraise ValueErroren lugar deassert, y el comentario indica explícitamente que esto es para que también tenga efecto en entornos de producción📎 vllm/v1/kv_cache_spec_registry.py:165-174. Esta es una decisión de ingeniería importante: el flag-Ode Python elimina los assert, pero los errores de configuración en producción deben exponerse al arrancar, no provocar un fallo en tiempo de ejecución.

Resumen del capítulo

Este capítulo analiza las dos estructuras de datos centrales de vLLM v1.RequestEs el vehículo del ciclo de vida de la solicitud dentro del motor; mediante la lista doble de tokens, el contador de programación asíncrona y el mecanismo de block hash, sustenta las dos funcionalidades centrales: el procesamiento por lotes continuo y la caché de prefijos.KVCacheSpecy su jerarquía de herencia definen las especificaciones de diseño de memoria de la KV cache, desde el estándarFullAttentionSpechastaMLAAttentionSpec、MambaSpec, cubriendo necesidades diversas de arquitecturas de modelos. El patrón de registro permite añadir nuevos tipos de spec sin modificar el código central, garantizando la extensibilidad del sistema.

Hasta aquí, hemos visto cómo Request se transforma desde EngineCoreRequest y cómo sustenta las decisiones de programación mediante contadores de estado, block hash y otros mecanismos. Pero, ¿cómo atraviesa realmente una solicitud externa el API Server, el chat template y el procesamiento multimodal hasta convertirse finalmente en EngineCoreRequest? El siguiente capítulo entra en la capa de entrada de solicitudes, rastreando por completo esta ruta desde HTTP/CLI hasta EngineCore.

En el capítulo anterior analizamos las dos estructuras de datos centrales del motor: Request y KVCacheSpec, y comprendimos cómo se desacopla la secuencia lógica de los bloques físicos de memoria. Pero, ¿cómo atraviesa realmente un cuerpo de solicitud HTTP o una cadena de Python el API Server, el chat template y el procesamiento multimodal hasta convertirse finalmente en EngineCoreRequest? Este capítulo rastreará por completo esta ruta y revelará cómo las tres vías de entrada —CLI síncrona, API asíncrona y clase LLM offline— convergen en el mismo núcleo del motor.

3.1 El punto de convergencia de las tres vías de entrada: AsyncLLMEngine y LLMEngine

Antes de profundizar en el análisis de solicitudes, es necesario ver claramente la topología de las tres vías de entrada. vLLM ofrece tres formas de uso:vllm serveel servicio HTTP compatible con OpenAI iniciado mediante, la herramienta de línea de comandosvllm, y la instanciación directa en Python de la claseLLMpara inferencia offline. Parecen independientes, pero en realidad comparten el mismo núcleo del motor.

Veamos primero el mecanismo de alias de la vía de API asíncrona.

📎 vllm/engine/async_llm_engine.py:7-7

Este archivo es tan corto que casi no parece un módulo: solo hace una cosa: apuntar el aliasAsyncLLMEnginehaciavllm.v1.engine.async_llm.AsyncLLM. Esta es una huella típica de migración arquitectónica. En la era de vLLM v0,AsyncLLMEngineera una clase enorme y compleja; tras la reescritura de la arquitectura v1, la nuevaAsyncLLMasumió las mismas responsabilidades. Para no romper el código de usuario existente, vLLM conservó la ruta del módulo antiguo como capa de compatibilidad.

Veamos ahora la entrada de la vía offline.

📎 vllm/entrypoints/llm.py:344-346

LLM.__init__finalmente llama aLLMEngine.from_engine_args, pasandoUsageContext.LLM_CLASS. Esta enumeraciónUsageContextes clave para distinguir las vías de entrada: permite al motor saber si se ejecuta en modo de procesamiento por lotes offline o en modo de servicio en línea, ajustando así las estrategias de registro, métricas y gestión de recursos.

📎 vllm/entrypoints/llm.py:357-359

Nótese aquí la asignación deself.renderer = self.llm_engine.rendereryself.input_processor = self.llm_engine.input_processor. La clase offlineLLMno implementa por sí misma el renderizado del chat template, sino que reutiliza elrendererinterno del motor. Esto significa que la lógica de análisis del chat template es el mismo código tanto en la vía offline como en la online, solo cambia el momento de invocación.

La relación de convergencia de las tres vías puede representarse con el siguiente diagrama de flujo de datos.

flowchart LR
    subgraph entry["入口层"]
        http["HTTP 请求体<br/>ChatCompletionRequest"]
        cli["CLI 参数<br/>vllm serve / vllm chat"]
        offline["Python 调用<br/>LLM.chat(messages)"]
    end

    subgraph parse["解析层"]
        chat_utils["chat_utils.parse_chat_messages<br/>-> ConversationMessage + mm_data"]
        renderer["renderer<br/>apply_chat_template -> token_ids"]
    end

    subgraph engine["引擎层"]
        async_llm["AsyncLLM<br/>add_request()"]
        llm_engine["LLMEngine<br/>add_request()"]
        core["EngineCore<br/>input_queue"]
    end

    http --> chat_utils
    cli --> chat_utils
    offline --> chat_utils
    chat_utils --> renderer
    renderer --> async_llm
    renderer --> llm_engine
    async_llm --> core
    llm_engine --> core

Este diagrama revela un diseño clave: sin importar si la solicitud proviene de HTTP, CLI o Python,chat_utilses el único punto de entrada para el procesamiento multimodal y del chat template. Unifica los formatos de entrada heterogéneos en una listaConversationMessagemásMultiModalDataDict, y luego los entrega al renderer para generar la secuencia de tokens.

3.2 chat_utils: de mensajes heterogéneos a estructura de conversación unificada

chat_utils.pyes el módulo más complejo de toda la capa de entrada de solicitudes; sus 2264 líneas de código manejan el formato compatible con OpenAI, extensiones personalizadas, incrustaciones multimodales, llamadas a herramientas y todas las formas de entrada. Su responsabilidad central puede resumirse en una frase: normalizar cualquier lista de mensajes enviada por el usuario en una listaConversationMessageque el chat template pueda entender, extrayendo simultáneamente los datos multimodales a unMultiModalDataDictindependiente.

Modelo intuitivo: traductor y clasificador de equipaje

Imaginachat_utilscomo el traductor y clasificador de equipaje de un aeropuerto. Los pasajeros (usuarios) vienen de diferentes países (formato OpenAI, formato personalizado, formato Harmony) y hablan idiomas distintos. El traductor primero traduce lo que dicen todos a un idioma de trabajo común (ConversationMessage), al mismo tiempo, clasificar el equipaje facturado de los pasajeros (imágenes, audio, video) en cintas transportadoras independientes (MultiModalDataDict), pegar etiquetas (UUID), y finalmente enviar tanto a las personas como al equipaje al mismo avión (motor).

Sin esta capa, el motor tendría que comprender los detalles de cada formato de entrada, la lógica de extracción de datos multimodales se dispersaría en cada punto de entrada, y cualquier nuevo formato requeriría modificar el núcleo del motor.

Estructura de datos: colaboración dual entre rastreador y analizador

chat_utilsEl núcleo de consiste en la colaboración de dos grupos de clases:BaseMultiModalItemTrackery sus subclases se encargan de «rastrear» los elementos multimodales,BaseMultiModalContentParsery sus subclases se encargan de «analizar» las partes de contenido.

Primero veamos la disposición de campos del rastreador.

📎 vllm/entrypoints/chat_utils.py:598-601

_items_by_modalityes undefaultdict[str, list[_T]], que almacena los elementos pendientes agrupados por modalidad (image, audio, video, etc.)._modality_orderse dedica específicamente avision_chunkregistrar la modalidad original de cada chunk (image o video), porque el modelo unificado de chunk visual mapea ambos avision_chunk, pero el procesamiento posterior necesita conocer el tipo original.

📎 vllm/entrypoints/chat_utils.py:613-615

use_unified_vision_chunk_modalityes uncached_property, que lee el indicadoruse_unified_vision_chunkdesde la configuración de HuggingFace. Se usacached_propertyen lugar de una propiedad normal porque esta verificación se activa en cada llamada aadd, y el caché evita la sobrecarga repetida degetattr.

El métodoadddel rastreador es el punto de entrada principal.

📎 vllm/entrypoints/chat_utils.py:656-684

addEl método primero llama a_validate_addpara realizar la validación, y luego almacena los elementos bajo diferentes claves según si se usa la modalidad unificada de chunk visual. Nótese el manejo especial deprompt_embeds: se añade directamente a_items_by_modality["prompt_embeds"]y devuelveNone, porque los embeddings precalculados no pasan por el HF processor y no tienen cadena de marcador de posición.

_validate_addLa lógica de validación en

📎 vllm/entrypoints/chat_utils.py:686-721

merece un examen detallado. Aquí hay una rama sutil: cuandoenable_mm_embeds=Truey el límite por prompt de esa modalidad es 0 y la modalidad original termina en_embeds, se omite la validación de cantidad. Esto es para permitir que las entradas de embeddings eludan el límite de cantidad de la modalidad original — los embeddings son precalculados y no consumen recursos de procesamiento de la modalidad original.

Impulsado por escenarios: cómo se analiza una solicitud de chat con imágenes

Supongamos que el usuario envía una solicitud de chat que contiene una URL de imagen y texto.parse_chat_messageses el punto de entrada de la ruta síncrona.

📎 vllm/entrypoints/chat_utils.py:2161-2197

parse_chat_messagescreaMultiModalItemTracker, recorre cada mensaje llamando a_parse_chat_message_content, y finalmente llama a_postprocess_messagespara procesar los parámetros de llamada de herramientas, y luego materializa los datos multimodales a través demm_tracker.resolve_items().

_parse_chat_message_contentse encarga del análisis de un solo mensaje.

📎 vllm/entrypoints/chat_utils.py:2007-2029

Primero normaliza el content:Nonese convierte en una lista vacía, una cadena se convierte en una sola parte de texto. Luego llama a_parse_chat_message_content_parts, donde el parámetrowrap_dictses determinado porcontent_format == "openai"— esto decide si la salida es una lista de diccionarios estructurados o una cadena concatenada.

_parse_chat_message_content_partsrecorre cada part.

📎 vllm/entrypoints/chat_utils.py:1814-1853

Cada part pasa por el procesamiento de_parse_chat_message_content_part. Siwrap_dicts=False, finalmente concatena el texto y los marcadores de posición en una sola cadena; siwrap_dicts=True, devuelve una lista de diccionarios estructurados.

_parse_chat_message_content_partes el núcleo de la distribución.

📎 vllm/entrypoints/chat_utils.py:1875-1884

Para parts de texto puro, primero se realiza la verificación de retención de marcadores de posición, y luego se decide el formato de retorno segúnwrap_dicts. Para parts estructurados, se llama a_parse_chat_message_content_mm_partpara extraer el tipo y el contenido.

📎 vllm/entrypoints/chat_utils.py:1690-1723

_parse_chat_message_content_mm_partbusca la función de análisis correspondiente a través deMM_PARSER_MAP. Nótese la condición deuuid is None— si el usuario proporcionó un UUID, significa que los datos multimedia pueden no estar en el cuerpo de la solicitud (ya se subieron por otros medios), en cuyo caso se toma la rama de campo de URL directa a continuación.

📎 vllm/entrypoints/chat_utils.py:1731-1733

Cuandopart_type is Noneouuid is not None, el código intenta extraer el campo URL directamente del part. Este «análisis permisivo» es para compatibilidad con clientes que no siguen estrictamente el formato OpenAI.

Volviendo a_parse_chat_message_content_part, los parts de tipo multimedia se distribuyen a los métodosmm_parsercorrespondientes.

📎 vllm/entrypoints/chat_utils.py:1923-1968

Cada tipo de medio llama al métodoparse_*correspondiente, y estos métodos internamente llaman atracker.addpara añadir el elemento al rastreador y devuelven una cadena de marcador de posición. Finalmente, segúninterleave_stringsse decide devolver el marcador de posición oNone。

📎 vllm/entrypoints/chat_utils.py:1984-1999

prompt_embeds. El procesamiento de es especial: independientemente deinterleave_strings, siempre devuelvePROMPT_EMBEDS_PLACEHOLDER_TOKEN. El comentario explica la razón — prompt_embeds se concatena en el desplazamiento de tokens, la posición es importante, y si se pasa pormissing_placeholdersla lógica de relleno previo desordenaría el orden.

Diferencias de la ruta asíncrona

La ruta asíncrona usaAsyncMultiModalItemTrackeryAsyncMultiModalContentParser. La diferencia principal está enresolve_items。

📎 vllm/entrypoints/chat_utils.py:906-952

. La versión asíncrona usaasyncio.gatherpara esperar concurrentemente todos los elementos multimodales. El comentario señala explícitamente: cada elemento rastreado ya es un awaitable independiente, y el conector asíncrono descarga el trabajo de decodificación bloqueante al pool de hilos, por lo que esperar secuencialmente una modalidad y luego otra aumentaría la latencia innecesariamente.return_exceptions=Truepermite que todas las tareas se completen o fallen antes de lanzar la excepción de forma unificada, evitando que el primer fallo abandone las solicitudes de red aún en curso.

Reflexión de diseño: por qué separar el rastreador del analizador

3.3 De mensaje a token: la transición entre renderer y EngineCore

chat_utilsLa listaConversationMessageyMultiModalDataDictproducidas aún necesitan pasar por el renderizado del chat template para convertirse en una secuencia de tokens. Este paso lo realiza el renderer, y después la solicitud entra realmente en el motor.

Impulsado por escenarios: renderizado del chat template y envío de la solicitud

parse_chat_messagesTras regresar, el llamador (comoOpenAIServingChat) pasaráconversationymm_dataal renderer. El renderer aplica la plantilla de chat, renderiza la listaConversationMessagea texto y luego la tokeniza en una secuencia de IDs de token. Los marcadores de posición multimodales (como<##IMAGE##>) se reemplazan tras la tokenización por tokens de marcador de posición específicos del modelo.

Una vez completado el renderizado, la solicitud se encapsula comoEngineCoreRequesty se entrega a la cola de entrada de EngineCore medianteAsyncLLM.add_request()oLLMEngine.add_request().

📎 vllm/entrypoints/llm.py:420-484

El métodoLLM.generateoffline muestra esta cadena: primero validarunner_type, obtiene los parámetros de muestreo predeterminados y luego llama a_run_completion。_run_completion. Internamente se llama al renderer para renderizar el prompt y después se entrega la solicitud mediantellm_engine.

📎 vllm/entrypoints/llm.py:615-708

LLM.chatEl método muestra la ruta de chat: recibe la listamessages, llama a_run_chat, que internamente llama aparse_chat_messagesy al renderer.

Reflexión de diseño: por qué el renderer está dentro del motor

Recuperación de errores y trampas en producción

_postprocess_messagesEl manejo de parámetros de llamadas a herramientas en es un típico obstáculo del entorno de producción.

📎 vllm/entrypoints/chat_utils.py:2118-2158

Cuando un mensaje del assistant contienetool_calls, el campoargumentspuede ser una cadena JSON, un diccionario o JSON inválido. El código intenta analizar la cadena JSON; si falla, registra una advertencia y la convierte forzosamente en un objeto vacío. El comentario explica el motivo: losargumentscon formato incorrecto existen en el historial de conversación, y si se hace fallar la solicitud aquí, cada ronda posterior fallará y la conversación no podrá recuperarse. Este es un diseño de tolerancia a fallos bien pensado: es preferible que el modelo vea parámetros de herramienta vacíos antes que bloquear toda la conversación.

Otra trampa es la protección contra inyección de marcadores de posición reservados.

📎 vllm/entrypoints/chat_utils.py:1856-1872

Cuandoenable_prompt_embedsestá activado,PROMPT_EMBEDS_PLACEHOLDER_TOKENse registra como un token especial indivisible. Si el texto del usuario contiene exactamente esta secuencia literal, el tokenizer la codificará como el mismo ID de token, y el renderer creerá erróneamente que este es el punto de concatenación, permitiendo al llamador mover o inyectar la posición de concatenación mediante contenido de texto plano._reject_reserved_placeholder_in_textrechaza este tipo de entrada durante el análisis de la parte de texto, cerrando esta vulnerabilidad de seguridad.

📎 vllm/entrypoints/chat_utils.py:1889-1892

Nótese que esta comprobación se invoca tanto en la ramaisinstance(part, str)como en la rama de texto estructurado, asegurando que todas las rutas de texto pasen por la protección.

Resumen del capítulo

Este capítulo ha trazado el primer tramo del recorrido de una solicitud desde el exterior hacia el sistema. Las tres rutas de entrada —API HTTP, CLI y la claseLLMoffline— convergen finalmente en la capa de análisis multimodal dechat_utils.BaseMultiModalItemTrackerse encarga de la gestión de estado,BaseMultiModalContentParserse encarga de la extracción de contenido, y su separación permite que las rutas síncronas y asíncronas compartan la lógica de trazado.parse_chat_messagesnormaliza mensajes heterogéneos en una listaConversationMessageyMultiModalDataDict, y luego los entrega al renderer interno del motor para completar el renderizado de la plantilla de chat y la tokenización. Finalmente, la solicitud se encapsula comoEngineCoreRequesty se entrega a la cola de entrada de EngineCore.

Reflexión y autoevaluación del capítulo

El próximo capítulo entrará en el planificador para ver cómo EngineCore orquesta estas solicitudes con procesamiento por lotes continuo y estrategias conscientes de la memoria de video.

Hasta aquí, la solicitud ha completado la transformación normalizada desde la entrada externa hasta EngineCoreRequest y ha llegado a la entrada del núcleo del motor. Pero una vez que la solicitud ingresa, no se ejecuta inmediatamente: el motor necesita decidir qué solicitudes procesar en cada paso y cómo asignar los recursos limitados de memoria de video. El próximo capítulo profundizará en el bucle de planificación de EngineCore, analizando cómo el Scheduler equilibra el rendimiento y la latencia en el procesamiento por lotes continuo, y cómo el chunked prefill, el prefix caching y la asignación de bloques KV trabajan juntos.

Después de que la solicitud ingresa a la cola de entrada de EngineCore, no se ejecuta inmediatamente. Qué solicitudes procesar en cada paso, cuánto presupuesto de tokens asignar a cada solicitud, a quién sacrificar primero cuando la memoria de video es insuficiente, todas estas decisiones se concentran en el métodoScheduler.schedule(). Este capítulo comienza con las estructuras de datos del planificador y rastrea cómo una llamada aschedule()organiza la cola waiting, la lista running y el pool de KV cache en un lote ejecutable.

4.1 Estructuras de datos del planificador: tres colas y un pool de memoria de video

La pregunta central que debe responder el planificador es:Bajo un presupuesto limitado de tokens y de bloques KV, ¿qué solicitudes deben avanzar cuántos tokens en este paso?Para entenderlo, primero hay que ver claramente qué estados tiene en sus manos.

El planificador mantiene tres tipos de contenedores de solicitudes.self.requestsEs un diccionario global,req_id -> Request, la única fuente de verdad para todas las solicitudes activas📎 vllm/v1/core/sched/scheduler.py:208-209。self.waitingyself.skipped_waitingson dos colas de prioridad, la primera contiene solicitudes en espera normal de planificación, la segunda contiene solicitudes que temporalmente no pueden planificarse debido a dependencias asíncronas o restricciones (como esperar KV remoto, esperar la compilación de la gramática de salida estructurada)📎 vllm/v1/core/sched/scheduler.py:208-209。self.runningEs una lista normal que almacena solicitudes que ya han entrado en estado de ejecución y poseen bloques KV📎 vllm/v1/core/sched/scheduler.py:208-209。

Aquí hay un diseño fácil de pasar por alto:max_num_running_reqsymax_num_active_reqsson dos límites superiores diferentes. El primero proviene demax_num_seqs, determina el número de ranuras del model runner; el segundo proviene demax_num_active_seqs, solo limita el número de solicitudes que pueden entrar en RUNNING, por defecto igual al primero📎 vllm/v1/core/sched/scheduler.py:123-131. Esta separación permite reducir el tamaño real del lote de decodificación concurrente sin reducir la capacidad de captura del CUDA graph.

El lado de la memoria de video está gestionado de manera unificada porKVCacheManager, que internamente poseeBlockPool。BlockPoolEl núcleo deself.blocksesKVCacheBlock(lista de todosfree_block_queue) y📎 vllm/v1/core/block_pool.py:171-177(una lista doblemente enlazada de bloques libres ordenada por orden de expulsión)null_block. Nótese la existencia deis_null=True: es el primer bloque extraído de la cabeza de la cola de libres,📎 vllm/v1/core/block_pool.py:183-187, el conteo de referencias no participa en el mantenimiento regular, se usa específicamente como marcador de posición

. Cuando una posición de token de una solicitud no necesita un bloque KV real (por ejemplo, una posición omitida por la ventana deslizante), en la block table se coloca este null block.BlockHashToBlockMapLa estructura de índice de la caché de prefijos esBlockHashWithGroupId, que mapeaKVCacheBlocka un{block_id: KVCacheBlock}o a un diccionario📎 vllm/v1/core/block_pool.py:56-59. ¿Por qué usar un tipo unión? El comentario da la respuesta: la mayoría de los hashes corresponden a un solo bloque, y usar un diccionario generaría una sobrecarga innecesaria de GC; solo cuando el mismo hash es compartido por múltiples bloques se promueve a diccionario📎 vllm/v1/core/block_pool.py:56-59. Esta es una típica compensación de complejidad de tipos por sobrecarga en tiempo de ejecución.

KVCacheBlockses el objeto de interfaz entre el planificador y el gestor de KV cache, que oculta las estructuras de datos internas. Sublockscampo estuple[Sequence[KVCacheBlock], ...], la dimensión externa es el grupo de KV cache, la interna es la secuencia de bloques📎 vllm/v1/core/kv_cache_manager.py:41-54. El comentario explica claramente por qué no se usa el bloque como dimensión externa: eso asumiría que todos los grupos tienen el mismo número de bloques, mientras que en el futuro podrían configurarse diferentes block sizes para distintos grupos📎 vllm/v1/core/kv_cache_manager.py:43-48。

flowchart LR
    subgraph Sched["Scheduler 状态"]
        W["waiting<br/>RequestQueue"]
        SW["skipped_waiting<br/>RequestQueue"]
        R["running<br/>list[Request]"]
        REQ["requests<br/>dict[str, Request]"]
    end
    subgraph KV["KVCacheManager"]
        BP["BlockPool.blocks<br/>list[KVCacheBlock]"]
        FQ["free_block_queue<br/>FreeKVCacheBlockQueue"]
        MAP["cached_block_hash_to_block<br/>BlockHashToBlockMap"]
    end
    W -->|"admit + allocate_slots"| R
    R -->|"preempt"| W
    R -->|"free / pop_blocks_for_free"| FQ
    FQ -->|"get_new_blocks"| BP
    BP -->|"cache_full_blocks"| MAP
    MAP -->|"get_cached_block"| W

Este diagrama ancla el flujo de datos entre el planificador y el pool de memoria de video: las solicitudes de la cola waiting entran en running a través deallocate_slots, las solicitudes de running vuelven a waiting cuando son expropiadas, los bloques liberados vuelven a la cola libre, y la tabla hash de caché de prefijos es la entrada para que las solicitudes en waiting acierten en la caché.

4.2 Flujo principal de schedule(): prioridad a running, complemento de waiting, expropiación como respaldo

schedule()es el método central de todo el planificador, devuelve unSchedulerOutput, que describe qué se va a ejecutar en este paso. El comentario al inicio del método señala la filosofía de diseño: en el planificador no hay distinción entre "fase de decodificación" y "fase de prefill", cada solicitud solo tienenum_computed_tokensynum_tokens_with_spec, y la tarea del planificador es hacer que la primera alcance a la segunda📎 vllm/v1/core/sched/scheduler.py:559-568. Esta perspectiva unificada es la base para que chunked prefill, prefix caching y decodificación especulativa puedan coexistir.

4.2.1 Inicialización de presupuesto y cálculo de umbrales

Antes de entrar en el bucle principal, el planificador establece dos presupuestos:token_budgetse inicializa enmax_num_scheduled_tokens，input_budgetse inicializa enmax_num_batched_tokens 📎 vllm/v1/core/sched/scheduler.py:577-580. Ambos suelen ser iguales, pero cuando el modelo puede añadir tokens en el lote (como en decodificación especulativa),max_num_scheduled_tokensserá menor quemax_num_batched_tokens, y la diferencia es el espacio reservado para draft tokens.

long_prefill_token_thresholdEl manejo de merece una mirada aparte. Su función es evitar que un prefill largo mate de hambre a otras solicitudes, pero si actualmente solo hay una solicitud, nadie puede morir de hambre, así que el umbral se pone a cero📎 vllm/v1/core/sched/scheduler.py:606-616. Cuandoadaptive_long_prefill_thresholdestá activado, el umbral también se eleva ainput_budget // num_eligible_reqs, garantizando que no se reduzca el presupuesto de una sola solicitud por debajo de su cuota justa📎 vllm/v1/core/sched/scheduler.py:617-622。

4.2.2 Bucle de planificación de solicitudes running

El bucle principal recorre desde la cabeza deself.running,req_indexes el cursor📎 vllm/v1/core/sched/scheduler.py:624-627. Para cada solicitud, primero se hacen una serie de comprobaciones de omisión:

• Bajo planificación asíncrona, si el marcador de posición de salida de la solicitud indica que ya alcanzómax_tokens, se omite para evitar ejecutar un paso de más📎 vllm/v1/core/sched/scheduler.py:631-645。
• En el escenario V2 + PP + asíncrono, si el paso actual aún no ha llegado anext_decode_eligible_step, se omite para coincidir con el ritmo de difusión de tokens de muestreo del lado del worker📎 vllm/v1/core/sched/scheduler.py:647-651。
• Cuando el balanceo de prefill de DP está activado, los chunks de prefill en pasos no alineados al ritmo se posponen📎 vllm/v1/core/sched/scheduler.py:653-657。

Tras pasar las comprobaciones de omisión, se calcula cuántos tokens puede avanzar esta solicitud en este paso:

num_new_tokens = request.num_tokens_with_spec
               + request.num_output_placeholders
               - request.num_computed_tokens

Luego se restringe sucesivamente porlong_prefill_token_threshold、token_budget、input_budget - draft_slotsymax_model_len. Si la solicitud lleva entrada de codificador, también pasa por el ajuste de📎 vllm/v1/core/sched/scheduler.py:670-688_try_schedule_encoder_inputsA continuación viene el paso más crítico: asignar KV blocks.📎 vllm/v1/core/sched/scheduler.py:700-712。

se envuelve en un bucleallocate_slotswhile True. Si devuelve📎 vllm/v1/core/sched/scheduler.py:742-747, significa que no hay suficiente memoria de video, y el planificador comienza la expropiación: selecciona víctimas según la política (la estrategia PRIORITY elige la de menor prioridad, la estrategia FCFS elige la del final de la lista running)None, llama a📎 vllm/v1/core/sched/scheduler.py:761-767para devolverla a la cola waiting, y luego reintenta la asignación_preempt_request. Si la víctima es la propia solicitud actual, significa que ya no hay objetos expropiables, se sale del bucle y la solicitud actual tampoco puede planificarse📎 vllm/v1/core/sched/scheduler.py:801-806Hay un detalle sutil en la lógica de expropiación: bajo la estrategia PRIORITY, si la solicitud expropiada ya está en📎 vllm/v1/core/sched/scheduler.py:807-813。

(es decir, en este paso ya se le asignaron recursos), hay que devolver por completo su presupuesto de tokens, blocks, tokens especulativos y presupuesto de codificadorscheduled_running_reqs. Esto garantiza la consistencia del libro de presupuestos.📎 vllm/v1/core/sched/scheduler.py:779-797Tras una asignación exitosa, la solicitud se añade a

, se registran los blocks y el número de tokens, y se deduce del presupuestoscheduled_running_reqs. Los tokens relacionados con decodificación especulativa se recortan y registran aquí📎 vllm/v1/core/sched/scheduler.py:815-8234.2.3 Admisión de solicitudes waiting📎 vllm/v1/core/sched/scheduler.py:825-841。

Tras terminar el bucle de running, si en este paso no hubo expropiación y el planificador no está pausado, se empieza a procesar la cola waiting

. Antes de la admisión se comprueban dos límites:📎 vllm/v1/core/sched/scheduler.py:868-872ymax_num_active_reqsLa planificación de solicitudes waiting tiene un paso más que la de running: la búsqueda en la caché de prefijos. Cuandoinput_budget 📎 vllm/v1/core/sched/scheduler.py:873-879。

, se llama arequest.num_computed_tokens == 0para buscar aciertos en la caché local_get_local_prefix_cache_hit. Si se ha configurado un KV connector, también se consultan los aciertos en la caché remota📎 vllm/v1/core/sched/scheduler.py:932-939Aquí hay una lógica refinada para manejar conflictos entre aciertos locales y remotos. Un acierto local puede no estar alineado a bloques (📎 vllm/v1/core/sched/scheduler.py:942-954。

), y si un acierto remoto supera estrictamente al acierto local completo, se descarta la cola del subbloque local para que la carga remota la sobrescriba, evitando copy-on-writepartial_tail. En caso contrario, se conserva la cola local y no se carga nada externo📎 vllm/v1/core/sched/scheduler.py:977-988Tras una admisión exitosa, la solicitud se saca de la cola waiting, su estado se establece en RUNNING y se añade a la lista running📎 vllm/v1/core/sched/scheduler.py:989-995。

. Si después de este paso sigue en prefill (📎 vllm/v1/core/sched/scheduler.py:1263-1319), se añade al conjuntonum_computed_tokens + num_new_tokens < request.num_tokens_inflight_prefillsCopiar📎 vllm/v1/core/sched/scheduler.py:1326-1328。

flowchart TD
    start["schedule() 开始"] --> init["初始化 token_budget / input_budget"]
    init --> run_loop{"running 循环<br/>req_index < len(running)<br/>且 token_budget > 0?"}
    run_loop -->|是| skip_check{"跳过条件?<br/>max_tokens 已达 /<br/>decode_eligible / defer_prefills"}
    skip_check -->|跳过| run_inc["req_index += 1"]
    run_inc --> run_loop
    skip_check -->|不跳过| calc["计算 num_new_tokens<br/>受多约束裁剪"]
    calc --> alloc{"allocate_slots<br/>返回 None?"}
    alloc -->|成功| admit_run["加入 scheduled_running_reqs<br/>扣减预算"]
    admit_run --> run_inc
    alloc -->|失败| can_preempt{"有可抢占请求?<br/>_request_blocks_can_be_freed"}
    can_preempt -->|否| break_run["跳出 running 循环"]
    can_preempt -->|是| preempt["_preempt_request<br/>踢回 waiting"]
    preempt --> alloc
    break_run --> wait_loop{"无抢占且未暂停?<br/>waiting 非空且 token_budget > 0?"}
    run_loop -->|否| wait_loop
    wait_loop -->|是| blocked{"blocked 状态?<br/>_is_blocked_waiting_status"}
    blocked -->|是且无法提升| skip_wait["移入 skipped_waiting"]
    skip_wait --> wait_loop
    blocked -->|否| prefix{"num_computed_tokens == 0?<br/>查找前缀缓存"}
    prefix -->|命中| alloc_wait["allocate_slots<br/>带 new_computed_blocks"]
    prefix -->|未命中| alloc_wait
    alloc_wait --> wait_ok{"分配成功?"}
    wait_ok -->|是| admit_wait["加入 running<br/>状态设为 RUNNING"]
    admit_wait --> wait_loop
    wait_ok -->|否| break_wait["跳出 waiting 循环"]
    wait_loop -->|否| build["构建 SchedulerOutput"]
    break_wait --> build

y la rama de expropiación. Obsérvese la ruta de reintento de expropiación tras el fallo deschedule()en el bucle running, y cómo las solicitudes en estado blocked del bucle waiting se mueven aallocate_slots 失败后的抢占重试路径，以及 waiting 循环中 blocked 状态请求被移入 skipped_waitingdel bypass.

4.3 El núcleo consciente de la memoria de video: allocate_slots y la apropiación

allocate_slotses la compuerta entre el planificador y la memoria de video. Su lista de parámetros es en sí misma un libro contable de memoria de video:num_new_tokenses el número de tokens que se van a calcular nuevos,num_new_computed_tokenses el número de tokens que aciertan nuevos en la caché de prefijos,num_external_computed_tokenses el número de aciertos externos proporcionados por el connector,num_lookahead_tokensson las ranuras reservadas para la decodificación especulativa📎 vllm/v1/core/kv_cache_manager.py:371-383。

El comentario al inicio del método describe con precisión el diseño de bloques mediante un diagrama ASCII📎 vllm/v1/core/kv_cache_manager.py:417-438：

| < comp > | < new_comp > | < ext_comp >  | < new >  | < lookahead > |
                                          |   < to be computed >     |
                        |            < to be allocated >           |

compson los tokens ya calculados,new_compes un acierto en la caché de prefijos,ext_compes un acierto externo,newes el cálculo nuevo de este paso,lookaheades la reserva especulativa. La asignación se divide en tres fases: primero se liberan los bloques innecesarios y se comprueba si hay suficientes bloques libres, luego se procesan los tokens de prefijo y, por último, se asignan bloques para los tokens de cálculo nuevo📎 vllm/v1/core/kv_cache_manager.py:458-461。

4.3.1 Línea de nivel de agua y control de admisión

allocate_slotshay dos compuertas de admisión. La primera esfull_sequence_must_fit: cuando está activada, primero se comprueba si toda la secuencia de solicitud (no solo el primer chunk) cabe; si no cabe, se devuelve directamenteNone 📎 vllm/v1/core/kv_cache_manager.py:515-531. Esto evita que, bajo chunked prefill, una admisión excesiva provoque oscilaciones en la KV cache.

La segunda es la línea de nivel de agua.watermark_blockssolo entra en vigor cuando el estado de la solicitud es WAITING o PREEMPTED y ya hay solicitudes planificadas📎 vllm/v1/core/kv_cache_manager.py:506-513. Exige que, tras la asignación, se conserve al menos una cierta proporción de bloques libres, evitando expulsiones y apropiaciones frecuentes.reserved_blocksse utiliza en escenarios de carga asíncrona de KV, para garantizar que los bloques reservados para prefill en curso no sean consumidos por nuevas solicitudes📎 vllm/v1/core/kv_cache_manager.py:564-570。

4.3.2 El coste y la recuperación de la apropiación

La apropiación también aborda el problema de las "salidas obsoletas" bajo planificación asíncrona.num_stale_output_tokensse establece ennum_in_flight_tokens, marcando todas las salidas en curso como obsoletas📎 vllm/v1/core/sched/scheduler.py:1571-1574. Estos tokens se seguirán entregando (descartarlos perturbaría la tasa de aceptación de la decodificación especulativa), pero no modificarán los contadores tras el restablecimiento.drop_stale_outputel indicador determina si se descarta o se entrega📎 vllm/v1/core/sched/scheduler.py:1539-1547。

4.3.3 Liberación diferida: el riesgo de lectura tras escritura en conectores asíncronos

Cuando se usa un KV connector y hay varios lotes en curso,defer_block_freese establece enTrue 📎 vllm/v1/core/sched/scheduler.py:175-181. La razón es que un paso puede seguir escribiendo en los KV block de una solicitud ya liberada, mientras que el connector consumidor podría reasignar y rellenar esos bloques mediante una carga no ordenada respecto a esa escritura.

La liberación diferida se implementa mediantedeferred_freesuna cola doble, donde cada entrada es(fence_seq, blocks) 📎 vllm/v1/core/sched/scheduler.py:388-390。_free_request_blocksComprobar_request_blocks_can_be_freed, si el último paso de planificación de la solicitud aún no se ha terminado de procesar, se ponen los bloques en la cola diferida📎 vllm/v1/core/sched/scheduler.py:2679-2688。_drain_deferred_freesenupdate_from_outputse avanzaprocessed_step_seqse llama después, liberando los bloques cuyo fence ya se ha satisfecho📎 vllm/v1/core/sched/scheduler.py:2701-2706。

4.4 Determinación de aciertos en la caché de prefijos y ciclo de vida de los bloques

La entrada de búsqueda de la caché de prefijos esKVCacheManager.get_computed_blocks. Primero comprueba si la caché está habilitada y si la solicitud no está marcada para omitir la lectura📎 vllm/v1/core/kv_cache_manager.py:286-287. Luego llama acoordinator.find_longest_cache_hit, pasandorequest.block_hashesymax_cache_hit_length = request.num_tokens - 1 📎 vllm/v1/core/kv_cache_manager.py:295-300。

¿Por quénum_tokens - 1? El comentario explica: cuando todos los tokens aciertan en la caché, es obligatorio recalcular el último token para obtener los logits📎 vllm/v1/core/kv_cache_manager.py:289-294. Este es un límite fácil de pasar por alto: incluso si el prefijo acierta por completo, hay que calcular al menos un token.

El ciclo de vida de los bloques lo gestionaBlockPool.get_new_blocksextrae bloques del frente de la cola de libres; si la caché está habilitada, primero llama a_maybe_evict_cached_blockpara borrar sus metadatos de hash y luego incrementa el contador de referencias📎 vllm/v1/core/block_pool.py:683-702。free_blocksdecide si devolver el bloque al frente o al final de la cola según si tiene hash: los bloques sin hash se reutilizan LIFO (mejor localidad de GPU), los bloques con hash se reutilizan FIFO (comportamiento de expulsión LRU)📎 vllm/v1/core/block_pool.py:785-805。

cache_full_blockses el momento en que un bloque se escribe en la tabla hash de la caché de prefijos. Recorre los bloques recién llenados, omite los bloques null y los bloques enmascarados, calcula el hash de cada bloque y lo inserta encached_block_hash_to_block 📎 vllm/v1/core/block_pool.py:272-300. Si el bloque ya tiene hash (escenario en que un bloque parcial se actualiza a bloque lleno), primero se elimina el hash antiguo y luego se inserta el nuevo📎 vllm/v1/core/block_pool.py:285-293。

touchel método gestiona el contador de referencias cuando hay un acierto en la caché: si el bloque está en la cola de libres (ref_cnt == 0), primero se saca de la cola y luego se incrementa el contador de referencias📎 vllm/v1/core/block_pool.py:754-770. Esto garantiza que los bloques acertados no sean expulsados.

Reflexiones de diseño

Resumen del capítulo

El núcleo del planificador esschedule()dos bucles en el método: el bucle running prioriza garantizar el avance de las solicitudes ya en ejecución, mientras que el bucle waiting admite nuevas solicitudes cuando el presupuesto lo permite. Cuando la memoria de video es insuficiente, se libera espacio mediante la preempción de la solicitud de menor prioridad en la lista running; la solicitud expropiada tiene sunum_computed_tokensrestablecido a 0, pero la caché de prefijos puede compensar parte del costo de recálculo.allocate_slotses la compuerta de memoria de video, que mediantefull_sequence_must_fit, la línea de nivel de agua yreserved_blockstres niveles de control de admisión evitan la sobreasignación. La caché de prefijos logra compartir entre solicitudes mediante indexación por hash de bloques, y la determinación de acierto tiene como límite superiornum_tokens - 1para garantizar que al menos se calcule un token y se obtengan los logits.

Reflexión y autoevaluación del capítulo

La salida del planificadorSchedulerOutputespecifica el contenido de ejecución de este paso: los IDs de bloque de las nuevas solicitudes, el número de tokens de las solicitudes en caché, los tokens especulativos, las entradas del codificador, etc. El siguiente capítulo rastreará cómo esta salida es consumida por ModelRunner, desdeSchedulerOutputhasta la propagación hacia adelante en la GPU.

En el capítulo anterior vimos que el Scheduler, en cada paso del bucle de programación, decide qué solicitudes entran en la cola running, cuáles son expropiadas y cuáles esperan por falta de memoria de video, y finalmente produce un SchedulerOutput, que describe qué debe calcularse en este paso: qué solicitudes, cuántos tokens para cada una y qué bloques KV usar. Pero esta lista es solo una intención lógica; la GPU necesita tensores físicos. Este capítulo rastrea cómo SchedulerOutput es distribuido por el Executor a los Workers, y luego traducido por GPUModelRunner en entradas ejecutables por la GPU como input_ids, positions, slot_mapping y block table, para finalmente, a través de forward_context, inyectar la descripción del lote compartida entre capas en cada capa del modelo, completando el salto desde la decisión de programación hasta la propagación hacia adelante.

5.1 Executor: enviar el resultado de la programación a cada tarjeta

Modelo intuitivo

Executores el "mensajero" entre EngineCore y los GPU Workers. Sin él, EngineCore tendría que saber por sí mismo cuántas tarjetas hay en el clúster, en qué proceso está cada tarjeta y cómoSchedulerOutputSerializar el pasado — la lógica de scheduling quedaría entrelazada con la topología distribuida.ExecutorExtrae esta responsabilidad: EngineCore solo se encarga de invocarexecute_model(scheduler_output), y el resto — «a quién enviar, cómo enviar, cuántos resultados recibir» — lo decide el Executor.

Jerarquía de clases y campos

ExecutorEs una clase base abstracta cuyos campos a nivel de clase codifican directamente las capacidades del backend📎 vllm/v1/executor/abstract.py:48-49：

uses_ray: bool = False  # whether the executor uses Ray for orchestration.
supports_pp: bool = False  # whether the executor supports PP

Estos dos flags no son decorativos — el código de capas superiores los lee para decidir si habilitar ciertas rutas de optimización.__init__En se inicializansleeping_tags、kv_output_aggregator、ec_output_aggregatortres campos de estado📎 vllm/v1/executor/abstract.py:119-120, utilizados respectivamente para el seguimiento de etiquetas del modo sleep, la agregación de salidas del conector KV y la agregación de salidas del conector de encoder.

Selección de backend:get_classEnrutamiento por ramas de

get_classEs una fábrica estática que, según la configuración dedistributed_executor_backend, devuelve la clase Executor concreta📎 vllm/v1/executor/abstract.py:51-96. Su estructura de ramas merece un examen detallado:

• Si la configuración en sí es untype, tras validar si es una subclase deExecutor, se usa directamente📎 vllm/v1/executor/abstract.py:52-61；
• "ray"Bajo la rama de hay además subramas de segundo nivel:VLLM_USE_RAY_V2_EXECUTOR_BACKENDCuando es verdadero se usaRayExecutorV2, de lo contrario se usaRayDistributedExecutor 📎 vllm/v1/executor/abstract.py:64-72；
• "mp"se mapea aMultiprocExecutor，"uni"se mapea aUniProcExecutor 📎 vllm/v1/executor/abstract.py:73-80；
• Los backends personalizados en forma de cadena se resuelven dinámicamente medianteresolve_obj_by_qualname📎 vllm/v1/executor/abstract.py:85-90。

flowchart TD
    start["Executor.get_class(vllm_config)"] --> check_type{"backend 是 type?"}
    check_type -->|是| verify_sub{"issubclass(Executor)?"}
    verify_sub -->|否| err_type["raise TypeError"]
    verify_sub -->|是| use_direct["executor_class = backend"]
    check_type -->|否| check_ray{"backend == 'ray'?"}
    check_ray -->|是| ray_v2{"VLLM_USE_RAY_V2?"}
    ray_v2 -->|是| use_rayv2["RayExecutorV2"]
    ray_v2 -->|否| use_ray["RayDistributedExecutor"]
    check_ray -->|否| check_mp{"backend == 'mp'?"}
    check_mp -->|是| use_mp["MultiprocExecutor"]
    check_mp -->|否| check_uni{"backend == 'uni'?"}
    check_uni -->|是| use_uni["UniProcExecutor"]
    check_uni -->|否| check_ext{"backend == 'external_launcher'?"}
    check_ext -->|是| use_ext["ExecutorWithExternalLauncher"]
    check_ext -->|否| check_str{"backend 是 str?"}
    check_str -->|是| resolve["resolve_obj_by_qualname"]
    check_str -->|否| err_unknown["raise ValueError"]

Paso a paso: el flujo de invocación de unaexecute_model

Contextualizando: EngineCore completa un paso de scheduling, obtieneSchedulerOutput, e invocaexecutor.execute_model(scheduler_output)。

Executor.execute_modelLa implementación de es minimalista📎 vllm/v1/executor/abstract.py:237-238：

def execute_model(
    self, scheduler_output: SchedulerOutput, non_block: bool = False
) -> ModelRunnerOutput | None | Future[ModelRunnerOutput | None]:
    output = self.collective_rpc(
        "execute_model", args=(scheduler_output,), non_block=non_block
    )
    return output[0]

sample_tokensSigue el mismo patrón📎 vllm/v1/executor/abstract.py:257-258, pero el tipo de retorno no incluyeNone— el muestreo produce resultados inevitablemente. La división de tareas entre estos dos métodos corresponde al diseño de «separación ejecución-muestreo» de vLLM v1:execute_modelpuede devolverNone(indicando que el forward ya se envió pero el muestreo se pospone), en cuyo caso el estado se almacena temporalmente enExecuteModelState.

Reflexiones de diseño

collective_rpcSe declara como@abstractmethod 📎 vllm/v1/executor/abstract.py:186-192, lo que significa que cada backend debe implementar por su cuenta «cómo enviar el RPC al Worker».MultiprocExecutorusa colas de memoria compartida,RayDistributedExecutorusa llamadas a actores de Ray,UniProcExecutorrealiza llamadas locales directas. Esta abstracción hace que el código de capas superiores no necesite preocuparse en absoluto por los detalles distribuidos.

Un detalle fácil de pasar por alto:supported_tasksestá marcado como@cached_property 📎 vllm/v1/executor/abstract.py:306-309, y el comentario dice directamente «evitar llamadas RPC innecesarias». Porqueget_supported_tasksrequiere comunicación entre procesos, y la lista de tareas no cambia durante el ciclo de vida del modelo, el caché es una optimización correcta y necesaria.

5.2 GPUModelRunner: de SchedulerOutput a tensores de entrada

Modelo intuitivo

GPUModelRunneres un «traductor»: traduce las descripciones lógicas deSchedulerOutput(ID de solicitud, número de tokens, ID de bloque) a tensores físicos que la GPU puede consumir directamente. Sin él, la capa del modelo tendría que lidiar por sí misma con preguntas como «¿en qué ranura KV está el séptimo token de la tercera solicitud?» — esto sería una fuga de responsabilidades catastrófica.

Estado central y diseño de memoria

GPUModelRunnerHereda de tres Mixins📎 vllm/v1/worker/gpu_model_runner.py:479-480：LoRAModelRunnerMixin、KVConnectorModelRunnerMixin、ECConnectorModelRunnerMixin, que proporcionan respectivamente capacidades de adaptación LoRA, conector KV y conector de encoder.

__init__En se cachean todos los objetos de configuración📎 vllm/v1/worker/gpu_model_runner.py:488-498, y se inicializan varios flags clave:

• check_ep_fault: solo cuando el paralelismo de datos > 1 y es un modelo MoE, consulta si el gestor EP all2all soporta tolerancia a fallos📎 vllm/v1/worker/gpu_model_runner.py:507-509；
• is_pooling_model: determinado porrunner_type == "pooling"📎 vllm/v1/worker/gpu_model_runner.py:515；
• enable_prompt_embeds: si habilitar la entrada de prompt embedding📎 vllm/v1/worker/gpu_model_runner.py:516。

ExecuteModelStatees unNamedTuple, que porta el estado temporal entreexecute_model()ysample_tokens()📎 vllm/v1/worker/gpu_model_runner.py:463-476. El diseño de sus campos revela la esencia de la separación ejecución-muestreo:logits、hidden_states、sample_hidden_stateses el producto del forward,spec_decode_metadata、slot_mappingsson los metadatos aún necesarios en la fase de muestreo. El comentario dice explícitamente que este es «el estado de caché temporal que se pasa después de que execute_model() devuelve None»📎 vllm/v1/worker/gpu_model_runner.py:464-464。

Step-by-Step：_update_statesCómo sincronizar el estado de caché

Contextualizando: el scheduler decide que este paso procesa las solicitudes A (nueva solicitud), B (continuación del decode del paso anterior), C (recuperada tras ser expropiada), mientras que la solicitud D ya se completó.

Primer paso: limpiar las solicitudes completadas.Recorrefinished_req_ids, extrae el estado del diccionarioself.requests, elimina deinput_batch. Nótese el caso límite señalado por el comentario:📎 vllm/v1/worker/gpu_model_runner.py:1202-1217yfinished_req_idspueden solaparse — cuando una solicitud es abortada y luego reenviada con el mismo ID, se consideran dos solicitudes distintasscheduled_req_ids📎 vllm/v1/worker/gpu_model_runner.py:1211-1215。

Segundo paso: poner a cero los bloques KV recién asignados.Sinew_block_ids_to_zerono está vacío, se invoca_zero_block_idspara poner a cero la memoria de video, evitando que NaN obsoletos contaminen los cálculos de atención o SSM📎 vllm/v1/worker/gpu_model_runner.py:1219-1222. Este es el requisito de seguridad previo para la reutilización de bloques de PagedAttention.

Tercer paso: calcular el conjunto de solicitudes no programadas.Este es el paso más propenso a errores📎 vllm/v1/worker/gpu_model_runner.py:1238-1247：

scheduled_req_ids = scheduler_output.num_scheduled_tokens.keys()
cached_req_ids = self.input_batch.req_id_to_index.keys()
resumed_req_ids = scheduler_output.scheduled_cached_reqs.resumed_req_ids
unscheduled_req_ids = cached_req_ids - (scheduled_req_ids - resumed_req_ids)

El comentario explica por qué esscheduled_req_ids - resumed_req_idsen lugar de directamentescheduled_req_ids: normalmentecached_req_idsyresumed_req_idsno se intersecan, pero en escenarios de expropiación forzada desencadenados porreset_prefix_cache, las solicitudes recuperadas deben eliminarse primero del lote persistente y luego reincorporarse📎 vllm/v1/worker/gpu_model_runner.py:1241-1246。

Cuarto paso: procesar nuevas solicitudes.scheduled_new_reqsPara cadaCachedRequestState 📎 vllm/v1/worker/gpu_model_runner.py:1295-1308, se construyeRANDOM_SEED. Si el tipo de muestreo estorch.Generator 📎 vllm/v1/worker/gpu_model_runner.py:1277-1284, se crea un_init_mrope_positionscon semilla. Si el modelo usa M-RoPE, se invoca📎 vllm/v1/worker/gpu_model_runner.py:1319-1321。

para precalcular las posicionesscheduled_cached_reqsQuinto paso: actualizar las solicitudes en ejecución.num_computed_tokens 📎 vllm/v1/worker/gpu_model_runner.py:1402Para cada📎 vllm/v1/worker/gpu_model_runner.py:1437-1448, se actualizareq_index is None, se maneja la adición o reemplazo de IDs de bloquereqs_to_add 📎 vllm/v1/worker/gpu_model_runner.py:1450-1465。

. Si la solicitud no está en el lote persistente ( condense()Rellenar los huecos dejados por las solicitudes de eliminación📎 vllm/v1/worker/gpu_model_runner.py:1511-1512，_may_reorder_batchHacer que el backend de atención reordene según sea necesario📎 vllm/v1/worker/gpu_model_runner.py:1513-1514，refresh_metadata()Actualizar los metadatos del lote📎 vllm/v1/worker/gpu_model_runner.py:1515-1516。

Preparación de tensores de entrada:_prepare_input_idsRuta rápida asíncrona de

_prepare_input_idsManeja un problema sutil: bajo programación asíncrona, el token de muestreo del paso anterior aún está en la GPU, y losinput_idsdel paso actual necesitan rellenarlos📎 vllm/v1/worker/gpu_model_runner.py:1767-1772。

Ruta normal (prev_sampled_token_ids is None) copia directamente el tensor de CPU a la GPU📎 vllm/v1/worker/gpu_model_runner.py:1788-1794. La ruta asíncrona recorre las solicitudes, calcula el índice del último token de cada solicitud en elinput_idsaplanado📎 vllm/v1/worker/gpu_model_runner.py:1809-1836. Los comentarios dan un ejemplo concreto:cu_num_tokens = [2, 5, 8]、draft_tokens = [1, 2, 2]cuandosample_flattened_indices = [0, 2, 5]，spec_flattened_indices = [1, 3, 4, 6, 7] 📎 vllm/v1/worker/gpu_model_runner.py:1820-1822。

Hay una optimización clave📎 vllm/v1/worker/gpu_model_runner.py:1859-1868：

if common_indices_match and max_flattened_index == (num_common_tokens - 1):
    self.input_ids.gpu[:num_common_tokens].copy_(
        self.input_batch.prev_sampled_token_ids[:num_common_tokens, 0],
        non_blocking=True,
    )
    return

Cuando el lote no cambia y no hay reordenamiento, los índices son0..N-1la misma permutación, se puede usar directamente una copia por segmento único, evitando el costo de scatter. Esta es una manifestación directa de la optimización de lote persistente.

slot_mappingy la block table

_get_slot_mappingsDevuelve dos formatos📎 vllm/v1/worker/gpu_model_runner.py:4078-4078: indexado por KV cache groupdict[int, torch.Tensor]para uso de los metadatos de atención, indexado por nombre de capadict[str, torch.Tensor]paraForwardContextuso. Para un KV cache group encoder-only, el slot mapping es un tensor todo ceros📎 vllm/v1/worker/gpu_model_runner.py:4096-4115; de lo contrario se obtiene por segmento desdeblock_table.slot_mapping.gpu📎 vllm/v1/worker/gpu_model_runner.py:4107-4109. El relleno final no utilizado-1, el comentario explica que esto esreshape_and_cacheuna necesidad en modo CUDA graph completo📎 vllm/v1/worker/gpu_model_runner.py:4118-4122。

_get_block_tableObtener el tensor de dispositivo para cada KV cache group📎 vllm/v1/worker/gpu_model_runner.py:2319-2335, y usarNULL_BLOCK_IDpara rellenar las filas de CUDAGraph padding — el bloque 0 se reserva como padding📎 vllm/v1/worker/gpu_model_runner.py:2332-2334。

5.3 forward_context: descripción de lote compartida entre capas

Modelo intuitivo

forward_contextes el «tablón de anuncios unificado» pegado al frente del aula: cada capa del modelo puede levantar la vista y ver la disposición de asientos (attention metadata) y las reglas (slot mapping) de este examen, sin tener que preguntar cada una por su cuenta. Sin él, cada capa de atención tendría que recibir esta información desde los parámetros — pero laforwardfirma de la capa del modelo es fija y no permite pasar parámetros individualmente por capa.

Estructura de datos

ForwardContextes un@dataclass 📎 vllm/forward_context.py:141-202, campos principales:

• no_compile_layers: desdestatic_forward_contextcopia, marca las capas que no participan en la compilación📎 vllm/forward_context.py:132-137；
• attn_metadata: mapeo de nombre de capa a metadatos de atención, en modo DBO es una lista de longitud 2 (uno por microbatch)📎 vllm/forward_context.py:144-152；
• slot_mapping: mapeo de nombre de capa a tensor de slot mapping📎 vllm/forward_context.py:145；
• cudagraph_runtime_mode: modo CUDA graph en tiempo de ejecución, por defectoNONE 📎 vllm/forward_context.py:155-157；
• batch_descriptor: descriptor de lote, usado para el despacho de CUDA graph📎 vllm/forward_context.py:158；
• is_padding: máscara booleana en el eje de tokens,Trueindica filas de padding📎 vllm/forward_context.py:162-165。

BatchDescriptores otro@dataclass(frozen=True) 📎 vllm/forward_context.py:30-57, el diseño de campos sigue el principio de «minimizar los elementos descriptivos»:num_tokens、num_reqs(puede ser None en modo PIECEWISE),uniform(todas las solicitudes tienen el mismo número de tokens),has_lora、num_active_loras. El comentario explicanum_active_lorasla razón de ser de: cuandocudagraph_specialize_lora_countestá habilitado, cada valor de cantidad de LoRA captura un CUDA graph independiente, porquefused_moe_lorael grid size de kernels como depende de este valor📎 vllm/forward_context.py:60-64。

Singleton global y gestión de contexto

_forward_contextes una variable global a nivel de módulo📎 vllm/forward_context.py:199-201, a través deoverride_forward_contextel gestor de contexto guarda el valor anterior al entrar y lo restaura al salir📎 vllm/forward_context.py:263-274。set_forward_contextes una envoltura de nivel superior📎 vllm/forward_context.py:277-394, que además maneja la construcción de metadatos DP, la creación automática de batch descriptor y la inyección de kwargs específicos de la plataforma.

Paso a paso: desdeexecute_modelhasta el forward del modelo

Escenario:GPUModelRunner.execute_modelya tiene preparados todos los tensores de entrada y está a punto de invocar el modelo.

Enexecute_model,set_forward_contextse invoca📎 vllm/v1/worker/gpu_model_runner.py:4408-4420：

with (
    set_forward_context(
        attn_metadata,
        self.vllm_config,
        num_tokens=num_tokens_padded,
        num_tokens_across_dp=num_tokens_across_dp,
        cudagraph_runtime_mode=cudagraph_mode,
        batch_descriptor=batch_desc,
        ubatch_slices=ubatch_slices_padded,
        slot_mapping=slot_mappings,
        skip_compiled=has_encoder_input,
        is_padding=is_padding,
    ),
    ...
):
    model_output = self._model_forward(...)

set_forward_contextinternamente primero construyeDPMetadata(si DP o MoE con paralelismo de secuencia está habilitado)📎 vllm/forward_context.py:299-328, luego invocacreate_forward_contextpara construir la instanciaForwardContext📎 vllm/forward_context.py:347-358, y finalmente medianteoverride_forward_contextestablece la variable global📎 vllm/forward_context.py:361-362。

La capa del modelo medianteget_forward_context()lee📎 vllm/forward_context.py:208-214. Si no está establecido, la aserción falla y sugiere usarset_forward_context。

sequenceDiagram
    participant EC as EngineCore
    participant EX as Executor
    participant W as Worker
    participant MR as GPUModelRunner
    participant FC as ForwardContext
    participant M as Model Layers

    EC->>EX: execute_model(SchedulerOutput)
    EX->>W: collective_rpc("execute_model", args)
    W->>MR: execute_model(scheduler_output)
    MR->>MR: _update_states(scheduler_output)
    MR->>MR: _prepare_inputs(...)
    MR->>MR: _get_slot_mappings(...)
    MR->>FC: set_forward_context(attn_metadata, slot_mapping, ...)
    FC-->>MR: context manager entered
    MR->>M: _model_forward(input_ids, positions, ...)
    M->>FC: get_forward_context()
    FC-->>M: ForwardContext
    M-->>MR: hidden_states
    MR->>MR: compute_logits(sample_hidden_states)
    MR-->>W: ExecuteModelState / None
    W-->>EX: ModelRunnerOutput
    EX-->>EC: output[0]

Reflexión de diseño

is_paddingEl diseño del campo merece atención📎 vllm/forward_context.py:162-165: el comentario dice «los consumidores pueden usarlo para omitir el trabajo de los padding token». Esta es una optimización en el escenario de CUDA graph — las filas de padding participan en la captura del grafo pero no deben producir cómputo real.

all_moe_layersymoe_layer_indexson un par de workarounds ingeniosos📎 vllm/forward_context.py:170-195. El comentario explica en detalle el problema:vllm.moe_forwardlos operadores personalizados codifican la cadena del nombre de capa directamente en el grafo, lo que provoca tiempos de arranque en frío de torch.compile excesivamente largos. La solución es almacenar la lista de nombres de capa enForwardContext, y los operadores personalizados extraen las cadenas en orden e incrementan un contador. El comentario también admite que esto depende del supuesto de que «los operadores personalizados se ejecutan en orden y torch.compile no reordena»📎 vllm/forward_context.py:182-184。

Reflexión de diseño y escollos en producción

Consistencia de estado en programación asíncrona. _update_statesbajo decodificación especulativa asíncrona adopta una estrategia de «suposición optimista»: asume que todos los draft token del paso anterior fueron aceptados, primero expandeoutput_token_ids, y luego registra una función de corrección diferida📎 vllm/v1/worker/gpu_model_runner.py:1376-1384. La función de corrección se invoca después de que el forward del modelo se inicia📎 vllm/v1/worker/gpu_model_runner.py:1509-1510, lee el número real de aceptados desde la GPU y reviertenum_computed_tokens 📎 vllm/v1/worker/gpu_model_runner.py:1547-1558. La sutileza de este diseño es que: la corrección ocurre después de que «el lote ya se ha lanzado», no bloquea el forward y mantiene la continuidad del pipeline asíncrono.

_may_reorder_batchLa condición de activación de.Este método primero verificakv_cache_groupssi está vacío📎 vllm/v1/worker/gpu_model_runner.py:1131-1132. El comentario explica por qué no se puede simplemente verificaris_attention_free: El modelo Mamba también es attention-free, pero utiliza KV cache para guardar su estado interno📎 vllm/v1/worker/gpu_model_runner.py:1116-1139. Solo los modelos que realmente no tienen KV cache group omiten el reordenamiento.

_prepare_input_idsla trampa del cálculo de índices.Cuando en el lote hay tanto solicitudes decode del paso anterior como solicitudes nuevas,num_common_tokens < total_without_spec, es necesario copiar primero el tensor de CPU y luego hacer scatter📎 vllm/v1/worker/gpu_model_runner.py:1849-1854. Sinum_common_tokens == 0, significa que ninguna solicitud se superpone con el paso anterior, se retorna directamente📎 vllm/v1/worker/gpu_model_runner.py:1855-1858. La distinción entre estas dos ramas es crucial — omitir cualquiera de ellas provocará queinput_idsparte quede sin inicializar.

AsyncGPUModelRunnerOutputla sincronización de streams.La copia de salida se realiza en un stream CUDA independiente📎 vllm/v1/worker/gpu_model_runner.py:308-328, usandoblocking=Trueel Event de para evitar el busy polling del lock del driver CUDA📎 vllm/v1/worker/gpu_model_runner.py:296-298。get_output()en primero synchronize y luego liberar la referencia del tensor de dispositivo📎 vllm/v1/worker/gpu_model_runner.py:336-340, el orden no puede invertirse — de lo contrario el tensor podría ser reciclado antes de que la copia se complete.

Resumen del capítulo

Este capítulo rastreóSchedulerOutputla ruta completa desde EngineCore hasta el forward en GPU.ExecutorMediantecollective_rpcse difunde el resultado de la programación a todos los Workers,GPUModelRunnerel_update_statessincroniza el estado de caché,_prepare_inputsconstruye los tensores de entrada,_get_slot_mappingsgenera el mapeo de slots de KV, y finalmenteset_forward_contextinyecta la descripción del lote en el contexto global para que las distintas capas del modelo la consuman. La ruta de programación asíncrona mantiene la continuidad del pipeline mediante suposición optimista + corrección diferida, mientras queForwardContextel diseño de singleton global resuelve la contradicción entre la firma fija de las capas del modelo y la inyección de metadatos entre capas.

Reflexión y autoevaluación del capítulo

← Capítulo anterior: Capítulo 4

En el capítulo anterior vimos cómo GPUModelRunner traduce los resultados de la programación en tensores físicos como input_ids, slot_mapping y block_table, y los inyecta en cada capa mediante forward_context. Pero el verdadero consumidor de tiempo de GPU —el cálculo de atención— aún queda en el aire. ¿Quién consume exactamente esos tensores en attn_metadata? ¿Cómo pueden FlashAttention, FlashInfer y Triton ser intercambiables bajo el mismo código de modelo? La respuesta está en la capa de abstracción AttentionBackend. Esta desacopla "cómo se calcula la atención" de "cómo la invoca el modelo": la capa del modelo solo mantiene una referencia a AttentionImpl y llama a la interfaz unificada forward(query, key, value, kv_cache, attn_metadata, output); mientras que el backend concreto se encarga de traducir block_table, slot_mapping, seq_lens en parámetros que su propio kernel pueda consumir. Este capítulo toma FlashAttentionBackend como hilo principal, porque cubre simultáneamente la semántica de gather de PagedAttention, la compatibilidad con CUDA Graph, la atención en cascada, el contexto distribuido DCP y las ramas más ricas. Si lo entiendes a fondo, los demás backends son solo variantes de mapeo de parámetros. La motivación de diseño de "registro de backend + interfaz unificada" es directa: los kernels de atención evolucionan muy rápido (FA2→FA3→FA4, iteraciones de FlashInfer, Triton propio), y si la capa del modelo dependiera directamente de un kernel concreto, cada actualización del kernel requeriría modificar el código del modelo. La capa de abstracción aísla los cambios detrás de un único método de fábrica: get_impl_cls().

Selección de backend: declaración de capacidades y construcción de metadatos

Modelo intuitivo

Piensa enAttentionBackendcomo un anuncio de empleo: no hace el trabajo, solo declara "qué dtypes, qué head_size, qué formatos de cuantización de KV cache y qué tipos de atención puedo manejar". El planificador toma la configuración del modelo para hacer coincidir, y si falla, pasa al siguiente candidato. Sin esta capa de declaración, el sistema descubriría en tiempo de ejecución que "este kernel no soporta este head_size" y colapsaría directamente.

Matriz de capacidades: los campos son el contrato

FlashAttentionBackendLos atributos de clase de son sus límites de capacidad.supported_dtypeslimita a fp16/bf16📎 vllm/v1/attention/backends/flash_attn.py:287-287；supported_kv_cache_dtypespermite adicionalmente la serie fp8📎 vllm/v1/attention/backends/flash_attn.py:298-299. Pero "declarar soporte" no equivale a "soporte incondicional"—supports_kv_cache_dtypepara KV cuantizado delega además enflash_attn_supports_kv_cache_dtypepara hacer juicios dependientes del dispositivo📎 vllm/v1/attention/backends/flash_attn.py:431-438。

Más fino aún essupports_combination: recibe un conjunto completo de parámetros combinados como head_size, dtype, block_size, use_mla, has_sink, y devuelveNonepara indicar disponibilidad, o una cadena para indicar el motivo del rechazo📎 vllm/v1/attention/backends/flash_attn.py:454-507. Por ejemplo, sink se rechaza en capacidades < 9.0📎 vllm/v1/attention/backends/flash_attn.py:467-468, y en SM90 FP8 KV con mm_prefix debe ir por Triton📎 vllm/v1/attention/backends/flash_attn.py:472-472. Este diseño de "devolver cadena de motivo" permite que la capa superior dé errores diagnosticables en lugar de retroceder silenciosamente.

La elección de block_size también está impulsada por las capacidades. Por defecto devuelveMultipleOf(16), pero SM90 FP8-KV fuerza 64📎 vllm/v1/attention/backends/flash_attn.py:297-324, y el kernel FA4 con head_size=256 fuerzaFA4_HD256_PAGE_SIZE 📎 vllm/v1/attention/backends/flash_attn.py:326-352. Esto explica por qué el tamaño de bloque del KV cache no se define al azar—está restringido inversamente por el tamaño de tile TMA del kernel.

Estructura de metadatos: diseño de campos de FlashAttentionMetadata

FlashAttentionMetadataes un dataclass, los campos se dividen en cuatro grupos📎 vllm/v1/attention/backends/flash_attn.py:511-566：

El primer grupo es la descripción básica del lote:num_actual_tokens(número real de tokens sin padding),max_query_len、query_start_loc(suma prefija, usada por el kernel varlen para localizar el inicio y fin de cada secuencia),seq_lens、block_table、slot_mapping 📎 vllm/v1/attention/backends/flash_attn.py:520-526. Nótese el diagrama ASCII en los comentarios del código fuente📎 vllm/v1/attention/backends/flash_attn.py:512-518, que distingue con precisióncontext_len(KV histórico),query_len(nuevo en esta iteración),seq_len(la suma de ambos)—esta es la clave para entender los parámetros del kernel varlen.

El segundo grupo son los campos de atención en cascada:use_cascade、common_prefix_len、cu_prefix_query_lensetc.📎 vllm/v1/attention/backends/flash_attn.py:528-533。

El tercer grupo son los campos de DCP (Decode Context Parallel):max_dcp_context_kv_len、dcp_context_kv_lens, y los contadores que distinguen el número de solicitudes decode/prefill📎 vllm/v1/attention/backends/flash_attn.py:535-544。

El cuarto grupo son la programación opcional y máscaras especiales:scheduler_metadata(usado por la programación AOT de FA3),causal(puede ser bool o tensor, soporta causalidad por secuencia),mm_prefix_query_range_tensor(rangos bidireccionales multimodales), campos relacionados con R-SWA📎 vllm/v1/attention/backends/flash_attn.py:546-566。

Paso a paso de build()

Escenario: un lote mixto, 3 secuencias decode + 2 secuencias prefill, sin cascada, sin DCP.

Primer paso, desdecommon_attn_metadatadesempaquetar los tensores base📎 vllm/v1/attention/backends/flash_attn.py:824-832. Segundo paso, decidir si habilitar la programación AOT:aot_schedule = self.aot_schedule and not fast_build and not envs.VLLM_BATCH_INVARIANT 📎 vllm/v1/attention/backends/flash_attn.py:836-838。self.aot_scheduleen__init__es determinado porget_flash_attn_version() == 3— solo FA3 admite metadatos de programación precalculados. Tercer paso, rellenar de forma perezosa en el primer build📎 vllm/v1/attention/backends/flash_attn.py:709-709: recorrer todas lasaot_sliding_windowcapas para recopilar la configuración de ventana deslizante; si la configuración es única, adoptarla; si hay más de una, desactivar AOTFlashAttentionImplCuarto paso, calcular📎 vllm/v1/attention/backends/flash_attn.py:848-851。

. Por defecto 0 (para que FA3 use la heurística); solo se establece enmax_num_splitscuando se habilita full CUDA graph y el número de tokens está dentro del rango de capturaself.max_num_splits 📎 vllm/v1/attention/backends/flash_attn.py:856-866. El comentario explica la razón:num_splits > 1asigna[num_splits, num_heads, num_tokens, head_size]búferes intermedios, con alto coste de memoria; solo vale la pena en escenarios de CUDA graph📎 vllm/v1/attention/backends/flash_attn.py:862-865。

Quinto paso, tomar la rama no en cascada y no DCP, llamar a_get_scheduler_metadatapara generar los metadatos de programación de FA3📎 vllm/v1/attention/backends/flash_attn.py:976-986. Sexto paso,_store_scheduler_metadatamaneja el escenario de CUDA graph: copiar los nuevos metadatos al búfer preasignado y poner a cero la parte restante📎 vllm/v1/attention/backends/flash_attn.py:671-684. Este paso de puesta a cero es crucial — el comentario señala explícitamente que, de lo contrario, algunos thread blocks leerían metadatos inválidos y sobrescribirían el búfer de salida📎 vllm/v1/attention/backends/flash_attn.py:671-672。

Séptimo paso, construirFlashAttentionMetadatay devolver📎 vllm/v1/attention/backends/flash_attn.py:992-1015。

flowchart TD
    start["build(common_prefix_len, common_attn_metadata)"] --> unpack["解包 query_start_loc / seq_lens / block_table / slot_mapping"]
    unpack --> aot{"aot_schedule 且非 fast_build 且非 BATCH_INVARIANT?"}
    aot -->|是| sw_check{"aot_sliding_window 已初始化?"}
    aot -->|否| maxsplit
    sw_check -->|否, 首次| collect["_get_sliding_window_configs 收集层滑窗"]
    collect --> sw_unique{"配置数量 == 1?"}
    sw_unique -->|是| set_sw["设置 aot_sliding_window"]
    sw_unique -->|否, >1| disable_aot["self.aot_schedule = False"]
    set_sw --> maxsplit
    disable_aot --> maxsplit
    sw_check -->|是| maxsplit["计算 max_num_splits"]
    maxsplit --> cg_check{"use_full_cuda_graph 且 tokens <= max_cudagraph_size?"}
    cg_check -->|是| set_splits["max_num_splits = self.max_num_splits"]
    cg_check -->|否| zero_splits["max_num_splits = 0"]
    set_splits --> branch
    zero_splits --> branch
    branch{"dcp_world_size > 1?"}
    branch -->|是| dcp_path["计算 dcp_context_kv_lens, 可能 skip"]
    branch -->|否| cascade_check{"common_prefix_len > 0?"}
    cascade_check -->|是| cascade_path["构造 prefix/suffix 双份 scheduler_metadata"]
    cascade_check -->|否| normal_path["_get_scheduler_metadata 单份"]
    dcp_path --> store
    cascade_path --> store
    normal_path --> store
    store["_store_scheduler_metadata: CUDA graph 时拷入预分配缓冲并清零尾部"] --> build_meta["构造 FlashAttentionMetadata"]
    build_meta --> mm_check{"mm_req_doc_ranges 非空?"}
    mm_check -->|是| fill_mm["fill_mm_prefix_query_ranges + 拷贝到 GPU"]
    mm_check -->|否| rswa_check
    fill_mm --> rswa_check{"rswa_window 非空?"}
    rswa_check -->|是| copy_rswa["拷贝 prefix_lens 到持久缓冲"]
    rswa_check -->|否| done
    copy_rswa --> done["返回 attn_metadata"]

---

forward(): la cadena completa desde los metadatos hasta la invocación del kernel

Modelo intuitivo

forward()es el "taller de ensamblaje final" del backend: recibe las Q/K/V calculadas por las capas del modelo, los tensores de KV cache y los metadatos construidos en el paso anterior, ajusta el diseño físico del KV cache a la forma esperada por el kernel y luego lo despacha al kernel concreto. Sin este paso, el kernel leería un diseño de memoria incorrecto y produciría errores silenciosos en la salida — más difíciles de depurar que un crash.

Transformación del diseño de memoria del KV cache

La forma física del KV cache de vLLM es[num_blocks, num_kv_heads, block_size, 2 * head_size]— K y V concatenados en la última dimensión📎 vllm/v1/attention/backends/flash_attn.py:1246-1247. Pero los kernels de FlashAttention esperan K y V separados, y con diseño[num_blocks, block_size, num_kv_heads, head_size]。

La transformación ocurre al inicio deforward():kv_cache.transpose(1, 2).split(self.head_size, dim=-1) 📎 vllm/v1/attention/backends/flash_attn.py:1310-1310。transpose(1,2)convierte[blocks, heads, block_size, 2D]en[blocks, block_size, heads, 2D]，splitcortando K y V a lo largo de la última dimensión. Nótese quetransposesolo cambia el stride sin mover datos, por lo que los kernels posteriores deben admitir acceso no contiguo.

Inmediatamente después vienecanonicalize_singleton_dim_strides 📎 vllm/v1/attention/backends/flash_attn.py:1310-1310. El comentario señala el motivo: cuandonum_kv_heads=1(común en escenarios TP), el stride de las dimensiones de tamaño 1 es degenerado, y FA3/FA4 en H100+ usan TMA, que exige que el stride esté alineado a al menos 16 bytes📎 vllm/v1/attention/backends/flash_attn.py:1310-1310. Esta es una trampa típica de "lógicamente equivalente, físicamente inválido".

Flujo de parámetros en la ruta no en cascada

Tras entrar en la ramaif not attn_metadata.use_cascade, los parámetros se mapean uno a uno📎 vllm/v1/attention/backends/flash_attn.py:1326-1342：cu_seqlens_q = query_start_loc，seqused_k = seq_lens，block_table = attn_metadata.block_table。descale_shapetoma(batch_size, num_kv_heads), usado para la difusión de escala de cuantización FP8 — el comentario indica que flash-attn espera que la forma de descale sea(num_sequences, num_kv_heads), usando.expand()para evitar copias📎 vllm/v1/attention/backends/flash_attn.py:1258-1258。

Luego viene el tratamiento de simetrización de la ventana deslizante._maybe_symmetrize_windowlógica: la ventana deslizante causal(w, 0)en escenarios no causales debe convertirse en(w, w), para que la query bidireccional pueda mirar en ambas direcciones📎 vllm/v1/attention/backends/flash_attn.py:587-589. El comentario también enfatiza que "la window de la propia capa tiene prioridad sobre la window del group", porque un KV cache group puede albergar simultáneamente capas con ventana y capas globales (como cuando Gemma-3 desactiva hybrid KV cache manager)📎 vllm/v1/attention/backends/flash_attn.py:1362-1365。

Rama de máscara: mm_prefix y R-SWA

Cuandomm_prefix_query_rangesno está vacío y se cumplen las condiciones de FA4 + causal estático, el código construye elmask_mod 📎 vllm/v1/attention/backends/flash_attn.py:1374-1407de CuTE-DSL. Las acciones clave soncausal = Falseysliding_window_size = None 📎 vllm/v1/attention/backends/flash_attn.py:1406-1407. El comentario explica la razón: la semántica de mm_prefix es(causal ∧ window) ∨ bidirectional-range, no un subconjunto de causal; tras FA #155, establecer mask_mod ya no limpia automáticamente causal/local, y el llamador debe desactivarlo explícitamente, de lo contrario la ruta causal integrada cortocircuitaría mask_mod📎 vllm/v1/attention/backends/flash_attn.py:1402-1405。

_make_mm_prefix_mask_modusafunctools.cachecaché📎 vllm/v1/attention/backends/flash_attn.py:1793-1802. El comentario da una razón contundente: elhash_callablede FA4 mezcla elrepr()de la unidad de cierre en la clave de compilación; el_load_q_rangeanidado tiene una dirección diferente en cada llamada, lo que provoca una recompilación JIT completa en cada forward📎 vllm/v1/attention/backends/flash_attn.py:1793-1802. Este es un ejemplo típico de trampa de rendimiento en entornos de producción.

Dentro de la máscara hay un detalle de conversión de coordenadas: FA4 pasa elq_idxlocal (0-based dentro del chunk de prefill actual), mientras quekv_idxes la posición absoluta. El código usaq_abs = q_idx + seqlen_k - seqlen_qpara restaurar la posición absoluta📎 vllm/v1/attention/backends/flash_attn.py:1859-1865。__vec_size__ = 1también tiene su razón de ser:_load_q_rangelee el lane 0; una llamada no puede abarcar filas de query📎 vllm/v1/attention/backends/flash_attn.py:1897-1897。

El mask_mod de R-SWA es similar, pero la semántica escausal & (in_prefix | in_window) 📎 vllm/v1/attention/backends/flash_attn.py:1945-1948, yuse_fast_sampling = Truehace que FA4 omita los bloques KV completamente enmascarados, sin cargar sus datos📎 vllm/v1/attention/backends/flash_attn.py:1950-1950。

Tratamiento especial de FA4 hd256

Cuandoself.fa4_hd256es verdadero, el código fuerza la alineación de página:num_pages = cdiv(max_seqlen_k, FA4_HD256_PAGE_SIZE)，max_seqlen_kredondea hacia arriba al límite de página,block_tabletrunca al número exacto de páginas,num_splits = 1 📎 vllm/v1/attention/backends/flash_attn.py:1442-1448. El comentario indica que el kernel hd256 requiere longitud alineada a página, block table de ancho exacto y no admite SplitKV.

Finalmente se llama a_FA4_DENSE_ATTENTION_KERNEL(...), pasando q, k, v, out, cu_seqlens_q, seqused_k, block_table, softcap, mask_mod, aux_tensors, etc.📎 vllm/v1/attention/backends/flash_attn.py:1450-1475。

Escritura en KV cache: do_kv_cache_update

forward()solo lee el KV cache; la escritura la realizado_kv_cache_update. Este llama areshape_and_cache_flash, usandoslot_mappingpara escribir de forma dispersa las K/V recién calculadas en el cache📎 vllm/v1/attention/backends/flash_attn.py:1532-1541. El comentario señala:key/valueestá padded mientras queslot_mappingNo, pero no se requiere segmentación manual, porque el op usaslot_mappingla shape de📎 vllm/v1/attention/backends/flash_attn.py:1527-1531para determinar el número real de tokens. Aquí no se realiza normalización de stride, porque no participa ningún kernel TMA📎 vllm/v1/attention/backends/flash_attn.py:1520-1521。

sequenceDiagram
    participant Model as 模型层 Attention
    participant Impl as FlashAttentionImpl
    participant KVC as kv_cache 张量
    participant Kernel as flash_attn_varlen_func
    Model->>Impl: forward(query, key, value, kv_cache, attn_metadata, output)
    Impl->>Impl: output_scale 非空? 抛 NotImplementedError
    Impl->>Impl: attn_metadata is None? 返回 output.fill_(0)
    Impl->>KVC: transpose(1,2).split(head_size)
    KVC-->>Impl: key_cache, value_cache
    Impl->>Impl: canonicalize_singleton_dim_strides(key_cache)
    Impl->>Impl: use_cascade?
    alt 非级联
        Impl->>Impl: 映射 cu_seqlens_q / seqused_k / block_table
        Impl->>Impl: _maybe_symmetrize_window
        Impl->>Impl: mm_prefix 或 R-SWA? 构造 mask_mod
        Impl->>Kernel: _FA4_DENSE_ATTENTION_KERNEL(q, k, v, out, ...)
        Kernel-->>Impl: output 就地写入
    else 级联
        Impl->>Kernel: cascade_attention(prefix + suffix 两次调用)
        Kernel-->>Impl: merge_attn_states 合并
    end
    Impl-->>Model: output

---

Reflexión de diseño: por qué se escribió así

La compatibilidad con CUDA Graph es una restricción invisible del diseño de metadatos。_store_scheduler_metadataEl patrón de "copiar dentro + poner a cero la cola" de📎 vllm/v1/attention/backends/flash_attn.py:671-684aparece repetidamente en el búfer persistente de R-SWA📎 vllm/v1/attention/backends/flash_attn.py:787-798y en la zona temporal de mm_prefix📎 vllm/v1/attention/backends/flash_attn.py:800-813. El patrón común es: preasignar en__init__un búfer persistente del tamaño máximo,build()y en📎 vllm/v1/attention/backends/flash_attn.py:1044-1046。

solo copiar, sin asignar. La razón se señala en los comentarios: durante la captura de CUDA graph no puede haber operaciones de asignación。supports_draft_decode_metadata_update = self.dcp_world_size == 1 📎 vllm/v1/attention/backends/flash_attn.py:742-742Exclusión mutua entre DCP y fused draft decodeskip_dcp_context_attention(). El comentario explica: fused draft decode reutiliza entre pasos de draft los objetos de metadatos capturados, pero las decisiones del lado del host en tiempo de build de DCP (como📎 vllm/v1/attention/backends/flash_attn.py:736-741) cambian la forma de los metadatos, y estos campos de Python no se actualizan in situ entre replays del graph

. Esta es una compensación típica de "cuando el rendimiento y la corrección entran en conflicto, elegir la corrección".。use_cascade_attentionUmbral heurístico de la atención en cascada📎 vllm/v1/attention/backends/flash_attn.py:1967-1967Se filtran con una serie de umbrales: common_prefix_len < 256 se rechaza directamente📎 vllm/v1/attention/backends/flash_attn.py:1978-1979, alibi/sliding_window/local_attention no son compatibles📎 vllm/v1/attention/backends/flash_attn.py:1982-1984, número de solicitudes < 8 se rechaza📎 vllm/v1/attention/backends/flash_attn.py:1985-1987, escenario DCP deshabilitado📎 vllm/v1/attention/backends/flash_attn.py:2011-2029. Tras pasar, todavía hay que usar un modelo de rendimiento aproximado para comparar el número de CTA y el número de waves entre cascade y FlashDecoding📎 vllm/v1/attention/backends/flash_attn.py:2009-2010。

. El comentario admite que este modelo es "very rough"：forward()Puntos problemáticos en producciónview/sliceHay un comentario destacado que advierte que, bajo piece-wise CUDA graph, este método se ejecuta en modo eager,📎 vllm/v1/attention/backends/flash_attn.py:1277-1284y que métodos aparentemente sin operaciones de GPU como[:num_actual_tokens]son en realidad muy lentos; cualquier cambio debe ser benchmarkeado

---

. Esto explica por qué en el código se usa ampliamente

el slicing en lugar de formas más "elegantes": cada punto es el resultado de una compensación de rendimiento.FlashAttentionBackendResumen del capítulosupports_*Este capítulo recorrebuild()el ciclo de vida completo del backend de atención: declaración de capacidades (CommonAttentionMetadataserie) → construcción de metadatos (FlashAttentionMetadatatraduceforward()atranspose+split) → invocación del kernel (

transforma el layout del KV cache, construye máscaras, despacha al kernel FA). Los mecanismos centrales incluyen: la transformación de layout

del KV cache, la normalización de strides degenerados, el patrón de búfer persistente bajo CUDA graph, la construcción de máscaras CuTE-DSL para mm_prefix/R-SWA, y la decisión heurística de la atención en cascada.logitsPrincipio de diseño clave: separación entre declaración de capacidades e implementación, preasignación de metadatos impulsada por la compatibilidad con CUDA graph, y prioridad a la corrección cuando el rendimiento y la corrección entran en conflicto (DCP deshabilita fused draft decode).

El siguiente capítulo pasa a muestreo y salida:

Hasta aquí, la cadena completa desde la interfaz abstracta hasta la implementación del kernel del backend de atención ya está conectada: la capa del modelo llama de forma unificada a través de AttentionImpl, el backend se encarga de traducir metadatos como block_table y slot_mapping a parámetros concretos del kernel, y la implementación de PagedAttention de FlashAttentionBackend muestra la semántica de gather bajo KV Cache paginado y la estrategia de compatibilidad con CUDA Graph. Pero el cálculo de atención solo produce estados ocultos; lo que el modelo finalmente debe emitir es el siguiente token. ¿Cómo se convierten estos estados ocultos en logits, y cómo pasan los logits por muestreo y postprocesamiento hasta devolverse finalmente al cliente como texto en streaming? El siguiente capítulo seguirá este último tramo.

En el capítulo anterior seguimos cómo el backend de atención traduce el block table a parámetros del kernel y completa el cálculo de atención tipo gather sobre memoria no contigua. Pero la atención solo produce estados ocultos: lo que el modelo realmente debe entregar al usuario es el texto del siguiente token. Este capítulo sigue ese último tramo: una vez que los estados ocultos se proyectan a logits mediante lm_head, cómo atraviesan una cadena de procesadores cuidadosamente ordenada (temperatura, penalizaciones, top-k/top-p, restricciones estructuradas), se muestrean a token id y luego, mediante el detokenizer, se restauran a texto y se envían en streaming. Cualquier paso desordenado o fuga de estado en esta cadena degradará silenciosamente la calidad de la salida.

Sampler: el orden de la cadena de procesadores es la corrección

Modelo intuitivo: el Sampler es como una línea de ensamblaje, y los logits son la pieza en bruto por procesar. Cada estación de la línea (processor) modifica la pieza, y el orden de las estaciones determina directamente el producto final: primero cortar y luego pulir no da lo mismo que primero pulir y luego cortar. Sin esta cadena, el modelo solo podría emitir la distribución de probabilidad original, y el usuario obtendría un "muestreo desnudo" sin control de temperatura, sin supresión de repeticiones y sin restricciones de formato.

Estructuras de datos y diseño de memoria

El propio Sampler esnn.Module, pero su estado central es extremadamente delgado: solo contiene el submódulotopk_topp_sampler, la banderalogprobs_modeyuse_fp64_gumbel.📎 vllm/v1/sample/sampler.py:61-64. Todo el estado real a nivel de lote está encapsulado enSamplingMetadata, y se pasa mediante los parámetros de forward. Este diseño de "Sampler sin estado + metadatos externos" es deliberado: la instancia de Sampler se crea una sola vez durante el ciclo de vida del motor, mientras que la composición del lote cambia en cada decode step; externalizar el estado es lo que permite que el Sampler se reproduzca de forma segura tras ser capturado por CUDA Graph.

La constante clave es_SAMPLING_EPS = 1e-5 📎 vllm/v1/sample/sampler.py:18. Cumple simultáneamente dos semánticas: una temperatura inferior a este valor se considera greedy, yapply_temperaturesirve como respaldo para evitar la división por cero.

Step-by-Step Walkthrough

Escenario concreto: en un batch se mezclan solicitudes greedy con solicitudes de muestreo aleatorio, y algunas además tienen logprobs activados.

Primer paso, tomar una instantánea de los logprobs originales.Antes de aplicar cualquier penalización o temperatura, si la solicitud necesita logprobs, primero se decide el contenido de la instantánea segúnlogprobs_mode.📎 vllm/v1/sample/sampler.py:84-93. Nótese que el comentario señala explícitamente la diferencia con V0: V1 usalogits originales(antes de penalizaciones y temperatura) para calcular los top-k logprobs.📎 vllm/v1/sample/sampler.py:72-77. Este es el contrato semántico: el logprob que ve el usuario debe reflejar la distribución real del modelo, no una distribución distorsionada por penalizaciones.

Segundo paso, unificar a float32. 📎 vllm/v1/sample/sampler.py:95-96Independientemente de si la entrada es bf16 o fp16, se convierte a float32. La razón es que el log_softmax, top-k y la probabilidad acumulada posteriores acumulan errores en baja precisión, especialmente cuando el vocab alcanza los 150 mil.

Tercer paso, cadena de procesadores que no alteran el argmax. apply_logits_processorsSe aplican secuencialmente: máscara de lista blanca de allowed token, exclusión de bad words,non_argmax_invariantprocesadores, términos de penalización📎 vllm/v1/sample/sampler.py:391-404. La clasificación aquí es el diseño central——non_argmax_invariantse refiere a aquellosque cambian el resultado greedyprocesadores (como min_tokens, logit_bias), que deben aplicarse antes del muestreo greedy; mientras que losargmax_invariantprocesadores (como min_p) no cambian el argmax y pueden posponerse hasta después de la temperatura.

Cuarto paso, muestreo. sampleEl método primero determina si es completamente aleatorio📎 vllm/v1/sample/sampler.py:256-271: siall_greedy, retorna directamente argmax; de lo contrario, primero calcula el resultado greedy como respaldo, luego aplica temperatura, procesadores que no alteran el argmax, top-k/top-p📎 vllm/v1/sample/sampler.py:275-291. Finalmente usatorch.wherepara seleccionar entre el resultado greedy y el aleatorio según el umbral de temperatura📎 vllm/v1/sample/sampler.py:305-306, y reutiliza el tensorgreedy_sampledcomo búfer de salida, evitando asignaciones adicionales.

Quinto paso, recolectar logprobs y encapsular la salida.Segúnnum_logprobshay tres casos: None solo retorna los logprobs del token especificado; -1 retorna todos los logprobs sin ordenar; de lo contrario top-k📎 vllm/v1/sample/sampler.py:120-131. Finalmente, el token id se convierte a int32 para comprimir el tamaño, y se expande a un tensor bidimensional[num_requests, 1]de📎 vllm/v1/sample/sampler.py:138-148。

flowchart TD
    in_logits["logits (bf16/fp16)"] --> snap{"需要 logprobs?"}
    snap -->|是| raw["compute_logprobs / clone<br/>raw_logprobs 快照"]
    snap -->|否| f32
    raw --> f32["logits.to(float32)"]
    f32 --> proc["apply_logits_processors"]
    proc --> mask{"allowed_token_ids_mask?"}
    mask -->|是| fill["masked_fill_(-inf)"]
    mask -->|否| bad
    fill --> bad{"bad_words_token_ids?"}
    bad -->|是| apply_bad["apply_bad_words"]
    bad -->|否| noninv
    apply_bad --> noninv["non_argmax_invariant 处理器"]
    noninv --> pen["apply_penalties"]
    pen --> sample["sample()"]
    sample --> allg{"all_greedy?"}
    allg -->|是| greedy["greedy_sample (argmax)"]
    allg -->|否| temp["apply_temperature"]
    temp --> arginv["argmax_invariant 处理器"]
    arginv --> topp["topk_topp_sampler"]
    topp --> where["torch.where(temp < EPS)"]
    greedy --> out
    where --> out["SamplerOutput<br/>sampled_token_ids"]

Reflexiones de diseño y errores comunes

¿Por qué los términos de penalización deben ir antes de la temperatura?La temperatura es un escalado de la distribución, la penalización es una suma o resta de puntos a tokens específicos. Si se escala primero y luego se penaliza, la magnitud absoluta de la penalización se amplifica o reduce por la temperatura, causando que el mismo conjunto de parámetros de penalización se comporte de manera inconsistente a diferentes temperaturas. V1 fija la penalización antes de la temperatura, garantizando la estabilidad semántica de los parámetros.

mark_unbackedLa trampa de compilación deEngather_logprobs,batched_count_greater_thanse compila, y cuando la dimensión batch cambia de 1 a ≥2 se dispara una recompilación por especialización 0/1 de dynamo📎 vllm/v1/sample/sampler.py:345-348。mark_unbackedmarca esa dimensión como completamente simbólica, evitando esta recompilación. En producción, si se observa un bloqueo repentino tras la primera solicitud de decode, muy probablemente sea este tipo de recompilación.

gpu_sync_allowedEl límite de sincronización de batched_count_greater_thaninternamente puede disparar una sincronización de GPU, vLLM usa el contextogpu_sync_allowed(first_only=True)para declarar explícitamente "aquí se permite sincronización, pero solo la primera vez"📎 vllm/v1/sample/sampler.py:345-348. Si se sincroniza inesperadamente dentro de una región de captura de CUDA Graph, causará que la captura falle——esta es la pista clave para diagnosticar problemas de captura de grafos.

Salida estructurada: máquina de estados de doble vía con máscara de bits y gramática

Modelo intuitivo: la salida estructurada es como ponerle al muestreador un par de "gafas gramaticales"——en cada paso solo puede ver los tokens que cumplen con el JSON schema o la gramática. Sin ellas, el modelo podría generar JSON con errores de sintaxis y el parser downstream colapsaría directamente. La esencia de la implementación de vLLM es: la máquina de estados gramatical avanza en el lado de CPU, mientras que las restricciones se pasan al muestreo del lado de GPU en forma de máscara de bits.

Estructuras de datos y diseño de memoria

StructuredOutputManageres un singleton a nivel de motor, que poseebackend(uno de xgrammar/guidance/outlines/lm-format-enforcer),reasoner_clsy dos pools de hilos📎 vllm/v1/structured_output/__init__.py:39-98。

La máscara de bits es la estructura de datos central:_grammar_bitmaskes un tensor int32 de forma[max_batch_size * (1 + max_num_spec_tokens), vocab_size/32]. Cada bit corresponde a si un token es válido.📎 vllm/v1/structured_output/__init__.py:327-336representa "todo 1"——todos los tokens válidos_full_mask = torch.tensor(-1, dtype=torch.int32)Los dos pools de hilos tienen una división clara:📎 vllm/v1/structured_output/__init__.py:59。

se encarga de la compilación de gramática (intensivo en CPU, número de workers es la mitad de los núcleos de CPU)executorse encarga del llenado paralelo de máscaras de bits para batches grandes, solo se habilita cuando el batch supera 128📎 vllm/v1/structured_output/__init__.py:71-78；executor_for_fillmaskInicialización de gramática.📎 vllm/v1/structured_output/__init__.py:62-69。

Step-by-Step Walkthrough

Cuando una solicitud entra por primera vez,es invocadogrammar_init. Si el backend no está inicializado, se selecciona la implementación según la configuración📎 vllm/v1/structured_output/__init__.py:115-176. Luego se envía la tarea de compilación: por defecto va por la vía asíncrona📎 vllm/v1/structured_output/__init__.py:130-165, pero en el modoexecutor.submitdebe ser síncronaexternal_launcherGeneración de máscara de bits.📎 vllm/v1/structured_output/__init__.py:167-176。

En cada decode step,genera máscaras para todas las solicitudes estructuradas del batchgrammar_bitmask. Los batches grandes van por la ruta paralela: se envían al pool de hilos en grupos de 16📎 vllm/v1/structured_output/__init__.py:314-442. Los batches pequeños van por la ruta serial, avanzando el estado gramatical token por token📎 vllm/v1/structured_output/__init__.py:346-373Alineación de máscaras bajo decodificación especulativa.📎 vllm/v1/structured_output/__init__.py:374-433。

Esta es la parte más ingeniosa. Cuando hay draft tokens, cada solicitud necesitafilas de máscara. La ruta serial procesa token por token: si un draft token es rechazado por la gramática, se registra1 + max_num_spec_tokens, y las filas posteriores copian directamente la máscara de esa filafailed_index. Esto garantiza que "tras el rechazo de un draft, el estado de restricción de las posiciones posteriores retrocede al punto de rechazo".📎 vllm/v1/structured_output/__init__.py:396-418Reversión de estado.

Durante el llenado de la máscara de bits, el estado gramatical avanzópasos, pero los draft tokens aún no han sido realmente aceptados, por lo que se debestate_advancementsretrocedergrammar.rollback(state_advancements). La aceptación real ocurre en📎 vllm/v1/structured_output/__init__.py:422-430copiaaccept_tokens 📎 vllm/v1/structured_output/__init__.py:444-466。

sequenceDiagram
    participant Sched as Scheduler
    participant Mgr as StructuredOutputManager
    participant Pool as executor_for_fillmask
    participant Gram as StructuredOutputGrammar
    participant GPU as GPU Runner

    Sched->>Mgr: grammar_bitmask(requests, ids, spec_tokens)
    Mgr->>Mgr: allocate_token_bitmask(max_batch*(1+spec))
    alt batch > 128 且无投机
        Mgr->>Pool: _async_submit_fill_bitmask(batch)
        Pool->>Gram: fill_bitmask(bitmask, index)
        Gram-->>Pool: 写入合法 token 位
        Pool-->>Mgr: Future.result()
    else 小 batch 或含投机
        loop 每个 req 的每个 spec token
            Mgr->>Gram: fill_bitmask(bitmask, cumulative_index)
            Mgr->>Gram: accept_tokens(req_id, [token])
            Gram-->>Mgr: True/False
            Note over Mgr: 失败则记录 failed_index<br/>后续行复制该行
        end
        Mgr->>Gram: rollback(state_advancements)
    end
    Mgr-->>Sched: bitmask.numpy() (NDArray int32)
    Sched->>GPU: 传入采样内核

¿Por qué external_launcher debe compilar de forma síncrona?

El comentario da la razón precisa: la compilación asíncrona haría que las transiciones de estado deocurrieran en momentos diferentes en distintos TP ranks, rompiendo la suposición de determinismo de la que depende external_launcherWAITING_FOR_STRUCTURED_OUTPUT_GRAMMAR → WAITING 状态转换在不同 TP rank 上发生于不同时刻，破坏 external_launcher 依赖的确定性假设 📎 vllm/v1/structured_output/__init__.py:47-56. Este es un caso típico del conflicto entre determinismo distribuido y optimización asíncrona.

Punto de inicio de la restricción bajo el modelo de razonamiento. _get_constraint_startDetermina desde qué token comenzar a aplicar la restricción gramatical📎 vllm/v1/structured_output/__init__.py:220-292. Para modelos con cadena de pensamiento, la fase de reasoning no debe estar sujeta a restricciones JSON; solo se activa tras finalizar el reasoning.enable_in_reasoningCuando es True, devuelve directamente 0 (restricción durante todo el proceso)📎 vllm/v1/structured_output/__init__.py:235-236. Si el reasoner soportafind_reasoning_end_offset, úsalo para localizar con precisión📎 vllm/v1/structured_output/__init__.py:261-267; de lo contrario, recurre a la búsqueda de retroceso token por token📎 vllm/v1/structured_output/__init__.py:287-291。

validate_tokensla semántica de prefijo deEn decodificación especulativa, los draft tokens pueden violar la gramática,validate_tokensdevuelve el "prefijo legal más largo"📎 vllm/v1/structured_output/__init__.py:294-312. Nótese que primero elimina el relleno especulativo (-1), luego calcula el punto de inicio de la restricción y, finalmente, solo realiza la validación gramatical sobre los tokens dentro del intervalo restringido.

Detokenizer: el juego de fronteras entre decodificación incremental y stop string

Modelo intuitivo: el detokenizer es como un escriba que transcribe carácter por carácter, traduciendo token ids a texto legible por humanos. La dificultad radica en que: los tokens y los caracteres no tienen correspondencia uno a uno (un token puede corresponder solo a medio carácter UTF-8), y el stop string puede abarcar múltiples tokens. Sin decodificación incremental, cada paso requeriría decodificar toda la secuencia desde el principio, y el costo O(n²) degradaría el throughput.

Estructuras de datos y diseño de memoria

IncrementalDetokenizerLa clase base solo contienetoken_idslista📎 vllm/v1/engine/detokenizer.py:32-33。BaseIncrementalDetokenizerañade campos relacionados con stop:stoplista,min_tokens、include_stop_str_in_output、stop_buffer_lengthy_last_output_text_offset 📎 vllm/v1/engine/detokenizer.py:70-94。

stop_buffer_lengthson clave: cuando el stop string no está incluido en la salida, equivale a la longitud del stop string más largo menos uno📎 vllm/v1/engine/detokenizer.py:87-90. Este "búfer de retroceso" asegura que la salida en streaming no emita prematuramente caracteres que podrían ser prefijo de un stop string.

Dos rutas de implementación:FastIncrementalDetokenizerUsa la librería tokenizersDecodeStream 📎 vllm/v1/engine/detokenizer.py:166-246；SlowIncrementalDetokenizerUsa el lado de Pythondetokenize_incrementally 📎 vllm/v1/engine/detokenizer.py:249-305. El criterio de selección es que la versión de tokenizers sea ≥ 0.22.0 y el tipo de tokenizer coincida📎 vllm/v1/engine/detokenizer.py:32-33📎 vllm/v1/engine/detokenizer.py:61-63。

Step-by-Step Walkthrough

decodificación incremental. updateRecibe nuevos token ids ystop_terminatedflag📎 vllm/v1/engine/detokenizer.py:96-142. Si stop termina y no incluye el stop string, el último token queda excluido de la decodificación📎 vllm/v1/engine/detokenizer.py:107-111. Luego, token por token, llama adecode_nextacumula texto📎 vllm/v1/engine/detokenizer.py:117-122。

detección de stop string. check_stop_stringsSolo busca dentro del rango de caracteres nuevos📎 vllm/v1/engine/detokenizer.py:308-360. El punto de inicio de búsqueda es1 - new_char_count - stop_string_len 📎 vllm/v1/engine/detokenizer.py:338, este desplazamiento asegura que los stop strings que cruzan fronteras de tokens también sean capturados. Cuando múltiples stop strings coinciden simultáneamente, se eligeel que se complete primero📎 vllm/v1/engine/detokenizer.py:342-347。

segmentación de salida en streaming. get_next_output_textSegún el parámetrodeltadecide si devolver todo o solo el incremento📎 vllm/v1/engine/detokenizer.py:148-163. Si no está completo, retienestop_buffer_lengthcaracteres sin emitir📎 vllm/v1/engine/detokenizer.py:145-146, usa_last_output_text_offsetpara registrar la posición ya enviada📎 vllm/v1/engine/detokenizer.py:148-163。

recuperación de excepciones. FastIncrementalDetokenizer._protected_stepManeja dos tipos de excepciones: OverflowError/TypeError registra en log y devuelve None📎 vllm/v1/engine/detokenizer.py:225-229; el error "Invalid prefix" entoncesreconstruye DecodeStreamy reintenta📎 vllm/v1/engine/detokenizer.py:222-246. Este último aborda el caso límite en que el tokenizer produce salida UTF-8 no monótona.

Reflexiones de diseño y trampas

El equilibrio de stop_buffer_length.Cuanto más largo el búfer, mayor la latencia del streaming (el tiempo hasta que el usuario ve el texto se pospone), pero menos probable es pasar por alto stop strings que cruzan tokens. Tomar "la longitud del stop string más largo menos uno" es la cota inferior exacta: cualquier prefijo de un stop string tiene como máximo esa longitud.

min_tokens y stop_check_offset.Cuando el número de tokens de salida no alcanzamin_tokens,stop_check_offsetse empuja continuamente hasta el final del texto📎 vllm/v1/engine/detokenizer.py:120-122, lo que significa que este fragmento de texto no será sometido a detección de stop. Esto evita que el modelo choque con un stop string al inicio y produzca una salida vacía.

Caché de added_token_ids en la ruta Fast.Cuandospaces_between_special_tokenses False, es necesario suprimir los espacios entre tokens especiales📎 vllm/v1/engine/detokenizer.py:192-207. El código almacena en cachéadded_token_idsen el objeto tokenizer📎 vllm/v1/engine/detokenizer.py:195-200, evitando reconstruir el diccionario en cada decode.

Reflexiones de diseño

Los tres módulos comparten una filosofía de diseño:separar el avance de estado de la verificación de restricciones, dejando que el lado de la GPU solo realice operaciones tensoriales sin estado. El Sampler no tiene estado; el estado está enSamplingMetadata; la máquina de estados gramatical avanza en el lado de la CPU, y la GPU solo consume la máscara de bits; el_last_output_text_offsetdel detokenizer es el único cursor de streaming. Esta separación permite que cada componente del lado de la GPU sea capturado por CUDA Graph.

Otra línea principal esel orden es semántica. El orden de la cadena de procesadores del Sampler, el punto de inicio de la restricción de la salida estructurada, el desplazamiento de detección de stop del detokenizer: cualquier error de orden no provocará un crash, solo producirá silenciosamente resultados erróneos; esto es precisamente lo más difícil de depurar en este tipo de código.

Resumen del capítulo

• La cadena de procesadores del Sampler se ordena estrictamente: instantánea de logprobs originales → float32 → lista blanca/bad words → non-argmax-invariant → penalizaciones → temperatura → argmax-invariant → top-k/top-p.
• La salida estructurada usa máscaras de bits para pasar el estado sintáctico del lado de la CPU a la GPU; bajo decodificación especulativa, mediantefailed_indexcopia yrollbackse garantiza la consistencia del estado.
• El Detokenizer usastop_buffer_lengthun búfer de retroceso para equilibrar la latencia del streaming con la detección de stop strings entre tokens; la ruta Fast depende de tokenizers ≥ 0.22.0 deDecodeStream。

Reflexiones y autoevaluación de este capítulo

del diseño original cubre exactamente la ventana mínima necesaria de "caracteres nuevos + posible prefijo de stop string que cruza el límite", garantizando que no se omita ninguna detección y evitando falsas coincidencias con el historial.

Estado de verificación: líneas FACT ancladas a ubicaciones reales

En el capítulo anterior recorrimos el último tramo del ciclo de vida de una inferencia única, desde el muestreo de logits hasta la salida en streaming. Pero cuando el modelo es demasiado grande para caber en una sola tarjeta, este pipeline debe dividirse entre múltiples dispositivos para ejecutarse de forma coordinada. La cuestión primordial de la inferencia distribuida no es "cómo particionar el modelo", sino "una vez particionado, quién habla con quién y de qué manera". vLLM delega estas dos cuestiones respectivamente a la topología de grupos de procesos de parallel_state.py y a la implementación del comunicador de custom_all_reduce.py. Este capítulo sigue la cadena "crear grupos → particionar → comunicar → reequilibrar carga" para desglosar capa por capa las estrategias de paralelismo de TP, PP y EP y las primitivas de comunicación subyacentes.

8.1 Topología de grupos de procesos: cómo se divide una malla de ranks en TP/PP/DP/EP

Modelo intuitivonew_group, aparecerá un desajuste de comunicación del tipo "creía que estabas en el grupo TP, pero en realidad estás en el grupo DP" — una vez que falta un rank en la comunicación colectiva, NCCL se colgará directamente en lugar de reportar un error.

Estructura de datos y diseño de memoria

GroupCoordinatores el portador de todo esto. El diseño de sus campos corresponde directamente a la "múltiple identidad de un proceso en múltiples dimensiones paralelas":

• rankes el rank global,rankses la lista de ranks globales de los miembros del grupo,world_sizees el tamaño del grupo📎 vllm/distributed/parallel_state.py:434-436。
• local_rankse usa para vincular el dispositivo,rank_in_groupes el índice dentro del grupo — el código fuente usa una tabla para distinguir con precisión ambos: en un grupo de 4 GPUs que abarca dos nodos, el rank 2 tienelocal_rankes 0 (es la primera GPU en el nodo 1), perorank_in_groupes 2📎 vllm/distributed/parallel_state.py:437-445。
• cpu_groupydevice_groupexisten en pares: el primero usa gloo para comunicación de metadatos/objetos, el segundo usa NCCL para comunicación de tensores📎 vllm/distributed/parallel_state.py:446-447。

Aquí hay un diseño clave:¿Por qué cada grupo debe mantener un grupo de CPU?Porquebroadcast_object、send_objecteste tipo de operaciones transmiten objetos de Python (bytes serializados), usar NCCL desperdicia memoria de GPU y puede contaminar el dispositivo CUDA actual.barrier()Los comentarios de lo dicen de forma muy directa: el barrier de NCCL internamente es un broadcast, que crea tensores de GPU a escondidas, fácilmente desordena el dispositivo actual, por lo que es obligatorio usar el grupo de CPU📎 vllm/distributed/parallel_state.py:1355-1362。

Step-by-Step：initialize_model_parallelCómo dividir la malla

Tomemos un escenario concreto: 8 GPUs, TP=2, PP=4, DP=1. Lo esencial es reestructurar la secuencia unidimensional de ranks en una malla multidimensional, y luego dividir a lo largo de cada dimensión.

Primer paso, construir la malla de ranks. El orden de diseño se define explícitamente comoExternalDP x DP x PP x PCP x TP 📎 vllm/distributed/parallel_state.py:2045-2060：

all_ranks = torch.arange(world_size).reshape(
    -1, data_parallel_size, pipeline_model_parallel_size,
    prefill_context_model_parallel_size, tensor_model_parallel_size,
)

Segundo paso, dividir el grupo TP: ver la malla como(-1, tp_size)luego unbind, obteniendo[g0,g1],[g2,g3],... 📎 vllm/distributed/parallel_state.py:2065-2077. Nota que el grupo TP pasa adicionalmenteuse_message_queue_broadcaster=True, porque el grupo TP necesita compartir memoria broadcast para distribuir metadatos.

Tercer paso, dividir el grupo PP:all_ranks.transpose(2, 4)mover la dimensión PP a la última dimensión y luego dividir, obteniendo[g0,g2,g4,g6],[g1,g3,g5,g7] 📎 vllm/distributed/parallel_state.py:2175-2188. Este es exactamente el ejemplo dado en el docstring📎 vllm/distributed/parallel_state.py:1997-1997。

Cuarto paso, dividir el grupo DP:transpose(1, 4)luego dividir📎 vllm/distributed/parallel_state.py:2195-2202。

Quinto paso, dividir el grupo EP — aquí hay un detalle fácil de pasar por alto: el grupo EP solo se crea bajo modelos MoE, los modelos dense lo omiten directamente📎 vllm/distributed/parallel_state.py:2210-2241. El conjunto de ranks del grupo EP esDP x PCP x TPel producto de , lo que significa que EP reutiliza las GPUs físicas de DP y TP, en lugar de ser una dimensión independiente.

flowchart TD
    start["initialize_model_parallel()"] --> grid["all_ranks = arange(world_size).reshape(-1, DP, PP, PCP, TP)"]
    grid --> tp["TP: view(-1, tp_size).unbind(0)"]
    grid --> pp["PP: transpose(2,4).reshape(-1, pp_size)"]
    grid --> dp["DP: transpose(1,4).reshape(-1, dp_size)"]
    grid --> ep_check{"model_config.is_moe?"}
    ep_check -->|是| ep["EP: transpose(1,2).reshape(-1, DP*PCP*TP)"]
    ep_check -->|否| skip["_EP 保持 None"]
    ep --> eplb_check{"enable_eplb?"}
    eplb_check -->|是| eplb["EPLB: 与 EP 同 rank 集，独立 PG"]
    eplb_check -->|否| no_eplb["_EPLB 保持 None"]
    tp --> done["logger.info_once 打印各维度 rank"]
    pp --> done
    dp --> done
    ep --> done
    skip --> done
    eplb --> done
    no_eplb --> done

Reflexiones de diseño y trampas

¿Por qué EPLB necesita un grupo de procesos independiente?Los comentarios dan la respuesta: aislar la comunicación de EPLB de la comunicación colectiva del forward de MoE, para prevenir que "el torch.distributed en tiempo de ejecución" y "el torch.distributed de EPLB" se bloqueen mutuamente📎 vllm/distributed/parallel_state.py:2243-2246. Esta es una compensación típica de "intercambiar un dominio de comunicación independiente por determinismo" — el costo de memoria de un PG adicional, a cambio de no bloquear el forward durante la transferencia de pesos.

Restricción de sincronización del grupo DPes la trampa más común en entornos de producción: todos los ranks dentro del mismo grupo DP deben llamar simultáneamente agenerate, de lo contrario hay deadlock📎 vllm/distributed/parallel_state.py:2048-2051. Porque dentro del grupo DP se hace all-reduce de gradientes/resultados de muestreo, cualquier rank ausente bloqueará permanentemente la comunicación colectiva.

Orden de destruccióntambién tiene sus matices.destroy()Primero destruir el device communicator, luego destruir device_group y cpu_group📎 vllm/distributed/parallel_state.py:1380-1393. Los comentarios explican la razón: el device communicator puede mantener áreas de trabajo de comunicación colectiva que dependen de estos PG (como el barrier IPC PCIe de FlashInfer), deben liberarse primero📎 vllm/distributed/parallel_state.py:1377-1377。

8.2 Primitivas de comunicación: cómo el all-reduce personalizado evita NCCL

Modelo intuitivo

El all-reduce de NCCL es un "camión de carga general", puede llevar cualquier carga, ir por cualquier camino, pero el costo de arranque y el costo de protocolo son fijos. Cuando necesitas hacer repetidamente all-reduce de tensores pequeños en una máquina de 8 GPUs totalmente interconectadas por NVLink (cada capa de attention/MLP de TP lo necesita), el "peaje" del camión de carga general se vuelve innegable. El all-reduce personalizado es un "carrito especializado": solo se habilita en escenarios intra-nodo, con NVLink totalmente interconectado y tamaño de tensor adecuado, usando una vezcudaMemcpypara reemplazar el handshake y el costo de protocolo de NCCL.

Estructura de datos y diseño de memoria

CustomAllreduceLa inicialización de es una combinación de "sondeo de capacidades + preasignación de recursos". Campos clave:

• _SUPPORTED_WORLD_SIZES = [2, 4, 6, 8, 16]: solo soporta estos tamaños de grupo📎 vllm/distributed/device_communicators/custom_all_reduce.py:113-129。
• meta_ptrs: metadatos de sincronización + búfer de resultados intermedios, tamañoops.meta_size() + max_size 📎 vllm/distributed/device_communicators/custom_all_reduce.py:291-294。
• buffer_ptrs: búfer IPC pre-registrado, en modo eager el tensor de entrada se copia primero y luego se calcula📎 vllm/distributed/device_communicators/custom_all_reduce.py:298-305。
• rank_data: tensor uint8 de 8MB, almacena las tuplas de punteros de búfer IPC de todos los ranks📎 vllm/distributed/device_communicators/custom_all_reduce.py:309-315。

¿Por qué pre-registrar los búferes?Porque la captura de CUDA Graph requiere que todas las direcciones estén fijas en el momento de la captura.register_graph_buffersAl final de la captura, difunde todas las direcciones de búfer utilizadas a todos los ranks y las registra📎 vllm/distributed/device_communicators/custom_all_reduce.py:474-491。

Paso a paso: el flujo de decisión de un all-reduce

Escenario: la salida de una capa MLP dentro del grupo TP necesita all-reduce, la entrada es un tensor bf16 de 4MB.

Primer paso,custom_all_reduceverificar si está deshabilitado, si cumpleshould_custom_ar 📎 vllm/distributed/device_communicators/custom_all_reduce.py:529-533。

Segundo paso,should_custom_arFiltrado elemento por elemento: se rechaza si world_size > 8; dtype debe ser fp32/fp16/bf16; el número de bytes debe ser múltiplo de 16; debe ser débilmente contiguo; solo se continúa si world_size==2 o si hay interconexión total📎 vllm/distributed/device_communicators/custom_all_reduce.py:493-508。

Tercer paso, derivar según si se está en captura de CUDA Graph: durante la captura se usaregistered=True(dirección ya fijada), de lo contrarioregistered=False(es necesario hacer memcpy primero a un búfer pre-registrado)📎 vllm/distributed/device_communicators/custom_all_reduce.py:529-545。

Cuarto paso, llamar realmente aops.all_reduce, pasandobuffer_ptrs[rank]ymax_size 📎 vllm/distributed/device_communicators/custom_all_reduce.py:519-527。

flowchart TD
    call["custom_all_reduce(input)"] --> disabled{"self.disabled?"}
    disabled -->|是| ret_none["return None → 回退 NCCL"]
    disabled -->|否| should{"should_custom_ar(input)?"}
    should -->|否| ret_none
    should -->|是| capturing{"self._IS_CAPTURING?"}
    capturing -->|是| stream_cap{"is_current_stream_capturing()?"}
    stream_cap -->|是| reg["all_reduce(registered=True)"]
    stream_cap -->|否| mimic["return empty_like(input) 模拟分配"]
    capturing -->|否| eager["all_reduce(registered=False) 先 memcpy"]
    reg --> out["返回 out 张量"]
    eager --> out

Reflexiones de diseño y trampas encontradas

La ruta de degradación en escenarios multi-nodoes la parte más ingeniosa de este código.same_nodeCuando es falso,mnnvl_onlyse pone a verdadero📎 vllm/distributed/device_communicators/custom_all_reduce.py:198-199, luego se verifica la capacidad MNNVL (Multi-Node NVLink). Si no todas las tarjetas del grupo soportan MNNVL, se deshabilita directamente la comunicación colectiva personalizada📎 vllm/distributed/device_communicators/custom_all_reduce.py:228-233。_group_can_attempt_mnnvlSe usa un all-reduce de CPU (operación MIN) para asegurar que todos los ranks sigan el mismo flujo de control📎 vllm/distributed/device_communicators/custom_all_reduce.py:59-73—esta es la protección clave en clústeres heterogéneos para evitar que "algunos ranks entren en la ruta MNNVL y otros vayan por NCCL" y provoquen un cuelgue.

El coste de la comprobación P2P：_can_p2precorre todos los peers haciendogpu_p2p_access_check, el comentario dice que el primer cálculo es caro pero se cachea📎 vllm/distributed/device_communicators/custom_all_reduce.py:278-278. En producción, si se detecta un arranque lento, se puede configurarVLLM_SKIP_P2P_CHECKpara omitirlo y confiar directamente en el informe P2P del driver📎 vllm/distributed/device_communicators/custom_all_reduce.py:86-100。

Selección de backend en tres niveles para reduce-scattermerece una mirada aparte:_select_reduce_scatter_backenddevuelve por prioridadmnnvl_multimem > mnnvl_lamport > legacy 📎 vllm/distributed/device_communicators/custom_all_reduce.py:601-636. La ruta multimem requiere que world_size esté en(2,4,8)y que la capacidad del dispositivo sea (10,0) o (10,3) (nivel Blackwell)📎 vllm/distributed/device_communicators/custom_all_reduce.py:103-104. Nótese queVLLM_BATCH_INVARIANTdeshabilita la ruta multimem📎 vllm/distributed/device_communicators/custom_all_reduce.py:628—porque el orden de reducción de multimem es no determinista y rompería la invariancia de lote.

8.3 EPLB: lógica de planificación del reequilibrio de carga de expertos

Modelo intuitivo

En un modelo MoE, 256 expertos lógicos se reparten entre 32 tarjetas, 8 por tarjeta. Pero bajo tráfico real, algunos "expertos populares" (por ejemplo, los que procesan estructuras gramaticales comunes) reciben una gran cantidad de tokens enrutados, lo que convierte a la tarjeta que los aloja en cuello de botella mientras las demás quedan ociosas. EPLB (Expert Parallel Load Balancer) consiste precisamente en "añadir réplicas a los expertos populares": copiar los pesos de los expertos populares a tarjetas ociosas para que los tokens se distribuyan hacia allí. Sin él, el throughput real de MoE quedaría limitado por la tarjeta más lenta.

Estructuras de datos y disposición de memoria

EplbModelStateutiliza tres tablas de mapeo para describir la relación "experto lógico ↔ experto físico":

• physical_to_logical_map: forma(num_moe_layers, num_physical_experts), cada ranura física almacena el id del experto lógico que aloja📎 vllm/distributed/eplb/eplb_state.py:105-120。
• logical_to_physical_map: forma(num_moe_layers, num_logical_experts, max_replicas+1), matriz dispersa, -1 indica que no hay mapeo📎 vllm/distributed/eplb/eplb_state.py:123-146。
• logical_replica_count: cuántas réplicas tiene cada experto lógico📎 vllm/distributed/eplb/eplb_state.py:147-161。

expert_load_windowes una ventana deslizante, forma(window_size, num_moe_layers, num_physical_experts) 📎 vllm/distributed/eplb/eplb_state.py:180-187. El comentario señala especialmente que ahora se registra la carga de todos los expertos físicos y no solo la de los expertos locales, para garantizar que las estadísticas sean consistentes entre distintos métodos de dispatch (naive all-to-all, DeepEP); bajo naive all-to-all, cada rank de DP aporta el mismo conjunto de tokens, por lo que la carga se multiplica por dp_size📎 vllm/distributed/eplb/eplb_state.py:180-187。

Paso a paso: la cadena completa de una reorganización

Escenario de ejemplo:expert_rearrangement_stepalcanza el umbral, se dispararearrange()。

Primer paso, mapear la carga física de vuelta a los expertos lógicos. Se usascatter_add_para agregar segúnphysical_to_logical_map, las ranuras inválidas (<0) se rellenan en el bucketinvalid_idxy finalmente se descartan📎 vllm/distributed/eplb/eplb_state.py:794-816。

Segundo paso, all-reduce entre ranks para obtener la carga lógica global._allreduce_listse concatenan las cargas de varios modelos, se hace un solo all-reduce y luego se separan, evitando múltiples comunicaciones📎 vllm/distributed/eplb/eplb_state.py:1045-1068。

Tercer paso, llamar a la estrategia para calcular el nuevo mapeo.policy.rebalance_expertsse ejecuta en el host, así que tanto la ventana de carga como el mapeo actual deben copiarse de vuelta a la CPU📎 vllm/distributed/eplb/eplb_state.py:859-867。

Cuarto paso, juicio de "omitir reorganización" específico de ROCm: si la mejora en el desequilibrio de carga entre ranks que aporta el nuevo mapeo es inferior al 5%, se omite esta reorganización📎 vllm/distributed/eplb/eplb_state.py:869-923. Es una optimización pragmática: la reorganización en sí tiene coste de comunicación, y si el beneficio no es suficiente, no se hace.

Quinto paso, ejecutar el traslado de pesos y confirmar el nuevo mapeo📎 vllm/distributed/eplb/eplb_state.py:925-942。

sequenceDiagram
    participant Main as 主线程 step()
    participant Policy as DefaultEplbPolicy
    participant Comm as EplbCommunicator
    participant Async as async_worker 线程
    Main->>Main: expert_rearrangement_step >= interval
    Main->>Main: scatter_add_ 物理负载→逻辑负载
    Main->>Main: _allreduce_list 跨 rank 聚合
    Main->>Policy: rebalance_experts(load, replicas, groups, nodes, gpus, map)
    Policy-->>Main: new_physical_to_logical_map
    alt 同步模式
        Main->>Comm: rearrange_expert_weights_inplace()
        Comm-->>Main: 权重搬运完成
        Main->>Main: _commit_eplb_maps()
    else 异步模式
        Main->>Main: eplb_stats = EplbStats(...); rebalanced = True
        Main->>Async: rearrange_event.record()
        Async->>Comm: 后台搬运权重到 expert_buffer
        Async-->>Main: pending_result 就绪
        Main->>Main: _move_to_workspace() 提交
    end

Reflexiones de diseño y trampas encontradas

La primitiva de sincronización en modo asíncronoes el punto más delicado de este código.rebalancedEl flag depende del GIL para sincronizarse entre el hilo principal y el worker async📎 vllm/distributed/eplb/eplb_state.py:194-203. Pero el comentario advierte:rebalanceddebe mantenerse consistente en todos los ranks, de lo contrario el all-reduce dentro de_all_ranks_result_readyse colgará📎 vllm/distributed/eplb/eplb_state.py:664-665。_all_ranks_result_readySe prioriza usar el grupo de CPU para el all-reduce, porque el grupo de CPU es más fiable📎 vllm/distributed/eplb/eplb_state.py:1024-1043。

La optimización de "grabación anticipada" de la ventana deslizante：_should_record_current_stepsolo activa la grabación cuando faltan como máximowindow_sizepasos para la siguiente reorganización📎 vllm/distributed/eplb/eplb_state.py:689-709. El comentario explica: los datos de losstep_interval - window_sizepasos previos a cada ciclo de reorganización serán sobrescritos por la ventana deslizante, así que grabarlos es inútil y desperdicia cómputo de GPU📎 vllm/distributed/eplb/eplb_state.py:1196-1199。should_record_tensores el mismo tensor escalar compartido por todas las capas, una solafill_actualiza todas las capas📎 vllm/distributed/eplb/eplb_state.py:272-278。

Reserva de capacidad para EP elástico：enable_elastic_epcuando,physical_expert_capacityse reserva segúnelastic_ep_max_dp_size, la tabla de mapeo rellena las ranuras sobrantes con -1📎 vllm/distributed/eplb/eplb_state.py:375-386. Así, al escalar no hace falta reasignar memoria de vídeo, solo rellenar las ranuras con -1 con expertos reales.reconfigure_physical_expert_slotsse encarga de refrescar la vista al escalar hacia arriba/abajo📎 vllm/distributed/eplb/eplb_state.py:1135-1160。

_commit_eplb_mapsEl manejo de pin memory en: cuandoPIN_MEMORYestá activado y el origen está en CPU, primero se copia a memoria pinned y luegonon_blocking=Truese copia asíncronamente a la GPU📎 vllm/distributed/eplb/eplb_state.py:1392-1400. Esto evita que la copia H2D bloquee el hilo principal: la tabla de mapeo se actualiza en cada capa y en cada ronda, y una copia síncrona se convertiría en cuello de botella.

Reflexiones de diseño

Tres bloques de código comparten una filosofía de diseño:intercambiar detección de capacidades por degradación determinista。GroupCoordinatorenworld_size == 1se hace bypass directamente de toda comunicación colectiva📎 vllm/distributed/parallel_state.py:736-738；CustomAllreducesi alguna condición no se cumple, se devuelveNonepermitiendo que el llamador recurra a NCCL📎 vllm/distributed/device_communicators/custom_all_reduce.py:532-533; EPLB omite la redistribución cuando la mejora es inferior al 5%📎 vllm/distributed/eplb/eplb_state.py:916. Este patrón de "fallo rápido + degradación elegante" permite que el mismo código se ejecute en todo el espectro de hardware, desde una sola GPU hasta MNNVL multi-nodo, sin necesidad de escribir ramas para cada configuración.

Otra característica común esla consistencia del flujo de control por encima del rendimiento。_group_can_attempt_mnnvlse usa all-reduce de CPU para forzar que todos los ranks sigan la misma rama📎 vllm/distributed/device_communicators/custom_all_reduce.py:59-73，_all_ranks_result_readyde manera similar📎 vllm/distributed/eplb/eplb_state.py:1024-1043. En sistemas distribuidos, "algunos ranks toman la ruta rápida y otros la ruta lenta" es mucho más peligroso que "todos los ranks toman la ruta lenta": lo primero provoca un cuelgue, lo segundo solo es lento.

Resumen del capítulo

• GroupCoordinatorse remodela la secuencia unidimensional de ranks en unaExternalDP x DP x PP x PCP x TPcuadrícula, dividiendo a lo largo de cada dimensión los grupos de procesos TP/PP/DP/EP/EPLB; cada grupo mantiene simultáneamente dos PG: CPU (gloo) y device (NCCL).
• CustomAllreducemediante detección de capacidades (misma máquina, NVLink totalmente interconectado, tamaño del tensor, dtype, alineación a 16 bytes) se decide si se asume el all-reduce, degradando a MNNVL o NCCL en escenarios multi-nodo.
• EPLB utiliza tres tablas de mapeo para describir la relación entre expertos lógicos/físicos, mediante ventanas deslizantes se estadística la carga, la estrategia calcula el nuevo mapeo, el comunicador transporta los pesos, y soporta modos síncrono y asíncrono.
• El principio de diseño común a los tres: detección de capacidades + degradación determinista + consistencia del flujo de control como prioridad.

Reflexiones y autoevaluación del capítulo

Hasta aquí, hemos aclarado los mecanismos de creación de grupos, partición y reequilibrio de carga para la comunicación entre tarjetas. Pero el desafío de comunicación en la inferencia distribuida no se limita al interior de una sola instancia: cuando prefill y decode se separan en instancias distintas, el KV Cache necesita transferirse entre nodos. En el próximo capítulo dejaremos la "comunicación entre tarjetas" para entrar en la "comunicación entre instancias": cómo se transfiere el KV Cache entre instancias de prefill y decode en un despliegue desagregado, y cómo la abstracción KV Connector unifica backends de transferencia como NIXL, Mooncake, etc.

En el capítulo anterior fijamos la perspectiva dentro de una única instancia de inferencia: cómo se forman los grupos de procesos TP/PP/DP/EP, cómo se particionan los tensores entre tarjetas y cómo EPLB reequilibra expertos en la capa MoE. Pero todos estos mecanismos se basan en una misma premisa: prefill y decode se ejecutan en la misma instancia, y el KV Cache permanece en la memoria local de la GPU de principio a fin. El despliegue desagregado (Prefill-Decode Disaggregation, abreviado como separación PD) rompe esta premisa. Divide prefill y decode en dos instancias vLLM independientes: la instancia de prefill solo realiza el cálculo forward del prompt, produce el KV Cache y lo entrega a la instancia de decode; la instancia de decode toma este KV Cache y continúa la generación autorregresiva. La ventaja de esto es que los recursos pueden configurarse de forma independiente según las características de cada fase: prefill es intensivo en cómputo, adecuado para TP grande y batch grande; decode es intensivo en acceso a memoria, adecuado para batch pequeño y programación de baja latencia. Ambos dejan de obstaculizarse mutuamente. El costo es que el KV Cache debe transferirse entre instancias. Este es el protagonista de este capítulo: KV Connector. El comentario de cabecera del archivo vllm/distributed/kv_transfer/kv_connector/v1/base.py ya enumera las primitivas centrales de toda la abstracción: el lado del Scheduler se encarga de vincular metadatos, consultar aciertos de caché remota y decidir si liberar bloques de forma asíncrona; el lado del Worker se encarga de la carga y guardado real del KV. El objetivo de diseño de esta interfaz es desacoplar por completo la lógica de programación de la capa superior de los backends de transferencia subyacentes (NIXL, Mooncake, MoRIIO). Desde una perspectiva de ingeniería, el mayor riesgo de la separación PD no es la lentitud de la transferencia, sino la inconsistencia de estado: la instancia de prefill cree que el KV ya se envió, pero la instancia de decode no lo recibió; o la instancia de decode libera el bloque antes de tiempo mientras prefill todavía está escribiendo en él. Lo que este capítulo busca esclarecer es precisamente cómo este sistema de conectores utiliza protocolos de handshake, leases, heartbeats y mecanismos de recuperación ante fallos para cubrir estos casos límite.

I. KVConnectorBase_V1: abstracción de doble rol y contrato de metadatos

Modelo intuitivo

KV Connector es como un sistema de mensajería entre dos sucursales. La tienda de Prefill calcula el producto semiterminado (KV Cache), lo empaqueta y lo envía a la tienda de Decode para continuar el procesamiento. Pero un sistema de mensajería no puede tener solo la acción de "enviar": necesita una guía de envío (metadata) que indique qué se envía y a dónde; necesita un mecanismo de acuse de recibo que confirme que el destinatario lo recibió; y también necesita un conjunto de reglas de timeout para evitar que un paquete quede atascado en el camino ocupando espacio en el estante.

Sin esta abstracción, cada backend de transferencia (NIXL, Mooncake) tendría que implementar su propia lógica de programación, y el Scheduler de vLLM tendría que escribir código de adaptación para cada backend. El valor de KVConnectorBase_V1 es fijar este contrato.

Doble rol: lado del Scheduler y lado del Worker

📎 vllm/distributed/kv_transfer/kv_connector/v1/base.py:137-142Se definen los dos roles del conector:

class KVConnectorRole(enum.Enum):
    # Connector running in the scheduler process
    SCHEDULER = 0
    # Connector running in the worker process
    WORKER = 1

Esta división no es arbitraria. El proceso del Scheduler se encarga de las decisiones globales de programación: qué solicitudes necesitan transferirse, cuándo se puede liberar un bloque; el proceso del Worker se encarga del traslado real de datos. Ambos se comunican medianteKVConnectorMetadatacomunicación.

📎 vllm/distributed/kv_transfer/kv_connector/v1/base.py:153-158Se define la clase base de metadatos en la dirección del Scheduler al Worker:

class KVConnectorMetadata(ABC):  # noqa: B024
    """Abstract Metadata used to communicate
    Scheduler KVConnector -> Worker KVConnector.
    """
    pass

En la dirección inversa, del Worker al Scheduler,📎 vllm/distributed/kv_transfer/kv_connector/v1/base.py:161-176se defineKVConnectorWorkerMetadata, que requiere implementar elaggregatemétodo, porque en un engine step puede haber múltiples workers devolviendo metadatos cada uno, y es necesario agregarlos antes de entregarlos al Scheduler.

Estructura de datos central: KVConnectorTransferResults

📎 vllm/distributed/kv_transfer/kv_connector/v1/base.py:87-96Se define la estructura de instantánea de los resultados de transferencia:

@dataclass
class KVConnectorTransferResults:
    finished_sending: set[str] = field(default_factory=set)
    finished_recving: set[str] = field(default_factory=set)
    failed_recving: set[str] = field(default_factory=set)

Nótese el diseño clave en los comentarios:Las recepciones fallidas también aparecerán enfinished_recvingdentro de. Esto es para que el Scheduler pueda liberar la solicitud del estado de "esperando transferencia" — incluso si la transferencia falla, la solicitud no puede quedarse atascada para siempre. La información de fallo se transmite por separado mediantefailed_recving, y el Scheduler decide en función de ello si reintentar o degradar.

Hooks del ciclo de vida: de la solicitud a la liberación

Todo el ciclo de vida del conector gira en torno a varios hooks clave. Del lado del Scheduler:

• get_num_new_matched_tokens 📎 vllm/distributed/kv_transfer/kv_connector/v1/base.py:485-518: consulta cuántos tokens puede acertar la caché remota. El comentario enfatiza especialmente que "solo debe considerarse el prefijo máximo realmente disponible"; si algunos tokens no se pueden obtener por problemas de conexión o desalojo, no deben contarse.
• update_state_after_alloc 📎 vllm/distributed/kv_transfer/kv_connector/v1/base.py:520-544: actualiza el estado tras la asignación de bloques. En el comentario hay una trampa fácil de pisar — para determinar si se debe cargar hay que mirarnum_external_tokens, y no siblocksestá vacío, porque los subconectores no seleccionados de MultiConnector también reciben bloques reales.
• request_finished 📎 vllm/distributed/kv_transfer/kv_connector/v1/base.py:579-598: se llama cuando la solicitud se completa, devuelveTruepara indicar que el conector asume la responsabilidad de la liberación asíncrona del bloque.

Del lado del Worker:

• start_load_kv / wait_for_layer_load: carga capa por capa, soporta pipeline.
• save_kv_layer / wait_for_save: guarda capa por capa.
• get_transfer_results 📎 vllm/distributed/kv_transfer/kv_connector/v1/base.py:396-397: devuelve el estado de finalización de la transferencia asíncrona.

📎 vllm/distributed/kv_transfer/kv_connector/v1/base.py:192-201También hay un diseño fácil de pasar por alto pero crucial —requires_kv_deliveryatributo:

@property
def requires_kv_delivery(self) -> bool:
    """Whether this connector hands off KV that must be reliably delivered.
    ...
    """
    return self._kv_transfer_config.is_kv_producer

El comentario explica la motivación: si una solicitud es expropiada antes de que se complete el traspaso de KV, debe recalcularse en lugar de dejarla completarse y traspasar bloques que ya fueron liberados por la expropiación. Solo el rol de producer necesita entrega confiable; si la caché best-effort se pierde, solo será un cache miss futuro.

Metadatos de handshake

📎 vllm/distributed/kv_transfer/kv_connector/v1/base.py:145-150define la clase base de los metadatos de handshake:

class KVConnectorHandshakeMetadata(ABC):  # noqa: B024
    """Metadata used for out of band connector handshake between
    P/D workers. This needs to serializable.
    """
    pass

"out of band" significa que el handshake no sigue la ruta normal de la solicitud, sino que se comunica directamente entre los workers P/D. Esto prepara el terreno para el protocolo de handshake ZMQ de NIXL.

---

II. Conector NIXL: handshake, registro y construcción de descriptores

Modelo intuitivo

NIXL (NVIDIA Inference Xfer Library) es la biblioteca de transporte de bajo nivel proporcionada por NVIDIA, que soporta múltiples backends como UCX, GDS, etc. El rol de NixlBaseConnectorWorker es como el centro de clasificación de una empresa de mensajería — primero necesita establecer una línea dedicada con el centro de clasificación de la otra parte (handshake), registrar la disposición de sus estanterías (registrar las regiones de memoria de KV Cache), y solo entonces puede recoger y enviar mercancías eficientemente por dirección.

Sin este mecanismo, cada transferencia tendría que renegociar direcciones y restablecer conexiones, y la latencia sería inaceptablemente alta.

Diseño de memoria: Region y Descriptor

Los conceptos centrales de NIXL sonregion(región de memoria) ydescriptor(descriptor). Cada capa de KV Cache se registra en NIXL como una o más regions, y cada region tiene dirección base, longitud de bloque y stride de bloque.

📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:740-751enumera los campos centrales relacionados con region:

# Number of NIXL regions. Currently one region per cache
# (so 1 per layer for MLA, otherwise 2 per layer)
self.num_regions = 0
self.region_mem_types: list[str] = []
self.region_group_ids: list[int] = []
self._uses_region_group_mapping = False
self.region_names: list[str] = []
self.region_num_blocks: list[int] = []
self._mixed_mem_types = False

📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:897-900explica además el origen del stride de bloque:

# Per-region block stride in bytes. Taken from the registered tensor's
# stride(0) so it stays correct under layouts that interleave layers
# within a block (BLHNC/BHLNC), where stride > block_len.
self.block_stride_per_layer = list[int]()

La idea clave aquí es:block_stride no es igual a block_len. En diseños con intercalado entre capas como BLHNC/BHLNC, el span real de un bloque puede ser mayor que la longitud de sus datos válidos. Si se usa directamente block_len como stride, se leerán direcciones incorrectas.

Protocolo de handshake: ZMQ + hash de compatibilidad

El handshake es la parte más compleja del conector NIXL.📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:974-1128El método_nixl_handshakede

muestra este proceso por completo.📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:988-998El primer paso es establecer el contexto del dispositivo CUDA.

# the first time we connect to a remote agent.
# be careful, the handshake happens in a background thread.
# it does not have an active cuda context until any cuda runtime
# call is made. when UCX fails to find a valid cuda context, it will
# disable any cuda ipc communication, essentially disabling any NVLink
# communication.
if not self.use_host_buffer:
    current_platform.set_device(self.device_id)

explica la razón:

Copiar📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:1029-1036：

msg = msgspec.msgpack.encode(
    (GET_META_MSG, remote_pp_rank, remote_rank)
)
# Set receive timeout to 5 seconds to avoid hanging on dead server
sock.setsockopt(zmq.RCVTIMEO, 5000)  # milliseconds
start_time = time.perf_counter()
sock.send(msg)
reply_parts = sock.recv_multipart()

El segundo paso es enviar la consulta de metadatos mediante ZMQ.📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:1042-1045Copiar

El timeout de 5 segundos es para evitar esperar indefinidamente si el par muere. Al mismo tiempo, el código usa RTT para estimar el desfase de reloj📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:1063-1080：

assert self.compat_hash is not None
if (
    self.enforce_compat_hash
    and handshake_payload.compatibility_hash != self.compat_hash
):
    raise RuntimeError(
        f"NIXL compatibility hash mismatch. "
        ...
    )

El tercer paso es la verificación de compatibilidad.📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:1372-1376Copiar

self.compat_hash = compute_nixl_compatibility_hash(
    self.vllm_config,
    self.backend_name,
    transfer_mode=self._TRANSFER_MODE,
)

:transfer_modeCopiar📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:163-166Nótese que

también participa en el hash —

el comentario de📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:824-835：

self._handshake_initiation_executor = ThreadPoolExecutor(
    # NIXL is not guaranteed to be thread-safe, limit 1 worker.
    max_workers=1,
    thread_name_prefix="vllm-nixl-handshake-initiator",
)
self._ready_requests = queue.Queue[tuple[ReqId, ReqMeta]]()
self._handshake_futures: dict[
    EngineId, Future[tuple[dict[tuple[int, int], str], float]]
] = {}
# Protects _handshake_futures and _remote_agents.
self._handshake_lock = threading.RLock()

max_workers=1Programación asíncrona del handshake_handshake_lockEl handshake es asíncrono y se ejecuta mediante un pool de hilos._handshake_futuresCopiar_remote_agentses porque NIXL no garantiza la seguridad de hilos.

_ensure_handshake 📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:1257-1317protege los dos diccionarios

y

.📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:172-310implementa el inicio idempotente del handshake: si ya se completó con éxito, devuelve None directamente; si está en proceso de handshake, devuelve el Future existente; de lo contrario, envía una nueva tarea y registra el callback._compute_desc_idsConstrucción de descriptores: de block ID a descriptor NIXL

Una vez completado el handshake, se deben construir descriptores para cada solicitud.📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:226-262El

# NOTE (NickLucche) With HMA, every kv group has the same number of layers
# and layers from different groups share the same kv tensor.
# eg block_ids=[[1, 2], [3]]->blocks [1, 2] need to be
# read across all regions, same for [3], but group0-group1 blocks will
# always differ (different areas). Therefore we can just flatten the
# block_ids and compute the descs ids for all groups at once.

es el núcleo.📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:285-304：

elif _is_ssm_spec(spec_type):
    # NOTE (NickLucche) SSM and Attention block regions can
    # be exchanged arbitrarily by manager.  Therefore, descs
    # are laid out as:
    #   [descs_fa (all regions) | descs_ssm (all regions)].
    # num_fa_descs offset must be computed per-engine since
    # P and D can have different num_blocks (and thus
    # different FA desc counts).

. El comentario explica el manejo en escenarios HMA:

Copiar📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:2130-2178Para modelos SSM híbridos, el diseño de descriptores es más complejoadd_remote_agentCopiar

Cuando D.world_size > P.world_size, múltiples workers D leen diferentes fragmentos de KV head desde el mismo worker P. La documentación proporciona un ejemplo concreto: D TP=4, P TP=2, tp_ratio=2. D-Worker0 lee la primera mitad de los KV head de P-Worker0, D-Worker1 lee la segunda mitad.

Para modelos MLA, el KV Cache se replica entre los workers TP, por lo que rank_offset siempre es 0.

Lease y heartbeat: evitar la liberación prematura de bloques

Este es uno de los diseños más ingeniosos del conector NIXL. Después de que la instancia Prefill envía el KV, no puede liberar el bloque inmediatamente, porque la instancia decode podría seguir leyendo. Pero si nunca se libera, la memoria de video se filtrará.

La solución es el lease.📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:528-528：

kv_lease_duration: int = vllm_config.kv_transfer_config.get_from_extra_config(
    "kv_lease_duration", 30
)
# NOTE (NickLucche): For now we use a hardcoded value for a simpler interface.
self._lease_extension = kv_lease_duration * 2 // 3

El lease predeterminado es de 30 segundos, y cada heartbeat lo extiende 20 segundos (2/3).

El manejo del heartbeat está en📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:3014-3034：

def _handle_heartbeat(self, payload: str) -> None:
    new_expiry = time.perf_counter() + self._lease_extension
    for req_id in payload.split(","):
        if req_id in self._reqs_to_send:
            old = self._reqs_to_send[req_id]
            self._reqs_to_send[req_id] = max(old, new_expiry)

Notamax(old, new_expiry)——el heartbeat solo puede extender el lease, no acortarlo.

La recuperación después de que expire el lease está en📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:2986-3012：

def _reap_expired_send_leases(self, done_sending: set[str]) -> None:
    """Reclaim expired send-side KV leases into ``done_sending``.

    ``_reqs_to_send`` is not ordered by expiry: heartbeats update the
    deadline in place, and mixed TTLs share the map, so a live head
    entry can sit in front of already-expired ones. Scan every entry
    rather than stopping at the first still-live request.
    """

El comentario señala un error fácil de cometer: no se puede detener el escaneo solo porque se encontró la primera solicitud no expirada, porque el heartbeat actualiza el tiempo de expiración in situ, lo que hace que el map no esté ordenado por tiempo de expiración.

Máquina de estados de transferencia y recuperación de fallos

El ciclo de vida de la transferencia se gestiona mediante_pop_done_transfers 📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:3036-3086:

for handle in handles:
    try:
        xfer_state = self.nixl_wrapper.check_xfer_state(handle)
        if xfer_state == "DONE":
            res = self.nixl_wrapper.get_xfer_telemetry(handle)
            self.xfer_stats.record_transfer(res)
            self.nixl_wrapper.release_xfer_handle(handle)
        elif xfer_state == "PROC":
            in_progress.append(handle)
        else:
            self._log_failure(
                failure_type="transfer_failed",
                req_id=req_id,
                xfer_state=xfer_state,
            )

La transferencia NIXL tiene tres estados:DONE(completado),PROC(en progreso), otros (fallo).

El manejo de fallos está en📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:3103-3127：

def _handle_failed_transfer(
    self,
    req_id: str,
    handle: int | None,
    failed_req_ids: set[str] | None = None,
    record_failed_transfer: bool = True,
) -> bool:
    if record_failed_transfer:
        self.xfer_stats.record_failed_transfer()
    if failed_req_ids is not None:
        failed_req_ids.add(req_id)
    return handle is None or self._try_release_xfer_handle(req_id, handle)

_try_release_xfer_handle 📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:3088-3101El comentario de es crucial:

except Exception as e:
    # A status error does not guarantee that the backend stopped DMA.
    self._log_failure(
        failure_type="transfer_release_failed",
        msg="Retaining handle and blocks until release succeeds",
        ...
    )
    return False

Un estado de error no garantiza que el backend haya detenido el DMA. Si la liberación falla, se deben conservar el handle y el bloque hasta que la liberación tenga éxito. Este es un diseño típico de "prefiero filtrar que usar incorrectamente".

Manejo de bloques de solicitudes fallidas

Cuando la recepción falla,📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:2876-2891muestra la lógica de manejo:

for req_id in done_recving:
    meta = self._recving_metadata.pop(req_id, None)
    assert meta is not None, f"{req_id} not found in recving_metadata list"

    # Skip KV sync and post-processing for failed requests
    if req_id in failed_recv_reqs:
        self._pending_recv_notifs.pop(req_id, None)
        # TODO (NickLucche) handle failed transfer for HMA.
        if not self._is_hma_required:
            self._invalid_block_ids.put(set(meta.local_block_ids[0]))
        logger.warning(
            "Skipping KV post-processing for failed request %s",
            req_id,
        )
        continue

El ID del bloque fallido se coloca en la cola_invalid_block_ids, y el Scheduler lo extrae medianteget_block_ids_with_load_errors 📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:3491-3504para decidir si reintentar.

Expulsión por TTL de motores remotos

Las instancias de larga duración encontrarán continuamente nuevos motores remotos; si no se limpian, la memoria crecerá indefinidamente.📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:3506-3532El de_evict_stale_enginesimplementa la expulsión por TTL:

def _evict_stale_engines(self) -> None:
    """Scan for and evict remote engines that have exceeded their TTL.

    Called from the main thread in when a new remote engine appears.
    We can only go OOM as we discover and register a new remote, therefore we make
    sure we clean up stale engine data structures before then.
    """
    if self._engine_ttl <= 0:
        return

    now = time.perf_counter()
    busy = self._engines_with_inflight_transfers()
    for eid, last_active in list(self._engine_last_active.items()):
        if now - last_active > self._engine_ttl and eid not in busy:
            self._cleanup_remote_engine(eid)

La restricción clave es el conjuntobusy——los motores con transferencias en progreso no pueden ser expulsados.📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:3534-3546El comentario de explica la razón:

"""Remote engines a transfer is still reading from.

The timestamp is stamped when a read is issued and not refreshed while
it runs, so a transfer that outlives the TTL leaves its engine looking
idle. A peer that has lost its NIC holds one indefinitely.
"""

Si la tarjeta de red del par se avería, la transferencia podría quedar colgada para siempre, la marca de tiempo no se actualizaría y el motor parecería estar inactivo.busyEl conjunto protege explícitamente esta situación.

Secuencia temporal del handshake y la transferencia

El siguiente diagrama de secuencia muestra la interacción central desde la solicitud hasta la finalización de la transferencia:

sequenceDiagram
    participant Sched as Scheduler
    participant Worker as NixlWorker
    participant BgThread as 握手后台线程
    participant Remote as 远程 NIXL Agent

    Sched->>Worker: build_connector_meta()
    Worker->>Worker: _ensure_handshake(engine_id)
    alt 已握手
        Worker->>Worker: 直接返回 None
    else 握手中
        Worker->>BgThread: 返回已有 Future
    else 新握手
        Worker->>BgThread: submit(_nixl_handshake)
        BgThread->>Remote: ZMQ GET_META_MSG
        Remote-->>BgThread: NixlHandshakePayload
        BgThread->>BgThread: 校验 compat_hash
        BgThread->>Remote: add_remote_agent()
        BgThread-->>Worker: done_callback 注册 _remote_agents
    end
    Worker->>Remote: prep_xfer_dlist + make_xfer_req
    Worker->>Worker: _recving_transfers[req_id] = handles
    Sched->>Worker: get_transfer_results()
    Worker->>Worker: _pop_done_transfers()
    alt xfer_state == DONE
        Worker->>Remote: release_xfer_handle
        Worker-->>Sched: finished_recving
    else xfer_state == PROC
        Worker->>Worker: 保留 handle 等待下一轮
    else 失败
        Worker->>Worker: _handle_failed_transfer
        Worker-->>Sched: failed_recving + invalid_block_ids
    end

---

III. Reflexiones de diseño: por qué se diseñó así

¿Por qué el handshake debe ser asíncrono?

El handshake implica un viaje de ida y vuelta por la red, que puede tardar decenas de milisegundos. Si se ejecuta de forma síncrona, bloquearía el bucle principal del Scheduler, afectando la programación de todas las solicitudes. El handshake asíncrono permite que el Scheduler procese primero otras solicitudes y notifique mediante callback cuando el handshake se complete.

Pero la asincronía también trae complejidad:_handshake_futuresel diccionario necesita protección con lock, en el callback hay que manejar tanto el éxito como el fallo, y además hay que evitar handshakes duplicados.

¿Por qué usar lease en lugar de conteo de referencias?

El conteo de referencias requiere que la instancia decode notifique explícitamente a prefill "ya terminé de leer". Pero si la instancia decode se cae, la notificación nunca llegará y el bloque de prefill se filtrará para siempre.

El lease es una solución más robusta: incluso si decode se cae, prefill recupera automáticamente el bloque cuando el lease expira. El mecanismo de heartbeat garantiza la renovación del lease en condiciones normales.

¿Por qué conservar el handle en caso de fallo?

📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:3088-3101El comentario de lo dice claramente: un estado de error no garantiza que el DMA se haya detenido. Si en ese momento se libera el handle, el DMA podría seguir escribiendo datos en memoria ya liberada, causando corrupción de datos o un crash. Prefiero filtrar temporalmente que asumir ese riesgo.

¿Por qué la expulsión por TTL debe verificar busy?

📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:3534-3546El comentario de revela un escenario de bug oculto: la marca de tiempo se registra al iniciar la lectura y no se actualiza durante la lectura. Si el tiempo de transferencia supera el TTL, el motor parecerá inactivo, pero en realidad sigue siendo leído. Si se expulsa en ese momento, la transferencia en curso fallará.

Puntos problemáticos en entornos de producción

1. Problema de contexto CUDA: el handshake se ejecuta en un hilo en segundo plano, se debe hacer explícitamenteset_device, de lo contrario UCX deshabilitará NVLink silenciosamente.

2. Incompatibilidad de hash de compatibilidad: la versión de vLLM, el modelo, el dtype, el KV layout y el attention backend de las instancias P/D deben ser completamente consistentes. Cuando no lo son, el handshake fallará y el mensaje de error indicará cómo deshabilitar la verificación (pero no se recomienda).

3. Expiración del lease: si la instancia decode tiene una carga muy alta, el heartbeat podría retrasarse, provocando que el lease expire. En los logs aparecerá la advertencia "Releasing expired KV blocks". Se puede aumentarkv_lease_duration。

4. Desajuste de TP: el TP heterogéneo requiere un layout block-contiguous (como LBHNC). Si se usa un layout no contiguo, el TP heterogéneo fallará.

5. Agotamiento de NIXL UAR：📎 vllm/distributed/kv_transfer/kv_connector/v1/nixl/base_worker.py:631-636advertencia en los comentarios de: cada hilo UCX asigna UAR (doorbell pages) a través de DevX; un uso excesivo de NIXL UAR agotará el espacio UAR de la NIC, provocando que NVSHMEM (usado por el kernel de DeepEP) falle durante la inicialización de RDMA.

---

Resumen del capítulo

Este capítulo profundizó en los mecanismos centrales del sistema KV Connector:

1. KVConnectorBase_V1definió la abstracción de doble rol del lado Scheduler y del lado Worker, medianteKVConnectorMetadatayKVConnectorTransferResultspara implementar el intercambio de metadatos y la retroalimentación de resultados de transferencia.

2. Conector NIXLes la implementación más madura; establece conexiones entre instancias P/D mediante un protocolo de handshake ZMQ, usa hash de compatibilidad para evitar desajustes de configuración y emplea un pool de hilos asíncronos para evitar bloquear el bucle principal.

3. Lease y heartbeatel mecanismo resolvió el problema de temporización en la liberación de blocks: prefill no libera inmediatamente después de enviar KV, sino que espera la renovación por heartbeat o la expiración del lease de decode.

4. Recuperación ante fallossigue el principio de "prefiero una fuga que un uso incorrecto": cuando la liberación falla, se conserva el handle; los block ID fallidos se reportan al Scheduler para decidir el reintento.

5. Expulsión por TTLevita el crecimiento ilimitado del estado de motores remotos durante ejecuciones prolongadas, pero debe proteger a los motores con transferencias en curso.

En el próximo capítulo pasaremos a otra dirección para eliminar sobrecarga: aceleración de compilación y CUDA Graph. Una vez que la separación PD resolvió el problema de utilización de recursos, el costo de arranque de un solo forward se convierte en el nuevo cuello de botella: cómo usar CUDA Graph para comprimir cientos o miles de lanzamientos de kernels en una sola reproducción.

Reflexión y autoevaluación de este capítulo

Hasta aquí, hemos visto cómo KV Connector establece un canal de datos confiable entre las instancias de prefill y decode, y cómo utiliza mecanismos de arrendamiento, latido y recuperación ante fallos para mantener la consistencia del estado. Pero la transferencia entre instancias es solo la mitad de la historia de la separación PD: una vez que KV Cache llega a la instancia de decode, el motor de inferencia aún debe ejecutar eficientemente cada paso de cómputo hacia adelante dentro de una sola instancia. Y la sobrecarga de programación de Python y el lanzamiento de kernels es precisamente el siguiente cuello de botella que limita la latencia de un solo paso. El próximo capítulo se centrará en la aceleración por compilación y CUDA Graph, para ver cómo vLLM utiliza torch.compile y el backend piecewise para eliminar estas sobrecargas, y cómo hace que CUDA Graph y las formas de procesamiento por lotes dinámico coexistan de manera coordinada.

En el capítulo anterior vimos que KV Connector, a través de conectores como NIXL y Mooncake, transfiere eficientemente el KV cache entre los motores de Prefill y Decode, permitiendo que la arquitectura desagregada reduzca el TTFT mientras mejora la utilización de recursos. Pero incluso si la transferencia es más rápida, en la decodificación autorregresiva todavía hay dos costos fijos que no pueden eliminarse mediante algoritmos: la sobrecarga de programación del intérprete de Python y la sobrecarga de lanzamiento de kernels de GPU. Cuando el forward del modelo se divide en cientos de operadores, y cada operador debe pasar por una llamada a función de Python y un lanzamiento de kernel CUDA, la sobrecarga del lado de la CPU es suficiente para que la GPU quede inactiva entre dos cómputos. Este capítulo analiza cómo vLLM utiliza torch.compile para fusionar operadores en un grafo estático, y luego usa CUDA Graph para grabar toda la secuencia de lanzamiento de kernels como una sola reproducción, reduciendo así estos dos tipos de sobrecarga a casi cero.

Caché de compilación y capa de adaptación del compilador: permitir la reutilización de resultados de compilación entre procesos

Modelo intuitivo

El beneficio de la aceleración por compilación es "compilar una vez, ejecutar muchas veces", pero el costo es que la primera compilación puede tardar varios minutos. Sin caché, cada reinicio del servicio requeriría recompilar, y el tiempo de arranque en frío sería inaceptable.CompilerInterfaceEsta capa debe resolver precisamente el problema de "cómo serializar el producto de compilación, cómo identificarlo con un hash y cómo acertarlo con precisión en el próximo inicio". Sin ella, el desastre al que se enfrenta el sistema no es un fallo, sino que cada reinicio degenera en una "primera ejecución"; en un entorno de producción con escalado automático, esto significa que las instancias escaladas no podrán proporcionar servicios de baja latencia durante varios minutos.

Estructuras de datos y contrato de interfaz

CompilerInterfaceSe define el contrato abstracto del adaptador del compilador, cuyo núcleo son cuatro métodos:initialize_cacheSe encarga de redirigir el directorio de caché del propio compilador al directorio de caché de vLLM📎 vllm/compilation/compiler_interface.py:36-51；compute_hashRecopila la información de configuración relacionada con el compilador para generar un hash📎 vllm/compilation/compiler_interface.py:53-62；compileEjecuta la compilación y devuelve el objeto invocable y el handle📎 vllm/compilation/compiler_interface.py:64-95；loadRestaura el producto de compilación a partir del handle📎 vllm/compilation/compiler_interface.py:97-103。

La decisión de diseño clave aquí escompileDevuelve una tupla de dos elementos(callable, handle)。callableEs el resultado de compilación directamente invocable dentro de este proceso;handleEs la credencial "utilizada para restaurar en el próximo inicio", y la documentación exige explícitamente que debe ser un "plain Python object, preferably a string or a file path"📎 vllm/compilation/compiler_interface.py:81-81. Esta separación permite que la ruta de acierto de caché y la ruta de primera compilación sigan códigos completamente diferentes; en caso de acierto, no se necesita en absolutocompile, solo se requiereload。

compile_rangeEl parámetro transporta la semántica de formas dinámicas. Los comentarios indican que "could be concrete size (if compile_sizes is provided), e.g. [4, 4] or a range [5, 8]", y que "Right now we only support one variable in ranges for all inputs, which is the batchsize (number of tokens) during inference"📎 vllm/compilation/compiler_interface.py:74-74. Esta es la restricción central de la estrategia de compilación de vLLM: todas las formas dinámicas se reducen a una única variable: el número de tokens.

Impulsado por escenarios: el flujo completo de una solicitud de compilación

Supongamos que el servicio se inicia por primera vez,InductorAdaptor.compileEs invocado. Primero incrementa el contador de compilación📎 vllm/compilation/compiler_interface.py:477-489, y luego entra en una pila de parches cuidadosamente construida.

El primer paso es una copia profunda del grafo. Los comentarios señalan que "inductor can inplace modify the graph, so we need to copy it"📎 vllm/compilation/compiler_interface.py:500-502, esto es un diseño defensivo: tras un fallo de compilación, el grafo original aún puede usarse para reintentar.

El segundo paso es instalar una serie de monkey-patches.hijacked_compile_fx_innerEnvuelve la función de compilación interna de Inductor, y tras completar la compilación extrae el hash deinductor_compiled_graph._fx_graph_cache_keyExtrae el hash📎 vllm/compilation/compiler_interface.py:512-536。hijack_compiled_fx_graph_hashIntercepta la propia función de cálculo de hash📎 vllm/compilation/compiler_interface.py:538-542. ¿Por qué "secuestrar" el hash? Porque vLLM necesita compilar por separado fuera del contexto de rastreo de Dynamo, y el cálculo de hash de Inductor depende de dicho contexto.

El tercer paso es_check_can_cacheparche, que retorna directamente sin realizar ninguna verificación📎 vllm/compilation/compiler_interface.py:544-551. El comentario explica la motivación: "Inductor refuses to cache the graph outside of Dynamo tracing context, and also disables caching for graphs with high-order ops. For vLLM, in either case, we want to cache the graph"📎 vllm/compilation/compiler_interface.py:544-551。

El cuarto paso es limpiar el contexto de trazado. Este es el punto más sutil: vLLM llama aPiecewiseCompileInterpreterdesde dentro decompile_fx, en este momento elFakeTensorModede Dynamo y elFakeTensorModede la entrada del subgrafo no son consistentes,detect_fake_mode()provocará un fallo de aserción📎 vllm/compilation/compiler_interface.py:615-622. El código guardaTracingContexty luego lo deja vacío, y registra un callback para restaurarlo al salir📎 vllm/compilation/compiler_interface.py:623-630。

flowchart TD
    start["InductorAdaptor.compile()"] --> deepcopy["copy.deepcopy(graph)"]
    deepcopy --> patch_stack["ExitStack 安装补丁"]
    patch_stack --> p1["patch compiled_fx_graph_hash"]
    patch_stack --> p2["patch FxGraphCache._get_shape_env"]
    patch_stack --> p3["patch _check_can_cache"]
    patch_stack --> p4["清空 TracingContext"]
    p4 --> call_fx["compile_fx(graph, example_inputs)"]
    call_fx --> check{"hash_str is None?"}
    check -->|"是"| err["RuntimeError: 编译失败<br/>建议删除 torch_compile_cache"]
    check -->|"否"| check2{"file_path is None?"}
    check2 -->|"是"| assert_err["AssertionError"]
    check2 -->|"否"| ret["return (compiled_graph, (hash_str, file_path))"]
    err --> cleanup["ExitStack 退出<br/>恢复 TracingContext"]
    assert_err --> cleanup
    ret --> cleanup

Reflexión de diseño: AlwaysHitShapeEnv y consistencia de caché

AlwaysHitShapeEnvEsta clase merece un análisis aparte. Su docstring expone directamente la motivación: vLLM solo ejecuta una vez la compilación de bytecode de Dynamo, pero necesita ejecutar múltiples veces la compilación de Inductor con diferentes formas más una forma genérica; la compilación para formas específicas ocurre fuera del contexto de Dynamo, donde no hay un shape environment disponible para Inductor, lo que provoca fallos en la búsqueda de la caché de código de Inductor📎 vllm/compilation/compiler_interface.py:114-131。

La solución es proporcionar un shape environment falso que "siempre acierta":evaluate_guards_expressionsiempre retornaTrue 📎 vllm/compilation/compiler_interface.py:144-145，get_pruned_guardsretorna una lista vacía📎 vllm/compilation/compiler_interface.py:144-145，produce_guards_expressionretorna una cadena vacía📎 vllm/compilation/compiler_interface.py:147-159. El comentario admite que estos métodos fueron "obtained by trial-and-error until it works"📎 vllm/compilation/compiler_interface.py:137-142——este es un punto frágil acoplado a la implementación interna de PyTorch, y también el lugar más propenso a problemas al actualizar PyTorch.

La composición del hash de caché también es clave.get_inductor_factorsrecopila tres tipos de factores: estado del sistemaCacheBase.get_system(), estado de PyTorchtorch_key(), y la configuración de Inductor y functorch📎 vllm/compilation/compiler_interface.py:165-185. Nótese que la configuración de functorch se recopila dentro del contexto depatch(_get_vllm_functorch_config()), lo que garantiza que "la configuración en tiempo de compilación y la clave de caché siempre sean consistentes"——el comentario dice explícitamente que esto es para mantener📎 vllm/compilation/compiler_interface.py:188-189yset_functorch_config()consistentesget_inductor_factors(). Si estos dos lugares no son consistentes, ocurrirá un desajuste de "se usó la configuración A en tiempo de compilación, pero la clave de caché se calculó según la configuración B", lo que provocará que se cargue un artefacto incorrecto aunque haya un acierto de caché.📎 vllm/compilation/compiler_interface.py:147-159Problemas en producción:

es un backport para torch < 2.10.0_patch_standalone_compile_atomic_save. Cambia📎 vllm/compilation/compiler_interface.py:205-243para usarCompiledArtifact.save()escritura en formato binario, el comentario indica que el propósito es "preventing corrupt cache files when multiple processes compile concurrently"write_atomic. En escenarios de arranque en frío simultáneo de múltiples réplicas, múltiples procesos escribirán concurrentemente el mismo archivo de caché; una escritura no atómica producirá archivos truncados, y los procesos posteriores leerán artefactos corruptos con comportamiento impredecible.📎 vllm/compilation/compiler_interface.py:208-210PiecewiseBackend: compilación por niveles de forma y despacho en tiempo de ejecución

Modelo intuitivo

es el centro de coordinación entre compilación y ejecución. Compila "un subgrafo FX" en "múltiples objetos invocables por nivel de forma", y en tiempo de ejecución selecciona el más adecuado según el número real de tokens. Sin él, o todas las formas usarían la misma compilación genérica (rendimiento subóptimo), o cada forma se compilaría por separado (explosión del tiempo de compilación).

PiecewiseBackendEstructura de datos: RangeEntry y rango de compilación

La estructura de datos central es

, que vincula la banderaRangeEntryycompile_range、compiledjuntosrunnablemantiene un📎 vllm/compilation/piecewise_backend.py:80-83。PiecewiseBackendLa construcción del rango de compilación se divide en dos pasos. Primero se procesarange_entries: dict[Range, RangeEntry] 📎 vllm/compilation/piecewise_backend.py:166-171。

(tamaños exactos), cada tamaño genera un intervalo de punto únicocompile_sizesdeRange(start=size, end=size). Nótese que aquí para la cadena📎 vllm/compilation/piecewise_backend.py:166-171se lanza directamente"cudagraph_capture_sizes", y se explica que "should be handled inNotImplementedError——esta es una declaración explícita de límite de responsabilidad. Luego se procesapost_init_cudagraph_sizes" 📎 vllm/compilation/piecewise_backend.py:166-171(intervalos), cada intervalo genera una entrycompile_rangesadmite dos modos mutuamente excluyentes, y el constructor fuerza esto con una aserción XOR📎 vllm/compilation/piecewise_backend.py:173-173。

PiecewiseBackend: el modo de compilación (con graph, sin compiled_runnables) usa📎 vllm/compilation/piecewise_backend.py:117-119; el modo precompilado (sin graph, con compiled_runnables) usacompile_all_ranges() 📎 vllm/compilation/piecewise_backend.py:193-194. Este diseño permite que el arranque en frío y el arranque en caliente compartan la misma clase, solo con diferente origen de datos.load_all_ranges() 📎 vllm/compilation/piecewise_backend.py:193-194Orientado a escenarios: del despacho en compilación al tiempo de ejecución

Fase de compilación

recorre todas las range entry, y para cada entry no compilada llama a：compile_all_rangesregistra el evento de trazado_log_compile_start. La rama clave está en la construcción de parámetros: si es un tamaño de punto único, se llama a📎 vllm/compilation/piecewise_backend.py:252-256para generar un FakeTensor de forma concretacreate_concrete_args; de lo contrario, se llama a📎 vllm/compilation/piecewise_backend.py:258-261para reutilizar directamente los metadatos del placeholder en el grafoget_fake_args_from_graphLa implementación de📎 vllm/compilation/piecewise_backend.py:262-263。

create_concrete_argsrevela los detalles de la concretización de formas simbólicas. Construye unShapeEnvconFakeTensorMode 📎 vllm/compilation/piecewise_backend.py:54, y luego recorre los nodos placeholder. Para entradas de tipoSymInt, usaconcretizepara reemplazar todos los símbolos libres consize 📎 vllm/compilation/piecewise_backend.py:47-52; para el tipoTensor, debe concretizar simultáneamente shape, stride, storage_offset, y usarcompute_required_storage_lengthpara calcular la longitud de almacenamiento requerida, y luego reconstruir el tensor medianteas_strided 重建张量 📎 vllm/compilation/piecewise_backend.py:64-73. ¿Por qué no se puede cambiar solo la shape? Porque stride y storage_offset también pueden contener signos, y los tres deben ser coherentes entre sí; de lo contrario,as_stridedprovocará un acceso fuera de límites.

Despacho en tiempo de ejecución：__call__es una ruta crítica. Si existesym_shape_indices, se extrae la forma en tiempo de ejecuciónargsdesde📎 vllm/compilation/piecewise_backend.py:357-362, y luego se llama a_find_range_for_shapepara buscar. La lógica de búsqueda tiene prioridad: primero se comprueba si coincide con uncompile_sizesexacto; si coincide, se devuelve ese intervalo de punto único📎 vllm/compilation/piecewise_backend.py:342-355; de lo contrario, se recorrecompile_rangespara encontrar el intervalo que contiene esa forma📎 vllm/compilation/piecewise_backend.py:342-355。

flowchart TD
    call["PiecewiseBackend.__call__(*args)"] --> has_sym{"sym_shape_indices 非空?"}
    has_sym -->|"是"| get_shape["runtime_shape = args[sym_shape_indices[0]]"]
    get_shape --> find["_find_range_for_shape(runtime_shape)"]
    find --> exact{"runtime_shape in compile_sizes?"}
    exact -->|"是"| exact_entry["返回 Range(start=shape, end=shape) 的 entry"]
    exact -->|"否"| scan["遍历 compile_ranges 找包含区间"]
    scan --> found{"找到?"}
    found -->|"否"| assert_fail["AssertionError: 形状超出编译范围"]
    found -->|"是"| entry_ok["返回对应 entry"]
    has_sym -->|"否"| static["取唯一已编译 entry"]
    static --> check_count{"compiled_entries 数量 == 1?"}
    check_count -->|"否"| count_err["AssertionError"]
    check_count -->|"是"| entry_ok
    exact_entry --> run["range_entry.runnable(*args)"]
    entry_ok --> run

Reflexión de diseño: serialización y manejo especial de CachingAutotuner

Durante la serialización también se habilita temporalmentebundled_autograd_cache 📎 vllm/compilation/piecewise_backend.py:222, lo cual resuena con la lógica en_get_vllm_functorch_config—cuandoVLLM_USE_MEGA_AOT_ARTIFACTno está habilitado, esa configuración esFalse 📎 vllm/compilation/compiler_interface.py:160-161, y al serializar se fuerza aTrue, asegurando que los artefactos se empaqueten.

load_all_rangeses la ruta de arranque en caliente; afirma que cada range puede encontrarse encompiled_runnablescon su key correspondiente; de lo contrario, lanza un error que incluye la lista de keys disponibles📎 vllm/compilation/piecewise_backend.py:329-339. Este mensaje de error está diseñado de forma muy práctica: enumera directamente las keys disponibles, lo que facilita diagnosticar desajustes de versión de caché.

Wrapper de CUDA Graph: captura, reproducción y despacho anidado

Modelo intuitivo

CUDA Graph graba "una secuencia de lanzamientos de kernels" como un grafo estático; después, cada reproducción solo requiere una llamada a la API.CUDAGraphWrapperes el ejecutor de la grabación y la reproducción. Su dificultad central es: el tamaño de batch de vLLM es dinámico, mientras que CUDA Graph exige direcciones de entrada fijas. La solución es "capturar por niveles según batch descriptor": se graba un grafo por cada nivel de forma, y en tiempo de ejecución se consulta la tabla por descriptor para reproducir.

Estructura de datos: CUDAGraphEntry y contrato de despacho

CUDAGraphEntrycontiene tres campos clave:batch_descriptorcomo clave de despacho📎 vllm/compilation/cuda_graph.py:128-135、cudagraphes el objeto de grafo capturado📎 vllm/compilation/cuda_graph.py:128-135、outputes la salida en el momento de la captura (se guarda con referencia débil para ahorrar memoria)📎 vllm/compilation/cuda_graph.py:128-135。input_addressessolo se usa en modo de depuración para validar que las direcciones de entrada coincidan al reproducir📎 vllm/compilation/cuda_graph.py:128-135。

CUDAGraphWrapperLa documentación de clase de describe con precisión el contrato de despacho: al inicializar se asigna un runtime mode (FULL o PIECEWISE)📎 vllm/compilation/cuda_graph.py:158-158; en tiempo de ejecución recibe runtime_mode y batch_descriptor desde el forward context y "blindly trust them"📎 vllm/compilation/cuda_graph.py:158-158; si runtime_mode es NONE o no coincide, llama directamente a📎 vllm/compilation/cuda_graph.py:158-158; de lo contrario, ejecuta captura o reproducción📎 vllm/compilation/cuda_graph.py:158-158。

La documentación también declara explícitamente un límite: "CUDAGraphWrapper does not store persistent buffers or copy any runtime inputs into that buffers for replay"📎 vllm/compilation/cuda_graph.py:164-164. Esto significa que la gestión de los búferes de entrada es responsabilidad del llamador: el wrapper solo se encarga del grafo en sí.

Guiado por escenarios: una captura y una reproducción

Ruta de captura: cuando__call__se activa y runtime_mode coincide, primero se comprueba si el forward context está disponible. Si no lo está (como en el forward del codificador visual), se llama directamente a la función subyacente📎 vllm/compilation/cuda_graph.py:232-233. Esta es la rama clave del escenario multimodal: el forward de ViT no pasa por CUDA Graph.

Luego se tomanbatch_descriptorycudagraph_runtime_mode 📎 vllm/compilation/cuda_graph.py:242-244. Si mode es NONE o no coincide, se llama directamente a📎 vllm/compilation/cuda_graph.py:246-256. Este diseño de "si no coincide, pasa directo" permite que los wrappers anidados coexistan: el wrapper FULL en la capa externa y el wrapper PIECEWISE en la capa interna; en tiempo de ejecución solo uno se activará.

Si elcudagraphde la entry es None, se entra en captura. Primero se llama avalidate_cudagraph_capturing_enabled()para validar la legalidad📎 vllm/compilation/cuda_graph.py:279, luego se registran las direcciones de entrada📎 vllm/compilation/cuda_graph.py:281-284, se creatorch.cuda.CUDAGraph() 📎 vllm/compilation/cuda_graph.py:285。

En el contexto de captura hay varias operaciones clave. Sigc_disableestá habilitado, se hace patch degc.collectytorch.accelerator.empty_cache 📎 vllm/compilation/cuda_graph.py:288-303. El comentario explica el motivo: en modo piecewise cada capa debe capturar un grafo, y ejecutar GC repetidamente haría la captura extremadamente lenta, así que "only run gc for the first graph, and disable gc for the rest"📎 vllm/compilation/cuda_graph.py:289-294. Luego se establece el graph pool id📎 vllm/compilation/cuda_graph.py:305-308, y se sincroniza el stream de copia del offloader📎 vllm/compilation/cuda_graph.py:310-312。

La captura real se ejecuta dentro del contextotorch.cuda.graph(cudagraph, pool=..., stream=...)self.runnable(*args, **kwargs) 📎 vllm/compilation/cuda_graph.py:315-321. Después de capturar, se llama aget_offloader().join_after_forward()para evitar errores de streams no unidos📎 vllm/compilation/cuda_graph.py:322-326. Siweak_ref_outputestá habilitado, se convierte la output a referencia débil para ahorrar memoria📎 vllm/compilation/cuda_graph.py:327-334. Finalmente, la entry guarda la output como referencia débil y el objeto de grafo📎 vllm/compilation/cuda_graph.py:338-339, perolo que se devuelve es la output original, no la referencia débil—el comentario enfatiza que esto es para que PyTorch gestione correctamente la memoria durante la captura📎 vllm/compilation/cuda_graph.py:343-346。

Ruta de reproducción: si la entry ya tiene un grafo, en modo de depuración se valida que las direcciones de entrada coincidan📎 vllm/compilation/cuda_graph.py:348-357, luego se sincroniza el offloader📎 vllm/compilation/cuda_graph.py:359-361, se llama aentry.cudagraph.replay()y se devuelveentry.output 📎 vllm/compilation/cuda_graph.py:362-363。

Reflexión de diseño: por qué la salida debe ser una referencia débil, mientras que el retorno debe ser una referencia fuerte

Esto esCUDAGraphWrapperlo más contraintuitivo enoutputestá gestionado por el cudagraph pool de PyTorch📎 vllm/compilation/cuda_graph.py:320. Si la entrada hace una referencia fuerte a output, entonces la memoria de video ocupada por este grafo nunca podrá liberarse; pero si se convierte a referencia débil durante la captura, PyTorch podría reclamar la memoria antes de que finalice la captura, provocando que la captura falle. Por eso el código usa una referencia débil dentro del bloque de captura📎 vllm/compilation/cuda_graph.py:334, almacena una referencia débil en la entrada📎 vllm/compilation/cuda_graph.py:338, pero el valor de retorno de la función es una referencia fuerte📎 vllm/compilation/cuda_graph.py:346. Este "estado de triple referencia" es un equilibrio preciso entre seguridad de memoria y eficiencia de memoria de video.

Otro diseño digno de mención es_all_instancesesteWeakSet 📎 vllm/compilation/cuda_graph.py:173-176. Permite queclear_all_graphspueda vaciar de una sola vez los grafos de todos los wrappers📎 vllm/compilation/cuda_graph.py:173-176, para reclamación de emergencia cuando la memoria de video está ajustada. Se usaWeakSeten lugar de un conjunto normal para no impedir que el wrapper sea recolectado por GC; de lo contrario, el propio wrapper se filtraría.

Errores en producción:__getattr__la implementación de lanza errores con contexto para atributos inexistentes en modo de depuración📎 vllm/compilation/cuda_graph.py:211-217. Esto parece trivial, pero al investigar "por qué falla cierta llamada a un método", poder ver la descripción de cadena del runnable envuelto por el wrapper es mucho más útil que unAttributeErrordesnudo.

Reflexión de diseño: desacoplamiento entre compilación y CUDA Graph

El documento de diseño registra explícitamente la motivación de esta refactorización. La compilación piecewise temprana era para soportar la captura piecewise de CUDA Graph, excluyendo los operadores que no soportan CUDA Graph (principalmente attention)📎 docs/design/cuda_graphs.md:25. Posteriormente se añadió soporte para full CUDA Graph, pero "this tight coupling between compilation and cudagraph capture led to an all-or-nothing experience with little flexibility"📎 docs/design/cuda_graphs.md:25。

Tras la refactorización, los objetivos son cuatro: distinguir explícitamente entre lotes prefill/mixed y uniform-decode y capturarlos por separado📎 docs/design/cuda_graphs.md:25-25; desacoplar la lógica de captura de CUDA Graph de la compilación, de modo que "capturing piecewise and full cudagraphs using the same compiled graph"📎 docs/design/cuda_graphs.md:25-25; despachar en tiempo de ejecución según la composición del lote📎 docs/design/cuda_graphs.md:25-25; y control centralizado para reducir la complejidad📎 docs/design/cuda_graphs.md:25-25。

BatchDescriptores la estructura central de la clave de despacho, que contienenum_tokens、num_reqs、uniform、has_loracuatro campos📎 docs/design/cuda_graphs.md:86-93。uniformEl flag es especialmente crítico: muchos backends de attention solo soportan full CUDA Graph cuando el lote es uniforme📎 docs/design/cuda_graphs.md:95-95. El documento también anticipa que esta estructura podría extenderse, por ejemplo añadiendouniform_query_lenpara soportar múltiples longitudes de uniform decode📎 docs/design/cuda_graphs.md:95-95。

La prioridad de despacho esFULL > PIECEWISE > None, y si la clave de despacho no existe, se recurre al modo NONE para ejecución eager📎 docs/design/cuda_graphs.md:112-115. Esta estrategia de "degradar en lugar de reportar error" garantiza que cualquier combinación de lotes pueda ejecutarse, solo que con rendimiento diferente.

AttentionCGSupportEl enum cuantifica la capacidad de CUDA Graph del backend, con valoresALWAYS=3 > UNIFORM_BATCH=2 > UNIFORM_SINGLE_TOKEN_DECODE=1 > NEVER=0 📎 docs/design/cuda_graphs.md:153-162. Los modelos de attention híbrida (como mamba mixer) toman el mínimo de las capacidades de todos los backends y degradan el modo CUDA Graph en consecuencia📎 docs/design/cuda_graphs.md:173-175. Este diseño desacopla la "declaración de capacidad" de la "selección de modo": añadir un nuevo backend solo requiere declarar su capacidad, y la estrategia de degradación se aplica automáticamente.

Resumen del capítulo

Reflexión y autoevaluación del capítulo

El siguiente capítulo se centrará en cuantización y kernels personalizados, para ver cómo vLLM interviene en el control de precisión desde la etapa de carga de pesos, y utiliza operadores altamente especializados para convertir realmente los beneficios de la cuantización en mejoras de throughput.

Este capítulo analizó los dos niveles del mecanismo de aceleración por compilación de vLLM. El primer nivel es CompilerInterface y PiecewiseBackend: el primero define el contrato de adaptación del compilador y la estrategia de hash de caché, usando AlwaysHitShapeEnv para eludir el problema de la falta de contexto de Dynamo; el segundo compila un único subgrafo FX en múltiples niveles de forma, despachando en tiempo de ejecución según el número de tokens. El segundo nivel es CUDAGraphWrapper: captura CUDA Graph por niveles según BatchDescriptor, logrando despacho anidado mediante coincidencia de runtime mode, permitiendo que los modos FULL y PIECEWISE coexistan en el mismo grafo compilado. El desacoplamiento de ambos es el núcleo de esta refactorización: los artefactos de compilación pueden ser reutilizados por ambos modos de CUDA Graph, y CUDA Graph también puede funcionar independientemente de la compilación. Sin embargo, la compilación y la captura de grafos resuelven la sobrecarga de programación; la precisión de los pesos del modelo y la eficiencia de los operadores siguen siendo otra línea principal de optimización. El siguiente capítulo se centrará en cuantización y kernels personalizados, para ver cómo vLLM analiza la configuración de cuantización, completa la conversión de formatos como FP8/INT4/AWQ/GPTQ durante la carga de pesos, y aprovecha _custom_ops y los kernels de Triton para exprimir aún más el rendimiento del hardware.

En el capítulo anterior vimos cómo torch.compile y CUDA Graph llevan al extremo la reducción de la sobrecarga de programación de Python y de lanzamiento de kernels. Pero por muy rápido que sea el despacho, si los pesos en sí son FP16 y la multiplicación de matrices usa GEMM genérico, el poder de cómputo del hardware sigue limitado por el ancho de banda de memoria y operadores ineficientes. La cuantización y los kernels personalizados son otra línea principal de optimización ortogonal: la primera reduce la precisión ya en la etapa de carga de pesos, y los segundos convierten realmente los beneficios de la cuantización en throughput. Este capítulo parte desde el punto de entrada de análisis de la configuración de cuantización y llega hasta el registro de operadores de _custom_ops y la programación de kernels de Triton.

11.1 Configuración de cuantización: de la cadena CLI a QuantKey

Modelo intuitivo

El rol del módulo de configuración de cuantización es como el traductor del menú de un restaurante. El usuario dice en recepción "quiero fp8_per_tensor" (cadena CLI), pero la cocina necesita el número exacto de receta (QuantKey). El traductor debe manejar tres tipos de entrada: abreviatura pura de CLI, metadatos de cuantización propios del checkpoint, y escenarios combinados de ambos. Sin esta capa de traducción, la cocina recibiría un montón de cadenas ambiguas y no podría decidir qué kernel invocar.

Estructuras de datos y diseño de memoria

Las estructuras de datos centrales sonQuantSpecyQuantizationConfigArgs. La primera describe las claves de cuantización de pesos y activaciones de un tipo de capa (linear o MoE), la segunda es la configuración de nivel superior visible al usuario.

📎 vllm/config/quantization.py:73-99

@config
class QuantSpec:
    weight: QuantKeyField = None
    activation: QuantKeyField = None

    def __str__(self) -> str:
        def quant_key_str(quant_key: QuantKey | None) -> str:
            if quant_key is None:
                return "None"
            return next(
                (
                    name
                    for name, known_quant_key in QUANT_KEY_NAMES.items()
                    if known_quant_key == quant_key
                ),
                str(quant_key),
            )
        return quant_key_str(self.weight)

weightyactivationson ambos opcionalesQuantKey。Nonela semántica es "revertir al valor predeterminado de la propia clase del método" — normalmente heredado del checkpoint; en escenarios de cuantización en línea, significa no cuantizar📎 vllm/config/quantization.py:74-74。QuantKeyen sí es un tipo complejo que contiene declaraciones deNamedTupleyClassVar[GroupShape], pydantic no puede introspectarlo directamente, por lo que el autor inyectó un validador personalizado medianteGetPydanticSchema_coerce_quant_key, que normaliza de forma unificada cadenas oQuantKey📎 vllm/config/quantization.py:60-69。

QuantizationConfigArgsel diseño de campos merece atención📎 vllm/config/quantization.py:102-126：

• linear / moe: actúan respectivamente sobre las capasLinearBaseyFusedMoEFactory;
• ignore: lista de nombres de capas a omitir en la cuantización; la cuantización en línea también admite comodines fnmatch;
• targets: sobrescritura de cuantización en línea capa por capa; la clave puede ser un nombre de capa exacto, una expresión regular con prefijore:, o un patrón fnmatch; el valor es mutuamente excluyente conlinear/moe.

targetsylinear/moeson mutuamente excluyentes, forzado pormodel_validator📎 vllm/config/quantization.py:172-179. Esta restricción no es formalismo:targetssigue la ruta de sobrescritura capa por capa,linear/moesigue la ruta predeterminada global; si ambas coexisten, "qué spec usa realmente cierta capa" se vuelve indeterminable.

Paso a paso: una resolución de--quantization fp8_per_tensor

Escenario: el usuario pasa por línea de comandos--quantization fp8_per_tensor, y al mismo tiempo especifica mediante--quantization-configla cuantización de activaciones de la capa MoE.

Primer paso,resolve_quantization_configes invocado con los argumentos: cadena de CLI y diccionario de configuración📎 vllm/config/quantization.py:233-235. Primero comprueba siquantizationestá enONLINE_QUANT_SHORTHAND_NAMES— esta tupla contiene todos los nombres abreviados más un"online" 📎 vllm/config/quantization.py:216-222。

Segundo paso,fp8_per_tensorcoincide con la tabla de abreviaturas,basese resuelve como_ONLINE_SHORTHANDS["fp8_per_tensor"], es decir, tanto linear como moe usankFp8StaticTensorSym 📎 vllm/config/quantization.py:188-190。

Tercer paso,quantization_configno está vacío, se construye como objetoQuantizationConfigArgs. Luego entra en la lógica de fusión📎 vllm/config/quantization.py:267-268: cada campo se decide conquantization_config.xxx or base.xxx— los campos establecidos explícitamente por el usuario tienen prioridad; los no establecidos heredan el valor predeterminado de la abreviatura. Aquí se usaoren lugar deif is not Nonede forma intencional:QuantSpecy la lista vacía son falsy; semánticamente, "no establecido" y "vacío" son equivalentes.

Cuarto paso, siquantizationno está en la tabla de abreviaturas (por ejemplo, es unawqpropio del checkpoint), yquantization_configesNone, la función devuelve directamenteNone 📎 vllm/config/quantization.py:256-257. Esto significa "no superponer cuantización en línea"; el método de cuantización del checkpoint sigue siendo dominante.

Hay una rama fácil de pasar por alto:_DEFERRED_ONLINE_SHORTHANDScontienemxfp4ymxfp8 📎 vllm/config/quantization.py:233-235. Estos dos nombres son tanto abreviaturas de CLI como nombres de métodos de cuantización de checkpoint. Cuando el usuario solo pasa--quantization mxfp4y noquantization_config, la función devuelveNoneen lugar debase 📎 vllm/config/quantization.py:267-268, posponiendo la decisión a los metadatos del checkpoint — solo cuando el checkpoint no tiene información de cuantización se recurre a la abreviatura en línea.

flowchart TD
    start["resolve_quantization_config(quantization, quantization_config)"]
    check_shorthand{"quantization in ONLINE_QUANT_SHORTHAND_NAMES?"}
    checkpoint_path{"quantization_config is None?"}
    return_none1["return None (checkpoint 主导)"]
    build_args["QuantizationConfigArgs(**quantization_config)"]
    get_base["base = _ONLINE_SHORTHANDS.get(quantization)"]
    cfg_none{"quantization_config is None?"}
    deferred{"quantization in _DEFERRED_ONLINE_SHORTHANDS?"}
    return_none2["return None (推迟到 checkpoint)"]
    return_base["return base"]
    merge["逐字段合并: cfg.xxx or base.xxx"]
    return_merged["return 合并后的 QuantizationConfigArgs"]

    start --> check_shorthand
    check_shorthand -->|否| checkpoint_path
    checkpoint_path -->|是| return_none1
    checkpoint_path -->|否| build_args
    check_shorthand -->|是| get_base
    get_base --> cfg_none
    cfg_none -->|是| deferred
    deferred -->|是| return_none2
    deferred -->|否| return_base
    cfg_none -->|否| merge
    merge --> return_merged

Consideraciones de diseño y trampas

_coerce_specel validador maneja un escenario sutil: cuandolinearomoereciben una cadena, primero consulta_ONLINE_SHORTHANDS; si coincide, extrae el spec del campo correspondiente; si no coincide, lo trata como un único nombre deQuantKey📎 vllm/config/quantization.py:130-139. Esto significa quelinear="fp8_per_tensor"ylinear="fp8_per_tensor_static"siguen dos rutas diferentes — la primera es una abreviatura de configuración completa, la segunda es una única clave de cuantización. Si en la abreviatura ese campo esNone(por ejemplo,int8_per_channel_weight_onlyno tiene campolinear), se lanza unValueErrorexplícito en lugar de devolver silenciosamenteNone 📎 vllm/config/quantization.py:130-139。

Una trampa común en producción:targetslas claves de expresión regular se precompilan y validan en_validate_targets📎 vllm/config/quantization.py:166-167, pero las claves de patrón fnmatch no se validan. Si el usuario escribe un patrón fnmatch que nunca coincide con ninguna capa, no se produce error; simplemente esa capa permanece sin cuantizar — al diagnosticar, hay que comprobar si el nombre de capa realmente coincide.

11.2 _custom_ops: registro de operadores e implementación fake

Modelo intuitivo

_custom_ops.pyes la capa de adaptación entre vLLM y los operadores subyacentes de CUDA/C++, como una aduana. En el espacio de nombrestorch.ops._Cde PyTorch se registran operadores C++ compilados, pero llamarlos directamente tiene tres problemas: el conjunto de operadores difiere entre plataformas (CUDA/ROCm/CPU/XPU),torch.compilenecesita implementaciones fake para inferir formas de salida, y algunos operadores requieren preprocesamiento de parámetros en el lado de Python._custom_opsencapsula estos problemas de forma unificada.

Estructuras de datos y mecanismo de registro

Al cargar el módulo, primero se llama acurrent_platform.import_kernels() 📎 vllm/_custom_ops.py:25-26, dando a la capa de plataforma la oportunidad de importar su propia biblioteca de operadores. Luego se defineregister_fake— bajoTYPE_CHECKINGes un decorador vacío; en tiempo de ejecución se importa desdetorch.library📎 vllm/_custom_ops.py:25-26。

La función principal de la implementación fake es hacer quetorch.compileconozca la forma de salida y el dtype del operador durante la fase de trazado, sin ejecutarlo realmente. Tomandoscaled_fp4_quantcomo ejemplo:

📎 vllm/_custom_ops.py:90-100

if hasattr(torch.ops, "_C") and hasattr(torch.ops._C, "scaled_fp4_quant"):

    @register_fake("_C::scaled_fp4_quant")
    def _scaled_fp4_quant_fake(
        input: torch.Tensor,
        input_scale: torch.Tensor,
        is_sf_swizzled_layout: bool,
    ) -> tuple[torch.Tensor, torch.Tensor]:
        n = input.shape[-1]
        m = input.numel() // n
        return create_fp4_output_tensors(m, n, input.device, is_sf_swizzled_layout)

Nótese la guardahasattr: la implementación fake solo se define cuando la plataforma realmente ha registrado_C::scaled_fp4_quant. Esto garantiza que importar el módulo en CPU o GPU antiguas no falle por falta de operadores.

create_fp4_output_tensorsmuestra los detalles del diseño de memoria de la salida de cuantización FP4📎 vllm/_custom_ops.py:69-87. Cuandois_sf_swizzled_layout=True, el tensor de escala debe organizarse en tiles de 128x4 según lo exigen los Tensor Core: el número de filas se redondea hacia arriba a múltiplos de 128, el número de columnas (n // 16) se redondea hacia arriba a múltiplos de 4, y cada 4 float8_e4m3 se empaquetan en un int32📎 vllm/_custom_ops.py:55-64. El comentario señala explícitamente que el kernel de cuantización NVFP4 pone a cero explícitamente todas las entradas de escala de padding, por lo que no se necesita un kernel de inicialización a cero por separado📎 vllm/_custom_ops.py:60-61。

Paso a paso: flujo de llamada de un AWQ GEMM

Escenario: el modelo ha cargado pesos cuantizados AWQ; en la propagación hacia adelante se debe hacer una multiplicación matricial entre activaciones y pesos cuantizados.

Primer paso, llamar aawq_gemm 📎 vllm/_custom_ops.py:587-592. La función primero comprueba la variable de entornoVLLM_USE_TRITON_AWQ. Si es verdadera, importa de forma diferidaawq_gemm_tritony llama — esta es una ruta de implementación puramente Triton, para plataformas que no admiten operadores CUDA o escenarios de depuración.

Segundo paso, la ruta predeterminada llama atorch.ops._C.awq_gemm, pasando input, qweight, scales, qzeros ysplit_k_iters 📎 vllm/_custom_ops.py:598-598。

Tercer paso, sitorch.ops._C.awq_gemmExiste, la implementación fake está registrada📎 vllm/_custom_ops.py:601-616. La forma que devuelve fake es(split_k_iters, num_in_feats, qweight.size(1) * 8)luego.sum(0)——esto simula con precisión la forma del resultado intermedio de split-K y la forma final tras la reducción.qweight.size(1) * 8Proviene del empaquetado de AWQ: cada int32 almacena 8 pesos de 4 bits.

Cuarto paso,awq_dequantizesigue una ruta similar📎 vllm/_custom_ops.py:553-559, pero la derivación de formas de la implementación fake es diferente:out_c = qout_c * 8, porque tras la desquantización el número de columnas se expande 8 veces📎 vllm/_custom_ops.py:587-592。

La función repack de la serie Marlin muestra otro patrón.gptq_marlin_repackLa implementación fake de calculapack_factor = 32 // num_bits, la forma de salida es(size_k // 16, size_n * 16 // pack_factor) 📎 vllm/_custom_ops.py:1103-1119. Aquí16es el Marlin tile size,size_k // 16indica que la dimensión K se divide por tiles. La versión MoE degptq_marlin_moe_repackitera en la capa de Python sobre cada expert llamando al repack de un solo expert📎 vllm/_custom_ops.py:1154-1172, y afirmasize_k % 16 == 0——esta es una restricción dura del formato Marlin.

flowchart LR
    input["input: torch.Tensor (FP16/BF16)"]
    qweight["qweight: torch.Tensor (INT32 packed)"]
    scales["scales: torch.Tensor"]
    qzeros["qzeros: torch.Tensor"]
    check_env{"VLLM_USE_TRITON_AWQ?"}
    triton_path["awq_gemm_triton(input, qweight, scales, qzeros, split_k_iters)"]
    cuda_path["torch.ops._C.awq_gemm(...)"]
    output["output: torch.Tensor (FP16/BF16)"]

    input --> check_env
    qweight --> check_env
    scales --> check_env
    qzeros --> check_env
    check_env -->|是| triton_path
    check_env -->|否| cuda_path
    triton_path --> output
    cuda_path --> output

Reflexiones de diseño y trampas

La implementación fake debe coincidir exactamente con la forma de salida del operador real, de lo contrariotorch.compileel grafo trazado tendrá formas que no coincidirán en tiempo de ejecución.create_fp4_output_tensorsEl comentario de enfatiza especialmente "Must match the C++ scaled_fp4_quant_func allocation exactly when padded_n is None"📎 vllm/_custom_ops.py:69-74. Este es un punto propenso a errores: si el lado de C++ cambia la lógica de asignación y el fake no se sincroniza, el grafo compilado fallará al reproducirse en CUDA Graph.

Otra trampa estorch.library.custom_opla regla de alias de .safeFusedQuantizeNvEl comentario de señala que torch 2.12+ no permite que la salida de un operador personalizado haga alias con ninguna entrada, por lo que el autor cambió el tensor de retorno a un parámetro in-place📎 vllm/_custom_ops.py:4650-4655. Esta práctica de "cambiar la forma de la API para sortear limitaciones del framework" es común en la capa de adaptación de operadores; al depurar hay que prestar atención a simutates_argsla declaración es consistente con el comportamiento real.

CPUDNNLGEMMHandlerMuestra otro patrón de gestión de recursos: el puntero del handler se guarda en un tensor int64,__del__al llamarrelease_dnnl_matmul_handlerlibera📎 vllm/_custom_ops.py:3708-3717. Guardar el puntero en un tensor es para evitar que la optimización de inline de enteros de Python lo elimine——esta es una técnica clásica de binding de bajo nivel.

11.3 Programación de kernels Triton:KernelOverridey re-binding entre módulos

Modelo intuitivo

El rol del programador de kernels Triton es como un sistema de reemplazo de puestos en una empresa. Cuando una plataforma (por ejemplo ROCm) necesita reemplazar con su propia implementación un kernel Triton del núcleo de vLLM, no puede modificar directamente el código del núcleo——eso contaminaría el upstream.dispatcherPermite que una plataforma registre un sustituto y luego cambia silenciosamente todas las referencias al kernel original por el sustituto. Sin este mecanismo, cada plataforma tendría que mantener un fork, con conflictos constantes al fusionar cambios del upstream.

Estructuras de datos y diseño de memoria

La estructura de datos central es_registryel diccionario yKernelOverridela clase📎 vllm/triton_utils/dispatcher.py:29-36。

KernelOverridecampos clave de📎 vllm/triton_utils/dispatcher.py:50-61：

• _impl: función de implementación de la plataforma;
• arg_names: tupla con los nombres de parámetros que refleja el kernel original, usada para el binding por palabra clave en el launch;
• constexprs: declaraciones constexpr heredadas del kernel original;
• func: apunta a la función de implementación, para introspección en warmup;
• _forward_by_name: flag booleano que decide si al hacer launch se reenvían los parámetros por palabra clave o por posición.

_forward_by_nameLa lógica de cálculo de es: compararinspect.signature(impl).parameterscon el del kernel originalarg_namessi son exactamente iguales📎 vllm/triton_utils/dispatcher.py:50-61. Si son iguales, significa que los nombres de parámetros de la implementación coinciden con los del kernel y se puede reenviar de forma segura por palabra clave; de lo contrario, hay que reenviar por posición según el orden de parámetros del kernel original.

Paso a paso: unregister_kernelsre-binding de

Escenario: la plataforma ROCm llama en la inicialización aregister_kernels({"vllm.v1.sample.rejection_sampler.expand_kernel": my_expand_impl})。

Primer paso,register_kernelsrecorre los overrides y para cada nombre llama a_resolve_kernel 📎 vllm/triton_utils/dispatcher.py:162-166。_resolve_kerneldivide el nombre por el último.en nombre de módulo y nombre de atributo📎 vllm/triton_utils/dispatcher.py:83-94. Si la primera letra del último segmento del nombre del módulo es mayúscula, significa que el kernel pertenece a alguna clase (JIT warmup owner); hay que importar primero el módulo padre y luegogetattrobtener la clase, devolver(类, 属性名); de lo contrario, importar el módulo mismo y devolver(模块, 属性名)。

Segundo paso, tras obtener el objeto del kernel original, construirKernelOverrideel wrapper y registrarlo en_registry 📎 vllm/triton_utils/dispatcher.py:167-169。

Tercer paso,_rebind_kernelsejecuta un escaneo de todos los módulos📎 vllm/triton_utils/dispatcher.py:97-144. Recorresys.modulesde todos los módulos en__dict__, y para cada valor de atributo hace comparación de identidad——ojo,isy no==, porque algunos valores de atributo (comoPlaceholderModuleel centinela) al hacer hash/eq pueden disparar importaciones o excepciones📎 vllm/triton_utils/dispatcher.py:116-123。

Cuarto paso, para los atributos que coinciden con el kernel original, directamentesetattrreemplazar por el wrapper📎 vllm/triton_utils/dispatcher.py:125-135. Para el JIT warmup owner (objetos cuya propiedad de instanciakernelapunta al kernel original), reemplazarvalue.kernely limpiar el caché de_kernel_arg_names, para que el binding del launch se derive de nuevo desde el wrapper📎 vllm/triton_utils/dispatcher.py:138-139。

Quinto paso,_rebind_kernelstras completarse, recién entonces reemplazar también el atributo en el sitio de definición por el wrapper📎 vllm/triton_utils/dispatcher.py:170-174. El comentario explica la importancia del orden: si se reemplaza primero el sitio de definición, al escanear ya no se encontrará el kernel original📎 vllm/triton_utils/dispatcher.py:170-171。

sequenceDiagram
    participant Platform as "ROCm 平台"
    participant Dispatcher as "register_kernels"
    participant Resolver as "_resolve_kernel"
    participant Scanner as "_rebind_kernels"
    participant Modules as "sys.modules"

    Platform->>Dispatcher: register_kernels({"vllm...expand_kernel": my_impl})
    Dispatcher->>Resolver: _resolve_kernel("vllm...expand_kernel")
    Resolver-->>Dispatcher: (module, "expand_kernel")
    Dispatcher->>Dispatcher: KernelOverride(original, my_impl)
    Dispatcher->>Scanner: _rebind_kernels([(original, wrapper)])
    Scanner->>Modules: 遍历所有模块 __dict__
    Modules-->>Scanner: 属性值列表
    Scanner->>Scanner: lookup(value) 身份比较
    Scanner->>Modules: setattr(module, attr, wrapper)
    Scanner->>Modules: value.kernel = wrapper (JIT owner)
    Scanner-->>Dispatcher: 重绑定完成
    Dispatcher->>Modules: setattr(host, attr, wrapper)
    Dispatcher-->>Platform: 注册完成

Reflexiones de diseño y trampas

KernelOverride.__getitem__devuelveself._launch, lo que hace quekernel[grid](**kwargs)esta sintaxis estándar de launch de Triton sea transparente para el wrapper📎 vllm/triton_utils/dispatcher.py:63-74。_launchLa lógica de reenvío de tiene tres casos📎 vllm/triton_utils/dispatcher.py:63-74: si hay argumentos posicionales, se pasan directamente;_forward_by_namesi es verdadero, se reenvía por palabra clave; de lo contrario, se comprueba si en kwargs hay nombres de parámetros que el kernel original no reconoce; si los hay, se lanzaRuntimeError, y si no, se extraen los valores en el orden de parámetros del kernel original y se reenvían por posición.

EsteRuntimeErrorEs una defensa importante: si los nombres de parámetros implementados por la plataforma no coinciden con los del kernel, y quien llama pasa parámetros que la implementación no reconoce, ignorarlos silenciosamente provocaría resultados erróneos difíciles de diagnosticar. Un error explícito expone el problema ya en la fase de registro.

Una trampa en producción:_rebind_kernelsEl escaneo de es de O(número de módulos × número de atributos × número de kernels). Para modelos grandes,sys.modulespuede haber miles de módulos, cada uno con cientos de atributos. Aunque solo se ejecuta una vez durante la inicialización, si hay muchos kernels registrados, el tiempo de arranque aumentará notablemente.lookupLa función usa un escaneo lineal en lugar de búsqueda por hash, y el comentario explica por qué: algunos valores de atributos no son hashables📎 vllm/triton_utils/dispatcher.py:116-123. Es un compromiso típico de "corrección antes que rendimiento".

Otra trampa:_resolve_kernelDetermina si es un atributo de clase mediante "la primera letra de la última sección del nombre del módulo en mayúscula"📎 vllm/triton_utils/dispatcher.py:83-94. Si el nombre de un módulo comienza con mayúscula (lo cual no sigue las convenciones de nomenclatura de Python pero es sintácticamente válido), se clasificará erróneamente como clase. Es un diseño de convención sobre configuración que depende de las normas de nomenclatura internas de vLLM.

Reflexiones de diseño

Los dos mecanismos, la configuración de cuantización y el registro de operadores, constituyen conjuntamente la superficie de ajuste "precisión-rendimiento" de vLLM.QuantizationConfigArgsEl diseño de refleja la separación entre "intención del usuario" y "valores predeterminados del método":NoneNo significa "no cuantizar", sino "dejar que la clase del método decida por sí misma". Esta decisión diferida permite que una misma configuración se adapte tanto a la cuantización de checkpoint como a la cuantización en línea.

_custom_opsEl patrón de implementación fake de es eltorch.compileestándar del ecosistema, pero lo distintivo de vLLM es elhasattruso generalizado de guardas. Esto permite que un mismo módulo se importe en CUDA, ROCm, CPU y XPU sin fallar, a costa de que cada operador requiere tres piezas de código: envoltorio de Python, implementación fake y guarda de plataforma.

El reenlace entre módulos del dispatcher de Triton es una solución agresiva. No depende de los hooks de importación de Python ni de__getattr__, sino que escanea y reemplaza directamente todas las referencias. La ventaja de este enfoque es que es exhaustivo: sin importar en cuántos lugares sefrom mod import kernelcopie el kernel, puede ser reemplazado; la desventaja es que es frágil: cualquier nueva forma de mantener una referencia al kernel (como la captura por closure) podría escapar del escaneo.

Resumen del capítulo

Reflexiones y autoevaluación del capítulo

下一章将转向高级推理特性，看前缀缓存如何复用 KV block、投机解码如何用小模型加速大模型、以及 LoRA 如何在不改基座权重的前提下动态切换适配器。

Este capítulo analiza las dos capas de infraestructura de cuantización y kernels personalizados de vLLM. La primera capa es el análisis de la configuración de cuantización: QuantSpec y QuantizationConfigArgs normalizan de forma unificada las cadenas de CLI, los metadatos del checkpoint y las sobrescrituras por capa en una QuantKey; resolve_quantization_config gestiona la expansión de abreviaturas y la fusión de campos, y _DEFERRED_ONLINE_SHORTHANDS resuelve escenarios de conflicto de nombres. La segunda capa es la adaptación de operadores: _custom_ops implementa el registro de operadores multiplataforma mediante guardas hasattr y register_fake; la implementación fake replica con precisión las formas de salida de los operadores reales para soportar torch.compile; el dispatcher implementa el reemplazo de plataforma de kernels Triton mediante KernelOverride y un escaneo de todo el módulo. Ambas capas sostienen conjuntamente la materialización de las ganancias de cuantización desde la carga de pesos hasta el cálculo forward. A continuación, pasaremos a las características avanzadas de inferencia que mejoran el throughput y reducen la latencia: cómo el caché de prefijos automático reutiliza los KV entre solicitudes, cómo la decodificación especulativa acelera la generación con un modelo borrador y cómo LoRA cambia dinámicamente de adaptador.

En el capítulo anterior profundizamos en el sistema de cuantización y la infraestructura de operadores personalizados de vLLM, vimos cómo se analiza la configuración de cuantización y se selecciona el kernel correspondiente, y cómo esquemas como FP8, INT4, AWQ y GPTQ completan la conversión durante la carga de pesos. Además, averiguamos cómo _custom_ops registra operadores CUDA, el mecanismo de programación de kernels Triton y cómo los kernels fusionados de MoE reducen los viajes de ida y vuelta a memoria. Estas capacidades de bajo nivel allanaron el camino para optimizaciones de inferencia más avanzadas. Este capítulo se centrará en tres características avanzadas de inferencia de vLLM: el caché de prefijos automático (APC), la decodificación especulativa y LoRA. Aunque parecen independientes, en realidad comparten el mismo conjunto de infraestructura subyacente: el hash de los bloques KV, la asignación de slots del planificador y la inyección dinámica de pesos durante la ejecución del modelo. La clave para entenderlas es comprender cómo llevan la «reutilización» al extremo sin romper la semántica de paginación de PagedAttention.

12.1 Caché de prefijos: cómo el block hash toma la huella digital de un prefijo

Modelo intuitivo

El caché de prefijos es como el «cuaderno de extractos de párrafos comunes» de una biblioteca: dos estudiantes escriben una redacción y ambos comienzan citando el mismo pasaje clásico; el profesor solo necesita corregir ese pasaje una vez, y luego revisar por separado las partes diferentes de cada uno. Sin él, cada solicitud tendría que hacer prefill de todo el prompt desde el principio, y en escenarios de preguntas y respuestas sobre documentos largos la potencia de cálculo se consumiría repetidamente varias veces.

Estructura de datos: del token al mapeo de block hash

El núcleo del caché de prefijos es «cómo determinar que los prefijos de dos solicitudes son iguales». La respuesta de vLLM es: dividir la secuencia de tokens en bloques y calcular un hash encadenado para cada bloque. Encadenado significa que el hash del N-ésimo bloque contiene el hash de los N-1 bloques anteriores, por lo que un block hash toma la huella digital única de todo el prefijo «desde el inicio de la secuencia hasta el final de ese bloque».

El portador del hash esBlockHash, que se define comobytesdeNewType, y no como unbytesdesnudo, con el objetivo de evitar a nivel de tipos el uso indebido de📎 vllm/v1/core/kv_cache_utils.py:59-62. Cuando es necesario combinar el block hash con el KV cache group id para formar una clave de diccionario, vLLM no usa una tupla, sino que concatena directamente el group id de 4 bytes en big-endian al final de los bytes del hash📎 vllm/v1/core/kv_cache_utils.py:75-76：

def make_block_hash_with_group_id(block_hash, group_id):
    return BlockHashWithGroupId(block_hash + group_id.to_bytes(4, "big", signed=False))

La función de hash en sí corre a cargo dehash_block_tokens, que alimenta a la función de hash el hash del bloque padre, la tupla de token ids del bloque actual y las claves adicionales juntos📎 vllm/v1/core/kv_cache_utils.py:650-680. Nótese que el hash padre del primer bloque no esNone, sino el globalNONE_HASH：

if not parent_block_hash:
    parent_block_hash = NONE_HASH

📎 vllm/v1/core/kv_cache_utils.py:674-675。NONE_HASHLa elección de la semilla de"vllm-none-hash"esconde un diseño de seguridad: para hashes criptográficos como SHA-256, la semilla es fija📎 vllm/v1/core/kv_cache_utils.py:105-126。resolve_none_hash_seed, lo que hace que distintos procesos de vLLM calculen el mismo hash para el mismo contenido y así compartan el caché de prefijos entre nodos; en cambio, para hashes no criptográficos como xxhash, la semilla es aleatoria por proceso, porque una semilla predecible permitiría a un atacante precalcular offline bloques en colisiónPYTHONHASHSEEDimplementa esta bifurcación:os.urandom(32) 📎 vllm/v1/core/kv_cache_utils.py:132-145。

la variable de entorno tiene prioridad; en caso contrario, los hashes criptográficos usan una semilla fija y los no criptográficos usan

Supongamos que una solicitud entra con 128 tokens y el block size es 16.get_request_block_hasherEl cierre devuelto se encarga del cálculo incremental📎 vllm/v1/core/kv_cache_utils.py:802-861：

Primer paso, determinar desde dónde empezar a calcular.start_token_idx = len(request.block_hashes) * hash_block_size 📎 vllm/v1/core/kv_cache_utils.py:812-812, es decir, el número de bloques ya calculados multiplicado por el tamaño de bloque. Si los tokens restantes no alcanzan un bloque, se devuelve vacío directamente📎 vllm/v1/core/kv_cache_utils.py:812-812。

Segundo paso, manejar el desplazamiento multimodal. Si la posición inicial cae dentro de alguna entrada multimodal, se necesita usarget_mm_features_in_windowpara reposicionarcurr_mm_idx 📎 vllm/v1/core/kv_cache_utils.py:823-832. Esto se debe a que el token placeholder de la entrada multimodal en sí no lleva semántica; es necesario incorporar el identificador de característica mm y su desplazamiento dentro del bloque como claves adicionales en el hash.

Tercer paso, calcular cada bloque en bucle.generate_block_hash_extra_keysRecopilar todas las claves adicionales📎 vllm/v1/core/kv_cache_utils.py:611-647, incluyendo nombre de LoRA, clave multimodal, cache salt, hash de prompt embeds. De estos, cache salt solo tiene efecto en el primer bloque📎 vllm/v1/core/kv_cache_utils.py:633-635, esto es intencional: la función de salt es aislar todo el espacio de nombres de caché, solo necesita inyectarse una vez en el punto de inicio de la cadena.

Cuarto paso,hash_block_tokenshashear juntos el hash padre, la tupla de tokens y las claves adicionales, y el resultado se usa como hash padre del siguiente bloque📎 vllm/v1/core/kv_cache_utils.py:851-857. La estructura en cadena se forma así.

Conversión de granularidad con múltiples block sizes

Cuando el modelo tiene múltiples grupos de KV cache y los block sizes son diferentes, la granularidad del hash y la granularidad de bloque del grupo pueden no coincidir.BlockHashListWithBlockSizeResolver este problema: no recalcula el hash, sino que aprovecha la propiedad del hash en cadena — el hash de un target block es el hash de su último hash block interno📎 vllm/v1/core/kv_cache_utils.py:2781-2851. Por ejemplo, cuando el hash block es 16 y el target block es 32, el hash de los tokens 0-31 es el segundo hash de tamaño 16 (que ya cubre en cadena 0-31)📎 vllm/v1/core/kv_cache_utils.py:2794-2806。_get_value_atLa implementación esself.block_hashes[(idx + 1) * self.scale_factor - 1] 📎 vllm/v1/core/kv_cache_utils.py:2848-2851。

flowchart TD
    req["Request 到达"] --> check{"剩余 token >= hash_block_size?"}
    check -->|否| empty["返回空列表"]
    check -->|是| mm{"起始位置在多模态窗口内?"}
    mm -->|是| reloc["get_mm_features_in_window 重定位 curr_mm_idx"]
    mm -->|否| extra
    reloc --> extra["generate_block_hash_extra_keys 收集 LoRA/MM/salt/embeds 键"]
    extra --> hash["hash_block_tokens 链式哈希"]
    hash --> append["追加到 new_block_hashes"]
    append --> advance["start_token_idx += hash_block_size"]
    advance --> check

Reflexiones de diseño y trampas encontradas

¿Por qué usar hash en cadena en lugar de hash independiente?El hash independiente no puede distinguir el caso de «el mismo bloque aparece en diferentes posiciones de prefijo». El hash en cadena hace que el block hash identifique de forma única todo el prefijo, lo cual es precisamentefind_longest_cache_hitla premisa para poder reutilizar KV de forma segura.

Trampa entre procesos del hash no criptográfico.Si se usa xxhash y no se configuraPYTHONHASHSEED, elNONE_HASHde cada proceso es diferente, lo que provoca que el caché de prefijos entre instancias falle por completo.init_none_hashSe imprimirá una advertencia📎 vllm/v1/core/kv_cache_utils.py:161-169. En producción, si se despliegan múltiples instancias compartiendo caché, se debe configurar explícitamentePYTHONHASHSEEDo cambiar a sha256.

La sutileza del desplazamiento multimodal. _gen_mm_extra_hash_keysUsar(mm_identifier, offset - start_token_idx)como clave adicional📎 vllm/v1/core/kv_cache_utils.py:552. El desplazamiento es relativo al inicio del bloque, de modo que el mismo ítem mm al aparecer en diferentes posiciones de bloque produce hashes distintos, evitando falsos aciertos.

12.2 Decodificación especulativa: la colaboración entre borrador y verificación

Modelo intuitivo

La decodificación especulativa es como si una secretaria redactara primero varias versiones de respuesta para el jefe, y el jefe solo tuviera que marcar rápidamente cuál sirve. El modelo borrador (drafter) predice múltiples tokens candidatos con un costo extremadamente bajo, y el modelo objetivo (target) verifica en paralelo estos candidatos en un solo forward, aceptando la parte que coincide. Sin esto, el modelo objetivo solo podría generar token por token de forma serial, y la utilización de GPU en la fase de decode sería extremadamente baja.

Estructura de datos: anotación del EAGLE group

El problema central de la decodificación especulativa en la gestión de KV cache es: ¿cómo se agrupan las capas KV del modelo borrador con las capas KV del modelo objetivo?_annotate_eagle_groupsSe usan dos reglas para identificar el grupo borrador📎 vllm/v1/core/kv_cache_utils.py:2134-2189：

Regla uno, impulsada por spec:non_causal_multi_token_decodeEl flag se declara enMLAAttentionSpec, lo establece la capa de atención borrador que ejecuta decode multi-token no causal, y puede sobrevivir a la operaciónmergede📎 vllm/v1/core/kv_cache_utils.py:2175-2177。

Regla dos, retroceso por posición: los borradores MTP (como DeepseekV4/V4.1 DSpark) reutilizan las propias capas decoder del modelo objetivo, no tienen marca en spec, pero sus capas de atención borrador siempre se registran después de todas las capas objetivo, por lo que se anota el grupo que contiene la última capa registrada📎 vllm/v1/core/kv_cache_utils.py:2183-2184. Esta regla solo tiene efecto cuando el grupo divide exactamentekv_cache_spectodas las capas📎 vllm/v1/core/kv_cache_utils.py:2183-2184。

Impulsado por escenario: asignación de KV en decodificación especulativa

Cuandospeculative_configestá habilitado yuse_eagle_block_drop()es verdadero,_annotate_eagle_groupsse invoca📎 vllm/v1/core/kv_cache_utils.py:2175-2177. El resultado de la anotaciónis_eagle_groupafecta la estrategia posterior de asignación de bloques — los bloques del grupo borrador pueden descartarse tras la verificación.

En la ruta principal deget_kv_cache_groups, la anotación ocurre después de la agrupación📎 vllm/v1/core/kv_cache_utils.py:2364-2365. Si ningún grupo es anotado como grupo borrador,_warn_if_unannotated_eagle_mambaemitirá una advertencia📎 vllm/v1/core/kv_cache_utils.py:2192-2222。

sequenceDiagram
    participant Sched as Scheduler
    participant Drafter as 草稿模型
    participant Target as 目标模型
    participant KV as KV Cache Manager
    Sched->>Drafter: 请求生成 k 个候选 token
    Drafter->>KV: 分配草稿组 block (is_eagle_group=True)
    Drafter-->>Sched: 返回候选 token 序列
    Sched->>Target: 并行验证候选 (一次前向)
    Target->>KV: 读取目标组 block
    Target-->>Sched: 返回接受/拒绝掩码
    Sched->>KV: 丢弃被拒绝的草稿 block

Reflexiones de diseño y trampas encontradas

¿Por qué el grupo borrador necesita anotación separada?Los tokens generados por el modelo borrador pueden ser rechazados tras la verificación, y el KV correspondiente debe descartarse. Si el KV borrador y el KV objetivo se mezclan en el mismo grupo, la operación de descarte afectaría por error al KV objetivo. La anotación permite al planificador recuperarlos con precisión.

Fragilidad de la regla de retroceso por posición.La regla dos depende de la convención de que «la capa borrador se registra al final»; en los comentarios se marca explícitamente como hacky check y se deja un FIXME📎 vllm/v1/core/kv_cache_utils.py:2158-2159. Cuando el caché de cola del borrador abarca múltiples grupos, esta regla solo anota el grupo que contiene la última capa, y necesita generalizarse.

Restricciones adicionales de los modelos Mamba.Si se habilita la decodificación especulativa pero ningún grupo se identifica como grupo borrador, y existe un grupo Mamba, se activa una advertencia📎 vllm/v1/core/kv_cache_utils.py:2211-2213. Esto normalmente significa que la spec de la capa borrador no se puede distinguir de la capa objetivo, y es necesario verificar el orden de registro del modelo.

12.3 LoRA: adaptadores dinámicos sin recargar la base

Modelo intuitivo

LoRA es como cambiarle la funda a un mismo teléfono: el cuerpo del teléfono (modelo base) no cambia, y al cambiar la funda (adaptador) se convierte en un estilo diferente. Sin esto, cada tarea de ajuste fino tendría que cargar una copia completa de los pesos, y la memoria de video no podría soportarlo.

Estructura de datos: caché LRU doble y arreglo de slots

LoRAModelManagerSe usan dos cachés LRU para gestionar el ciclo de vida de los adaptadores📎 vllm/lora/model_manager.py:115-120：

self._registered_adapters: AdapterLRUCache[LoRAModel] = AdapterLRUCache(
    self.capacity, self.deactivate_adapter
)
self._active_adapters: AdapterLRUCache[None] = AdapterLRUCache(
    self.lora_slots, self._deactivate_adapter
)

capacityEs el número total de adaptadores que se pueden almacenar en caché del lado de CPU (max_cpu_loras）📎 vllm/lora/model_manager.py:340-342，lora_slotsEs el número de adaptadores que se pueden activar simultáneamente del lado de GPU (max_loras）📎 vllm/lora/model_manager.py:345-346。_registered_adaptersCuando se elimina, se activadeactivate_adaptercallback📎 vllm/lora/model_manager.py:71-74, lo que asegura que al desalojar la caché de CPU también se limpien las copias en GPU.

lora_index_to_idEs un arreglo de longitudlora_slotsque mapea índices de slot de GPU a id de adaptador📎 vllm/lora/model_manager.py:122. Este arreglo es el índice central cuando el punica wrapper realiza cálculos LoRA por lotes.

Guiado por escenarios: activación de adaptadores

Cuando una solicitud llega con un adaptador LoRA,activate_adapterse invoca📎 vllm/lora/model_manager.py:352-409：

Primer paso, verificar si ya está activado; si lo está, retornar directamente📎 vllm/lora/model_manager.py:352-354。

Segundo paso, buscar un slot libre. Recorrerlora_index_to_idpara encontrar el primerNone 📎 vllm/lora/model_manager.py:362-362. Si no hay slot libre, lanzarValueError("No free lora slots") 📎 vllm/lora/model_manager.py:368-368。

Tercer paso, actualizar el estado y recorrer todos los módulos envueltos, invocandomodule.set_lora(index, lora_a, lora_b)para copiar los pesos al stacked buffer de GPU📎 vllm/lora/model_manager.py:377-401. Si algún módulo no tiene pesos LoRA correspondientes, invocarreset_lora(index)para poner a cero📎 vllm/lora/model_manager.py:378-385。

Cuarto paso, si no se aplicó ningún peso, imprimir un log de depuración único📎 vllm/lora/model_manager.py:411-416. Esto es el comportamiento esperado bajo paralelismo de pipeline o paralelismo de expertos: algunos ranks no poseen las capas adaptadas.

Envoltura de módulos: de nn.Linear a BaseLayerWithLoRA

_create_lora_modulesRecorrer todos los módulos nombrados del modelo📎 vllm/lora/model_manager.py:462-606. Lógica clave:

• OmitirPPMissingLayer 📎 vllm/lora/model_manager.py:473-474。
• Filtrar segúntarget_modules: si no se especifica, usaris_supported_lora_modulepara determinar; de lo contrario usar_match_target_modules 📎 vllm/lora/model_manager.py:479-493。
• Manejar módulos alias: un mismo módulo subyacente puede ser accedido por múltiples rutas (por ejemplo, el gate de MoE está tanto en el block como dentro del runner). En este caso, redirigir el atributo alias al mismo wrapper, pero no registrarlo de nuevo, de lo contrarioactivate_adapterinvocará sobre el aliasreset_loray borrará los pesos recién establecidos📎 vllm/lora/model_manager.py:512-527。
• Usarfrom_layerpara crear el wrapper y reemplazar el módulo original📎 vllm/lora/model_manager.py:546-553。

Reflexiones de diseño y trampas

Los cambios en el diseño de slots activan la actualización del mapeo. set_adapter_mappingNo solo compara si el mapping cambió, sino que también comparalora_index_to_idla instantánea de tupla de📎 vllm/lora/model_manager.py:1323-1331. La razón está claramente explicada en los comentarios: unadd_lora()fuera de banda puede activar el desalojo LRU y reasignar slots, mientras que el batch en ejecución y su mapping no cambiaron📎 vllm/lora/model_manager.py:1323-1331. Si solo se mira el mapping, los metadatos de punica usarán un diseño de slots obsoleto.

Segmentación EP de MoE.Cuando se habilita el paralelismo de expertos, el checkpoint posee los pesos de todos los expertos globales, pero cada rank solo poseelocal_num_experts._stack_moe_lora_weightsPrimero segúnglobal_num_expertsreshape, luego segmentar[expert_start:expert_end] 📎 vllm/lora/model_manager.py:966-977. Cuando no es EP, la segmentación es no-op.

Momento de pin_memory.El empaquetado de pesos (comopack_moe) puede invalidar la asignación de pin_memory, por lo que pin_memory se ejecuta después de fusionar todos los pesos📎 vllm/lora/model_manager.py:916-934. Los comentarios señalan explícitamente dos razones: los modelos MoE tienen una gran cantidad de pesos LoRA, y hacer pin demasiado pronto tiene un costo notable; el empaquetado puede invalidar la asignación📎 vllm/lora/model_manager.py:916-921。

Reflexión de diseño: el punto de sinergia de los tres

Las tres características convergen en la capa de gestión de KV cache. La caché de prefijos reutiliza KV mediante block hash; la decodificación especulativa medianteis_eagle_groupanotaciones para distinguir KV borrador; LoRA mediante_gen_lora_extra_hash_keysincorpora el nombre del adaptador al block hash📎 vllm/v1/core/kv_cache_utils.py:568-581, asegurando que secuencias de tokens idénticas con adaptadores diferentes no colisionen erróneamente entre sus KV.

generate_block_hash_extra_keysColoca la clave LoRA al principio de la lista de claves adicionales📎 vllm/v1/core/kv_cache_utils.py:640-642, junto con las claves multimodales, cache salt y prompt embeds, formando la entrada hash completa. Esto garantiza que: incluso si dos solicitudes tienen tokens completamente idénticos, siempre que sus adaptadores LoRA sean diferentes, sus block hash serán diferentes y los KV no se mezclarán.

Resumen del capítulo

Reflexiones y autoevaluación del capítulo

El siguiente capítulo se centrará en el sistema de plugins y la extensibilidad, para ver cómo vLLM admite formas de despliegue diversas mediante la abstracción de plataformas, los procesadores de IO y la extensión de endpoints.

Este capítulo analiza los mecanismos subyacentes de las tres características avanzadas de inferencia de vLLM. El núcleo de la caché de prefijos es el hash de bloques encadenado: hash_block_tokens hashea juntos el hash padre, la tupla de tokens y las claves adicionales, y la estrategia de semilla de NONE_HASH equilibra el uso compartido entre procesos y la seguridad frente a colisiones. La decodificación especulativa distingue los grupos de KV de borrador mediante la anotación is_eagle_group. LoRA gestiona el ciclo de vida de los adaptadores mediante una caché LRU doble y un arreglo de slots, y mezcla el nombre del adaptador en el hash de bloques para lograr el aislamiento de caché. Estas características muestran en conjunto la profundidad y flexibilidad de vLLM en la optimización de inferencia. A continuación, nos centraremos en el sistema de plugins y la extensibilidad de vLLM, para ver cómo los plugins de plataforma se adaptan a nuevo hardware, cómo los plugins de IO processor intervienen en el procesamiento de entradas multimodales y cómo los plugins de endpoints inyectan rutas de API personalizadas. Comprender el orden de carga del registro y descubrimiento de plugins revelará cómo ampliar las capacidades de vLLM sin modificar el código central.

En el capítulo anterior vimos que características avanzadas como la caché de prefijos, la decodificación especulativa y LoRA están profundamente acopladas en las rutas centrales del planificador, la gestión de KV y la ejecución del modelo. Pero para que un motor de inferencia llegue realmente a producción, no basta con el rendimiento: debe responder a una pregunta más espinosa: cuando la comunidad quiere integrar nuevo hardware, un nuevo formato de entrada multimodal o una ruta HTTP personalizada, ¿cómo hacerlo sin forkear el código central? Este es precisamente el sentido de la existencia del sistema de plugins. La arquitectura de vLLM es naturalmente multiproceso: el proceso frontend del API Server, el proceso EngineCore y el proceso Worker correspondiente a cada rango TP/PP. Si el mecanismo de plugins simplemente «ejecutara un fragmento de código al importar», entonces o bien se ejecutaría repetidamente en cada proceso provocando una acumulación de efectos secundarios, o bien se ejecutaría solo en el proceso principal y los Workers no obtendrían la extensión. Lo que este capítulo va a desglosar es cómo vLLM utiliza el mecanismo estándar de entry_points de Python, junto con la triple restricción de grupo + límite de proceso + momento de carga, para construir un sistema de plugins que cubra todos los procesos y a la vez controle con precisión la superficie expuesta. Nos centramos en tres líneas principales: plugins de plataforma (adaptación a nuevo hardware), plugins de IO processor (intervención en el procesamiento de entradas multimodales) y plugins de endpoints (inyección de rutas de API personalizadas). Las estrategias de carga de los tres son completamente distintas; comprender esta diferencia es comprender la filosofía de equilibrio de vLLM entre «capacidad de extensión» y «límite de seguridad».

I. Descubrimiento y carga de plugins: el contrato de agrupación de entry_points

Modelo intuitivo: los «canales de difusión» de los plugins

Imagina el sistema de plugins de vLLM como un conjunto de canales de difusión. Cada paquete de plugin, al instalarse, mediantesetup.pydeentry_points«registra» en algún canal su indicativo (plugin name) y su función de respuesta (plugin value). vLLM escanea estos canales al arrancar y decide qué canales se «escuchan» en qué procesos.

Sin este mecanismo, extender vLLM solo sería posible modificando el código fuente: cada vez que la comunidad añadiera hardware, habría que mantener un fork, lo que acabaría fragmentando las versiones. El valor del mecanismo de agrupación radica en:Un mismo paquete de plugin puede registrarse solo en un canal específico, quedando limitado a cargarse en un proceso determinado。

Estructura de datos: cinco constantes de grupo y una bandera global

vLLM define envllm/plugins/__init__.pyla parte superior cinco constantes de entry point group, cada una correspondiente a una estrategia de carga:

📎 vllm/plugins/__init__.py:16-30

DEFAULT_PLUGINS_GROUP = "vllm.general_plugins"
IO_PROCESSOR_PLUGINS_GROUP = "vllm.io_processor_plugins"
PLATFORM_PLUGINS_GROUP = "vllm.platform_plugins"
STAT_LOGGER_PLUGINS_GROUP = "vllm.stat_logger_plugins"
ENDPOINT_PLUGINS_GROUP = "vllm.endpoint_plugins"

En los comentarios se esconde información clave:DEFAULT_PLUGINS_GROUPentodos los procesoscargar (process0, engine core, worker);IO_PROCESSOR_PLUGINS_GROUP solo en process0；PLATFORM_PLUGINS_GROUPse carga en todos los procesos, pero el momento de activación escurrent_platformla primera vez que se accede;STAT_LOGGER_PLUGINS_GROUPsolo en process0 y en modo asíncrono;ENDPOINT_PLUGINS_GROUPsolo en el proceso frontend del API Server.

Inmediatamente después hay una variable global a nivel de móduloplugins_loaded = False 📎 vllm/plugins/__init__.py:32-33que actúa como guarda para la carga idempotente — el comentario dice explícitamente "make sure one process only loads plugins once".

Paso a paso: unaload_plugins_by_groupsecuencia completa de llamadas

Supongamos el escenario: el usuario registró ensetup.pyvllm.general_pluginsbajoregister_dummy_modely ahora vLLM arranca, algún proceso llama aload_general_plugins()。

Primer paso: guarda de idempotencia. load_general_pluginsPrimero compruebaplugins_loadedsi ya esTruedevuelve directamente📎 vllm/plugins/__init__.py:77-90. Aquí hay una sutileza: la guarda se activaantesde cargar, lo que significa que aunque la carga posterior lance una excepción, no se reintentará. Esto es intencional — un fallo en la carga de un plugin no debería provocar que el proceso lo intente repetidamente.

Segundo paso: descubrimiento.Entra enload_plugins_by_groupy medianteimportlib.metadata.entry_points(group=group)obtiene todos los entry points instalados bajo ese grupo📎 vllm/plugins/__init__.py:36-45. Si está vacío, registra un log de debug y devuelve un diccionario vacío.

Tercer paso: niveles de log.El código fuente distingue el nivel de log entre grupos por defecto y no por defecto:is_default_groupcuando es verdadero usalogger.debug, de lo contrario usalogger.info 📎 vllm/plugins/__init__.py:47-54. La motivación es práctica —vllm.general_pluginsnormalmente agrupa una gran cantidad de plugins de registro de modelos, y usar INFO saturaría los logs; en cambio, los plugins de plataforma/endpoint son pocos e importantes, y merecen ser visibles con INFO.

Cuarto paso: filtrado por lista blanca.Leeenvs.VLLM_PLUGINS, si esNonecarga todos, de lo contrario solo carga los plugins cuyos nombres estén en la lista📎 vllm/plugins/__init__.py:62-70. Nótese queplugin.load()está envuelto en try/except, de modo que el fallo de carga de un solo plugin solo registra un log de exception y no afecta a los demás plugins📎 vllm/plugins/__init__.py:68-72。

Quinto paso: ejecución.Vuelve aload_general_pluginsy para cada función cargada llama directamente afunc() 📎 vllm/plugins/__init__.py:77-90. Por eso la documentación insiste en que las funciones de plugin deben serreentrantes (re-entrant)— pueden ser llamadas múltiples veces en múltiples procesos.

El siguiente diagrama de flujo describe laload_plugins_by_groupruta de decisión completa de

flowchart TD
    start["load_plugins_by_group(group)"] --> discover["entry_points(group=group)"]
    discover --> empty{"len(discovered) == 0?"}
    empty -->|是| ret_empty["返回 {}"]
    empty -->|否| log["按 is_default_group 选 log_level"]
    log --> loop["遍历 discovered_plugins"]
    loop --> check{"allowed_plugins is None<br/>或 plugin.name in allowed?"}
    check -->|否| skip["跳过该插件"]
    check -->|是| load["func = plugin.load()"]
    load --> load_ok{"加载成功?"}
    load_ok -->|否| log_exc["logger.exception 记录"]
    load_ok -->|是| add["plugins[name] = func"]
    skip --> next["下一个插件"]
    log_exc --> next
    add --> next
    next --> loop
    loop --> ret["返回 plugins 字典"]

Reflexión de diseño: por qué usar entry_points en lugar de un archivo de configuración

---

II. Plugins de plataforma: la capa de abstracción para la adaptación de hardware

Modelo intuitivo: la plataforma es el "traductor de dialectos de hardware"

PlatformLa clasees el único traductorde toda la conversación entre vLLM y el hardware. El código del modelo solo llama acurrent_platform.get_attn_backend_cls()、current_platform.is_cuda_alike()métodos abstractos comoimport torch.cuda, nunca directamente aif device == "xpu". Sin esta capa de abstracción, cada vez que se soportara un nuevo hardware habría que añadir ramas de

en el código del modelo, lo que acabaría convirtiéndose en espagueti.

PlatformEstructura de datos: disposición de campos de la clase base Platformvllm/platforms/interface.pyes una clase pura (no se usa instanciada), y sus atributos de clase clave se definen al inicio de📎 vllm/platforms/interface.py:135-179：

class Platform:
    _enum: PlatformEnum
    device_name: str
    device_type: str
    dispatch_key: str = "CPU"
    ray_device_key: str = ""
    device_control_env_var: str = "VLLM_DEVICE_CONTROL_ENV_VAR_PLACEHOLDER"
    ray_noset_device_env_vars: list[str] = []
    simple_compile_backend: str = "inductor"
    dist_backend: str = ""
    supported_quantization: list[str] = []
    additional_env_vars: list[str] = []
    _global_graph_pool: Any | None = None

_enumes el valor de enumeración dePlatformEnumque determinais_cuda()、is_rocm()y otras comprobaciones📎 vllm/platforms/interface.py:69-78。device_control_env_vares la abstracción de "variable de entorno de visibilidad de dispositivo" independiente de la plataforma — CUDA esCUDA_VISIBLE_DEVICES, y otras plataformas definen cada una su📎 vllm/platforms/interface.py:151-152。_global_graph_pooles la caché de memoria de CUDA graph a nivel de clase, inicializada de forma perezosa medianteget_global_graph_pool📎 vllm/platforms/interface.py:1210-1215。

Cabe destacar la lógica de respaldo de__getattr__📎 vllm/platforms/interface.py:1189-1208: cuando se accede a un atributo que no existe en Platform, intenta reenviarlo desde el espacio de nombrestorch.<device_type>. Esto permite que el código de plataforma escribacurrent_platform.memory_allocated()mientras en realidad llama atorch.cuda.memory_allocated(). Pero el código fuente excluye deliberadamente los métodos dunder — de lo contrario, al comprobar pickle__getstate__obtendríaNonee intentaría llamarlo📎 vllm/platforms/interface.py:1182-1185。

Paso a paso: la conversión de device ID entre tres espacios de nombres

Lo más propenso a errores en la abstracción de plataforma es elespacio de nombres de device ID. Los comentarios del código fuente enumeran explícitamente tres📎 vllm/platforms/interface.py:275-283：

• logical: el local rank interno de vLLM, que indexa_assigned_physical_gpu_ids
• visible: el número de torch/CUDA tras el remapeo deCUDA_VISIBLE_DEVICESen el proceso actual
• physical: el GPU ID global usado por APIs de topología como NVML, no afectado por variables de entorno

Supongamos el escenario: a un proceso Worker se le asigna la GPU física[4, 5], la variable de entornoCUDA_VISIBLE_DEVICES=4,5, y ahora hay que convertir el local rank 0 atorch.device("cuda:0")。

Primer paso: logical → physical. device_id_to_physical_device_id(0)Primero consulta_assigned_physical_gpu_ids, si ya está establecido lo indexa y devuelve directamente4 📎 vllm/platforms/interface.py:296-297. Si no está establecido, divide la lista separada por comas dedevice_control_env_vary toma el elemento 0📎 vllm/platforms/interface.py:305-311. Nótese que el código fuente trata deliberadamente lacadena vacíacomo no establecida — esta es una configuración válida cuando Ray arranca el motor sobre un placement group puramente de CPU📎 vllm/platforms/interface.py:296-297。

Segundo paso: physical → visible. logical_device_id_to_visible_device_id(0)Una vez obtenido physical4, se descompone la variable de entorno en[4, 5], se busca4el índice de0y se devuelve📎 vllm/platforms/interface.py:316-339. Si el physical ID no está en la lista visible, se lanzaRuntimeError——esta es una protección estricta para evitar el uso indebido de dispositivos no visibles entre procesos.

set_assigned_physical_gpu_idsEl diseño idempotente deRuntimeError 📎 vllm/platforms/interface.py:38-56también merece atención: establecer el mismo valor repetidamente es una no-op, mientras que establecer un valor diferente lanza

. Esto evita que el mapeo de dispositivos sea sobrescrito accidentalmente en entornos multihilo.

Registro de plugins de plataforma e inyección de configuraciónvllm.platform_pluginsLos plugins de plataforma se registran mediante el grupoNone, y la función del plugin devuelve el nombre completamente cualificado de la clase de plataforma (o📎 docs/design/plugin_system.md:50-50para indicar que el entorno actual no es compatible)📎 docs/design/plugin_system.md:100-100：

• _enum. La implementación mínima proporcionada por la documentación requiere quePlatformEnum.OOT（out-of-tree）
• device_typenormalmente se establezca en
• check_and_update_configy devuelva la cadena de tipo de dispositivo que PyTorch reconocese invoca temprano durante la inicialización de vLLM,worker_cls
• get_attn_backend_clsy se debe establecer aquí
• get_device_communicator_clsdevuelve el nombre de la clase del backend de atención

check_and_update_configdevuelve el nombre de la clase del comunicador📎 vllm/platforms/interface.py:583-592es el hook más crítico del plugin de plataformaVllmConfig. Recibe una referencia a📎 docs/design/plugin_system.md:105-105y la modifica in situ, pudiendo ajustar block size, graph mode, etc. La documentación enfatiza que "lo más importante es que worker_cls debe establecerse aquí"

——porque vLLM necesita saber qué clase Worker usar para instanciar el proceso de trabajo.

Reflexión de diseño: estrategia de tres fases para la alineación del block sizeupdate_block_size_for_backend 📎 vllm/platforms/interface.py:666-708La lógica más compleja en la interfaz de plataforma es

Phase 1. Se divide en tres fases para garantizar que el block size sea compatible con el backend de atención:--block-size: si el usuario no especifica explícitamente_preferred_block_size_for_backends, se llama a📎 vllm/platforms/interface.py:687-697para seleccionar el block size mínimo compatible con todos los backends📎 vllm/platforms/interface.py:622-663。

Phase 2. Esta función enumera valores candidatos usando LCM (mínimo común múltiplo), porque algunos backends (como CPU_MLA) solo aceptan tamaños exactos y no múltiplos📎 vllm/platforms/interface.py:699-702。

Phase 3: los modelos híbridos (attention + mamba) necesitan alinear el block con el mamba page size📎 vllm/platforms/interface.py:704-708。

---

Este diseño por fases refleja la realidad que enfrenta vLLM: las restricciones sobre el block size de diferentes hardware, esquemas de cuantización y arquitecturas de modelos entran en conflicto entre sí, y no pueden resolverse con una única fórmula. Dividirlo en fases permite que cada restricción se maneje de forma independiente y, finalmente, se tome la solución que satisface todas las restricciones.

III. IO Processor y plugins de endpoint: procesamiento de entrada y extensión de API

Modelo intuitivo: IO Processor es una "capa de traducción multimodal"

La entrada de un modelo multimodal (como LLaVA) no es texto puro, sino una mezcla de texto + imágenes. El plugin IO Processor se encarga de convertir los datos multimodales originales en tensores que el modelo puede consumir, y luego convertir la salida del modelo de vuelta a un formato legible por humanos. Es como un traductor de aduanas: el idioma extranjero que entra (imagen/audio) se traduce al idioma nativo del modelo, y el idioma nativo del modelo que sale se traduce de vuelta al idioma extranjero.

Paso a paso: descubrimiento e instanciación de IO Processorio_processor_pluginEscenario: cargar un modelo con un HF config que tiene el campo

. get_io_processorPrimer paso: determinar el nombre del plugin.plugin_from_initSe prioriza elhf_configpasado explícitamente; de lo contrario, se leeio_processor_plugindesde el campo📎 vllm/plugins/io_processors/__init__.py:42-50deNone. Si ambos están vacíos, se devuelve📎 vllm/plugins/io_processors/__init__.py:52-54。

——lo que indica que el modelo no necesita IO processorSegundo paso: cargar todos los plugins instalados.load_plugins_by_group(IO_PROCESSOR_PLUGINS_GROUP)Se llama a📎 vllm/plugins/io_processors/__init__.py:59-61。

para obtener todos los plugins de ese grupoTercer paso: construir el mapeo cargable.processor_cls_qualnameSe recorre cada plugin, se llama a su función para obtenerNone, y si no esloadable_plugins 📎 vllm/plugins/io_processors/__init__.py:66-76se registra en

. Nótese que la llamada a la función de cada plugin también está envuelta en try/except, de modo que un fallo individual no afecta a los demás.Cuarto paso: validación e instanciación.ValueErrorSi el número de plugins cargables es 0, se lanza📎 vllm/plugins/io_processors/__init__.py:66-76indicando "se requiere un plugin IOProcessor pero no hay ninguno instalado"ValueError. Si el nombre del plugin requerido por el modelo no está en la lista de cargables, se lanza📎 vllm/plugins/io_processors/__init__.py:80-81y se listan todos los nombres de plugins disponiblesresolve_obj_by_qualname. Finalmente, mediante📎 vllm/plugins/io_processors/__init__.py:80-81。

se resuelve el nombre de la clase y se instancia

Plugins de endpoint: postura de seguridad de denegación por defectoLos plugins de endpoint son la categoría más especial de este capítulo, porque。load_endpoint_pluginsno se cargan por defectoload_plugins_by_group. La cadena de documentación de📎 vllm/plugins/__init__.py:93-94。

explica claramente la razón: los plugins de endpoint añaden rutas HTTP al API Server, ampliando la superficie de exposición de red, por lo que adoptan una postura de "denegación por defecto" más estricta queLa regla concreta es: solo cuando el nombre del pluginVLLM_PLUGINSaparece explícitamente en, y surequired_tasksesNoneo tiene intersección con las tasks soportadas por el servidor, se carga📎 vllm/plugins/__init__.py:108-108。

Escenario: el usuario instaló un plugin de endpoint pero olvidó establecerVLLM_PLUGINS。

Primer paso: comprobar si VLLM_PLUGINS no está establecido.Sienvs.VLLM_PLUGINS is None, primero se descubren los plugins de ese grupo y, si los hay, se registra un warning indicando "debe estar explícitamente en la allowlist"📎 vllm/plugins/__init__.py:126-126. Nótese que los comentarios del código fuente señalan especialmente:VLLM_PLUGINS=""se interpreta como[""]en lugar deNone, por lo que se considera una "allowlist que no coincide con ningún plugin", en vez de "no establecido"📎 vllm/plugins/__init__.py:108-108. Esta distinción de límites es importante——una cadena vacía es un "no cargar nada" explícito, mientras queNonees "no configurado".

Segundo paso: cargar e instanciar.Tras obtener la función de fábrica medianteload_plugins_by_group, se llama una por una afactory()para instanciar📎 vllm/plugins/__init__.py:133-141. Si la instanciación falla, se registra la excepción y se continúa.

Tercer paso: control por task.Se compruebaplugin.required_tasks, y si no esNoney tiene intersección consupported_tasksSin intersección, se omite este plugin📎 vllm/plugins/__init__.py:144-145. Esto permite que el mismo paquete de plugin registre diferentes endpoints para distintas tareas (como embedding vs generation).

El siguiente diagrama de secuencia describe la interacción completa del plugin de endpoint desde el descubrimiento hasta la carga:

sequenceDiagram
    participant App as "API Server 前端进程"
    participant Loader as "load_endpoint_plugins()"
    participant Env as "envs.VLLM_PLUGINS"
    participant EP as "entry_points(ENDPOINT_PLUGINS_GROUP)"
    participant Factory as "plugin factory()"

    App->>Loader: load_endpoint_plugins(supported_tasks)
    Loader->>Env: 读取 VLLM_PLUGINS
    alt VLLM_PLUGINS is None
        Loader->>EP: entry_points(group)
        EP-->>Loader: discovered plugins
        Loader-->>App: 返回 [] (记 warning)
    else VLLM_PLUGINS 已设置
        Loader->>EP: load_plugins_by_group(group)
        EP-->>Loader: factories 字典
        loop 每个 factory
            Loader->>Factory: factory()
            Factory-->>Loader: EndpointPlugin 实例
            Loader->>Loader: 检查 required_tasks 交集
            alt tasks 不匹配
                Loader->>Loader: 跳过 (记 info)
            else tasks 匹配
                Loader->>Loader: append 到结果列表
            end
        end
        Loader-->>App: 返回 endpoint_plugins 列表
    end

Reflexión de diseño: el límite de proceso determina la estrategia de carga

La diferencia en las estrategias de carga de los tres tipos de plugins es, en esencia, un mapeo dellímite de proceso:

Trampa en producción: degradación silenciosa por fallo de carga de plugins

load_plugins_by_groupPara cada plugin,plugin.load()se envuelve con try/except, y en caso de fallo solo se registra la exception📎 vllm/plugins/__init__.py:68-72. Esto significa queun plugin corrupto no impedirá que vLLM se inicie, pero tampoco dará un error explícito—el usuario puede confundirse preguntándose "¿por qué mi plugin no funciona?".

Sugerencia de diagnóstico: ajusta el nivel de log a DEBUG, busca"Failed to load plugin". Si el plugin está bajo el grupovllm.general_plugins, el nivel de log predeterminado es DEBUG, y se necesita habilitarlo explícitamente para ver los detalles de carga📎 vllm/plugins/__init__.py:49-50。

Otra trampa es el momento de activación del guardiánplugins_loaded📎 vllm/plugins/__init__.py:77-90: se establece antes de la cargaTrue. Si la primera carga falla por alguna razón (como una excepción en el escaneo de entry_points), las llamadas posteriores retornarán directamente sin reintentar. Esto puede causar el extraño fenómeno de "el plugin funciona a veces sí y a veces no" en entornos de prueba.

---

Resumen del capítulo

El sistema de plugins de vLLM se construye sobre Pythonentry_points, mediantecinco constantes de grupose dividen los tipos de extensión, medianteel límite de procesose determina el alcance de carga, medianteVLLM_PLUGINSla lista blancase controla el conjunto de carga. Los plugins de plataforma usan la clase basePlatformpara abstraer las diferencias de hardware, y su conversión de tres espacios de nombres de device ID (logical/visible/physical) es el núcleo de la gestión de dispositivos entre procesos; los plugins de IO processor se activan mediante el campoio_processor_plugindel HF config, y se encargan de la traducción de entradas multimodales; los plugins de endpoint adoptan una postura de "rechazo predeterminado", y solo se cargan cuando hay un allowlist explícito y el task coincide, para controlar la superficie de exposición de red.

Las tres líneas principales comparten el mismo mecanismo de descubrimiento, pero las diferencias en las estrategias de carga reflejan el equilibrio de vLLM entre la "conveniencia de extensión" y el "límite de seguridad": los plugins que no exponen red se cargan por defecto, los plugins que exponen red deben ser opt-in.

Reflexiones y autoevaluación de este capítulo

---

El siguiente capítulo girará hacia los compromisos arquitectónicos, los escollos en producción y la evolución futura; reuniremos los mecanismos desglosados en los trece capítulos anteriores, examinaremos las concesiones de vLLM entre rendimiento, mantenibilidad y extensibilidad, y vislumbraremos la dirección de evolución de los motores de inferencia.

Hasta aquí, hemos visto con claridad cómo vLLM, mediante el mecanismo de agrupación de entry_points, el momento de carga consciente de los límites de proceso y las estrategias diferenciadas para los tres tipos de plugins (plataforma, IO processor y endpoint), abre la superficie de extensión mientras mantiene estable el código central. Este sistema de plugins permite que nuevo hardware, nuevos formatos de entrada y nuevas rutas de API se integren de forma no invasiva, pero la extensibilidad en sí misma también implica más dimensiones que deben sopesarse. El siguiente capítulo cerrará el libro, sistematizando las tensiones en las decisiones clave de diseño de vLLM——procesamiento por lotes continuo frente a fragmentación de memoria de video, CUDA Graph frente a formas dinámicas, despliegue separado frente a sobrecarga de red——y ofrecerá una lista de escollos en entornos de producción y una ruta de diagnóstico, además de vislumbrar las tendencias de evolución hacia el frontend en Rust, la capa IR y el hardware heterogéneo.

En el capítulo anterior desglosamos el mecanismo de extensión por plugins de vLLM y vimos cómo los plugins de plataforma, los plugins de IO processor y los plugins de endpoint permiten que el motor se adapte a nuevo hardware, nuevas modalidades y nuevas API sin modificar el código central. Esta extensibilidad permite que vLLM abrace rápidamente los cambios, pero cuantos más puntos de extensión haya, más complejas serán las rutas de interacción en entornos de producción. Cuando problemas reales como la fragmentación de la memoria de video, el fallo del handshake de NCCL, la invalidación de la caché de compilación y la fluctuación de red aparecen simultáneamente, los mecanismos presentados en los trece capítulos anteriores se tensionan entre sí y exponen tensiones que no se manifestaban en entornos ideales. Este capítulo no introduce nuevos mecanismos centrales, sino que reúne estos mecanismos, tomando como ancla la documentación oficial de troubleshooting, combinándolos con el diseño de la herramienta bench del frontend en Rust, para examinar las concesiones entre rendimiento y operabilidad, y ofrecer una ruta de diagnóstico accionable.

I. Niveles de optimización: un contrato explícito entre tiempo de arranque y rendimiento en ejecución

Modelo intuitivo

Los niveles de optimización son como los "modos de escena" de una cámara: el modo automático (-O2) sirve para la mayoría de escenarios, pero cuando necesitas una captura rápida (depuración), cambiar al modo manual (-O0) responde de inmediato, a costa de una caída en la calidad de imagen (rendimiento). vLLM convierte este compromiso en un contrato explícito de cuatro niveles, en lugar de esconderlo entre decenas de flags booleanos para que el usuario los combine por su cuenta.

Distribución de campos de los cuatro niveles

vLLM ofrece-O0hasta-O3cuatro niveles📎 docs/design/optimization_levels.md:5-5. El principio de diseño central es:los flags establecidos explícitamente por el usuario tienen prioridad sobre los valores predeterminados del nivel de optimización 📎 docs/design/optimization_levels.md:5-5. Esto significa que el nivel de optimización es solo un conjunto de valores predeterminados, no una restricción rígida.

-O0desactiva todo: sin autotuning, sin compilación, sin cudagraph📎 docs/design/optimization_levels.md:32-33. En concreto, se reduce a cuatro interruptores:cudagraph_mode=NONE、mode=NONE, todas las fusiones desactivadas,enable_flashinfer_autotune=False 📎 docs/design/optimization_levels.md:37-40。

-O1es el punto de equilibrio para escenarios de desarrollo: habilitaPIECEWISEcudagraph y el modoVLLM_COMPILE. Nótese un detalle sutil:📎 docs/design/optimization_levels.md:50-51yfuse_norm_quantsolo se habilitan cuando uno de los operadores usa un kernel personalizado; de lo contrario, la fusión automática de Inductor funciona mejorfuse_act_quant. Esta es una decisión de diseño típica de "no quitarle trabajo al compilador".📎 docs/design/optimization_levels.md:61es el valor predeterminado, orientado a producción

-O2. Sobre la base de📎 docs/design/optimization_levels.md:66-67añade-O1cudagraph yFULL_AND_PIECEWISEactualmente equivale afuse_allreduce_rms 📎 docs/design/optimization_levels.md:72-73。-O3, reservando-O2para optimizaciones experimentales más agresivas en el futuro📎 docs/design/optimization_levels.md:80-81。

Flujo de selección guiado por escenarios

Cuando un usuario ejecutavllm serve model -O1, ¿qué ocurre internamente? El siguiente diagrama de flujo muestra cómo interactúan los niveles de optimización con los flags del usuario:

flowchart TD
    start["用户启动 vllm serve -O1"] --> parse["解析 optimization_level=1"]
    parse --> load_defaults["加载 O1 默认值集合"]
    load_defaults --> check_user{"用户是否显式设置了<br/>cudagraph_mode?"}
    check_user -->|是| user_wins["使用用户值<br/>覆盖 O1 默认"]
    check_user -->|否| use_default["使用 O1 默认<br/>PIECEWISE"]
    user_wins --> check_fusion{"fuse_norm_quant<br/>是否涉及自定义 kernel?"}
    use_default --> check_fusion
    check_fusion -->|是| enable_fuse["启用该 fusion"]
    check_fusion -->|否| skip_fuse["跳过，交给 Inductor"]
    enable_fuse --> done["配置完成，进入引擎初始化"]
    skip_fuse --> done

La clave de este flujo está en la ramacheck_user: lo que el usuario establece explícitamente siempre tiene prioridad📎 docs/design/optimization_levels.md:5-5. Esto evita problemas difíciles de diagnosticar como "el nivel de optimización sobrescribió silenciosamente mi flag de depuración".

Reflexiones de diseño y escollos

La trampa de producción más común de los niveles de optimización esun tiempo de arranque demasiado largo. La documentación recomienda explícitamente: cuando el tiempo de arranque sea demasiado largo, usar-O0o-O1 📎 docs/design/optimization_levels.md:87. Pero aquí hay un coste oculto——-O0sin cudagraph, el coste de lanzamiento por CPU de cada kernel queda al descubierto, y en escenarios de alta concurrencia el throughput puede caer varias veces.

Otra trampa esel error de compilación。-O2: elFULL_AND_PIECEWISEcudagraph de-O2tiene suposiciones más fuertes sobre la estructura del modelo; ciertos modelos personalizados fallan al compilar con-O1pero funcionan condebug_dump_path. La documentación recomienda usar📎 docs/design/optimization_levels.md:88para obtener más información de depuración-O0. La ruta de diagnóstico debería ser: primero usar-O1、-O2para confirmar que la funcionalidad es correcta, luego subir gradualmente a

---

II. Lista de trampas en producción: ruta de diagnóstico desde el síntoma hasta la causa raíz

Modelo intuitivo

La resolución de fallos en producción es como el triaje en urgencias: no puedes hacer un chequeo completo a todos los pacientes, primero debes reducir el alcance rápidamente según los síntomas (OOM, hang, crash) y luego profundizar de forma dirigida. La documentación de troubleshooting de vLLM es esencialmente un manual de triaje.

Clasificación de síntomas y herramientas de diagnóstico

La documentación divide los problemas comunes en varias categorías principales; las revisaremos en orden de dificultad de diagnóstico progresiva.

Primera categoría: descarga/carga del modelo colgada.El síntoma es una larga falta de respuesta tras el arranque. La causa raíz suele ser una red lenta o un sistema de archivos compartido lento.📎 docs/usage/troubleshooting.md:11-11El medio de diagnóstico es--load-format dummyomitir la carga de pesos, aislando si lo lento es la descarga o la carga.📎 docs/usage/troubleshooting.md:23-23Esta es una técnica típica de "aislamiento por bisección".

Segunda categoría: OOM de memoria de video.La documentación apunta directamente al documento de configuración conserving_memory.📎 docs/usage/troubleshooting.md:23Pero el OOM en producción a menudo no se debe a que el modelo sea demasiado grande, sino a la fragmentación del KV cache o a un número de solicitudes concurrentes superior al esperado.

Tercera categoría: cambios en la calidad de generación.Esta es una trampa fácil de pasar por alto. v0.8.0 cambió el origen de los parámetros de muestreo por defecto: de los valores neutros por defecto de vLLM a los del autor del modelo.generation_config.json 📎 docs/usage/troubleshooting.md:23-23En la mayoría de los casos esto mejora la calidad, pero la configuración de algunos modelos resulta peor.📎 docs/usage/troubleshooting.md:23-23El método de diagnóstico es revertir a--generation-config vllmcomparar📎 docs/usage/troubleshooting.md:23-23。

Cuarta categoría: cuelgue (hang).Esta es la categoría más difícil de diagnosticar. La documentación ofrece un conjunto de variables de entorno de depuración progresivas📎 docs/usage/troubleshooting.md:41-41：

• VLLM_LOGGING_LEVEL=DEBUG: activar logs detallados
• VLLM_LOG_STATS_INTERVAL=1.: salida de alta frecuencia del estado de la cola y de los aciertos de caché
• CUDA_LAUNCH_BLOCKING=1: localizar qué kernel de CUDA está fallando
• NCCL_DEBUG=TRACE: activar logs detallados de NCCL
• VLLM_TRACE_FUNCTION=1: registrar todas las llamadas a funciones, pero ralentiza más de 100 veces📎 docs/usage/troubleshooting.md:41

Aquí hay una disciplina operativa importante: tras depurar hay que desactivar estas variables de entorno, o abrir directamente un shell nuevo, de lo contrario la configuración de depuración residual seguirá ralentizando el sistema.📎 docs/usage/troubleshooting.md:11-11。

La trampa de los límites de proceso en la depuración con breakpoints

La arquitectura multiproceso de vLLM hace que los breakpoints convencionalespdbdejen de funcionar: si el breakpoint se ejecuta en un subproceso, lanzaráBdbQuit 📎 docs/usage/troubleshooting.md:45-54. Dos soluciones: usarforked-pdb 📎 docs/usage/troubleshooting.md:57-61, o establecerVLLM_ENABLE_V1_MULTIPROCESSING=0para mantener el planificador en el mismo proceso📎 docs/usage/troubleshooting.md:63-68。

Diagnóstico de comunicación distribuida

El despliegue distribuido tiene documentación de diagnóstico específica. La recomendación central es:establecer las variables de entorno al crear el clúster, porque las variables se propagan a todos los nodos; mientras que establecerlas en el shell solo afecta al nodo local.📎 docs/serving/distributed_troubleshooting.md:16-16。

Un problema frecuente esNo available node types can fulfill resource request, que aparece incluso cuando el clúster tiene suficientes GPU.📎 docs/serving/distributed_troubleshooting.md:16-16La causa raíz suele ser que el nodo tiene múltiples IP y vLLM eligió la incorrecta. La solución es usarVLLM_HOST_IPpara especificarla explícitamente, yray statuspara verificar📎 docs/serving/distributed_troubleshooting.md:16-16。

Script de diagnóstico para fallos de inicialización de NCCL

La documentación proporciona un script de diagnóstico completo que verifica la pila de comunicación capa por capa.📎 docs/usage/troubleshooting.md:89-150Su diseño es muy jerárquico:

flowchart TD
    start["运行诊断脚本"] --> nccl_test["测试 PyTorch NCCL<br/>dist.all_reduce"]
    nccl_test --> nccl_ok{"value == world_size?"}
    nccl_ok -->|否| hw_broken["硬件/驱动故障<br/>联系系统管理员"]
    nccl_ok -->|是| gloo_test["测试 PyTorch GLOO<br/>CPU 通信"]
    gloo_test --> gloo_ok{"value == world_size?"}
    gloo_ok -->|否| gloo_fail["GLOO 配置问题<br/>检查网络接口"]
    gloo_ok -->|是| pynccl_test["测试 vLLM PyNcclCommunicator"]
    pynccl_test --> pynccl_ok{"all_reduce 正确?"}
    pynccl_ok -->|否| pynccl_fail["vLLM NCCL 封装问题"]
    pynccl_ok -->|是| graph_test["测试 CUDA Graph 内 all_reduce"]
    graph_test --> graph_ok{"g.replay() 后正确?"}
    graph_ok -->|否| graph_fail["CUDA Graph 捕获问题<br/>检查 stream 语义"]
    graph_ok -->|是| success["sanity check 成功"]

Lo ingenioso de este script es que aísla capa por capa: primero verifica el PyTorch NCCL de más bajo nivel, luego el GLOO del lado de CPU, después el propio envoltorio PyNcclCommunicator de vLLM, y finalmente la comunicación dentro de CUDA Graph.📎 docs/usage/troubleshooting.md:90-146. Cada fallo de capa apunta a una causa raíz diferente.

Un detalle digno de mención en el script:pynccl.disabled = Falsees por compatibilidad hacia atrás con la versión 0.6.4 y anteriores.📎 docs/usage/troubleshooting.md:121-125. En 0.6.5+ está habilitado por defecto, pero se conserva esta línea para que los usuarios que lean la documentación más reciente no se confundan.

En las pruebas multinodo, la documentación usa deliberadamente--rdzv_backend=staticen lugar dec10d, porquec10den multinodo fallará por un fallo de resolución DNS.📎 docs/usage/troubleshooting.md:168-168. Esta es una configuración típica de "solo lo sabes tras haber pisado el pozo".

Reflexiones de diseño y trampas

Fallo de inicialización de NCCL（ncclCommInitRank(reporta unhandled system error) normalmente apunta a dos causas raíz: falta deIPC_LOCKcapability o/dev/shmno montado📎 docs/usage/troubleshooting.md:311-311. Ambas son trampas clásicas del despliegue en contenedores.

Desajuste de la cadena de herramientas CUDA PTX（the provided PTX was compiled with an unsupported toolchain) indica que el PTX dentro del wheel fue compilado con una versión superior del CUDA toolkit.📎 docs/usage/troubleshooting.md:325-327. La solución es habilitar la compatibilidad hacia adelante de CUDA: en Docker añadir-e VLLM_ENABLE_CUDA_COMPATIBILITY=1 📎 docs/usage/troubleshooting.md:325-327, en bare metal instalar elcuda-compatpaquete y establecerVLLM_CUDA_COMPATIBILITY_PATH 📎 docs/usage/troubleshooting.md:325-327。

Problema conocido de sobrecarga de memoria de NCCL：vLLM >= 0.4.3, <= 0.10.1.1estableceNCCL_CUMEM_ENABLE=0para evitar un bug de NCCL; los procesos externos que se conectan a vLLM también deben establecer esta variable, de lo contrario se colgarán o fallarán.📎 docs/usage/troubleshooting.md:375. Tras la corrección en NCCL 2.22.3, las versiones nuevas eliminaron esta sobrescritura para permitir optimizaciones de rendimiento.📎 docs/usage/troubleshooting.md:375. Este caso demuestra que:el contrato de variables de entorno entre procesos es una dependencia implícita de los sistemas distribuidos, y debe sincronizarse al actualizar.

---

III. Frontend en Rust: la filosofía de diseño de cero copias de la herramienta bench

Modelo intuitivo

Si el frontend en Python es una navaja suiza "completa pero pesada", la herramienta bench en Rust es un bisturí "hecho solo para pruebas de carga". Su objetivo de diseño no es la cobertura funcional, sino minimizar la sobrecarga del propio cliente bajo alta concurrencia, para que las cifras medidas reflejen de verdad el rendimiento del servidor.

Estructuras de datos y diseño de memoria

La estructura de datos central de la herramienta bench esRequestFuncInput 📎 rust/src/bench/src/backends/mod.rs:59-89. Utiliza ampliamenteArc<str>yArc<[u32]>en lugar deString/Vec, que es el núcleo del diseño de copia cero.

Veamos algunos campos clave:prompt: Arc<str> 📎 rust/src/bench/src/backends/mod.rs:50-52——múltiples solicitudes concurrentes pueden compartir la misma cadena de prompt, evitando clonar una copia por cada solicitud.prompt_token_ids: Option<Arc<[u32]>> 📎 rust/src/bench/src/backends/mod.rs:77——los token ID precalculados se envían directamente al servidor, omitiendo la tokenización del lado del servidor📎 rust/src/bench/src/backends/mod.rs:74-76。

Lo más ingenioso esmulti_modal_content: Option<Arc<[Arc<str>]>> 📎 rust/src/bench/src/backends/mod.rs:81. El comentario explica: el contenido multimodal se trata como fragmentos JSON preserializados, y el backend de chat los concatena directamente en el flujo de bytes del payload, evitando cualquier análisis o copia profunda de datos de imagen base64📎 rust/src/bench/src/backends/mod.rs:78-80. Esta es una estructura de doble capaArc: la capa externaArc<[...]>comparte todo el arreglo, la capa internaArc<str>comparte un solo fragmento.

chat_messages_json: Option<Arc<str>>tiene la prioridad más alta, se concatena directamente tal cual en el payload📎 rust/src/bench/src/backends/mod.rs:82-85。

Deserialización sin asignaciones

El análisis de respuestas en streaming SSE es otro punto clave de rendimiento. El comentario señala explícitamente: usar deserialización tipada para evitar construir un árbol completo deserde_json::Value, extrayendo solo los campos necesarios📎 rust/src/bench/src/backends/mod.rs:20-24。

CompletionChunkconservando solochoicesyusagedos campos📎 rust/src/bench/src/backends/mod.rs:20-24，ChatChunkDe manera similar📎 rust/src/bench/src/backends/mod.rs:33-37。#[serde(default)]hace que el campo faltantechoicespor defecto sea un arreglo vacío📎 rust/src/bench/src/backends/mod.rs:20-24, que es el caso común en respuestas en streaming.

Flujo de solicitudes orientado a escenarios

Cuando se emite una solicitud de prueba de carga, ¿cómo fluyen los datos? El siguiente diagrama de flujo de datos muestra la transformación desde la entrada hasta la salida:

flowchart LR
    input["RequestFuncInput<br/>Arc&lt;str&gt; prompt"] --> build["build_headers<br/>+ payload 拼接"]
    build --> send["reqwest::Client<br/>send_request"]
    send --> sse["SSE 流式响应<br/>字节流"]
    sse --> parse["CompletionChunk<br/>类型化反序列化"]
    parse --> output["RequestFuncOutput<br/>ttft/itl/tpot"]

BackendLa enumeración usa despacho estático para evitar el problema de los objetos trait async📎 rust/src/bench/src/backends/mod.rs:150-154。send_requestA través dematchse despacha a la implementación concreta📎 rust/src/bench/src/backends/mod.rs:158-168。get_backendSegúnBackendKinddevuelve el backend correspondiente📎 rust/src/bench/src/backends/mod.rs:172-181。

Un detalle:API_KEYusaOnceLockcaché, evitando hacer una llamada al sistema de variables de entorno por cada solicitud📎 rust/src/bench/src/backends/mod.rs:186-188。build_headersinserta secuencialmente Content-Type, Authorization, extra headers, request-id📎 rust/src/bench/src/backends/mod.rs:191-215。

Reflexiones de diseño y errores comunes

RequestFuncOutputEl diseño de campos dettft（time to first token）、itltambién merece atención:tpot（time per output token）📎 rust/src/bench/src/backends/mod.rs:93-105(arreglo de latencia entre tokens),

---

. Estos tres indicadores corresponden a diferentes dimensiones de rendimiento: TTFT refleja el prefill y la latencia de cola, ITL refleja la estabilidad del decode, TPOT refleja el rendimiento general. Si en la prueba de carga solo se mira la latencia promedio, se oculta la fluctuación de ITL.

Reflexión de diseño: la lógica subyacente de las compensaciones arquitectónicas

El procesamiento por lotes continuo permite que el lote se reorganice en cada paso, mejorando enormemente el rendimiento, pero a costa de una asignación y liberación extremadamente frecuentes de la caché KV. El mecanismo de tabla de bloques de PagedAttention está diseñado precisamente para hacer frente a esta asignación de alta frecuencia: los bloques de tamaño fijo eliminan la fragmentación externa, pero introducen la sobrecarga de direccionamiento indirecto de la tabla de bloques y la fragmentación interna (el último bloque puede no estar lleno). Esta es una compensación típica de "intercambiar tasa de fragmentación por una capa de indirección", la misma idea que la paginación de memoria virtual de los sistemas operativos.CUDA Graph vs formas dinámicas.PIECEWISECUDA Graph requiere formas estáticas, pero el tamaño de lote del procesamiento por lotes continuo cambia en cada paso. La solución de vLLM esFULL_AND_PIECEWISEy📎 docs/design/optimization_levels.md:50,72modo-O0——capturar como grafo la parte que puede volverse estática, manteniendo la parte dinámica en modo eager.-O2Desactivar completamente cudagraph es para depuración,-O1activarlo todo es para producción, el

intermedio es el compromiso.Despliegue separado vs sobrecarga de red.IPC_LOCK、/dev/shm）📎 docs/usage/troubleshooting.md:311-311KV Connector permite separar prefill y decode en diferentes instancias, pero la transferencia de caché KV entre instancias introduce latencia de red. Los requisitos de configuración de GPUDirect RDMA en la documentación (

indican que esta ruta tiene requisitos estrictos de infraestructura. La fluctuación de red provoca tiempos de espera en la transferencia de KV, lo que a su vez desencadena reintentos o degradación.Operabilidad vs rendimiento.VLLM_TRACE_FUNCTION=1Los niveles de optimización, las variables de entorno de depuración y los scripts de diagnóstico son costos pagados por la operabilidad.📎 docs/usage/troubleshooting.md:41puede ralentizar 100 veces

---

, pero es el último recurso para localizar problemas de cuelgue. Un motor maduro debe proporcionar estas herramientas "lentas pero que permiten ver con claridad".

Resumen de este capítulo

Este capítulo cierra el libro, reexaminando los mecanismos de los trece capítulos anteriores desde la perspectiva de producción.-O0Los niveles de optimización (-O3a📎 docs/design/optimization_levels.md:5-5) son un contrato explícito entre el tiempo de arranque y el rendimiento en ejecución; los flags del usuario siempre tienen prioridad sobre los valores predeterminados del nivelArc. La lista de errores comunes en producción cubre la ruta completa de diagnóstico desde la carga del modelo, OOM de memoria de video, cambios en la calidad de generación hasta fallos de comunicación distribuida, con una metodología central de "aislamiento por bisección" y "verificación capa por capa". La herramienta bench en Rust usa

Tres líneas centrales de compensación recorren todo el libro: procesamiento por lotes continuo frente a fragmentación de memoria de video, CUDA Graph frente a formas dinámicas, y despliegue desagregado frente a sobrecarga de red. Comprender estas tensiones es más importante que memorizar cualquier mecanismo individual, porque cada ajuste en un entorno de producción consiste, en esencia, en encontrar un punto de equilibrio entre estas tensiones.

Reflexiones y autoevaluación de este capítulo

---

para compartir texto, la segunda usa

para compartir la secuencia de tokens.