candle-kernels 0.9.2 fails to build CUDA kernels for sm75; build script unconditionally compiles MoE WMMA kernels that instantiate BF16 WMMA fragments.

Hi everyone!

I am in the mid of translating python code to candle, and I'm facing issue with cuda capability (no problem on python side)

The downstream project does not use MoE or BF16, but enabling `candle-core/cuda` still pulls and builds these kernels, causing the whole CUDA build to fail.

An unreviewed AI-made patch is at the bottom. But it seems clunky and introduces CANDLE_BUILD_MOE_KERNELS. I think it would it be better with a feature for the crate? Happy to work on a PR for this if considered for inclusion

/J

## Environment

- GPU: NVIDIA Quadro RTX 5000
- Compute capability: 7.5 / `sm_75`
- Driver: 570.133.20
- CUDA toolkit tested: 12.8
- `CUDA_COMPUTE_CAP=75`
- Crate: `candle-kernels 0.9.2`

## Error

Build fails in `candle-kernels` while compiling `src/moe/moe_wmma.cu` and `src/moe/moe_wmma_gguf.cu`:

```text
src/moe/moe_wmma.cu(172): error: incomplete type
"nvcuda::wmma::fragment" is not allowed

src/moe/moe_wmma.cu(173): error: incomplete type
"nvcuda::wmma::fragment" is not allowed

The nvcc invocation targets sm_75:

nvcc --gpu-architecture=sm_75 ... src/moe/moe_wmma.cu

## Why this seems problematic

Turing / sm75 does not support BF16 WMMA fragments, so these kernels cannot compile for this target. However, they are built unconditionally as part of candle-kernels, even when downstream users only need normal F32 ops such as conv2d.

In my downstream use case, StarDist inference uses F32 tensors and does not need MoE or BF16 kernels.

## Local workaround

I patched candle-kernels/build.rs locally to avoid including the MoE WMMA files in the default PTX build and to skip building libmoe.a unless explicitly requested.

That allowed CUDA F32 inference to build and run correctly on sm75.

## Suggested fix

Please add a feature flag or environment variable to disable MoE/WMMA kernels during candle-kernels build, for example:

- Cargo feature: moe-kernels
- or env var: CANDLE_BUILD_MOE_KERNELS=1

Suggested behavior:

- Default build should avoid compiling MoE WMMA kernels unless they are needed, or
- Build script should skip BF16 WMMA kernels when CUDA_COMPUTE_CAP < 80, or
- MoE kernels should be behind an opt-in feature/env flag.

This would let downstream crates use Candle CUDA F32 operations on sm75/Turing GPUs without requiring unused BF16 WMMA support.

## Suggested patch

--- original build.rs
+++ patched build.rs
@@ -10,16 +10,36 @@
// Build for PTX
let out_dir = PathBuf::from(env::var("OUT_DIR").unwrap());
let ptx_path = out_dir.join("ptx.rs");
+ let kernel_paths = vec![
+ "src/affine.cu",
+ "src/binary.cu",
+ "src/cast.cu",
+ "src/conv.cu",
+ "src/fill.cu",
+ "src/indexing.cu",
+ "src/quantized.cu",
+ "src/reduce.cu",
+ "src/sort.cu",
+ "src/ternary.cu",
+ "src/unary.cu",
+ ];
+
let builder = bindgen_cuda::Builder::default()
.arg("--expt-relaxed-constexpr")
.arg("-std=c++17")
- .arg("-O3");
+ .arg("-O3")
+ .kernel_paths(kernel_paths);
let bindings = builder.build_ptx().unwrap();
bindings.write(&ptx_path).unwrap();

// Remove unwanted MOE PTX constants from ptx.rs
remove_lines(&ptx_path, &["MOE_GGUF", "MOE_WMMA", "MOE_WMMA_GGUF"]);

+ if env::var("CANDLE_BUILD_MOE_KERNELS").as_deref() != Ok("1") {
+ println!("cargo:rustc-link-lib=dylib=cudart");
+ return;
+ }
+
let mut moe_builder = bindgen_cuda::Builder::default()

Start in candle-kernels/build.rs and inspect how the PTX and MoE libraries are selected, then reproduce the build with CUDA_COMPUTE_CAP=75 using the reported sm_75 target. Done means the default CUDA build no longer compiles unsupported or unused MoE WMMA kernels while regular F32 operations still build successfully.

Written by the indexing model from the issue text.