This document is relevant for: Trn2, Trn3

Block Packing Vision Attention#

Overview#

Block Packing Vision Attention eliminates wasted computation in multi-image attention by packing variable-sized image token sequences into fixed-size blocks using First-Fit-Decreasing (FFD) bin packing. Each block is processed as an independent batch element, removing all cross-block attention and delivering 2x-2.5x efficiency gains over standard approaches.

Problem Statement#

Vision understanding models apply all-to-all attention independently within each image — every visual token in one image attends to every other token in that same image. When a request contains multiple images at varying resolutions, the token sequences differ in length and must be organized for batched computation. The two standard approaches both incur significant waste:

Standard approach 1: Batch-dimension packing#

Each image occupies a separate batch element, padded to the length of the largest image:

Image A: 1024 tokens ████████████████
Image B:   64 tokens ██░░░░░░░░░░░░░░  (93% padding)
Image C:  384 tokens ██████░░░░░░░░░░  (62% padding)
                     ↑ padded to 1024

Short sequences carry massive padding overhead and require two-dimensional bucketing (number of images × max image size).

Standard approach 2: Sequence-dimension concatenation#

All image tokens are concatenated into a single sequence with minimal trailing padding:

[Image A: 1024 tokens][Image B: 64 tokens][Image C: 384 tokens][pad: 64]
 ← total: 1536 tokens (next bucket) →

This minimizes padding but computes a full seq_len × seq_len attention matrix. A non-causal mask prevents cross-image attention, but the kernel still executes the masked-out regions. For the example of two images (1024 + 2048 tokens):

  • Useful compute: 1024² + 2048² = 5.2M

  • Total compute: 3072² = 9.4M

  • Wasted: 44%

The waste grows quadratically with more images. For 24 small images of 128 tokens each:

  • Useful compute: 24 × 128² = 393K

  • Total compute: 3072² = 9.4M

  • Wasted: 95.8%

Solution: Block Packing with FFD#

Block packing restructures the computation by grouping images into fixed-size blocks rather than one flat sequence or one-per-batch-element:

  1. Pack image token sequences into blocks of a compile-time-fixed size using First-Fit-Decreasing (FFD).

  2. Process each block as an independent batch element — the attention kernel computes only the block_size × block_size matrix per block with no cross-block computation.

First-Fit-Decreasing (FFD) Algorithm#

FFD is a bin-packing heuristic that operates in two steps:

  1. Sort images by token count in descending order.

  2. Place each image into the first block that has sufficient remaining capacity.

FFD guarantees at most ⌈11/9 · OPT⌉ + 1 blocks, where OPT is the theoretical minimum — worst-case within 22% of optimal. In practice, production image-size distributions (heterogeneous resolutions with natural size gaps) achieve optimal or OPT+1 packing.

Worked Example#

Consider 5 images: one 1024-token, one 512-token, one 256-token, and two 128-token images, with block_size = 1024:

Block 1: [Image 1K]                    → 1024/1024 used
Block 2: [Image 512][Image 256][2×128] → 1024/1024 used

Compute comparison:

Approach

Total compute

Useful compute

Efficiency

Batch packing (pad to 1024)

5 × 1024²

1024² + 512² + 256² + 2×128²

~28%

Sequence packing (flat)

2048²

same

~35%

Block packing (1024)

2 × 1024²

same

~70%

Block packing achieves 2.5x better efficiency than batch packing and 2x better than sequence packing in this scenario.

Relationship to Standard Approaches#

Block packing generalizes both standard approaches:

  • When block_size = total sequence length: a single block spans the entire attention matrix, reducing to standard sequence-dimension concatenation.

  • When all images have equal size: setting block_size to that size (rounded to the next multiple of 128) gives each image its own block, reducing to batch-dimension parallelism.

  • When images vary in size (the common case): block packing is strictly more efficient than both alternatives.

How It Works#

Step 1: Preprocessing (Framework Level)#

The framework performs FFD packing before kernel dispatch:

  1. Sort images by token count (descending).

  2. Assign each to the first block with remaining capacity ≥ image token count.

  3. Shuffle vision token position IDs to reflect the image-to-block mapping — the result is a [num_blocks × block_size] tensor with padded positions at each block’s tail.

  4. Compute bound_min and bound_max tensors per block, encoding the local start/end indices of each image within that block.

