[Engine] strata-dsv41 does not build for Pascal (sm_61): nvcuda::wmma used without an sm_70 guard in ds4_gguf.cu / ds4_iq.cu #4173

Area

[Engine]

Desktop app, Terminal or source?

Source (a build)

What happened

Building strata-dsv41 for Pascal (-DCMAKE_CUDA_ARCHITECTURES=61 -DSTRATA_EXPERIMENTAL_SM60=ON -DSTRATA_BUILD_DSV41=ON) fails in the sm_61 device pass: CUDA's <mma.h> declares nvcuda::wmma only for __CUDA_ARCH__ >= 700, and src/dsv41/ds4_gguf.cu (mx_set_tc_kernel) and src/dsv41/ds4_iq.cu (set_tc_kernel) use it unguarded:

src/dsv41/ds4_gguf.cu(713): error: name must be a namespace name
src/dsv41/ds4_iq.cu(812): error: name must be a namespace name
src/dsv41/ds4_gguf.cu(742): error: name followed by "::" must be a class or namespace name
src/dsv41/ds4_gguf.cu(742): error: identifier "fg" is undefined
...

(line numbers from 7dd5cf1; the same code is on main at 4b9e05a.)

Both kernels are never launched below compute capability 7 (tc_ok() checks major >= 7 in both files), so they only need to compile. The Qwen engine had the same problem and fixed it in a6bdc3f (WMMA kernel bodies compiled only for sm_70+ device passes). Locally I added a stub header included after <mma.h> in both files, and the engine then builds and runs on 2x Tesla P40 (smoke test passes, output below):

// wmma_sm6x_stub.cuh: CUDA declares nvcuda::wmma only for sm_70+; tc_ok() never launches the tensor-core kernels
// below compute capability 7, so these stubs only let them compile (and trap if one ever runs)
#pragma once
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 700 && !defined(STRATA_USE_HIP)
namespace nvcuda { namespace wmma {
struct matrix_a {}; struct matrix_b {}; struct accumulator {}; struct row_major {}; struct col_major {};
enum layout_t { mem_row_major, mem_col_major };
template <typename Use, int M, int N, int K, typename T, typename L = void> struct fragment { T x[1]; };
template <typename F, typename V> __device__ inline void fill_fragment(F&, V) { __trap(); }
template <typename F, typename P> __device__ inline void load_matrix_sync(F&, const P*, unsigned) { __trap(); }
template <typename D, typename A, typename B, typename C> __device__ inline void mma_sync(D&, const A&, const B&, const C&) { __trap(); }
template <typename P, typename F> __device__ inline void store_matrix_sync(P*, const F&, unsigned, layout_t) { __trap(); }
} }
#endif

Guarding the two kernel bodies the way a6bdc3f does would be the cleaner fix.

Separately, for anyone trying the same: on a CPU without AVX2 the build needs -DSTRATA_ISA_FLOOR=avx instead of -DSTRATA_PORTABLE=ON (as in packaging/build_dsv41_cuda.sh); with PORTABLE the engine dies with SIGILL at setup's smoke test (exit -4).

Smallest steps or command that reproduce it

cmake -S . -B build -DCMAKE_BUILD_TYPE=Release -DSTRATA_ENABLE_CUDA=ON -DSTRATA_BUILD_TESTS=OFF -DSTRATA_BUILD_DSV41=ON \
  -DSTRATA_NATIVE_EXPERTS=ON -DSTRATA_ISA_FLOOR=avx -DSTRATA_EXPERIMENTAL_SM60=ON -DCMAKE_CUDA_ARCHITECTURES=61 \
  -DCMAKE_CUDA_COMPILER=/usr/local/cuda-12.3/bin/nvcc -DSTRATA_GGML_DIR=<llama.cpp at 3cf03257>
cmake --build build --target strata-dsv41

Reproduction status

Reproduced locally

Strata version

Source 7dd5cf1 (0.2.4.2), engine strata-dsv41 0.1.49.2; checked the same code on main 4b9e05a

Graphics card, memory and system

2x Tesla P40 24 GB (sm_61, PCIe 3.0 x8) · Xeon E5-2651 v2 (AVX, no AVX2) · 128 GB DDR3 · Ubuntu 24.04 · CUDA 12.3, g++-12

Model, size and settings

DeepSeek-V4.1-Flash IST-DAS-2.00bpw (ISTA-DASLab IQ2_XXS) · --gguf-dir setup · both GPUs

Logs

With the stub: ./setup.sh --family deepseek41 --gguf-dir ... --gpus 0,1 --cuda 12 packs 48,488 tensors and passes the test question ("The capital of France is Paris."). 96 greedy tokens of a coding answer decode correctly.