deepseek-ai/DeepSeek-V3 · error · AssertionError

Input tensor must be contiguous

Error message

Input tensor must be contiguous

What it means

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.

Source

Thrown at inference/kernel.py:51

    y = y.to(y_ptr.dtype.element_ty)
    tl.store(y_ptr + offs, y)
    tl.store(s_ptr + pid, s)


def act_quant(x: torch.Tensor, block_size: int = 128, scale_fmt: Optional[str] = None) -> Tuple[torch.Tensor, torch.Tensor]:
    """
    Quantizes the input tensor `x` using block-wise quantization.

    Args:
        x (torch.Tensor): The input tensor to be quantized. Must be contiguous and its last dimension size must be divisible by `block_size`.
        block_size (int, optional): The size of the blocks to be used for quantization. Default is 128.
        scale_fmt (Optional[str], optional): The format of the scale. Default is None.
    Returns:
        Tuple[torch.Tensor, torch.Tensor]: A tuple containing:
            - The quantized tensor with dtype `torch.float8_e4m3fn`.
            - A tensor of scaling factors with dtype `torch.float32`.
    """
    assert x.is_contiguous(), 'Input tensor must be contiguous'
    assert x.size(-1) % block_size == 0, f'Last dimension size must be divisible by block_size (block_size={block_size})'
    y = torch.empty_like(x, dtype=torch.float8_e4m3fn)
    s = x.new_empty(*x.size()[:-1], x.size(-1) // block_size, dtype=torch.float32)
    grid = lambda meta: (triton.cdiv(x.numel(), meta['BLOCK_SIZE']), )
    act_quant_kernel[grid](x, y, s, BLOCK_SIZE=block_size, scale_fmt=scale_fmt)
    return y, s


@triton.jit
def weight_dequant_kernel(x_ptr, s_ptr, y_ptr, M, N, BLOCK_SIZE: tl.constexpr):
    """
    Dequantizes weights using the provided scaling factors and stores the result.

    Args:
        x_ptr (tl.pointer): Pointer to the quantized weights.
        s_ptr (tl.pointer): Pointer to the scaling factors.
        y_ptr (tl.pointer): Pointer to the output buffer for dequantized weights.
        M (int): Number of rows in the weight matrix.

View on GitHub (pinned to 9b4e9788e4)

Solutions

1. Call .contiguous() on the tensor right before act_quant: y, s = act_quant(x.contiguous())
2. Better: restructure upstream so the producer already yields contiguous memory (avoid the extra copy)
3. If this fires inside the model's own forward, check for a mismatched repo version where a transpose was added/removed

Example fix

# before
q, qs = act_quant(x.transpose(-1, -2))  # view, not contiguous

# after
t = x.transpose(-1, -2).contiguous()
q, qs = act_quant(t)

if not x.is_contiguous():
    x = x.contiguous()  # or raise, if the copy is unacceptable
y, s = act_quant(x)

def contiguous_or_bust(t: torch.Tensor) -> torch.Tensor:
    return t if t.is_contiguous() else t.contiguous()

When it happens

Trigger: 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().

Common situations: 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.

Related errors

• Last dimension size must be divisible by block_size (block_s
• Input tensors must be contiguous
• Scaling factor tensors must be contiguous
• Input tensors must have 2 dimensions
• Warning: Missing scale_inv tensor for ${weight_name}, skippi

AI-assisted analysis of deepseek-ai/DeepSeek-V3@9b4e9788e4 (2026-08-14). Data as JSON: /api/errors/25387273c5ef752a. Report an issue: GitHub.