huggingface/candle · error

input has to be contiguous

Error message

input has to be contiguous

What it means

candle's cuda_fwd custom op requires a contiguous CUDA tensor. It calls layout.contiguous_offsets(); if the layout is strided/non-contiguous it returns None and the op bails, because the CUDA kernel indexes raw f32 memory assuming a dense layout.

Source

Thrown at candle-core/src/custom_op.rs:782

        let group_dims = candle_metal_kernels::utils::get_block_dims(b, 1, 1);
        let encoder: &candle_metal_kernels::metal::ComputeCommandEncoder = encoder.as_ref();
        encoder.set_output_buffer(0, Some(sto.buffer()), 0);
        encoder.dispatch_threads(grid_dims, group_dims);

        Ok(())
    }

    #[cfg(feature = "cuda")]
    fn cuda_fwd(&self, sto: &mut CudaStorage, layout: &Layout) -> Result<()> {
        use crate::cuda_backend::WrapErr;
        use cudarc::driver::PushKernelArg;

        let elem_count = layout.shape().elem_count();
        let stream = sto.device.cuda_stream();
        // TODO: support more dtypes.
        let sto = sto.as_cuda_slice::<f32>()?;
        let sto = match layout.contiguous_offsets() {
            None => crate::bail!("input has to be contiguous"),
            Some((o1, o2)) => sto.slice(o1..o2),
        };
        let (g, b) = if elem_count % 32 == 0 {
            (elem_count / 32, 32)
        } else {
            (elem_count, 1)
        };
        let cfg = cudarc::driver::LaunchConfig {
            grid_dim: (g as u32, 1, 1),
            block_dim: (b as u32, 1, 1),
            shared_mem_bytes: 0,
        };
        let mut builder = stream.launch_builder(&self.func);
        builder.arg(&sto);
        unsafe { builder.launch(cfg) }.w()?;
        Ok(())
    }
}

View on GitHub (pinned to d5fee525bf)

Solutions

1. Call .contiguous()? on the tensor before applying the custom op.
2. Restructure the preceding ops (e.g. avoid transpose) so the tensor is naturally contiguous.
3. If you own the op, use strided indexing in the kernel to support non-contiguous inputs.

Example fix

// before
let out = t.apply(&custom_op)?;
// after
let out = t.contiguous()?.apply(&custom_op)?;

let t = if t.is_contiguous() { t } else { t.contiguous()? };

fn is_contiguous_f32(t: &candle_core::Tensor) -> bool {
    t.dtype() == candle_core::DType::F32 && t.is_contiguous()
}

When it happens

Trigger: Applying a CudaCustomOp to a tensor produced by slicing, transpose, permute, narrow, or other strided views without calling .contiguous() first.

Common situations: Passing a transposed or sliced view into a custom CUDA kernel; chaining ops that leave the tensor non-contiguous on GPU.

Related errors

• input is not a f32 tensor
• quantized embedding requires contiguous ids
• seqlens_q has to be contiguous
• seqlens_k has to be contiguous
• seqlens_q has to be contiguous

AI-assisted analysis of huggingface/candle@d5fee525bf (2026-09-02). Data as JSON: /api/errors/d829efd1c04148f6. Report an issue: GitHub.