Forem: Mayank Ketkar

One Week in Ray: 21 Bugs Between Us and a Production ML Pipeline

Mayank Ketkar — Sun, 15 Feb 2026 17:48:09 +0000

The pipeline finished. All four checkpoints evaluated. Metrics written to cloud storage. Zero errors in the logs. And 68% of the results were empty strings.

Not errors. Not exceptions. Not null values with a helpful stack trace. Empty strings. The pipeline had done everything it was supposed to do -- read images, distribute them to GPU workers, run inference, collect results, write output -- and two-thirds of the results were blank. The metrics dashboard showed accuracy numbers. They were plausible. They were also computed on one-third of the actual data.

We had built a system that was confidently wrong.

This is the story of how we found and fixed 21 bugs in 9 days, turning a silently broken GPU inference pipeline into one that survives actor crashes, CUDA deadlocks, and stalled workers with zero data loss. The pipeline evaluates driving scenes through a fine-tuned vision-language model -- classifying road conditions and predicting vehicle waypoints. Safety-critical decisions. Full dataset coverage is not optional.

The tools were fine. Ray Data, vLLM, Kubernetes, experiment tracking -- none of them were broken. Every single one of the 21 bugs was in how they were integrated.

Act I: "The Pipeline Works!"

The False Sense of Security

The pipeline reported success because it was configured to tolerate failure. Somewhere in the early development, someone had set max_errored_blocks to a generous value. The reasoning was sound: if a handful of samples fail due to transient cloud storage errors, don't crash the entire 8,600-sample evaluation. Keep going. Report what you have.

The problem is that "keep going" and "silently discard 68% of your data" are the same instruction when the errors are systemic rather than transient.

The first clue was a boolean argument that wasn't being parsed correctly. The evaluation config had a flag -- something like --use_detailed_metrics -- that was being read as a string. In Python, bool("False") is True. Every string except the empty string is truthy. So the flag was always on, regardless of what you passed. This wasn't causing the 68% empty results directly, but it was a symptom of a larger disease: nobody was validating that configurations were doing what they claimed.

# The bug: bool("False") == True
use_detailed = bool(args.use_detailed_metrics)  # Always True

# The fix: explicit parsing
use_detailed = args.use_detailed_metrics.lower() in ("true", "1", "yes")

A TypeError guard in the inference wrapper was catching exceptions from malformed inputs and returning empty dictionaries. Silently. No log line. No counter. Just an empty result that looked indistinguishable from a successful inference that happened to produce no output.

# The silent swallower
try:
    result = engine.generate(inputs)
except TypeError:
    result = {}  # Looks fine. Is not fine.

This is the student-turning-in-blank-pages problem. The teacher collects 30 papers, sees 30 papers in the stack, and assumes everyone answered the questions. It isn't until you actually read them that you realize 20 are blank. Our pipeline was counting papers, not reading them.

The core insight: error tolerance without error visibility is just data loss with extra steps. The pipeline was designed to be resilient. It was resilient to the wrong things. It tolerated failures that should have been hard crashes, and it reported success when it should have reported catastrophe.

Act II: "The Machine Runs Out of Room"

Three Places Data Piles Up

In a Ray pipeline, data accumulates in three places: the driver node (the coordinator), the workers (the GPU actors), and the object store (shared memory between them). Understanding which one is drowning is the difference between a targeted fix and a week of thrashing.

Our pipeline had the architecture equivalent of an hourglass. All data flowed through a single point -- the driver -- before fanning out to workers. Here's what that looked like:

Cloud Storage (8,600 samples x 1-2 MB = ~17 GB)
         │
         ▼
  ┌──────────────┐
  │  Ray Driver   │ ← deserializes ALL images here
  │  (1 node)     │ ← holds ~17 GB in Python heap
  └──────┬───────┘
         │ ray.data.from_items() serializes to object store
    ┌────┴────┬────────┐
    ▼         ▼        ▼
 ┌──────┐ ┌──────┐ ┌──────┐
 │ GPU  │ │ GPU  │ │ GPU  │   6 actors
 │ W-1  │ │ W-2  │ │ W-3  │   (VLM inference)
 └──┬───┘ └──┬───┘ └──┬───┘
    │         │        │
    └────┬────┴────────┘
         ▼
  ┌──────────────┐
  │  Results      │ ← take_all() pulls everything back to driver
  └──────────────┘

The driver was doing all the heavy lifting: reading every sample from cloud storage, decoding every base64 image into memory, building Python dictionaries with ~2 MB of image data each, then handing the entire collection to ray.data.from_items(). That call serializes everything into the object store. For 8,600 samples at 2 MB each, that's roughly 17 GB of Python objects materialized on a single machine before any GPU touches a single pixel.

The Before: Driver-Side Materialization

def load_dataset(self, dataset_path, num_samples):
    mds_dataset = StreamingDataset(streams=[stream], shuffle=False)
    samples = []
    for idx in range(len(mds_dataset)):
        raw = mds_dataset[idx]
        prompt_blocks = self._parse_images_to_base64(raw)
        samples.append({
            "sample_id": idx,
            "prompt_blocks": prompt_blocks,  # ~2 MB per sample
            "target": raw["target"],
            "metadata": raw["metadata"],
        })
    # All 17 GB lives in driver memory at this point
    dataset = ray.data.from_items(samples)  # 8,600 x 2 MB = OOM
    return dataset

Under load, the driver ran out of memory. Python's garbage collector fought with Ray's object store for the same physical RAM. Workers received empty batches -- not because they failed, but because the objects they were trying to read had been evicted from the object store to make room for newer ones. No error. No exception. Just empty data.

The After: Worker-Side Shard Reading

def load_dataset(self, dataset_path, num_samples):
    return ray.data.read_datasource(
        MDSShardedDatasource(dataset_path, num_samples),
        parallelism=64,
    )

class MDSShardedDatasource(Datasource):
    """Each worker reads its own shard directly from cloud storage."""

    def get_read_tasks(self, parallelism):
        # Driver only reads index.json (~1 KB)
        shard_paths = self._compute_shard_assignments(parallelism)

        for shard in shard_paths:
            yield ReadTask(
                read_fn=lambda p=shard: self._read_mds_shard(p),
                metadata=BlockMetadata(
                    num_rows=shard.num_samples,
                    size_bytes=shard.estimated_size,
                ),
            )

The driver now reads a single index file -- roughly 1 KB -- and tells each of 64 workers which shard to pull from cloud storage. Each worker deserializes only its own portion. Memory usage on the driver drops from 17 GB to effectively nothing.

The same antipattern existed in reverse for output. results.materialize() followed by take_all() pulled every result back to the driver before writing to disk. We replaced it with write_parquet() which streams results directly from workers to cloud storage. The driver never touches the result data.

The Object Store Spill Cascade

Even after fixing the driver bottleneck, we hit a second memory wall when running multiple evaluation pipelines concurrently. Four pipelines, each preprocessing 8,600 samples with PIL Images and PyTorch tensors (not Arrow-safe, so Ray falls back to pickle serialization -- full copy per consumer):

4 pipelines x 8,600 samples x ~25 MB preprocessed = ~860 GB object store pressure

Our node had 154 GB of object store (30% of 512 GB RAM). The cascade:

Object store hits 100%
Ray spills to /tmp (10-50x slower)
/tmp fills up
Ray evicts unconsumed objects
max_errored_blocks absorbs the errors
Pipeline "succeeds" with missing data

The fix was architectural: preprocess once, fan out to N inference pools via ray.put(). One copy in the object store, N readers. Object store usage dropped from hundreds of GB to single-digit GB.

Act III: "Everything Fails All the Time" -- The Eight-Fix Monday

This was the centerpiece. Eight fixes shipped in a single day, each addressing a different failure mode. Together, they form a defense-in-depth pattern -- six layers of protection between a failure and data loss.

Think of it like a substitute teacher managing a classroom of 32 students (GPU workers). Any individual student might fall asleep (stall), have a meltdown (crash), or turn in a blank worksheet (empty result). The substitute can't prevent any of these. But they can notice them, respond to them, and make sure the class still finishes the assignment.

Layer 1: Error Containment -- The _error Column

The first layer is the simplest. When any stage encounters an error, it tags the row instead of crashing.

def preprocess(self, row):
    try:
        return self._process(row)
    except Exception as e:
        row["_error"] = f"preprocess: {e}"
        return row  # Row survives with error tag

def postprocess(self, row):
    if row.get("_error") or row.get("generated_text") is None:
        row["_error"] = row.get("_error", "inference: empty response")
        return row
    return self._normal_postprocess(row)

Once _error is set, every downstream stage checks for it and passes the row through untouched. The row shows up in the final output with a clear failure reason. No data loss. No silent swallowing. You can grep the output for _error and see exactly which samples failed and why.

Layer 2: Per-Sample Fallback

Batch inference is fast -- 64 samples at once through the GPU. But one corrupted image in a batch of 64 can crash the entire batch. Without isolation, that one bad apple kills 63 good results.

def __call__(self, batch):
    try:
        return self._generate_batch(batch)              # Fast path: full batch
    except Exception as batch_err:
        logger.warning("Batch of %d failed (%s), retrying per-sample",
                       len(batch["prompt"]), batch_err)
        results = []
        for i in range(len(batch["prompt"])):
            try:
                result = self._generate_single(batch, i)
                results.append(result)
            except Exception as sample_err:
                results.append({
                    "generated_text": None,
                    "_error": f"inference: {sample_err}",
                })
        return self._merge_results(results)

The happy path tries the full batch. On failure, it falls back to per-sample inference. Only the broken sample gets an error tag. The other 63 flow through normally. The cost is one extra inference call per bad sample. The benefit is 63 results that would otherwise be lost.

Layer 3: Stall Watchdog

This is the one that saved us at 98% progress. A vLLM actor hit a multimodal cache bug and froze -- not crashed, not errored, frozen. A futex deadlock in the CUDA runtime. No error signal. No exit code. Just silence.

class VLMInferenceEngine:
    def __init__(self, config):
        self._last_progress = time.time()
        self._watchdog = threading.Thread(
            target=self._watchdog_loop, daemon=True
        )
        self._watchdog.start()

    def _watchdog_loop(self):
        """Dead man's switch: if no batch completes in 120s, force-kill."""
        while True:
            time.sleep(30)  # Check every 30 seconds
            stall_seconds = time.time() - self._last_progress
            if stall_seconds > 120:
                logger.error(
                    "WATCHDOG: No progress for %ds. Force-killing actor.",
                    stall_seconds
                )
                os._exit(1)  # Not sys.exit(). Not an exception. Hard kill.

    def __call__(self, batch):
        outputs = self.llm.generate(inputs, sampling_params)
        self._last_progress = time.time()  # Pet the watchdog
        return self._format_outputs(outputs, batch)

Why os._exit(1) instead of sys.exit() or raising an exception? Because when you're in a CUDA deadlock, polite shutdown is not an option. sys.exit() raises SystemExit, which Python's exit handlers try to process. Those handlers can themselves deadlock on the same futex. os._exit() calls the C runtime's _exit() directly -- no handlers, no cleanup, no chance to hang. The process dies. Ray detects the death. Ray restarts the actor on a fresh node.