Step 2: Execution (Kernel)#

The packed tensor is dispatched to the attention_cte prefill kernel:

  • num_blocks maps to the kernel’s batch dimension.

  • block_size maps to the kernel’s sequence length (seqlen_q).

  • The kernel processes blocks sequentially, using bound_min/bound_max to restrict attention to valid image-local regions within each block.

Bucketing occurs on two dimensions:

  1. Block size — determined at compile time based on maximum image resolution.

  2. Number of blocks — computed as max_seq_len // block_size.

Choosing the Right Block Size#

Block size is the key tuning parameter. The right value depends on your workload characteristics:

Guidelines#

In general, we want to balance intra-block padding with the wasted compute of fully padded blocks. Efficient packing will produce fewer blocks with less intra-padding within each block, which reduce both.

Workload

Recommended block size

Rationale

Single large image (e.g., 4K×4K)

Match the image token count (rounded to 128)

Single block, no packing overhead

Many small images (e.g., thumbnails)

1024

Packs many images per block, minimizes block count — fewer blocks means less wasted compute from fully padded trailing blocks and better NeuronCore utilization along the sequence dimension

Mixed resolutions

1024–2048

Balances packing density against per-block padding

Video workload

Multiple of num_tokens_per_frame (rounded to 128)

Aligns block boundaries to frame boundaries, avoiding mid-frame splits (e.g., 16 tokens/frame × 64 frames = 1024)

Max resolution known at deploy time

max_num_tokens_per_image (rounded to 128)

Guarantees any single image fits in one block

Key Considerations#

  1. Block size must be ≥ the largest single image’s token count. If an image exceeds the block size, it cannot be packed and the kernel will fail. Set block size to at least max_num_tokens_per_image for your model’s configured resolution.

  2. Larger blocks = fewer blocks but more intra-block padding. A 4096-token block can absorb any combination of smaller images, but each block carries more unused slots when images don’t fill it completely.

  3. Smaller blocks = more blocks but better packing density. A 1024-token block forces large images into dedicated blocks while letting small images share blocks efficiently. However, very small blocks increase the batch dimension and may underutilize compute for workloads dominated by large images.

  4. Alignment constraint: block size must be a multiple of 128 to satisfy NKI kernel tile alignment requirements.

Decision Flowchart#

What is your max single-image token count?
│
├─ ≤ 1024  → block_size = 1024
│             (good packing density)
│
├─ ≤ 2048  → block_size = 2048
│             (accommodates large images, still packable)
│
└─ > 2048  → block_size = ceil(max_tokens / 128) × 128
              (ensures single-image fit, packing for smaller images)

Example: Estimating Efficiency#

For a request with N images producing token counts [t₁, t₂, ..., tₙ] and a chosen block_size B:

Useful compute  = Σ tᵢ²
Block compute   = num_blocks × B²
Efficiency      = Useful compute / Block compute

Where num_blocks is determined by FFD packing. Higher efficiency means less wasted compute.

Performance Impact#

The efficiency gain depends on workload heterogeneity:

  • Homogeneous images (all same resolution): Block packing matches batch-dimension parallelism — no gain, no regression.

  • Heterogeneous images (mixed resolutions): 2x–2.5x efficiency improvement over the best standard approach.

  • Many small images: Up to 10x+ improvement over flat sequence packing (e.g., 24 thumbnails: ~4% → ~50%+ efficiency).

Key Files#

File

Purpose

vllm_neuron/model/qwen3_vl/utils/vision_block_packing.py

FFD packing, block scatter, bounds computation

vllm_neuron/model/qwen3_vl/utils/vision_preprocessing.py

CPU preprocessing (RoPE, attention bounds from grid_thw)

vllm_neuron/model/qwen3_vl/vision_encoder_bf16.py

Vision encoder (block-level forward with bound_min/bound_max)

vllm_neuron/model/qwen3_vl/model_bf16.py

Model integration (packing invocation, position shuffling)

vllm_neuron/functional/attention/attention_cte.py

NKI attention kernel (attention_cte with sequence bounds)

This document is relevant for: Trn2, Trn3