{"record":{"id":"25387273c5ef752a","repo":"deepseek-ai/DeepSeek-V3","slug":"input-tensor-must-be-contiguous","errorCode":null,"errorMessage":"Input tensor must be contiguous","messagePattern":"Input tensor must be contiguous","errorType":"validation","errorClass":"AssertionError","httpStatus":null,"severity":"error","filePath":"inference/kernel.py","lineNumber":51,"sourceCode":"    y = y.to(y_ptr.dtype.element_ty)\n    tl.store(y_ptr + offs, y)\n    tl.store(s_ptr + pid, s)\n\n\ndef act_quant(x: torch.Tensor, block_size: int = 128, scale_fmt: Optional[str] = None) -> Tuple[torch.Tensor, torch.Tensor]:\n    \"\"\"\n    Quantizes the input tensor `x` using block-wise quantization.\n\n    Args:\n        x (torch.Tensor): The input tensor to be quantized. Must be contiguous and its last dimension size must be divisible by `block_size`.\n        block_size (int, optional): The size of the blocks to be used for quantization. Default is 128.\n        scale_fmt (Optional[str], optional): The format of the scale. Default is None.\n    Returns:\n        Tuple[torch.Tensor, torch.Tensor]: A tuple containing:\n            - The quantized tensor with dtype `torch.float8_e4m3fn`.\n            - A tensor of scaling factors with dtype `torch.float32`.\n    \"\"\"\n    assert x.is_contiguous(), 'Input tensor must be contiguous'\n    assert x.size(-1) % block_size == 0, f'Last dimension size must be divisible by block_size (block_size={block_size})'\n    y = torch.empty_like(x, dtype=torch.float8_e4m3fn)\n    s = x.new_empty(*x.size()[:-1], x.size(-1) // block_size, dtype=torch.float32)\n    grid = lambda meta: (triton.cdiv(x.numel(), meta['BLOCK_SIZE']), )\n    act_quant_kernel[grid](x, y, s, BLOCK_SIZE=block_size, scale_fmt=scale_fmt)\n    return y, s\n\n\n@triton.jit\ndef weight_dequant_kernel(x_ptr, s_ptr, y_ptr, M, N, BLOCK_SIZE: tl.constexpr):\n    \"\"\"\n    Dequantizes weights using the provided scaling factors and stores the result.\n\n    Args:\n        x_ptr (tl.pointer): Pointer to the quantized weights.\n        s_ptr (tl.pointer): Pointer to the scaling factors.\n        y_ptr (tl.pointer): Pointer to the output buffer for dequantized weights.\n        M (int): Number of rows in the weight matrix.","sourceCodeStart":33,"sourceCodeEnd":69,"githubUrl":"https://github.com/deepseek-ai/DeepSeek-V3/blob/9b4e9788e4a3a731f7567338ed15d3ec549ce03b/inference/kernel.py#L33-L69","documentation":"Thrown by act_quant (inference/kernel.py:51), the Triton block-wise FP8 quantization kernel launcher: Triton kernels index raw memory, so the input tensor x must be contiguous (no stride gaps from slicing/transposing) before its data pointer is handed to act_quant_kernel. Applies to activations being quantized to float8_e4m3fn ahead of fp8_gemm.","triggerScenarios":"Calling act_quant(x) where x came from a non-contiguous op result — e.g. x.transpose(-1,-2), a tensor sliced on the last dim (x[..., ::2]), or an attention output reshaped from a transposed view. Layout-holding ops like .transpose() return views that fail is_contiguous().","commonSituations":"Custom model code that quantizes a hidden_states view after permute/transpose; feeding a weight loaded with a stride-preserving narrow; versions of the model code where wkv outputs were reshaped before quant.","solutions":["Call .contiguous() on the tensor right before act_quant: y, s = act_quant(x.contiguous())","Better: restructure upstream so the producer already yields contiguous memory (avoid the extra copy)","If this fires inside the model's own forward, check for a mismatched repo version where a transpose was added/removed"],"exampleFix":"# before\nq, qs = act_quant(x.transpose(-1, -2))  # view, not contiguous\n\n# after\nt = x.transpose(-1, -2).contiguous()\nq, qs = act_quant(t)","handlingStrategy":"type-guard","validationCode":"if not x.is_contiguous():\n    x = x.contiguous()  # or raise, if the copy is unacceptable\ny, s = act_quant(x)","typeGuard":"def contiguous_or_bust(t: torch.Tensor) -> torch.Tensor:\n    return t if t.is_contiguous() else t.contiguous()","tryCatchPattern":null,"preventionTips":["Never pass transpose/slice views into Triton kernels — call .contiguous() first","Prefer reordering upstream ops so producers already emit contiguous tensors","Add is_contiguous() checks in unit tests for custom quantization paths"],"tags":["triton","fp8","quantization","tensor-layout","deepseek"],"backgroundTag":null,"analyzedSha":"9b4e9788e4a3a731f7567338ed15d3ec549ce03b","analyzedAt":"2026-08-14T19:02:32.748Z","schemaVersion":2},"datasetVersion":"2026-08-15T17:31:12.345Z"}