A train's dead man's switch requires the engineer to press a button every 30 seconds. If the engineer is incapacitated, the switch is released and the train brakes. Our watchdog requires the actor to complete a batch every 120 seconds. If it doesn't, the actor is killed and restarted. The system doesn't diagnose why the actor stopped responding. It just knows that silence is dangerous.

Layer 4: Engine Retry with Exponential Backoff

Initializing a GPU inference engine can fail for transient reasons: stale CUDA state from a previously crashed process, fragmented GPU memory, a race condition during model loading. Without retry logic, a single init failure means a permanently broken actor.

def _initialize_engine(self):
    """Retry engine init with exponential backoff and GPU cleanup."""
    for attempt in range(3):
        try:
            self.llm = LLM(
                model=self.model_path,
                gpu_memory_utilization=0.85,  # 10% headroom for restarts
            )
            self._initialized = True
            logger.info("Engine initialized on attempt %d/3", attempt + 1)
            return
        except Exception as e:
            logger.warning("Init attempt %d/3 failed: %s", attempt + 1, e)
            # Aggressive GPU cleanup between attempts
            if torch.cuda.is_available():
                torch.cuda.empty_cache()
                torch.cuda.synchronize()
            gc.collect()
            time.sleep(10 * (2 ** attempt))  # 10s, 20s, 40s

    # All retries exhausted: mark permanently dead
    self._init_error = True
    logger.error("Engine init failed after 3 attempts. Marking actor as dead.")

The gpu_memory_utilization=0.85 is deliberate. We dropped it from 0.95 to leave 10% headroom for cold starts after a crash. That 10% sounds wasteful until you realize it's the difference between a restart succeeding and entering a death spiral.

Layer 5: Dead Engine Fast-Path (Circuit Breaker)

Once an engine is permanently dead, every batch that arrives should be immediately returned as an error, not retried.

def __call__(self, batch):
    if self._init_error:
        # Don't waste time. Don't waste GPU. Fast-return errors.
        return {
            col: batch[col] for col in batch
        } | {
            "generated_text": [None] * len(batch["prompt"]),
            "_error": ["engine_init_failed"] * len(batch["prompt"]),
        }

    if not self._initialized:
        self._initialize_engine()
        if self._init_error:
            return self.__call__(batch)  # Will hit fast-path above

    return self._run_inference(batch)

This is a circuit breaker pattern. Once the circuit opens (engine dead), all requests fast-fail. Ray detects that the actor is producing only errors and stops routing new work to it. Without this, the dead actor would keep accepting batches, running through the retry logic on every single one, and returning empty results after a 70-second delay each time.

Layer 6: Pipeline-Level Fault Tolerance

All five layers above protect individual actors and samples. The sixth layer protects the pipeline itself.

ctx = ray.data.DataContext.get_current()
ctx.max_errored_blocks = 10         # Tolerate up to 10 failed blocks
ctx.actor_task_retry_on_errors = True
ctx.actor_init_retry_on_errors = True

processed = dataset.map_batches(
    VLMInferenceEngine,
    batch_size=64,
    num_gpus=1,
    compute=ActorPoolStrategy(size=6, max_tasks_in_flight_per_actor=4),
    max_restarts=3,          # Actor dies 3 times -> permanent failure
    max_task_retries=3,      # Block retried 3 times -> marked errored
)

max_restarts=3 is actor-level: a GPU that crashes three times is considered permanently unhealthy. max_task_retries=3 is block-level: a batch of data that fails on three different actors is considered poison. Together, they create a two-tier retry policy. A transiently sick GPU gets three chances to recover. A poison batch gets three tries across different actors. The combination of these six layers meant that when our actor froze at 98% progress, the watchdog killed it, Ray restarted it on a fresh node, the remaining 178 samples were processed, and the final output contained every single one of the 8,600 input samples. Zero dropped.

Act IV: "The Fleet"

With the pipeline itself hardened, we turned to the operational layer -- running dozens of evaluations across multiple checkpoints on a shared cluster. This introduced a new category of problems: not "does the pipeline work?" but "does the fleet of pipelines play nice together?"

The Metrics Deadlock

Our experiment tracking SDK had a method -- experiment.end() -- that blocks until its internal upload queue drains. Under normal conditions, this takes a few seconds. Under throttled conditions (the SDK rate-limits itself when you log too many metrics too fast), the queue never drains. The call blocks forever.

This was happening inside Ray actors. A deadlocked actor holds its GPU allocation, accepts no new work, and never exits. We were losing GPUs to a metrics library.

# Before: blocks forever under throttling
def close(self):
    self.experiment.end()  # Can hang for hours

# After: bounded timeout in a daemon thread
def close(self, timeout=60):
    if not self.experiment:
        return
    self._flush_metrics_queue()
    thread = threading.Thread(target=self.experiment.end, daemon=True)
    thread.start()
    thread.join(timeout=timeout)
    if thread.is_alive():
        logger.warning(
            "Experiment end() did not complete in %ds. "
            "Server will auto-close the orphaned experiment.", timeout
        )

The daemon thread ensures that if end() hangs, the actor can still exit cleanly. The tracking server auto-closes orphaned experiments after a timeout, so no data is permanently lost. We also batched all log_metrics() calls into a single call per batch (was 10 separate network round-trips), disabled auto-logging of environment info and git state, and set explicit timeouts of 60 seconds instead of the default 3,600.

Job Queuing: The Bouncer Pattern

When you submit 80 evaluation jobs to a shared Ray cluster, you need to control how many run concurrently. Without resource gating, all 80 try to start simultaneously, each requesting 6 GPUs. The cluster has 24 GPUs. Four jobs fit. The other 76 sit in PENDING state, holding driver-node resources (CPU, memory for the coordinator process) while waiting for GPUs that won't be free for hours.

# entrypoint_resources acts as a bouncer:
# each job "holds" 6 GPUs worth of admission tickets
ray.job_submission.JobSubmissionClient(
    address=cluster_address
).submit_job(
    entrypoint="python run_eval.py --checkpoint $CKPT",
    entrypoint_resources={"GPU": 6},  # Job won't start until 6 GPUs free
    runtime_env={"env_vars": {"CHECKPOINT": ckpt_path}},
)

entrypoint_resources is the bouncer at the door. Each job declares upfront how many GPUs it needs. Ray's scheduler won't start the job until those GPUs are available. Jobs queue cleanly instead of stampeding. Four jobs run concurrently on 24 GPUs. As each finishes, the next one starts. No manual orchestration.

The Version Mismatch

We lost half a day to a Ray version mismatch. The cluster was running Ray 2.38. Our container image had Ray 2.35. The job submitted successfully, the actors started, and then inference produced subtly wrong results. No crash. No version check error. Just a different code path in the data serialization layer that handled edge cases differently.

The fix was pinning the Ray version in the container image to match the cluster exactly. But the real lesson was adding a version check to the job entrypoint:

import ray
expected = "2.38.0"
actual = ray.__version__
if actual != expected:
    raise RuntimeError(
        f"Ray version mismatch: expected {expected}, got {actual}. "
        f"Update your container image."
    )

Force-Stopping Stuck Jobs

Even with all the resilience logic, some jobs get stuck in PENDING state and never start -- usually because a previous job didn't release its resources cleanly. The Ray dashboard shows the job as PENDING, but there's no built-in "force stop" button that works reliably for jobs that haven't started yet.

from ray.job_submission import JobSubmissionClient

client = JobSubmissionClient(address=cluster_address)
for job_info in client.list_jobs():
    if job_info.status in ("PENDING", "RUNNING"):
        elapsed = time.time() - job_info.start_time
        if elapsed > 7200:  # 2 hours without progress
            logger.warning("Force-stopping stuck job %s", job_info.job_id)
            client.stop_job(job_info.job_id)

We wrapped this in a cron job that runs every 30 minutes. Aggressive, but the alternative was waking up to find 6 GPUs held hostage by a zombie job.

Act V: "The Result"

Here is where we ended up, 9 days after finding the 68% empty results:

Metric	Day 1 (broken)	Day 9 (fixed)
Results present	2,752 / 8,600 (32%)	8,600 / 8,600 (100%)
Results empty	5,848 (68%)	0
Crash recovery	None (manual restart)	Automatic (watchdog + Ray retry)
Actor stall detection	None	120s watchdog, `os._exit(1)`
Per-sample isolation	None (one bad sample kills batch)	Batch fallback to per-sample
Concurrent checkpoints	1 (manual)	4 concurrent, 80+ queued
Job control	Cannot kill stuck jobs	REST API + cron reaper
Metrics deadlocks	Frequent (hours-long hangs)	Zero (60s timeout)

We ran 80+ checkpoint evaluations unattended over the following week. Every single one produced a complete result set. Two actors crashed during that period -- one to a CUDA OOM, one to the same multimodal cache bug. Both were detected by the watchdog within 150 seconds, restarted by Ray, and the affected batches were reprocessed. Zero data loss in both cases.

What We Learned

The 21 fixes fell into a clear pattern. They weren't individually difficult. Most were 5-20 lines of code. The hard part was finding them, because the pipeline's error-tolerance features were actively hiding the problems.

Three principles emerged:

Reconcile relentlessly. Input count must equal output count. If they don't match, the run failed, regardless of what the logs say. We added a hard assertion at the end of every pipeline run: assert output_count == input_count.

Fail loud, recover quiet. Every failure should produce a visible signal -- a log line, an error tag, a metric increment. Recovery should be automatic and silent. The worst failure mode is one that is both invisible and unrecoverable.

Defense in depth. No single layer prevents all failures. The watchdog doesn't prevent CUDA deadlocks -- it detects and recovers from them. The per-sample fallback doesn't prevent corrupted images -- it isolates them. The circuit breaker doesn't prevent init failures -- it stops wasting time on dead actors. Each layer handles the failures that slip past the layer above it.

Production is not a feature you ship. It is a place your code lives. The weather there is harsh and unpredictable. You don't make code "production-ready" by adding a feature flag. You make it production-ready by assuming everything will fail, and building the systems to notice, contain, and recover -- automatically, silently, completely.

Our pipeline doesn't crash less than it used to. It crashes exactly as often. It just recovers now.

All code examples are simplified for clarity. The actual implementations include additional error handling, logging, and configuration.

The Ghost in the Batch: How vLLM Silently Switches Algorithms

Mayank Ketkar — Sun, 15 Feb 2026 17:47:55 +0000

You run Qwen3-VL on a single prompt. You record the output logprobs to full precision. Then you run the exact same prompt again, batched with 15 others. Same model, same weights, same GPU, same code. The logprobs are different.

Not catastrophically -- the top token usually agrees -- but the numbers have shifted at the seventh decimal place, and in an autoregressive loop, that hairline fracture propagates. By output position 8, 29% of your tokens have diverged.

You are not going crazy. vLLM silently changed the algorithm.

The Crime Scene

Setup: Qwen3-VL 2B on an NVIDIA H200. Identical prompts at BS=1 (one at a time) vs BS=16 (sixteen at once). VLLM_BATCH_INVARIANT=1 is enabled -- all GEMMs are deterministic via persistent Triton kernels. Yet:

