sgl-project/sglang · error · ValueError

renorm kernels require a CUDA/HIP tensor

Error message

renorm kernels require a CUDA/HIP tensor

What it means

The renormalization Triton kernels are GPU-only (CUDA/HIP); _prepare_probs checks probs.is_cuda and rejects CPU tensors because no CPU fallback is compiled and launching a Triton kernel on CPU data would crash.

Source

Thrown at python/sglang/kernels/ops/sampling/renorm_triton.py:58

    out_ptr,
    row_sums_ptr,
    numel,
    vocab_size: tl.constexpr,
    BLOCK_SIZE: tl.constexpr,
):
    offsets = tl.program_id(0).to(tl.int64) * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
    mask = offsets < numel
    row = offsets // vocab_size
    values = tl.load(out_ptr + offsets, mask=mask, other=0.0).to(tl.float32)
    denominator = tl.load(row_sums_ptr + row, mask=mask, other=1.0)
    tl.store(out_ptr + offsets, values / denominator, mask=mask)


def _prepare_probs(probs: torch.Tensor) -> torch.Tensor:
    if probs.ndim != 2:
        raise ValueError(f"probs must be 2D, got shape={tuple(probs.shape)}")
    if not probs.is_cuda:
        raise ValueError("renorm kernels require a CUDA/HIP tensor")
    return probs.float().contiguous()


def _renorm_from_pivots(probs_fp32: torch.Tensor, pivots: torch.Tensor) -> torch.Tensor:
    batch_size, vocab_size = probs_fp32.shape
    num_chunks = triton.cdiv(vocab_size, _BLOCK_SIZE)
    out = torch.empty_like(probs_fp32)
    partial_sums = torch.empty(
        (batch_size, num_chunks), device=probs_fp32.device, dtype=torch.float32
    )
    _mask_and_partial_sum_kernel[(batch_size, num_chunks)](
        probs_fp32,
        pivots,
        out,
        partial_sums,
        vocab_size=vocab_size,
        num_chunks=num_chunks,
        BLOCK_SIZE=_BLOCK_SIZE,

View on GitHub (pinned to 0132848349)

Solutions

1. Move probs to the GPU: probs = probs.cuda() before calling
2. Use torch-native renormalization on CPU (sort + cumsum mask) when no GPU is available

Example fix

# before
out = top_p_renorm_probs_triton(probs_cpu, 0.9)
# after
out = top_p_renorm_probs_triton(probs_cpu.cuda(), 0.9)

assert probs.is_cuda, 'renorm requires GPU tensors'
if not probs.is_cuda: probs = probs.cuda()

def on_gpu_for_renorm(p): return p.is_cuda

When it happens

Trigger: Calling top_p/top_k renorm with a CPU tensor, or a CUDA tensor that was moved with .cpu() for logging and then reused.

Common situations: Notebook experimentation on CPU before moving to GPU, or unit tests that run sampling logic without a GPU device.

Related errors

• Input probs contains NaN.
• cos and sin must be CUDA and last-dim contiguous
• nvfp4_gemm_swiglu_nvfp4_quant requires CUDA tensors
• probs must be 2D, got shape={tuple(probs.shape)}
• top_p must be scalar or have one value per row, got {top_ps.

AI-assisted analysis of sgl-project/sglang@0132848349 (2026-08-28). Data as JSON: /api/errors/926e17d3dfe158b7. Report an issue: GitHub.