Windows: zig build native fails on main (c9fab624): 21 CUDA kernels' explicit instantiations omit __restrict__, and GLM's setenv doesn't link #1210

On main (c9fab624), the native Windows build from docs/WINDOWS.md fails at two places. Both came in after #1010 was checked on 6ee828c2, and the Windows CI job doesn't compile kernels or link tensorfold-native, so neither shows there.

Host: x86_64 Windows 11, RTX 5090 (sm_120), driver 610.88, Zig 0.17.0, MSVC 19.44 (VS 2022 Build Tools), CUDA 13.0.2 (V13.0.88) and 13.4 (V13.4.59).

zig build native install -Dsm=120 -Dstrip=true -Dnvcc="%CUDA_PATH%\bin\nvcc.exe"

1. nvcc with cl.exe: "no instance of function template ... matches the specified type"

21 kernel files fail with 272 errors of that one kind, identically on nvcc 13.0 and 13.4:

• exl3_block/: dense3, exl3_experts, lin2, linear, pfdense_k8, pfdense_k10, pfdense_k12, pfdense_k16, x3gm, x3ld, x3ld_epi (through its includes), x3pf
• experts, experts_rows, flashnext_attn, flashnext_fact, flashnext_fexp, flashnext_hcf, prefill_attention, qmm, window_attention

Each failing explicit instantiation names a kernel whose definition marks pointer parameters __restrict__, while the instantiation's parameter list doesn't. With cl.exe as the host compiler, nvcc counts __restrict__ in the parameter types, so the instantiation matches no template; with GCC it doesn't. For example qmm.cu:228:

template __global__ void tf_qmm::qmm_kernel<32, 16, 64, 1, 4, 4, false, false, false>(
    const __nv_bfloat16*, const float*, const uint32_t*, const __nv_bfloat16*, const __nv_bfloat16*,
    void*, float*, int, int, int, int, int, int, int);

against a definition whose seven pointers are all __restrict__. Instantiations of kernels without __restrict__ in the same files compile (dense3's lanes_linear_kernel takes a LinTable by value and passes; its lanes_unpack_kernel fails). prefill_attention.cu at 6ee828c2 already had instantiations without __restrict__ and still compiles on 13.4, because its definitions had none then.

Writing each instantiation's pointers with the definition's __restrict__ (103 instantiations in 20 files, some inside #defines) makes all 21 files compile on 13.0 and 13.4.

2. lld-link: undefined symbol setenv in cuda_native.openRank

zig/src/families/glm/cuda_native.zig:106 declares extern "c" fn setenv, used by openRank for the three NCCL defaults. The Windows C runtime has no setenv, so tensorfold-native.exe doesn't link. The pattern heat.zig and flashnext/cuda_lanes.zig already use, @import("sys").c on Windows (win_clock.zig's _putenv_s wrapper) and the extern elsewhere, links.

With both changes

tensorfold-native.exe and tf-cuda-test.exe build for sm_120 with nvcc 13.0, and on the RTX 5090:

mlx-community/Qwen3.5-2B-MLX-4bit@93760be4 serves (--context 8192 --no-thinking, greedy, drafts on). For one prose and one code prompt, a repeated request and a "draft": false request return the same text as the drafted one.

Not done yet: zig build test test-kernel-contract test-golden fuzz-cortheon, lean_check.py against the merge base, kernel-contract on sm_120, and a Linux check that the changed kernels' fatbins are byte-identical before and after. I'd run those for the PR.

Notes

• With CUDA 13, cublasLt64_13.dll is in bin\x64, not bin. tf-cuda-test cublaslt reports CublasLtUnavailable until that directory is on PATH. docs/WINDOWS.md says no toolkit is needed to run; a line about the cuBLAS runtime for --vision and cuBLASLt projections could help.
• A Windows CI step that compiles the fatbins with nvcc (or only the kernels with explicit instantiations) would catch the first kind of break.

Open pull requests with the same patterns

• qwen CUDA: long-prompt prefill 5.4x faster (4096-row chunks, tensor-core attention, the delta-rule chain) #1080, qwen CUDA: decode attention on tensor cores (fp8 and bf16 caches), decode after long prompts 1.5-2.3x faster #1168 and qwen CUDA: --prompt-precision int8, prompt matmuls on int8 activations (cold prefill 20-27% shorter) #1176 add instantiations of tf_gdn_prefill::chain_kernel, tf_prefill_attention::hpattn_kernel, tf_prefill_attention::hpattn8_kernel and tf_qmm_prefill::prefill_kernel without the __restrict__ their definitions have, so they'd break the Windows build the same way when they land.
• glm/cuda: KV pages from the core pool, kept states sized by their position (#1142) #1163 and glm cuda: rank 0's NCCL id goes over the ranks' lines, so a /32 ring's followers bootstrap on NCCL_SOCKET_IFNAME #1178 add extern "c" fn unsetenv next to setenv in glm/cuda_native.zig; Windows has no unsetenv either (sys.c wraps it). Either would also conflict textually with a change to that declaration.

A rule for both might help: an explicit instantiation repeats its definition's __restrict__, and C-runtime environment calls go through sys.c.

I'd like to send the two fixes as one PR. Is adding __restrict__ to the instantiations the change you'd want, or would you rather the definitions changed?