Metric	BS=1 vs BS=16
Bitwise logprob match	6.1% (30/490)
Top-1 token match	78.6%
Semantic agreement	100%

The profiler gives the first clue: flash attention takes 5.15x longer per call in BS=16 (6.17ms vs 1.20ms). Same kernel name. Same call count (392). If it were the same algorithm processing more data, you'd expect it to scale with tokens -- not blow up 5x per call.

This is not a scaling problem. This is a different recipe.

What Attention Does (60 Seconds)

For each new token, attention looks back at every previous token, computes a relevance score for each, normalizes them (softmax), and takes a weighted average:

output = softmax(Q @ K^T / sqrt(d)) @ V

The critical observation: softmax involves summing over all previous tokens. If you have 640 tokens, that is 640 numbers being added. The order of that summation will matter shortly.

The Optimization That Changes Everything

Imagine 16 requests sharing the same 400-token system prompt. Without optimization, each scans those 400 tokens independently -- 6,400 redundant KV reads (80% waste).

vLLM's cascade attention splits the work:

Prefix (done ONCE): Flash attention over the shared 400 tokens for all 16 queries. Produces partial output + Log-Sum-Exp (LSE) statistic.
Suffix (per-request): Flash attention over each request's unique tokens (~100 each). Produces partial output + LSE.
Merge: Combine using LSE-weighted rebalancing.

# flash_attn.py:1040 -- Cascade Implementation
def cascade_attention(output, query, key_cache, value_cache, ...):
    # Step 1: Process shared prefix ONCE
    prefix_output, prefix_lse = flash_attn_varlen_func(
        q=query, k=key_cache, v=value_cache,
        max_seqlen_k=common_prefix_len,
        block_table=block_table[:1],
        causal=False,
        return_softmax_lse=True,
    )

    # Step 2: Process each request's unique suffix
    suffix_output, suffix_lse = flash_attn_varlen_func(
        q=query, k=key_cache, v=value_cache,
        max_seqlen_k=max_kv_len - common_prefix_len,
        block_table=block_table[:, num_common_kv_blocks:],
        causal=True,
        return_softmax_lse=True,
    )

    # Step 3: Merge -- THIS is where determinism breaks
    merge_attn_states(output, prefix_output, prefix_lse,
                      suffix_output, suffix_lse)

Total KV reads: 400 + 16x100 = 2,000 (vs 8,000). A 4x bandwidth reduction.

Mathematically, this produces the identical result. But we don't live in the world of exact arithmetic.

Why 1 + 2 + 3 Does Not Equal 3 + 2 + 1

This is the section that explains everything.

IEEE 754 floating-point addition is not associative:

>>> a, b, c = 1e-7, 1.0, -1.0
>>> (a + b) + c    # 1.1920928955078125e-07
>>> a + (b + c)    # 1e-07

Same inputs. Same operations. Different answers. This is the IEEE 754 spec -- finite precision means rounding depends on order.

Connect this to attention:

Single-pass (BS=1): softmax([t1, t2, ..., t640]) @ V -- one summation, one rounding chain
Cascade (BS=16): merge(softmax([t1,...,t512]) @ V_prefix, softmax([t513,...,t640]) @ V_suffix) -- two summations + LSE merge

The merge math:

# Simplified merge_attn_states logic
max_lse = max(prefix_lse, suffix_lse)  # numerical stability

prefix_weight = exp(prefix_lse - max_lse)
suffix_weight = exp(suffix_lse - max_lse)

output = (prefix_weight * prefix_output
        + suffix_weight * suffix_output) \
        / (prefix_weight + suffix_weight)

# Math: IDENTICAL to single-pass
# IEEE 754: DIFFERENT by ~1e-7 (FP32) or ~1e-3 (FP16)

~1e-7 per element sounds negligible. But autoregressive generation feeds each output back as input through ~28 transformer layers. That 1e-7 compounds:

Position 5: 17% token divergence
Position 8: 29% token divergence

The Three Gates

When does vLLM activate cascade? Silently, based on three runtime conditions:

# flash_attn.py:962
def use_cascade_attention(common_prefix_len, query_lens, ...):
    # Gate 1: Is shared prefix long enough?
    if common_prefix_len < 256:
        return False  # "Not worth the overhead"

    # Gate 2: Are there enough requests?
    num_reqs = len(query_lens)
    if num_reqs < 8:
        return False  # "Too few to benefit"

    # Gate 3: Performance heuristic
    cascade_time = cascade_waves * num_prefix_tiles
    flash_decoding_time = cdiv(flash_decoding_ctas, num_sms)
    return cascade_time < flash_decoding_time

BS=1 always fails Gate 2. It uses single-pass attention.

BS=16 with a system prompt passes all three. It uses cascade attention.

Your BS=1 benchmark is literally running a different algorithm than your BS=16 production system.

The Smoking Gun: 30 Matching Samples

In our 490-pair experiment, exactly 30 matched bitwise -- always the last 3 per batch. As a batch of 50 processes, requests finish and leave. When fewer than 8 remain, Gate 2 closes. The last requests revert to single-pass and match BS=1 perfectly.

Batch Position	Active Requests	Cascade?	Matches BS=1?
1-47	50 down to 8	Yes	No
48	7	No	Yes
49	6	No	Yes
50	5	No	Yes

The Fork in the Code

The branch at flash_attn.py:673:

if not attn_metadata.use_cascade:
    # SINGLE PASS: one call over all KV tokens
    flash_attn_varlen_func(
        q=query[:num_actual_tokens],
        k=key_cache,
        v=value_cache,
        ...
        num_splits=attn_metadata.max_num_splits,
    )
    return output

# CASCADE: two calls + merge
cascade_attention(
    output[:num_actual_tokens],
    query[:num_actual_tokens],
    key_cache, value_cache,
    ...
    common_prefix_len=attn_metadata.common_prefix_len,
)

Same function signature. Completely different execution.

vLLM already knows these conflict. When VLLM_BATCH_INVARIANT=1 is set, it auto-disables cascade:

# vllm/config/vllm.py:994
if vllm_is_batch_invariant() and not self.model_config.disable_cascade_attn:
    self.model_config.disable_cascade_attn = True
    logger.warning_once(
        "Disabling cascade attention when VLLM_BATCH_INVARIANT is enabled."
    )

The Fix: Three Options

Option 1: Surgical (Recommended)

llm = LLM(model="your-model", disable_cascade_attn=True)

Forces single-pass for all batch sizes. 5-15% throughput loss for shared-prefix workloads.

Option 2: Remove Prefix Detection

llm = LLM(model="your-model", enable_prefix_caching=False)

Gate 1 never opens. Also disables other prefix caching benefits.

Option 3: Full Determinism

export VLLM_BATCH_INVARIANT=1

Replaces ALL cuBLAS GEMMs + auto-disables cascade. ~2.4x performance cost.

Option	Cascade	GEMM Determinism	Cost
Default	Enabled	Non-deterministic	Baseline
`disable_cascade_attn=True`	Disabled	Non-deterministic	~5-15%
`enable_prefix_caching=False`	Disabled	Non-deterministic	Medium
`VLLM_BATCH_INVARIANT=1`	Auto-disabled	Deterministic	~2.4x

The Broader Lesson

Cascade attention is not a bug. It is a well-engineered bandwidth optimization. The issue is the silence.

This pattern recurs across GPU inference:

Flash Decoding splits attention across thread blocks -- same associativity issue
cuBLAS GEMM selects different tile sizes by matrix shape -- same op, different rounding
torch.compile fuses differently between eager/compiled -- same model, different graph

Every time a framework says "mathematically equivalent," ask: equivalent in the reals, or in IEEE 754?

The ghost in the batch is not malicious. It is an optimization doing its job. But now you know it is there, you know when it activates, and you know how to control it.

Key Takeaways

BS=1 and BS>=8 run different attention algorithms in vLLM. Single-pass vs cascade, by design.
Cascade saves 4x memory bandwidth by processing shared prefixes once. The merge step introduces FP divergence.
Three silent gates control activation: prefix >= 256 tokens, num_reqs >= 8, perf heuristic.
One flag fixes it: disable_cascade_attn=True or VLLM_BATCH_INVARIANT=1.
"Mathematically equivalent" != "numerically identical." This applies across all GPU ML.

Key files: flash_attn.py:673 (fork), flash_attn.py:962 (gates), flash_attn.py:1040 (cascade), merge_attn_states.py (merge), vllm/config/vllm.py:994 (auto-disable)

How to Read GPU Profiling Logs: A Ground-Up Guide

Mayank Ketkar — Sun, 15 Feb 2026 17:47:38 +0000

You ran nsys profile, got a 2GB .nsys-rep file, exported it to SQLite, and found yourself staring at 88 tables with names like CUPTI_ACTIVITY_KIND_KERNEL and ENUM_WDDM_PAGING_QUEUE_TYPE. The kernel names are integer IDs. The timestamps are in nanoseconds. Nothing is human-readable. You closed the file and went back to guessing.

This post exists so you never have to guess again.

I'm going to teach you to read any nsys trace in under 10 minutes — using four tables, one SQL join pattern, and four queries. Then we'll use those tools to solve a real mystery: why does a model give different results at batch size 1 vs batch size 16, even though both traces show exactly 8,955 kernel launches?

What nsys actually records

Imagine you're standing in a factory watching an assembly line. You have a stopwatch, and your job is to write down: when each machine started, when it stopped, which machine it was, and what it was building.

That's what NVIDIA Nsight Systems does for your GPU. It records every kernel launch, every memory copy, every synchronization event — with nanosecond timestamps.

The output is two files:

profile.nsys-rep     ← visual report (open in the Nsight GUI)
profile.sqlite       ← raw data in a database (query with SQL)

The .sqlite file is the gold mine. Everything in the .nsys-rep is derived from it.

88 tables, but only 4 matter

Open any nsys SQLite export and you'll find 88 tables. Most are enum lookup tables or metadata. You need these four:

Table	What it records
`CUPTI_ACTIVITY_KIND_KERNEL`	Every GPU kernel execution: start, end, grid, block, name
`CUPTI_ACTIVITY_KIND_MEMCPY`	Every memory transfer: CPU→GPU, GPU→CPU, GPU→GPU
`CUPTI_ACTIVITY_KIND_MEMSET`	Every memory fill (zeroing buffers)
`NVTX_EVENTS`	Human-readable markers programmers add to their code

Plus one helper table: StringIds — the Rosetta Stone that maps integer IDs to actual names.

Here's how to discover them:

import sqlite3
# nsys export --type=sqlite profile.nsys-rep  -> produces profile.sqlite
conn = sqlite3.connect('profile.sqlite')
cursor = conn.cursor()
cursor.execute("SELECT name FROM sqlite_master WHERE type='table'")
tables = [r[0] for r in cursor.fetchall()]
print(f"Total tables: {len(tables)}")  # 88 tables!
# But only 4 matter:
#   CUPTI_ACTIVITY_KIND_KERNEL  ← GPU kernel executions
#   CUPTI_ACTIVITY_KIND_MEMCPY  ← memory transfers
#   CUPTI_ACTIVITY_KIND_MEMSET  ← memory fills
#   NVTX_EVENTS                 ← human-readable markers
#   + StringIds                 ← name lookup table

