vllm - 💡(How to fix) Fix [RFC]: Convert Triton kernels from raw pointers to block pointers [1 participants]

Motivation.

vLLM's core dense-inference Triton kernels currently use raw pointer arithmetic (tl.load(ptr + offset, mask=...)). Triton's recommended and forward-looking memory access API is block pointers (tl.make_block_ptr / tl.advance), which:

1. Enable hardware portability — Block pointers abstract away raw address arithmetic, allowing Triton backends targeting non-NVIDIA hardware (e.g., Intel XPU, IBM Spyre/AIU) to lower these kernels. Some accelerators with tiled memory architectures fundamentally cannot accept raw pointer arithmetic.

2. Leverage Hopper TMA — On NVIDIA Hopper (H100/H200), block pointers can be lowered to Tensor Memory Accelerator (TMA) instructions, potentially improving memory throughput.

3. Align with Triton's direction — The Triton project is moving toward structured memory access as the primary API.

4. Improve readability — Block pointer code separates memory layout (shape, strides, offsets) from computation (load, compute, store).

Some newer vLLM kernels (FLA, Mamba, LoRA) already use block pointers, but the core inference-path kernels do not.

Proposed Change.

Convert the following simple, element-wise kernels to use block pointers:

These are mechanical conversions — replacing tl.load(ptr + offset, mask=...) with tl.load(block_ptr, boundary_check=...) — and carry low risk of behavioral change.

Validation approach

• Per-kernel numerical equivalence tests comparing raw-pointer and block-pointer outputs across multiple shapes and dtypes (fp16, bf16, fp32).
• Performance benchmarking on A100/H100 to confirm no regression.
• Integration test via vllm.LLM.generate() to verify identical output tokens.

Prior work

We have working block-pointer conversions for some of these kernels with numerical equivalence tests, and have validated the conversion approach across kernels of varying complexity.

Future work

If there is interest, the same conversion could be extended to the more complex kernels on the dense inference path — KV cache reshape, decode softmax+reduceV, merge attention states, and the attention kernels (prefill, decode stage 1, GQA decode, chunked prefill/paged decode).

Questions for maintainers

1. Is there interest in accepting block-pointer conversions of these kernels?
2. Any preference on tl.make_block_ptr vs tl.make_tensor_descriptor? The former is more stable; the latter is newer and maps to TMA on Hopper.
3. Should the block-pointer versions replace the raw-pointer versions in-place, or coexist as an alternative backend (e.g., selectable via config)?
4. Any specific performance benchmarking requirements or hardware targets we should cover?

Feedback Period.

No response

CC List.

No response

Any Other Things.

No response

Before submitting a new issue...

• Make sure you already searched for relevant issues, and asked the chatbot living at the bottom right corner of the documentation page, which can answer lots of frequently asked questions.