{"record":{"id":"d829efd1c04148f6","repo":"huggingface/candle","slug":"input-has-to-be-contiguous","errorCode":null,"errorMessage":"input has to be contiguous","messagePattern":"input has to be contiguous","errorType":"exception","errorClass":null,"httpStatus":null,"severity":"error","filePath":"candle-core/src/custom_op.rs","lineNumber":782,"sourceCode":"        let group_dims = candle_metal_kernels::utils::get_block_dims(b, 1, 1);\n        let encoder: &candle_metal_kernels::metal::ComputeCommandEncoder = encoder.as_ref();\n        encoder.set_output_buffer(0, Some(sto.buffer()), 0);\n        encoder.dispatch_threads(grid_dims, group_dims);\n\n        Ok(())\n    }\n\n    #[cfg(feature = \"cuda\")]\n    fn cuda_fwd(&self, sto: &mut CudaStorage, layout: &Layout) -> Result<()> {\n        use crate::cuda_backend::WrapErr;\n        use cudarc::driver::PushKernelArg;\n\n        let elem_count = layout.shape().elem_count();\n        let stream = sto.device.cuda_stream();\n        // TODO: support more dtypes.\n        let sto = sto.as_cuda_slice::<f32>()?;\n        let sto = match layout.contiguous_offsets() {\n            None => crate::bail!(\"input has to be contiguous\"),\n            Some((o1, o2)) => sto.slice(o1..o2),\n        };\n        let (g, b) = if elem_count % 32 == 0 {\n            (elem_count / 32, 32)\n        } else {\n            (elem_count, 1)\n        };\n        let cfg = cudarc::driver::LaunchConfig {\n            grid_dim: (g as u32, 1, 1),\n            block_dim: (b as u32, 1, 1),\n            shared_mem_bytes: 0,\n        };\n        let mut builder = stream.launch_builder(&self.func);\n        builder.arg(&sto);\n        unsafe { builder.launch(cfg) }.w()?;\n        Ok(())\n    }\n}","sourceCodeStart":764,"sourceCodeEnd":800,"githubUrl":"https://github.com/huggingface/candle/blob/d5fee525bfde3273eb7c9b75fd2bc4937be867ca/candle-core/src/custom_op.rs#L764-L800","documentation":"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.","triggerScenarios":"Applying a CudaCustomOp to a tensor produced by slicing, transpose, permute, narrow, or other strided views without calling .contiguous() first.","commonSituations":"Passing a transposed or sliced view into a custom CUDA kernel; chaining ops that leave the tensor non-contiguous on GPU.","solutions":["Call .contiguous()? on the tensor before applying the custom op.","Restructure the preceding ops (e.g. avoid transpose) so the tensor is naturally contiguous.","If you own the op, use strided indexing in the kernel to support non-contiguous inputs."],"exampleFix":"// before\nlet out = t.apply(&custom_op)?;\n// after\nlet out = t.contiguous()?.apply(&custom_op)?;","handlingStrategy":"validation","validationCode":"let t = if t.is_contiguous() { t } else { t.contiguous()? };","typeGuard":"fn is_contiguous_f32(t: &candle_core::Tensor) -> bool {\n    t.dtype() == candle_core::DType::F32 && t.is_contiguous()\n}","tryCatchPattern":null,"preventionTips":["Call .contiguous()? after slicing/transposing before custom CUDA ops.","Avoid interleaving views (narrow, transpose) immediately before custom ops.","Add an is_contiguous assertion in debug builds around custom op calls."],"tags":["cuda","contiguity","gpu","custom-op"],"backgroundTag":"tensor-not-contiguous","analyzedSha":"d5fee525bfde3273eb7c9b75fd2bc4937be867ca","analyzedAt":"2026-09-02T00:15:47.023Z","contentChangedAt":null,"schemaVersion":2},"datasetVersion":"2026-09-09T06:17:21.866Z"}