The #1 gotcha: names are integers

This is the single most confusing thing when you first look at nsys data. Kernel names are not stored as strings. They're stored as integer foreign keys into the StringIds table.

When you query the kernel table, you'll see:

demangledName = 58
shortName     = 59

These are pointers. To get the actual name, you JOIN:

JOIN StringIds s ON k.shortName = s.id

And suddenly 59 becomes vectorized_elementwise_kernel.

Every query you write will include this JOIN. It becomes muscle memory after your second trace.

Decoding kernel names: the Rosetta Stone

Once you run that JOIN, you'll see names like nvjet_tst_192x192_64x4_2x1_v_bz_coopB_TNN. This isn't gibberish — every character means something:

nvjet_tst _ 192x192 _ 64x4 _ 2x1 _ v _ bz _ coopB _ TNN
   │          │        │      │     │    │     │       │
   │          │        │      │     │    │     │       └─ transpose: T=yes, N=no
   │          │        │      │     │    │     │          TNN = A^T × B → C
   │          │        │      │     │    │     │
   │          │        │      │     │    │     └─ cooperative mode
   │          │        │      │     │    │        coopA/coopB = how SMs share work
   │          │        │      │     │    │
   │          │        │      │     │    └─ block-zero init strategy
   │          │        │      │     │
   │          │        │      │     └─ layout: v=vertical, h=horizontal
   │          │        │      │        (how tiles map to SMs)
   │          │        │      │
   │          │        │      └─ warp tiling: 2 warps in M, 1 in N
   │          │        │
   │          │        └─ block tile: 64×4 threads per block
   │          │
   │          └─ output tile: 192×192 chunk per SM
   │
   └─ nvjet_tst = NVIDIA JIT persistent kernel (deterministic!)

Think of it like a car's VIN number. Once you know the format, you can read any GPU kernel at a glance.

Besides nvjet_tst_*, you'll encounter:

device_kernel — output of torch.compile. Opaque, but often 70%+ of GPU time.
vectorized_elementwise_kernel — PyTorch's generic ops (add, multiply, cast).
rms_norm_kernel / act_and_mul_kernel — normalization and activation.
*_splitK_* — Split-K GEMM with atomic reduction. Potential non-determinism source.

Grid and block: mapping kernels to hardware

The kernel table has gridX/Y/Z and blockX/Y/Z columns. These map to physical GPU hardware.

An H200 has 132 Streaming Multiprocessors (SMs) — 132 independent assembly lines. Each can process one thread block at a time.

  GPU with 132 SMs:
  ┌─────┬─────┬─────┬─────┬─── ── ──┬─────┐
  │SM 0 │SM 1 │SM 2 │SM 3 │  ...    │SM131│
  │blk 0│blk 1│blk 2│blk 3│         │blk  │
  │128  │128  │128  │128  │         │131  │
  │thrds│thrds│thrds│thrds│         │thrds│
  └─────┴─────┴─────┴─────┴─── ── ──┴─────┘

  grid=132 × block=128 = 16,896 total threads

gridX=1: One SM active, 131 idle. Tiny work.
gridX=132: Every SM busy. What persistent GEMM targets.
gridX=64,000: Blocks queue in waves. GPU stays saturated.

Four SQL queries that answer every question

Every GPU investigation follows the same four-step pattern: The Census, The Lineup, The Stakeout, and The Timeline.

Level 1: The Census — "What's slow?"

-- Where is the GPU spending its time?
SELECT
    s.value AS kernel_name,
    COUNT(*) AS call_count,
    ROUND(SUM(k.end - k.start) / 1e6, 2) AS total_ms,
    ROUND(AVG(k.end - k.start) / 1e3, 1) AS avg_us
FROM CUPTI_ACTIVITY_KIND_KERNEL k
JOIN StringIds s ON k.shortName = s.id
GROUP BY s.value
ORDER BY total_ms DESC
LIMIT 10;

On our H200 running vLLM inference:

Kernel	Calls	Total (ms)	%
device_kernel (torch.compile)	748	877.38	70.2%
vectorized_elementwise_kernel	883	60.98	4.9%
nvjet_tst_192x192_64x4_2x1...	28	51.94	4.2%

70% of GPU time in one kernel type — the compiled vision encoder.

Level 2: The Lineup — "What category?"

SELECT
  CASE
    WHEN LOWER(s.value) LIKE '%nvjet%' THEN 'Persistent GEMM'
    WHEN LOWER(s.value) LIKE '%cublas%' THEN 'cuBLAS (DANGER)'
    WHEN LOWER(s.value) LIKE '%device_kernel%' THEN 'torch.compile'
    WHEN LOWER(s.value) LIKE '%flash%' THEN 'Flash Attention'
    WHEN LOWER(s.value) LIKE '%norm%' THEN 'Normalization'
    ELSE 'Other'
  END AS category,
  COUNT(*) AS calls,
  ROUND(SUM(k.end-k.start)/1e6, 2) AS total_ms
FROM CUPTI_ACTIVITY_KIND_KERNEL k
JOIN StringIds s ON k.shortName = s.id
GROUP BY category ORDER BY total_ms DESC;

Results: 17.3% Persistent GEMM, 7.7% Elementwise, 2% Normalization, and... 0.1% Flash Attention. Just 1.2ms. Barely a blip.

Remember that 0.1%. It becomes 5.15x the smoking gun.

Level 3: The Stakeout — "What changed?"

-- Drill into every launch of a specific kernel
SELECT
    k.start,
    ROUND((k.end - k.start) / 1e3, 1) AS dur_us,
    k.gridX, k.gridY, k.gridZ,
    k.blockX, k.registersPerThread,
    k.dynamicSharedMemory
FROM CUPTI_ACTIVITY_KIND_KERNEL k
JOIN StringIds s ON k.shortName = s.id
WHERE s.value = 'reshape_and_cache_flash_kernel'
ORDER BY k.start
LIMIT 20;

Level 4: The Timeline — "What happened when?"

-- Unified GPU timeline: kernels + memory transfers
SELECT 'KERNEL' AS type,
       k.start,
       ROUND((k.end - k.start) / 1e3, 1) AS dur_us,
       s.value AS name
FROM CUPTI_ACTIVITY_KIND_KERNEL k
JOIN StringIds s ON k.shortName = s.id
UNION ALL
SELECT 'MEMCPY',
       m.start,
       ROUND((m.end - m.start) / 1e3, 1),
       'copyKind=' || m.copyKind
FROM CUPTI_ACTIVITY_KIND_MEMCPY m
ORDER BY start
LIMIT 50;

Case study: the cascade attention smoking gun

Here's the mystery: We're running vLLM inference with batch-invariant mode, which guarantees bitwise-identical results regardless of batch size. BS=1 works perfectly. BS=16 gives different results. Both traces: exactly 8,955 kernel launches. Where's the bug?

Step 1: The Census — nothing obvious

Metric	BS=1	BS=16	Ratio
Total kernels	8,955	8,955	1.00x
Total GPU time	1.25s	1.37s	1.10x
Memcpy ops	1,099	1,147	1.04x
Memset ops	345	429	1.24x

Step 2: The Lineup — the outlier appears

Most categories grow modestly. Persistent GEMM +35%, Normalization +52%. Expected.

But Flash Attention: 1.20ms → 6.17ms — 5.15x increase.

In absolute terms it's 6ms. But the ratio is an extreme statistical outlier.

Step 3: The Stakeout — same calls, more data

reshape_and_cache_flash_kernel runs 392 times in both profiles (same count!), but takes 5.15x longer per call in BS=16. More data per call, not more calls.

Memory: 83% more device-to-device copies (53 vs 29 ops).

GEMM: 7 new kernel variants in BS=16 (88ms) that don't exist in BS=1 (which had 7 different variants, only 18ms).

Step 4: The code — root cause

Every clue points to one place:

# vllm/v1/attention/backends/flash_attn.py
def use_cascade_attention(common_prefix_len, query_lens, ...):
    # Too short prefix — not worth splitting
    if common_prefix_len < 256:
        return False  # ← BS=1 exits here

    # Too few requests — not worth splitting
    num_reqs = len(query_lens)
    if num_reqs < 8:
        return False  # ← BS=1-7 exits here

    # BS=16 with shared system prompt (>256 tokens):
    # → cascade ON → split prefix/suffix → LSE merge
    # → mathematically equivalent, but FP-different
    return True  # ← determinism breaks here

Cascade attention splits attention into prefix and suffix passes, then merges with LSE arithmetic. Mathematically equivalent. Floating-point different. That's the determinism break.

The fix: disable_cascade_attn=True.

The Evidence Board
| Clue | Evidence | Verdict |
|------|----------|---------|
| 5.15x attention slowdown | reshape_and_cache_flash_kernel: 1.20ms → 6.17ms | More data per call |
| 83% more D-to-D copies | 29 → 53 ops, 81 → 163 MB | Internal tensor splitting |
| 7 new GEMM variants | 88ms new vs 18ms removed | Autotuner adapting |
| Identical kernel count | 8,955 both profiles | Same graph, different PATH |
| All clues → | CASCADE ATTENTION | flash_attn.py:673 |

The cheat sheet

┌─────────────────────────────────────────────────────────────┐
│  NSYS SQLITE CHEAT SHEET                                    │
├─────────────────────────────────────────────────────────────┤
│                                                             │
│  THE 4 TABLES:                                              │
│    CUPTI_ACTIVITY_KIND_KERNEL  → kernel executions          │
│    CUPTI_ACTIVITY_KIND_MEMCPY  → memory transfers           │
│    CUPTI_ACTIVITY_KIND_MEMSET  → memory fills               │
│    NVTX_EVENTS                 → human-added markers        │
│    + StringIds                 → name lookup                │
│                                                             │
│  THE JOIN (everywhere):                                     │
│    JOIN StringIds s ON k.shortName = s.id                   │
│                                                             │
│  TIMESTAMPS:  nanoseconds                                   │
│    / 1e3 = μs    / 1e6 = ms    / 1e9 = seconds             │
│                                                             │
│  COPYKIND:  1=CPU→GPU  2=GPU→CPU  8=GPU→GPU                │
│                                                             │
│  GRID/BLOCK:                                                │
│    grid × block = total threads                             │
│    grid=132 → all H200 SMs active                           │
│    grid=1 → one SM, 131 idle                                │
│                                                             │
│  DANGER ZONE:                                               │
│    *cublas* → non-deterministic GEMM                        │
│    *splitK* → non-deterministic reduction                   │
│    cascade/merge_attn → FP divergence                       │
│                                                             │
└─────────────────────────────────────────────────────────────┘

Your first 10 minutes

# 1. Capture
nsys profile -o my_trace --stats=true python my_script.py

# 2. Export to SQLite
nsys export --type=sqlite my_trace.nsys-rep

# 3. Find your bottleneck
python3 -c "
import sqlite3
conn = sqlite3.connect('my_trace.sqlite')
cur = conn.cursor()
cur.execute('''
    SELECT s.value, COUNT(*), ROUND(SUM(k.end-k.start)/1e6,2) AS ms
    FROM CUPTI_ACTIVITY_KIND_KERNEL k
    JOIN StringIds s ON k.shortName = s.id
    GROUP BY s.value ORDER BY ms DESC LIMIT 10
''')
for row in cur.fetchall():
    print(f'{row[0]:50s}  calls={row[1]:5d}  total={row[2]:8.2f} ms')
