[x-Pepsy]

I have a 2060 (sm_75) and tried to run Qwen3.6-35B-A3B on it. The CUDA backend doesn't get as far as loading the kernels there: NVRTC stops on cp.async, __reduce_max_sync/__reduce_min_sync and the m16n8k16/m16n8k32 shapes, all of which want sm_80.

This adds an sm_75 path. It's not the fastest thing on earth, but it runs:

• f16 m16n8k16 becomes two m16n8k8 mmas on the same accumulator. In the k16 layout a0/a1 and a2/a3 are the two k-halves of A, and b0/b1 the two k-halves of B, so the split is just two calls with the same register mapping.
• cp.async copies through registers instead. A predicated-off cp.async zero-fills its bytes, and the prefill attention relies on that for keys past the last one, so the fallback writes zeros too instead of leaving stale shared memory.
• The two _reduce*_sync calls in topk_warp are plain shuffles now.
• The int8 GEMM is off on Turing: there is no m16n8k16/m16n8k32 int8 mma there either. gemmi/gemmis aren't built and gemm_raw takes the f16 GEMM. I added Backend::int8_gemm so the kernel tests build the reference the backend really uses; Vulkan keeps what it had.
• sm_75 has 64 KB of shared memory per block. attn_prefill at head_dim 256 uses a 16-key tile instead of 32, and attn_decode_gqa at head_dim 256 skips its cp.async KV prefetch (it already has the register path for q8). The dynamic shared memory opt-in is now the device max minus the kernel's statics instead of a fixed 96 KB.

Numbers on this box (RTX 2060 6 GB, Ryzen 5 7600X, 31 GB RAM, UD-Q2_K_XL, ctx 8192, ~2.6 GB of experts resident), one run:

prefill 512 / 2048 / 4096:  236 / 1201 / 1443 tok/s
decode at the same depths:  66 / 63 / 61 tok/s

With --kv q8 --ctx 131072 a full 128k prefill takes 4m40s (467 tok/s) and decoding at that depth sits around 32 tok/s.

Tests: cargo test --release with STRATA_MODEL pointing at that model. Everything that runs on CUDA passes: backend, cpu_moe, gguf, gpu_model 10/10, kernels 20/20, model, ops, quant, serve, tokenizer. kv_q8 and llama_numerics only fail on their Vulkan half here, because the only Vulkan device on this laptop is an AMD iGPU that reports no int8 and gives NaN/different results; I skipped Vulkan locally to check the CUDA side.

What I couldn't check: this is the only GPU I have, so the sm_80 paths are tested by compiling them, not by running them. The suite's split test runs both halves on the same card and passes, but that's not the same as two real GPUs.

If Turing isn't something the project wants to carry, feel free to close this.

[Vasyl-Trefilov]

First: Nice username, I like it, and ty for contribute
Second:
This is a clean port, and the notes on the k16 split and the cp.async zero-fill made it easy to check.

What I ran on an RTX 3060 (sm_86), with this PR merged onto current main:

• Full suite passes.
• Output is bit-identical to main on CUDA and Vulkan: the perplexity logprobs for Qwen3-4B Q4_K_M, Qwen3-0.6B Q8_0 and Qwen3-30B-A3B UD-Q2_K_XL hash the same. (The 30B was compared at equal expert placement, because the resident count follows free VRAM.)
• Every kernel compiles for sm_75 through NVRTC and ptxas, with and without --kv q8, on current main too.

Two changes before merging:

1. Keep __reduce_max_sync / __reduce_min_sync on sm_80+. topk_warp runs on every MoE decode step, and the shuffles cost about 0.7% of 30B decode here (112.1 against 112.8 tok/s, three alternating runs each). With this guard, it is back to main's speed:

   #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 800
           unsigned mk = warp_max_u(bk);
           bi = warp_min_u(bk == mk ? bi : 0xffffffffu);
   #else
           unsigned mk = __reduce_max_sync(0xffffffff, bk);
           bi = __reduce_min_sync(0xffffffff, bk == mk ? bi : 0xffffffffu);
   #endif

2. Rebase onto main. src/cuda/mod.rs conflicts in four places. Main added an attn_decode_gqa4_128 kernel and a rope frequency table (fr). The resolution I tested:

   • keep the .chain(["attn_decode_gqa4_128".to_string()]) line next to your int8_mma filter;
   • add both int8_mma and fr to the constructor;
   • in attn_decode, use let cp = !self.kv_q8 && (hd == 128 || (hd == 256 && self.int8_mma));. gqa4_128 needs one 16 KB chunk, which fits in Turing's 64 KB.
   • src/vulkan/mod.rs only needs both new methods kept.

One interaction for whichever lands second: #9 adds prequant_gemm_x, and on CUDA its early return needs !self.int8_mma || as well, the same check as in gemm_raw. I can add that on merge.