{"record":{"id":"926e17d3dfe158b7","repo":"sgl-project/sglang","slug":"renorm-kernels-require-a-cuda-hip-tensor","errorCode":null,"errorMessage":"renorm kernels require a CUDA/HIP tensor","messagePattern":"renorm kernels require a CUDA/HIP tensor","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/sampling/renorm_triton.py","lineNumber":58,"sourceCode":"    out_ptr,\n    row_sums_ptr,\n    numel,\n    vocab_size: tl.constexpr,\n    BLOCK_SIZE: tl.constexpr,\n):\n    offsets = tl.program_id(0).to(tl.int64) * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)\n    mask = offsets < numel\n    row = offsets // vocab_size\n    values = tl.load(out_ptr + offsets, mask=mask, other=0.0).to(tl.float32)\n    denominator = tl.load(row_sums_ptr + row, mask=mask, other=1.0)\n    tl.store(out_ptr + offsets, values / denominator, mask=mask)\n\n\ndef _prepare_probs(probs: torch.Tensor) -> torch.Tensor:\n    if probs.ndim != 2:\n        raise ValueError(f\"probs must be 2D, got shape={tuple(probs.shape)}\")\n    if not probs.is_cuda:\n        raise ValueError(\"renorm kernels require a CUDA/HIP tensor\")\n    return probs.float().contiguous()\n\n\ndef _renorm_from_pivots(probs_fp32: torch.Tensor, pivots: torch.Tensor) -> torch.Tensor:\n    batch_size, vocab_size = probs_fp32.shape\n    num_chunks = triton.cdiv(vocab_size, _BLOCK_SIZE)\n    out = torch.empty_like(probs_fp32)\n    partial_sums = torch.empty(\n        (batch_size, num_chunks), device=probs_fp32.device, dtype=torch.float32\n    )\n    _mask_and_partial_sum_kernel[(batch_size, num_chunks)](\n        probs_fp32,\n        pivots,\n        out,\n        partial_sums,\n        vocab_size=vocab_size,\n        num_chunks=num_chunks,\n        BLOCK_SIZE=_BLOCK_SIZE,","sourceCodeStart":40,"sourceCodeEnd":76,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/sampling/renorm_triton.py#L40-L76","documentation":"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.","triggerScenarios":"Calling top_p/top_k renorm with a CPU tensor, or a CUDA tensor that was moved with .cpu() for logging and then reused.","commonSituations":"Notebook experimentation on CPU before moving to GPU, or unit tests that run sampling logic without a GPU device.","solutions":["Move probs to the GPU: probs = probs.cuda() before calling","Use torch-native renormalization on CPU (sort + cumsum mask) when no GPU is available"],"exampleFix":"# before\nout = top_p_renorm_probs_triton(probs_cpu, 0.9)\n# after\nout = top_p_renorm_probs_triton(probs_cpu.cuda(), 0.9)","handlingStrategy":"validation","validationCode":"assert probs.is_cuda, 'renorm requires GPU tensors'\nif not probs.is_cuda: probs = probs.cuda()","typeGuard":"def on_gpu_for_renorm(p): return p.is_cuda","tryCatchPattern":null,"preventionTips":["Keep sampling tensors on device end-to-end","Implement a CPU sort+cumsum fallback for tests"],"tags":["sampling","top-p","cuda","device-placement"],"backgroundTag":"tensor-device-mismatch","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}