"

In 10 minutes you'll know which kernel is your bottleneck. GPU profiling is not a dark art. It's four tables, one JOIN, and four queries.

The answer is in the trace. Go look.

All data from real H200 GPU traces: 8,955 kernel launches during Qwen3-VL-2B inference with vLLM 0.15.2.

Two Ways to Move Tensors Without Stopping: Inside vLLM's Async GPU Transfer Patterns

Mayank Ketkar — Wed, 11 Feb 2026 21:12:16 +0000

A single torch.cuda.synchronize() in the wrong place can erase every optimization you spent weeks building. Your GPU sits idle, your pipeline stalls, and your inference latency doubles. In vLLM's distributed serving stack, tensors move between GPUs constantly: billions of parameters shuffled during weight updates, and key-value cache blocks shipped between nodes during live inference. The codebase solves these two problems with two radically different async patterns -- and studying them side-by-side reveals a masterclass in GPU concurrency.

The Problem Is Waiting

Imagine you need to get 14GB of model weights (a 7B parameter model in BF16) from a trainer process onto your inference GPU. The naive approach: transfer everything, wait, then resume inference.

At NVLink bandwidth (~900 GB/s), that's ~15ms. Not terrible. But at PCIe 5.0 (~40 GB/s effective), it's 350ms of dead time. And during RLHF training, this happens every few hundred steps.

vLLM faces this problem in two distinct contexts:

Weight distribution: Getting updated model weights from a trainer to inference workers (bulk, periodic)
KV cache migration: Shipping key-value cache blocks between disaggregated prefill and decode nodes (streaming, continuous)

Each context demands a different async pattern. Let's build the understanding from scratch.

CUDA Streams: Separate Lanes on the Same Highway

Before we can overlap anything, we need the mechanism that makes overlap possible.

Think of a factory with multiple conveyor belts feeding the same set of assembly stations. Each belt keeps its items in order. The stations (GPU compute units) pull work from any belt that has items ready:

Stream 0 (default):  [kernel A] → [kernel B] → [kernel C]
Stream 1 (transfer):  [NCCL broadcast chunk 0] → [NCCL broadcast chunk 1]

A and the broadcast can run AT THE SAME TIME on the GPU.
There is NO ordering between streams unless you add one.

A CUDA stream is just an ordered queue of GPU operations. Operations within a stream execute sequentially. Operations across streams can overlap. That's the entire mental model.

stream = torch.cuda.Stream()
with torch.cuda.stream(stream):
    # Everything here goes on the new stream, not the default one
    do_transfer()
# Back on the default stream -- transfer and default work can overlap

The key primitive: stream.synchronize() blocks the CPU until all operations on that one stream finish. It does NOT wait for other streams. This surgical waiting is what makes pipelining possible.

Quick recap: A CUDA stream is an ordered queue. Two streams run independently. stream.synchronize() waits for just one. Everything that follows builds on this single idea.

The uint8 Trick: One Blob to Rule Them All

Before we can pipeline transfers, we need to solve a packaging problem. A model has weights in different dtypes -- bfloat16, float32, maybe int8 quantization scales. NCCL broadcasts a single contiguous buffer. You can't torch.cat tensors of different dtypes.

The solution: everything is just bytes in memory.

# packed_tensor.py, lines 62-69
# Get weight tensor (any dtype: bf16, fp32, int8...)
tensor = (
    post_iter_func(next(iterator))
    .contiguous()           # Ensure contiguous memory
    .view(torch.uint8)      # Reinterpret as raw bytes — NO COPY
    .view(-1)               # Flatten to 1D for heterogeneous cat
)
packing_tensor_list[buffer_idx].append(tensor)
packing_tensor_sizes[buffer_idx] += tensor.numel()

.view(torch.uint8) is NOT a cast or a copy. It reinterprets the same memory as raw bytes -- zero cost. A bfloat16 tensor of shape [4096, 4096] becomes 33,554,432 uint8 values. Now every weight looks the same, and you can torch.cat them into one contiguous blob.

The consumer reverses the process using stored metadata (name, shape, dtype):

# Consumer unpacks: raw bytes → original dtype → original shape
tensor.contiguous().view(dtype).view(*shape)

Both sides iterate in the same order. If the order mismatches, you get silent corruption -- more on this in the antipatterns section.

Double-Buffered Weight Transfer: The Assembly Line

This is the centerpiece. vLLM's packed_tensor.py defines two constants:

# packed_tensor.py, lines 13-14
DEFAULT_PACKED_BUFFER_SIZE_BYTES = 1024 * 1024 * 1024  # 1GB
DEFAULT_PACKED_NUM_BUFFERS = 2

1GB per buffer. Two buffers. The model weights get chunked into ~1GB pieces, and the two buffers alternate roles: one is being packed with fresh weights while the other's broadcast is still in flight.

Here's the actual producer loop:

# packed_tensor.py, lines 39-86 (simplified)
streams = [torch.cuda.Stream() for _ in range(num_buffers)]
buffer_idx = 0

while True:
    streams[buffer_idx].synchronize()  # GATE: wait for buffer
    with torch.cuda.stream(streams[buffer_idx]):
        try:
            packing_tensor_list[buffer_idx] = []
            packing_tensor_sizes[buffer_idx] = 0
            while True:
                tensor = (post_iter_func(next(iterator))
                    .contiguous().view(torch.uint8).view(-1))
                packing_tensor_list[buffer_idx].append(tensor)
                packing_tensor_sizes[buffer_idx] += tensor.numel()
                if packing_tensor_sizes[buffer_idx] > target_size:
                    break
            packed = torch.cat(packing_tensor_list[buffer_idx])
            group.broadcast(packed, src=src)  # ASYNC on GPU!
            buffer_idx = (buffer_idx + 1) % num_buffers  # ROTATE
        except StopIteration:
            # Flush final partial buffer
            if len(packing_tensor_list[buffer_idx]) > 0:
                packed = torch.cat(packing_tensor_list[buffer_idx])
                group.broadcast(packed, src=src)
            break

The timeline shows the overlap:

Time ──────────────────────────────────────────────────────►

Stream 0:  ┌──pack──┐┌──broadcast──┐          ┌──pack──┐┌──broadcast──┐
           │ buf[0] ││   buf[0]    │          │ buf[0] ││   buf[0]    │
           └────────┘└─────────────┘          └────────┘└─────────────┘
                                    ▲ sync                              ▲ sync
Stream 1:            ┌──pack──┐┌──broadcast──┐          ┌──pack──┐
                     │ buf[1] ││   buf[1]    │          │ buf[1] │
                     └────────┘└─────────────┘          └────────┘
                                              ▲ sync

Iteration 0 (buffer 0, Stream 0): Pack ~1GB of weights, broadcast via NCCL. The broadcast is GPU-async -- returns to CPU immediately.

Iteration 1 (buffer 1, Stream 1): streams[1].synchronize() returns instantly (empty stream). Pack into buffer 1, broadcast. Stream 0's broadcast may still be running. This is the overlap.

Iteration 2 (buffer 0, Stream 0): streams[0].synchronize() -- the critical call. We're about to reuse buffer 0. Must wait for Stream 0's previous broadcast to finish. Then safely overwrite.

Why Does the Broadcast Return Instantly?

Because NCCL operations are asynchronous on the GPU:

# pynccl.py, lines 342-366
def broadcast(self, tensor: torch.Tensor, src: int, stream=None):
    if stream is None:
        stream = current_stream()
    if src == self.rank:
        sendbuff = buffer_type(tensor.data_ptr())
        recvbuff = buffer_type(tensor.data_ptr())  # Required by NCCL
    else:
        sendbuff = buffer_type()        # Null — receiver doesn't send
        recvbuff = buffer_type(tensor.data_ptr())
    self.nccl.ncclBroadcast(
        sendbuff, recvbuff, tensor.numel(),
        ncclDataTypeEnum.from_torch(tensor.dtype),
        src, self.comm,
        cudaStream_t(stream.cuda_stream),
    )
    # No synchronize()! Returns to CPU immediately.
    # Caller synchronizes at TOP of next iteration.

No synchronize() after the call. The synchronization happens at the top of the next iteration, when we need to reuse the buffer.

Also critical: NCCL broadcast is a collective. Every process must call it the same number of times, in the same order. Mismatch = deadlock.

Checkpoint: We've covered streams (independent lanes), uint8 packing (uniform bytes), and double-buffering (overlap via alternating buffers). This pattern handles bulk, structured transfers -- but still blocks the caller. What if we need to not block at all?

KV Connector Background Threads: Transfer Without Stopping Inference

In disaggregated serving, prefill nodes compute KV caches and ship them to decode nodes. This must happen during live inference without stalling the decode engine.

vLLM's p2p_nccl_engine.py takes a fundamentally different approach: background Python threads with producer-consumer queues.

The main inference thread does this:

# p2p_nccl_engine.py, lines 254-258
if self.send_type == "PUT_ASYNC":
    with self.send_queue_cv:
        self.send_queue.append(item)
        self.send_queue_cv.notify()  # Wake background thread
    return True  # Returns IMMEDIATELY to inference loop

Five lines. Append to a deque, notify a condition variable, return. Zero blocking.

Meanwhile, a daemon thread drains the queue:

# p2p_nccl_engine.py, lines 476-484
def send_async(self):
    """Background daemon thread — drains send queue."""
    while True:
        with self.send_queue_cv:
            while not self.send_queue:
                self.send_queue_cv.wait()   # Releases GIL, sleeps
            item = self.send_queue.popleft()
            if not self.send_queue:
                self.send_queue_cv.notify() # Signal: queue drained
        self.send_sync(item)  # NCCL send + stream.synchronize()

The background thread sleeps on send_queue_cv.wait() (zero CPU cost), wakes when notified, and performs the NCCL P2P send on a dedicated CUDA stream (self.send_stream).

 Inference Thread                    Background Threads
 ────────────────                    ──────────────────
   decode step                       ┌───────────────────┐
     send_tensor()                   │   Send Thread      │
       .append() + .notify() ──────►│   cv.wait()        │
     return True (instantly!)        │   ncclSend(item)   │
   next decode step                  └───────────────────┘
     recv_tensor()                   ┌───────────────────┐
       cv.wait() ◄──────────────────│   Listener Thread  │
       tensor = recv_store[id]       │   zmq.poll()       │
   continue inference                │   ncclRecv(tensor)  │
                                     │   cv.notify() ─────┘
                                     └───────────────────┘

The GIL Question

Python's GIL means only one thread runs Python at a time. How does this overlap?

NCCL calls are C library calls that release the GIL. While the send thread is inside ncclSend(), the inference thread is free to run. The threads work because they spend 99% of their time in GIL-releasing C calls:

Main thread:  [Python][Python][Python][Python]
Send thread:  [cv.wait][ncclSend...........][cv.wait]
               GIL-free  GIL-free (C lib)    GIL-free

Implicit vs. Explicit Coordination

Dimension	Weight Transfer	KV Connector
Coordination	Implicit (NCCL collective)	Explicit (ZMQ + queues)
Overlap target	Pack ↔ Broadcast	Transfer ↔ Inference
Blocks inference?	Yes	No
Concurrency	2 CUDA streams	OS threads + dedicated streams
Failure mode	Hang (collective mismatch)	Timeout (missed ZMQ)
Best for	Bulk, periodic	Streaming, continuous

Weight transfer coordinates implicitly through collectives -- simple but rigid. KV connector coordinates explicitly through ZMQ messages -- more machinery but fully async.

Eight Antipatterns

Synchronization

1. Syncing the wrong stream. streams[1].synchronize() when you needed streams[0]. Buffer 0 gets corrupted mid-broadcast. Silent wrong answers.

2. Missing synchronize entirely. Works under light load, race condition under heavy load. Nondeterministic, passes tests, fails in production.

3. torch.cuda.synchronize() instead of stream.synchronize(). Waits for ALL streams. Correctness maintained, performance destroyed (2-5x slower).

Buffers

4. One buffer (num_buffers=1). Must sync after every broadcast. All complexity, zero overlap.

5. Too many buffers (num_buffers=8). 8GB of VRAM wasted. Near-zero benefit over 2 buffers -- bandwidth is the bottleneck, not packing speed.

Coordination

6. Mismatched iteration order. Producer: [layer0, layer1, layer2]. Consumer: [layer2, layer0, layer1]. Every collective matches. No hang. Wrong bytes in wrong layers. Silent catastrophe.

7. Forgetting NCCL is a collective. One process skips a broadcast via early-exit. Others block forever. Deadlock, no error message.

8. Default stream from background thread. Send thread doesn't specify a stream. NCCL work goes on the default stream. Serializes with inference. Zero overlap despite threading complexity. Fix: self.send_stream = torch.cuda.Stream().

Bonus: The GIL Illusion
Python threads only provide parallelism when threads spend time in C extensions that release the GIL. If your background thread does tensor slicing in Python between NCCL calls, it holds the GIL and blocks inference. Keep the thread body a tight loop of C calls.

The Takeaway

These two patterns form a reusable vocabulary for GPU concurrency:

Double-buffered CUDA streams: Overlap data packing with transfer. Use for bulk, periodic transfers.
Background threads + producer-consumer queues: Decouple transfer from the critical path. Use for streaming, continuous transfers.

The meta-lesson: the hardest part of GPU concurrency is not making things fast -- it's making things correct while fast. stream.synchronize() is one line of code that separates "works" from "silently corrupts." The antipatterns exist because the failure mode of a missing sync is not a crash but wrong answers.

The next time you see a stream.synchronize() call in GPU code, don't skip over it. That one line is the load-bearing wall.

Code from vLLM's main branch: vllm/distributed/weight_transfer/packed_tensor.py and vllm/distributed/kv_transfer/kv_connector/v1/p2p/p2p_nccl_engine.py.

Your Ray Data Pipeline Works at 10K Samples. Here's Why It Crashes at 1M.

Mayank Ketkar — Mon, 09 Feb 2026 14:25:51 +0000

There's a moment every ML infrastructure engineer knows: the evaluation pipeline that worked perfectly on 10,000 samples crashes catastrophically when you point it at a million.

The model didn't change. The GPUs are fine. The inference code is identical. The data pipeline is the bottleneck — and it fails in ways that are completely invisible at small scale.

I spent a week scaling a Ray Data pipeline from 8,600 samples to 965,000 multi-image samples for a vision-language model (Qwen3-VL). Every sample contained 5-10 video frames — so the real data volume was closer to 5-10 million images flowing through the system. Along the way, I hit five distinct distributed systems problems, each of which required a different fix.

This is the field guide.

The Pipeline

S3 (MDS shards) → Read → Preprocess (CPU) → Inference (GPU) → Results
   966 shards       download,    decode images,     Qwen3-VL        CSV +
   ~266MB each      base64       resize, format     via vLLM        metrics

Each stage has workers. Ray Data orchestrates passing items between stages automatically — like conveyor belts between factory stations. The key insight: this is a streaming pipeline. Data should flow through continuously, not accumulate at any stage.

At 10K samples, everything fits in memory. At >1M samples with multi-image inputs, nothing does.

Problem 1: The AllToAll Barrier That Defeats Streaming

The original pipeline had a repartition() call between data loading and preprocessing:

# Step 1: Load dataset (streaming from S3)
dataset = dataset_handler.load_dataset(
    dataset_path=config.dataset_s3_path,
    num_samples=config.num_samples,
)

# Step 2: Repartition for parallel preprocessing  <-- THE BARRIER
dataset = dataset.repartition(config.num_blocks)    # AllToAll!
# Nothing downstream starts until ALL of ReadMDS finishes

# Step 3: Parallel preprocessing
processed = dataset.map(preprocessor, concurrency=160)

# Step 4: GPU inference
processed = processed.map_batches(VLMEngine, num_gpus=1, ...)

Repartition is an AllToAll barrier. Think of it like a quality checkpoint at a car factory where every single chassis must arrive before any chassis can move to the paint shop. Even if the first 100 are ready, they sit idle while chassis #965,000 is still being welded.

At 10K samples, the repartition was instant. At 1M samples being streamed from S3, it meant: download all 256GB first, hold it in memory, then start processing. GPUs idle for the entire download.

The fix: Remove it. With a streaming datasource producing multiple blocks, preprocessing can pull rows directly:

# Step 1: Stream data from S3 (no full materialization)
dataset = dataset_handler.load_dataset(
    dataset_path=config.dataset_s3_path,
    num_samples=config.num_samples,
)

# Step 2: Parallel preprocessing (CPU-bound)
processed = dataset.map(
    SceneIQPreprocessor,
    fn_constructor_kwargs={"config": config},
    compute=ActorPoolStrategy(size=config.num_preprocessing_workers),
)

# Step 3: GPU inference — data flows here immediately
processed = processed.map_batches(
    VLMEngine,
    batch_size=config.batch_size,
    num_gpus=1,
    concurrency=config.num_vllm_engines,
)

# Step 4: Postprocess and collect results
results = processed.map(postprocess).materialize()

Problem 2: Ray Thinks Each Task Uses 500MB. It's Actually 39GB.

The custom datasource creates ReadTask objects that download MDS shards from S3. Each task reports expected memory via BlockMetadata.size_bytes.

The original code set this to the raw shard size on S3 (~266MB). But in memory:

On S3 (compressed):     266 MB
After base64 encoding:  364 MB  (1.37x expansion)
In PyArrow table:       ~500 MB (column overhead)
Total per-task (16 tasks, ~60 shards each): ~39 GB

The fix: Tell Ray the truth:

# Before: Ray thought each task used ~266MB (raw shard size)
# After:  Tell Ray the actual in-memory size

for shard_info in task_shards:
        task_bytes += shard_info["raw_data"]["bytes"]

# Base64 expands ~1.37x, plus PyArrow/dict overhead.
# Use 4x raw bytes as conservative in-memory estimate.
estimated_mem = task_bytes * 4

meta = BlockMetadata(
    num_rows=task_samples,
    size_bytes=estimated_mem,    # was just task_bytes
    input_files=input_files,
    exec_stats=None,
)

This single line was the difference between "workers crash every run" and "steady-state for 12 hours."

Problem 3: 16 Tasks for 966 Shards = 64GB Per Task

Even with correct memory estimation, 16 tasks was far too few. This constant changed four times. Each wrong value crashed the cluster:

# This constant changed 4 times. Each wrong value
# crashed the cluster.
#
# 16  -> ~60 shards/task -> 64 GB -> OOM
# 48  -> ~20 shards/task -> 21 GB -> OOM
# 128 -> ~8 shards/task  -> 8.5 GB -> OOM (barely)
# 512 -> ~2 shards/task  -> 2 GB   -> Stable
#
# More tasks = less memory per task.
# Ray schedules them across CPUs, not all at once.
_DEFAULT_MAX_TASKS = 512

Having 512 tasks doesn't mean 512 simultaneous downloads. It means 512 small, schedulable units. Ray runs a handful at a time based on available CPU.

Problem 4: CPU Oversubscription

The original config was designed for a large cluster (280 CPU). On a smaller cluster (56 CPU):

# BEFORE: CPU oversubscription (crashed)
# vLLM:         20 engines x 4 CPU = 80 CPU
# Preprocessing: 160 workers x 1 CPU = 160 CPU
# Total:                               240 CPU
# Available:                            56 CPU
# Headroom:                              0 CPU  <-- CRASH

# AFTER: Right-sized for actual hardware
num_vllm_engines: 6         # 6 x 4 CPU = 24 CPU
num_preprocessing_workers: 16  # 16 x 1 CPU = 16 CPU
# Total:                               40 CPU
# Available:                           56 CPU
# Headroom:                            16 CPU  <-- Safe

Preprocessing is fast — 16 workers can easily keep 6 GPUs fed at ~22 samples/s.

Problem 5: Two Engines on a 64GB Pod

8 vLLM engines across 6 worker nodes means some pods get 2 engines:

Worker pod: 64 GB RAM
2 vLLM engines: ~30-40 GB (model + KV cache)
+ Object store: ~10-15 GB
+ ReadMDS data: ~5-10 GB
= ~55-65 GB --> OOM

The fix: 1 engine per worker node. Physical constraint, not a tuning parameter.

Under the Hood: The Custom Datasource

The fix for Problems 2 and 3 lives in a custom Ray Datasource. Here's the actual code.

The architecture: the driver reads a lightweight index.json from S3 (under 1KB), groups shards into bounded ReadTasks, and each task independently downloads and decodes its shards on a Ray worker:

# Maximum number of ReadTasks to create.
# More tasks = less memory per task (avoids OOM),
# but too many adds scheduling overhead.
#
# With datasets up to ~1000 shards at ~266MB each,
# use a high limit so each task gets 1-2 shards
# (~1GB in memory), preventing worker OOM.
# Ray Data will schedule them across available CPUs,
# so only a few run concurrently.
_DEFAULT_MAX_TASKS = 512

The get_read_tasks() method is where memory estimation happens. This is the method Ray Data calls to plan its work — and where the 4x multiplier prevents OOM:

def get_read_tasks(self, parallelism, per_task_row_limit=None):
    # Collect shards up to effective_samples limit
    active_shards = []
    remaining = self._effective_samples
    for shard in self._shard_infos:
        if remaining <= 0:
            break
        n = min(shard["samples"], remaining)
        active_shards.append((shard, n))
        remaining -= n

    num_shards = len(active_shards)
    num_tasks = min(num_shards, self._max_tasks)

    # Distribute shards across tasks (contiguous groups)
    shards_per_task = (num_shards + num_tasks - 1) // num_tasks

    read_tasks = []
    for group in shard_groups:
        task_bytes = sum(s["raw_data"]["bytes"] for s in group)

        # Base64 expands ~1.37x, plus PyArrow/dict overhead.
        # Use 4x raw bytes as conservative in-memory estimate
        # so Ray can schedule tasks without OOM-killing workers.
        estimated_mem = task_bytes * 4

        meta = BlockMetadata(
            num_rows=task_samples,
            size_bytes=estimated_mem,  # was just task_bytes!
            input_files=input_files,
            exec_stats=None,
        )

        read_tasks.append(ReadTask(
            read_fn=_make_read_fn(),
            metadata=meta,
        ))

    return read_tasks

Each ReadTask runs this function on a Ray worker — it downloads 1-2 shards, decodes samples, and returns a PyArrow table:

def _read_mds_shards(remote_path, shard_group, offset, max_samples):
    """Download MDS shards from S3 and extract samples.
    Runs on Ray workers — each task handles 1-2 shards."""
    s3 = boto3.client("s3")
    rows = []

    for shard_info in shard_group:
        # Download shard from S3 to temp directory
        basename = shard_info["raw_data"]["basename"]
        bucket, key = _parse_s3_path(f"{remote_path}/{basename}")
        s3.download_file(bucket, key, local_path)

        # Decode samples via MDSReader
        reader = MDSReader.from_json(dirname=tmp_dir, obj=shard_info)
        for idx in range(shard_samples):
            raw_sample = reader.get_item(idx)
            row = _extract_sample(raw_sample, sample_id)
            if row is not None:
                rows.append(row)

    # Return as PyArrow table for Ray Data streaming
    return [pa.table({
        "sample_id": [r["sample_id"] for r in rows],
        "system_prompt": [r["system_prompt"] for r in rows],
        "prompt_blocks": [r["prompt_blocks"] for r in rows],
        "target": [r["target"] for r in rows],
    })]

The key design choices: each task is self-contained (downloads its own shards, no shared state), the memory estimate is conservative (4x raw bytes), and the PyArrow table output integrates directly with Ray Data's streaming execution.

The Expanding Suitcase

Here's what most people miss. Each sample isn't text — it's 5-10 video frames:

Per sample in ReadMDS:      ~5-20 MB  (base64 strings)
Per sample in preprocessing: ~50-100 MB (PIL Images, uncompressed!)
  A 1 MB JPEG -> ~10 MB as a PIL Image
  x 10 frames = 100 MB per sample

This is the "expanding suitcase" problem. You pack vacuum-sealed clothes (compressed images, ~266MB per shard). At the destination, you unseal them and they expand 4-10x.

The saving grace: vLLM is the bottleneck. At ~22 samples/s, it's slow enough that data doesn't pile up. Natural backpressure keeps the pipeline stable.

The Result

Parameter	Original	Final	Change
Repartition	AllToAll barrier	Removed	Eliminated
Memory estimation	1x (raw bytes)	4x multiplier	Realistic
ReadMDS max_tasks	16	512	+3100%
vLLM engines	20	6 (1/node)	-70%
Preprocessing workers	160	16	-90%
CPU utilization	100%	71%	Headroom

1M multi-image samples. ~22 samples/s. 12+ hours. Zero OOM crashes.

The model code didn't change. Every fix was in how data gets to the GPUs.

Five Takeaways

Repartition kills streaming. It's an AllToAll barrier that forces full materialization. Remove it if your datasource already produces multiple blocks.
BlockMetadata.size_bytes is your memory contract with Ray. For image/video data, in-memory size can be 4-10x on-disk size. Set it explicitly.
More tasks = less memory per task. The simplest OOM fix. 512 tasks doesn't mean 512 simultaneous downloads.
Right-size for physical topology. Count GPUs, CPUs, and RAM per node, not just totals. One engine per node avoids hidden RAM contention.
The data pipeline is always the bottleneck at scale. At 10K, GPU is the constraint. At 1M, the plumbing is. The inference code doesn't need to change.

Ray Data custom datasources: docs.ray.io | Performance tips: docs.ray.io

Compiling the Vision Encoder: Squeezing 3% More Throughput from Qwen3-VL on Hopper GPUs

Mayank Ketkar — Mon, 09 Feb 2026 02:08:38 +0000

When you run a vision-language model through vLLM, the framework does something clever: it compiles the LLM decoder with torch.compile, fuses operators, and captures CUDA graphs for maximum throughput. But there is a component it quietly leaves behind -- the Vision Transformer (ViT) encoder that processes your images. It runs in plain eager mode, every single time.

We changed that for Qwen3-VL. The result: 3.4% higher throughput on an NVIDIA H200, three previously unknown bugs discovered and fixed, and a one-flag change that any vLLM user can enable today.

This post walks through the engineering story -- why the encoder was left behind, how we ported compilation support from a sibling model, what broke along the way, and what the profiler actually says about where the time goes.

Why Does the Encoder Run Eager?

vLLM's compilation infrastructure is built around the LLM decoder. When you launch an inference server, the startup sequence compiles the decoder's forward pass with torch.compile, traces its graph, and captures CUDA graphs at various batch sizes. This eliminates Python overhead and enables kernel fusion across attention, LayerNorm, and MLP layers.

The multimodal encoder -- the ViT that converts raw image pixels into embedding vectors -- gets none of this treatment. The reason is a single boolean flag in vLLM's compilation config:

compile_mm_encoder: bool = False
"""Whether or not to compile the multimodal encoder.
Currently, this only works for Qwen2_5_vl and mLLaMa4
models on selected platforms. Disabled by default until
more models are supported/tested to work."""

The default is False, and for good reason. Vision encoders face a fundamental tension with compilation: variable input shapes. Different requests can carry images at different resolutions, producing different numbers of patches. CUDA graphs require fixed tensor shapes at capture time. A general-purpose serving framework cannot assume that every image will be the same size.

But for batch inference workloads with fixed-size images -- which is common in production pipelines processing standardized camera frames, satellite tiles, or document pages -- this conservatism leaves performance on the table. If your images are all the same resolution, the encoder always receives identically shaped tensors, and torch.compile can fully specialize.

There was a second, more specific problem: Qwen3-VL simply lacked the compilation decorators. Its sibling model, Qwen2.5-VL, already had full torch.compile support for its encoder. Qwen3-VL shared much of the same architecture (including the identical attention implementation), but the compilation wiring was never ported over.

The Pattern: Porting from Qwen2.5-VL

vLLM uses a decorator-based system for selective compilation. Rather than compiling an entire model's forward pass (which would break on Python control flow, NumPy calls, and dynamic branching), it compiles individual submodules whose forward() methods contain only clean tensor operations.

Qwen2.5-VL already had this wired up for three encoder submodules: VisionPatchEmbed, VisionBlock, and VisionPatchMerger. Our task was to replicate the exact same pattern in Qwen3-VL.

The Decorator

Each compilable submodule gets a @support_torch_compile decorator that declares which tensor dimensions are dynamic and provides a gating function:

@support_torch_compile(
    dynamic_arg_dims={"x": 0},
    enable_if=should_torch_compile_mm_vit,
)
class Qwen3_VisionPatchEmbed(nn.Module):
    ...

The dynamic_arg_dims={"x": 0} tells torch.compile that dimension 0 of the input tensor x can vary between calls (different numbers of patches), so it should not bake that shape into the compiled graph. The enable_if callback is a one-liner that checks whether the user opted in:

def should_torch_compile_mm_vit(vllm_config: VllmConfig) -> bool:
    return vllm_config.compilation_config.compile_mm_encoder

When compile_mm_encoder is False (the default), the decorator sets self.do_not_compile = True and the forward pass runs in eager mode -- zero overhead, zero behavior change. When it is True, the decorator wraps the module in torch.compile on first call and uses compiled execution from then on.

The Model Tags

The second piece of wiring is set_model_tag, a context manager that tells the compilation backend to use separate caches for encoder versus decoder components. Without tags, the encoder and decoder would share a single compile cache, causing shape mismatches when the compiler tries to reuse a graph compiled for decoder weight shapes on encoder weights.

In Qwen3_VisionTransformer.__init__():

# DO NOT MOVE THIS IMPORT
from vllm.compilation.backends import set_model_tag

with set_model_tag("Qwen3_VisionPatchEmbed", is_encoder=True):
    self.patch_embed = Qwen3_VisionPatchEmbed(...)

with set_model_tag("Qwen3_VisionPatchMerger", is_encoder=True):
    self.merger = Qwen3_VisionPatchMerger(...)

# Deepstack mergers need a separate tag (different weight shapes!)
with set_model_tag("Qwen3_VisionPatchMerger_deepstack", is_encoder=True):
    self.deepstack_merger_list = nn.ModuleList([...])

with set_model_tag("Qwen3_VisionBlock", is_encoder=True):
    self.blocks = nn.ModuleList([
        Qwen3_VisionBlock(...) for _ in range(depth)
    ])

That comment about DO NOT MOVE THIS IMPORT is not a joke -- it matches the exact pattern in Qwen2.5-VL and relates to import ordering constraints with the compilation backend (see vllm#27044).

Notice the deepstack mergers get their own tag, separate from the main merger. This was not in the original plan. It was the fix for Bug #2, which we will get to shortly.

What Gets Compiled

The Qwen3-VL vision encoder has three distinct compilable submodules:

Submodule	What It Does	Dynamic Dims
`Qwen3_VisionPatchEmbed`	Reshape + Conv3D + reshape (pixels to patch embeddings)	`x: dim 0`
`Qwen3_VisionBlock` (x24)	LayerNorm -> Attention -> Residual -> LayerNorm -> MLP -> Residual	`x, cu_seqlens, cos, sin: dim 0`
`Qwen3_VisionPatchMerger`	LayerNorm -> Linear -> GELU -> Linear (merge spatial patches)	`x: dim 0`

The outer VisionTransformer.forward() -- which orchestrates these submodules -- is deliberately not compiled. It contains NumPy operations (np.array, np.cumsum), Python control flow (isinstance, list comprehensions), and .tolist() calls that would cause graph breaks. The per-submodule pattern avoids all of this.

Zero Graph Breaks

The first compile attempt was the moment of truth. We enabled TORCH_LOGS=+dynamo and TORCH_COMPILE_DEBUG=1, loaded a handful of test images, and watched TorchDynamo trace through the encoder.

Result: zero graph breaks. The single COMPILING GRAPH event reported:

COMPILING GRAPH due to GraphCompileReason(
    reason='return_value',
    user_stack=[<FrameSummary file qwen3_vl.py, line 1169 in forward>],
    graph_break=False
)

This was expected but still satisfying. The per-submodule compilation pattern is specifically designed to isolate clean tensor operations from Python control flow. Each compiled forward method contains nothing but torch operations -- reshapes, linear projections, attention, LayerNorm, residual additions. No data-dependent control flow, no Python-side data structures, no calls that escape the Dynamo graph.

The key insight: if you tried to compile the entire VisionTransformer.forward() as one graph, you would hit graph breaks immediately on the NumPy calls that compute positional embeddings and cumulative sequence lengths. By compiling only the inner submodules, you get all the fusion benefits with none of the graph break headaches.

Three Bugs Found and Fixed

Zero graph breaks did not mean zero problems. The first full run crashed. Then it crashed differently. Then it crashed a third way. Here is what we found.

Bug 1: `AssertionError: Forward context is not set` in `profile_run()`

The crash:

AssertionError: Forward context is not set.
Please use `set_forward_context` to set the forward context.

What happened: When vLLM starts up, it runs a profiling pass (profile_run()) to determine memory usage. This calls self.model.embed_multimodal() to profile the encoder. In eager mode, this works fine -- the encoder's forward methods are just regular PyTorch calls.

But with @support_torch_compile, the compilation backend wraps each submodule in a CUDAGraphWrapper. The wrapper's __call__ method reads forward_context.cudagraph_runtime_mode to decide whether to execute via CUDA graph or fall through to eager. Without a forward context set, it crashes.

The fix: Wrap the profiling call in set_forward_context:

with set_forward_context(attn_metadata=None, vllm_config=self.vllm_config):
    dummy_encoder_outputs = self.model.embed_multimodal(
        **batched_dummy_mm_inputs
    )

Since attn_metadata=None, the wrapper sees CUDAGraphMode.NONE and falls through to eager execution -- exactly the behavior we want during profiling.

Bug 2: `AssertionError: expected size 1024==4096`

The crash:

AssertionError: expected size 1024==4096, stride 1==1 at dim=0

What happened: Qwen3-VL has two types of patch mergers. The main merger has a LayerNorm over context_dim=1024 (the per-patch hidden size before spatial merging). The deepstack mergers have a LayerNorm over hidden_size=4096 (the full hidden size, via use_postshuffle_norm=True). Both use the Qwen3_VisionPatchMerger class.

In our initial implementation, both mergers shared the same set_model_tag("Qwen3_VisionPatchMerger") context. This meant they shared a single compiled graph cache. When torch.compile traced through the main merger (norm weight shape (1024,)), it cached a graph with that shape baked in. When the deepstack merger tried to reuse the same cached graph with its (4096,) weights -- crash.

The fix: Separate model tags:

with set_model_tag("Qwen3_VisionPatchMerger", is_encoder=True):
    self.merger = ...          # LayerNorm over 1024

with set_model_tag("Qwen3_VisionPatchMerger_deepstack", is_encoder=True):
    self.deepstack_merger_list = ...  # LayerNorm over 4096

Same Python class, different compile caches. The tag system was designed exactly for this -- but you have to remember to use it when two instances of the same class have different weight shapes.

Bug 3: Same as Bug 1, but in `_execute_mm_encoder()`

The profiling fix (Bug 1) resolved the startup crash, but the same AssertionError appeared during actual inference. The encoder execution path in _execute_mm_encoder() also called embed_multimodal() without setting forward context.

The fix: Same pattern -- wrap the encoder execution loop in set_forward_context(attn_metadata=None, ...).

Defense-in-Depth

After fixing both call sites, we added a belt-and-suspenders guard in CUDAGraphWrapper.__call__ itself:

def __call__(self, *args, **kwargs):
    if not is_forward_context_available():
        return self.runnable(*args, **kwargs)  # Eager fallback
    forward_context = get_forward_context()
    ...

If any future code path calls a compiled encoder submodule without setting forward context, it gracefully falls through to eager execution instead of crashing. This is defense-in-depth -- the primary fix is ensuring all call sites set the context, but the guard protects against regressions.

Profiling: Where the Time Goes

With compilation working, we instrumented the encoder with torch.cuda.Event timing to measure exactly how much each component contributes and how much compilation helps.

The Encoder Is Only 13.5% of Total Inference Time

For Qwen3-VL-2B on our workload, the ViT encoder processes each image once to produce embedding tokens, then the LLM decoder generates the output sequence. The decoder dominates.

Component	Baseline (ms)	Compiled (ms)	Speedup
PatchEmbed	5.2	6.2	-19%
VisionBlocks (24)	352.5	330.2	+6.3%
PatchMerger	3.8	5.3	-39%
Total Encoder	450.3	430.5	+4.4%

VisionBlocks Win, Small Ops Lose

The 24 VisionBlocks are where compilation shines. Each block runs LayerNorm -> Attention -> Residual -> LayerNorm -> MLP -> Residual. The Inductor backend fuses these into fewer, more efficient kernels. Blocks 1-23 show a consistent 7-8% per-block speedup, accumulating to a 22.3ms reduction.

PatchEmbed and PatchMerger show the opposite: compilation makes them slower. These are tiny operations (~0.3ms per call). The @support_torch_compile decorator adds Python dispatch overhead on every call, and at this scale, the overhead exceeds the fusion benefit. It is a classic tradeoff -- compilation has a per-call dispatch cost that only pays off when the compiled operation is large enough.

A pragmatic optimization would be to remove the @support_torch_compile decorators from PatchEmbed and PatchMerger, compiling only VisionBlocks. The net encoder speedup would actually be slightly higher without the small-op regressions. But the dispatch overhead is small in absolute terms (a few milliseconds total), and having all submodules wired for compilation maintains consistency with the Qwen2.5-VL pattern.

Why 4.4% Encoder Speedup Becomes 3.4% End-to-End

With the encoder representing 13.5% of total inference time, even a 4.4% encoder speedup translates to only ~0.6% of total wall time through Amdahl's Law. The actual measured end-to-end improvement is larger than that simple calculation suggests, likely because the compilation also reduces Python overhead and improves memory access patterns in ways that benefit the surrounding orchestration code.

End-to-End Benchmark

We ran a full A/B comparison over ~8,000 samples on an NVIDIA H200, with 10-sample warmup excluded from measurements.

Metric	Baseline	Compiled	Delta
Throughput	32.33 samp/s	33.42 samp/s	+3.4%
Generate time	266.1s	257.4s	-8.7s
Per-sample latency	30.93ms	29.92ms	-1.0ms
Model load time	37.3s	50.2s	+12.9s

The 3.4% throughput improvement is consistent across scales. We saw similar relative gains at 100 samples (+0.9% -- noisier at smaller scale) and at the full dataset (+3.4%).

The model load time increase (+12.9s) is a one-time cost for Dynamo bytecode transforms and Inductor codegen on the encoder submodules. On subsequent runs, the compilation cache (~/.cache/vllm/torch_compile_cache/) eliminates recompilation entirely -- subsequent startups are only marginally slower than baseline. In a production serving context, this compilation happens once at server startup and all subsequent inference benefits from the speedup.

Break-Even Analysis

Parameter	Value
One-time compilation overhead	12.9s
Per-sample time saving	~1.0ms
Break-even point	~12,900 samples

For the first-ever run (cold compilation cache), you need to process approximately 13,000 samples before the compilation overhead is amortized. For any subsequent run with a warm cache, the benefit is immediate.

Output Correctness

An important caveat: compiled and baseline modes produce slightly different outputs on some inputs. This is expected behavior from torch.compile -- the Inductor backend may apply different operator fusion, reduction ordering, and kernel implementations that change floating-point rounding at the bit level. These tiny differences in intermediate activations can cascade through the encoder, shift logits by small amounts, and occasionally flip the argmax for borderline tokens during autoregressive decoding.

Concretely:

Both modes are individually deterministic -- the same mode always produces the same output for the same input, run after run.
They are not cross-compatible -- baseline and compiled may differ on some samples.
The differences are small in magnitude and affect only a fraction of samples.

This is a property of torch.compile itself, not of our changes. If your application requires bitwise reproducibility between compiled and non-compiled modes, this is worth knowing. If you only need consistency within a single mode (the more common requirement), both modes deliver it.

When Would This Matter More?

A 3.4% throughput improvement is real and free (once the cache is warm), but it is bounded by the encoder's share of total inference time. For Qwen3-VL-2B, the ViT encoder is small relative to the LLM decoder. Several scenarios would amplify the benefit:

Larger ViT encoders. Qwen3-VL-72B has a proportionally larger vision encoder. The same 7-8% per-block VisionBlock speedup applied to more expensive encoder blocks would yield a larger end-to-end improvement.

Video workloads. Video inputs require processing many frames, multiplying encoder invocations per request. The encoder's share of total time increases, and the compilation benefit compounds.

High-concurrency serving. When many requests arrive simultaneously, encoder latency adds up across the batch. Shaving 4.4% off each encoder call reduces queuing delay.

Bandwidth-bound GPUs. The H200 is a compute-rich Hopper GPU. On more bandwidth-constrained hardware like the L40S, the operator fusion from torch.compile (which reduces memory traffic by eliminating intermediate tensor materializations) would likely yield a larger percentage improvement.

Higher-resolution images. More patches per image means more work in the VisionBlocks, which are the primary beneficiaries of compilation.

How to Enable It

One flag:

from vllm import LLM

llm = LLM(
    model="Qwen/Qwen3-VL-2B-Instruct",
    compilation_config={"compile_mm_encoder": True},
    # ... other settings
)

Or via the CLI:

vllm serve Qwen/Qwen3-VL-2B-Instruct \
    --compilation-config '{"compile_mm_encoder": true}'

That is it. No model changes, no custom code, no configuration gymnastics. The flag tells vLLM to apply torch.compile to the ViT encoder submodules during model initialization. The first inference call that includes images will trigger compilation (or load from cache), and all subsequent calls use the compiled kernels.

First Run vs. Subsequent Runs

On the very first run with a new model or new vLLM version, you will see a longer model load time (~13s extra) as TorchDynamo traces and Inductor generates code for the encoder submodules. These artifacts are cached to ~/.cache/vllm/torch_compile_cache/.

On all subsequent runs, the cached artifacts load in seconds, and the throughput benefit is immediate.

Conclusion

This was a small change -- six modifications across two files for the core enablement, plus four files touched for bug fixes. The pattern was already established by Qwen2.5-VL; we just ported it to Qwen3-VL. But small changes can have disproportionate engineering value when they uncover latent bugs.

The three bugs we found -- missing set_forward_context in two encoder execution paths, and shared compile caches for mergers with different weight shapes -- are not specific to Qwen3-VL. They would affect any model that enables compile_mm_encoder. The fixes (including the defense-in-depth guard in CUDAGraphWrapper) benefit the entire vLLM multimodal compilation infrastructure.

The profiling results tell an honest story: the ViT encoder is a small fraction of end-to-end time for a 2B parameter model, so even a solid 4.4% encoder speedup translates to a modest 3.4% end-to-end gain. But it is a free 3.4% -- one flag, cached after the first run, no accuracy impact within a single mode. For larger models, video workloads, or bandwidth-constrained hardware, the benefit would be larger.

Sometimes the most useful engineering work is not building something new, but noticing that a capability already exists in the codebase and was never wired up for your model.

Summary of Changes

File	Change
`vllm/model_executor/models/qwen3_vl.py`	`@support_torch_compile` decorators on 3 encoder submodules + `set_model_tag` wiring
`vllm/config/compilation.py`	Updated `compile_mm_encoder` docstring to include Qwen3-VL
`vllm/v1/worker/gpu_model_runner.py`	`set_forward_context` wrapper in `_execute_mm_encoder()` and `profile_run()`
`vllm/compilation/cuda_graph.py`	`is_forward_context_available()` guard in `CUDAGraphWrapper.__call__`

Hardware and Software

GPU: NVIDIA H200 (141 GB HBM3e)
vLLM: 0.15.x (main branch)
PyTorch: 2.9.x
Model: Qwen3-VL-2B-Instruct (fine-tuned checkpoint)
Workload: ~8,000 fixed-resolution images, single GPU, temperature=0.0, max_tokens=128