diff --git a/zig/kernels/cuda/exl3_block/dense3.cu b/zig/kernels/cuda/exl3_block/dense3.cu
index cd908255da..7afcbf4e45 100644
--- a/zig/kernels/cuda/exl3_block/dense3.cu
+++ b/zig/kernels/cuda/exl3_block/dense3.cu
@@ -379,24 +379,29 @@ template __global__ void dsv41_dense3::lanes_linear_kernel<4, 2, 4, 3>(const dsv
 template __global__ void dsv41_dense3::lanes_linear_kernel<4, 2, 4, 4>(const dsv41_x3seg::LinTable, int, int);
 template __global__ void dsv41_dense3::lanes_linear_kernel<4, 2, 8, 1>(const dsv41_x3seg::LinTable, int, int);
 template __global__ void dsv41_dense3::lanes_linear_kernel<4, 2, 8, 2>(const dsv41_x3seg::LinTable, int, int);
-template __global__ void dsv41_dense3::lanes_unpack_kernel<4, 2>(const uint32_t*, half*, int, int64_t, int64_t);
+template __global__ void dsv41_dense3::lanes_unpack_kernel<4, 2>(const uint32_t* __restrict__, half* __restrict__, int,
+    int64_t, int64_t);
 template __global__ void dsv41_dense3::lanes_linear_kernel<6, 2, 4, 3>(const dsv41_x3seg::LinTable, int, int);
 template __global__ void dsv41_dense3::lanes_linear_kernel<6, 2, 4, 4>(const dsv41_x3seg::LinTable, int, int);
 template __global__ void dsv41_dense3::lanes_linear_kernel<6, 2, 8, 1>(const dsv41_x3seg::LinTable, int, int);
 template __global__ void dsv41_dense3::lanes_linear_kernel<6, 2, 8, 2>(const dsv41_x3seg::LinTable, int, int);
-template __global__ void dsv41_dense3::lanes_unpack_kernel<6, 2>(const uint32_t*, half*, int, int64_t, int64_t);
+template __global__ void dsv41_dense3::lanes_unpack_kernel<6, 2>(const uint32_t* __restrict__, half* __restrict__, int,
+    int64_t, int64_t);
 template __global__ void dsv41_dense3::lanes_linear_kernel<8, 2, 4, 3>(const dsv41_x3seg::LinTable, int, int);
 template __global__ void dsv41_dense3::lanes_linear_kernel<8, 2, 4, 4>(const dsv41_x3seg::LinTable, int, int);
 template __global__ void dsv41_dense3::lanes_linear_kernel<8, 2, 8, 1>(const dsv41_x3seg::LinTable, int, int);
 template __global__ void dsv41_dense3::lanes_linear_kernel<8, 2, 8, 2>(const dsv41_x3seg::LinTable, int, int);
-template __global__ void dsv41_dense3::lanes_unpack_kernel<8, 2>(const uint32_t*, half*, int, int64_t, int64_t);
+template __global__ void dsv41_dense3::lanes_unpack_kernel<8, 2>(const uint32_t* __restrict__, half* __restrict__, int,
+    int64_t, int64_t);
 template __global__ void dsv41_dense3::lanes_linear_kernel<10, 2, 4, 3>(const dsv41_x3seg::LinTable, int, int);
 template __global__ void dsv41_dense3::lanes_linear_kernel<10, 2, 4, 4>(const dsv41_x3seg::LinTable, int, int);
 template __global__ void dsv41_dense3::lanes_linear_kernel<10, 2, 8, 1>(const dsv41_x3seg::LinTable, int, int);
 template __global__ void dsv41_dense3::lanes_linear_kernel<10, 2, 8, 2>(const dsv41_x3seg::LinTable, int, int);
-template __global__ void dsv41_dense3::lanes_unpack_kernel<10, 2>(const uint32_t*, half*, int, int64_t, int64_t);
+template __global__ void dsv41_dense3::lanes_unpack_kernel<10, 2>(const uint32_t* __restrict__, half* __restrict__, int,
+    int64_t, int64_t);
 template __global__ void dsv41_dense3::lanes_linear_kernel<12, 2, 4, 3>(const dsv41_x3seg::LinTable, int, int);
 template __global__ void dsv41_dense3::lanes_linear_kernel<12, 2, 4, 4>(const dsv41_x3seg::LinTable, int, int);
 template __global__ void dsv41_dense3::lanes_linear_kernel<12, 2, 8, 1>(const dsv41_x3seg::LinTable, int, int);
 template __global__ void dsv41_dense3::lanes_linear_kernel<12, 2, 8, 2>(const dsv41_x3seg::LinTable, int, int);
-template __global__ void dsv41_dense3::lanes_unpack_kernel<12, 2>(const uint32_t*, half*, int, int64_t, int64_t);
+template __global__ void dsv41_dense3::lanes_unpack_kernel<12, 2>(const uint32_t* __restrict__, half* __restrict__, int,
+    int64_t, int64_t);
diff --git a/zig/kernels/cuda/exl3_block/exl3_experts.cu b/zig/kernels/cuda/exl3_block/exl3_experts.cu
index a0959e216b..d87056848d 100644
--- a/zig/kernels/cuda/exl3_block/exl3_experts.cu
+++ b/zig/kernels/cuda/exl3_block/exl3_experts.cu
@@ -245,8 +245,9 @@ __global__ void down_combine_kernel(const float* __restrict__ Z, const int* __re
 
 // upstream grouped_kernel at mul1 for every (nt, warps, pf) x TF_RANGES, and dequant_kernel<2, K2> (tests)
 #define TF_EXL3X_INST(NT, W, PF, LO, HI) template __global__ void tf_exl3x::grouped_kernel<2, NT, W, PF, LO, HI>( \
-    const half*, const half*, const int64_t*, const int64_t*, const int*, const int*, const int*, const int*, \
-    const int*, float*, int, int, int, int, int, int);
+    const half* __restrict__, const half* __restrict__, const int64_t* __restrict__, const int64_t* __restrict__, \
+    const int* __restrict__, const int* __restrict__, const int* __restrict__, const int* __restrict__, \
+    const int* __restrict__, float* __restrict__, int, int, int, int, int, int);
 TF_EXL3X_INST(8, 4, 1, 8, 8)
 TF_EXL3X_INST(8, 4, 1, 2, 10)
 TF_EXL3X_INST(8, 4, 1, 2, 16)
@@ -256,28 +257,31 @@ TF_EXL3X_INST(8, 4, 2, 2, 16)
 TF_EXL3X_INST(4, 4, 2, 8, 8)
 TF_EXL3X_INST(4, 4, 2, 2, 10)
 TF_EXL3X_INST(4, 4, 2, 2, 16)
-template __global__ void tf_exl3x::dequant_kernel<2, 2>(const uint32_t*, half*, int, int);
-template __global__ void tf_exl3x::dequant_kernel<2, 3>(const uint32_t*, half*, int, int);
-template __global__ void tf_exl3x::dequant_kernel<2, 4>(const uint32_t*, half*, int, int);
-template __global__ void tf_exl3x::dequant_kernel<2, 5>(const uint32_t*, half*, int, int);
-template __global__ void tf_exl3x::dequant_kernel<2, 6>(const uint32_t*, half*, int, int);
-template __global__ void tf_exl3x::dequant_kernel<2, 7>(const uint32_t*, half*, int, int);
-template __global__ void tf_exl3x::dequant_kernel<2, 8>(const uint32_t*, half*, int, int);
-template __global__ void tf_exl3x::dequant_kernel<2, 9>(const uint32_t*, half*, int, int);
-template __global__ void tf_exl3x::dequant_kernel<2, 10>(const uint32_t*, half*, int, int);
-template __global__ void tf_exl3x::dequant_kernel<2, 11>(const uint32_t*, half*, int, int);
-template __global__ void tf_exl3x::dequant_kernel<2, 12>(const uint32_t*, half*, int, int);
-template __global__ void tf_exl3x::dequant_kernel<2, 13>(const uint32_t*, half*, int, int);
-template __global__ void tf_exl3x::dequant_kernel<2, 14>(const uint32_t*, half*, int, int);
-template __global__ void tf_exl3x::dequant_kernel<2, 15>(const uint32_t*, half*, int, int);
-template __global__ void tf_exl3x::dequant_kernel<2, 16>(const uint32_t*, half*, int, int);
+template __global__ void tf_exl3x::dequant_kernel<2, 2>(const uint32_t* __restrict__, half* __restrict__, int, int);
+template __global__ void tf_exl3x::dequant_kernel<2, 3>(const uint32_t* __restrict__, half* __restrict__, int, int);
+template __global__ void tf_exl3x::dequant_kernel<2, 4>(const uint32_t* __restrict__, half* __restrict__, int, int);
+template __global__ void tf_exl3x::dequant_kernel<2, 5>(const uint32_t* __restrict__, half* __restrict__, int, int);
+template __global__ void tf_exl3x::dequant_kernel<2, 6>(const uint32_t* __restrict__, half* __restrict__, int, int);
+template __global__ void tf_exl3x::dequant_kernel<2, 7>(const uint32_t* __restrict__, half* __restrict__, int, int);
+template __global__ void tf_exl3x::dequant_kernel<2, 8>(const uint32_t* __restrict__, half* __restrict__, int, int);
+template __global__ void tf_exl3x::dequant_kernel<2, 9>(const uint32_t* __restrict__, half* __restrict__, int, int);
+template __global__ void tf_exl3x::dequant_kernel<2, 10>(const uint32_t* __restrict__, half* __restrict__, int, int);
+template __global__ void tf_exl3x::dequant_kernel<2, 11>(const uint32_t* __restrict__, half* __restrict__, int, int);
+template __global__ void tf_exl3x::dequant_kernel<2, 12>(const uint32_t* __restrict__, half* __restrict__, int, int);
+template __global__ void tf_exl3x::dequant_kernel<2, 13>(const uint32_t* __restrict__, half* __restrict__, int, int);
+template __global__ void tf_exl3x::dequant_kernel<2, 14>(const uint32_t* __restrict__, half* __restrict__, int, int);
+template __global__ void tf_exl3x::dequant_kernel<2, 15>(const uint32_t* __restrict__, half* __restrict__, int, int);
+template __global__ void tf_exl3x::dequant_kernel<2, 16>(const uint32_t* __restrict__, half* __restrict__, int, int);
 // down_combine_kernel at both outputs the binding takes (fp32, bf16: TF_DSV41_MOE_BF16)
-template __global__ void tf_exl3_experts::down_combine_kernel<float>(const float*, const int*, const half*, float*, \
-    const float*, float*, int, int, int, int, int);
-template __global__ void tf_exl3_experts::down_combine_kernel<__nv_bfloat16>(const float*, const int*, const half*, float*, \
-    const float*, __nv_bfloat16*, int, int, int, int, int);
+template __global__ void tf_exl3_experts::down_combine_kernel<float>(const float* __restrict__, \
+    const int* __restrict__, const half* __restrict__, float* __restrict__, const float* __restrict__, \
+    float* __restrict__, int, int, int, int, int);
+template __global__ void tf_exl3_experts::down_combine_kernel<__nv_bfloat16>(const float* __restrict__, \
+    const int* __restrict__, const half* __restrict__, float* __restrict__, const float* __restrict__, \
+    __nv_bfloat16* __restrict__, int, int, int, int, int);
 // rot_in_kernel for both inputs the binding takes
-template __global__ void tf_exl3_experts::rot_in_kernel<__nv_bfloat16>(const __nv_bfloat16*, int, const int*, const half*, const half*, half*, \
-    half*, int, int, int);
-template __global__ void tf_exl3_experts::rot_in_kernel<half>(const half*, int, const int*, const half*, const half*, half*, \
-    half*, int, int, int);
+template __global__ void tf_exl3_experts::rot_in_kernel<__nv_bfloat16>(const __nv_bfloat16* __restrict__, int, \
+    const int* __restrict__, const half* __restrict__, const half* __restrict__, half* __restrict__, \
+    half* __restrict__, int, int, int);
+template __global__ void tf_exl3_experts::rot_in_kernel<half>(const half* __restrict__, int, const int* __restrict__, \
+    const half* __restrict__, const half* __restrict__, half* __restrict__, half* __restrict__, int, int, int);
diff --git a/zig/kernels/cuda/exl3_block/lin2.cu b/zig/kernels/cuda/exl3_block/lin2.cu
index 1e41a2309c..34d3c4f9ee 100644
--- a/zig/kernels/cuda/exl3_block/lin2.cu
+++ b/zig/kernels/cuda/exl3_block/lin2.cu
@@ -467,10 +467,12 @@ __global__ void __launch_bounds__(WK * 32, MINB) lanesc_kernel(const dsv41_x3seg
     template __global__ void dsv41_lin2::lanes2_kernel<K2, 2, 8, 1>(const dsv41_x3seg::LinTable, int, int);
 LIN2_LANES(4) LIN2_LANES(6) LIN2_LANES(8) LIN2_LANES(10) LIN2_LANES(12)
 #define LIN2_LINEAR(K2, WK)                                                                                         \
-    template __global__ void dsv41_lin2::linear2_kernel<K2, 2, WK, true>(const half*, const uint32_t*, long long, \
-        long long, const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);          \
-    template __global__ void dsv41_lin2::linear2_kernel<K2, 2, WK, false>(const half*, const uint32_t*, long long, \
-        long long, const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
+    template __global__ void dsv41_lin2::linear2_kernel<K2, 2, WK, true>(const half* __restrict__, \
+        const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+        void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__); \
+    template __global__ void dsv41_lin2::linear2_kernel<K2, 2, WK, false>(const half* __restrict__, \
+        const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+        void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
 #define LIN2_LINEAR_W(K2) LIN2_LINEAR(K2, 2) LIN2_LINEAR(K2, 4) LIN2_LINEAR(K2, 8)
 LIN2_LINEAR_W(2) LIN2_LINEAR_W(3) LIN2_LINEAR_W(4) LIN2_LINEAR_W(5) LIN2_LINEAR_W(6) LIN2_LINEAR_W(7)
 LIN2_LINEAR_W(8) LIN2_LINEAR_W(10) LIN2_LINEAR_W(12) LIN2_LINEAR_W(14) LIN2_LINEAR_W(16)
diff --git a/zig/kernels/cuda/exl3_block/linear.cu b/zig/kernels/cuda/exl3_block/linear.cu
index 6ebb689d18..03702147a2 100644
--- a/zig/kernels/cuda/exl3_block/linear.cu
+++ b/zig/kernels/cuda/exl3_block/linear.cu
@@ -284,80 +284,124 @@ __global__ void __launch_bounds__(32) unpack_kernel(const uint32_t* __restrict__
 }  // namespace tf_exl3_linear
 
 // Launched at mul1 (TF_EXL3_WIDTHS(X, 2) + K2 3, 5, 7): linear_kernel<K2, 2, WK> at WK 2 / 4 / 8, unpack_kernel<K2, 2>.
-template __global__ void tf_exl3_linear::linear_kernel<2, 2, 2>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::linear_kernel<2, 2, 4>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::linear_kernel<2, 2, 8>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::unpack_kernel<2, 2>(const uint32_t*, half*, int, int64_t, int64_t);
-template __global__ void tf_exl3_linear::linear_kernel<3, 2, 2>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::linear_kernel<3, 2, 4>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::linear_kernel<3, 2, 8>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::unpack_kernel<3, 2>(const uint32_t*, half*, int, int64_t, int64_t);
-template __global__ void tf_exl3_linear::linear_kernel<4, 2, 2>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::linear_kernel<4, 2, 4>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::linear_kernel<4, 2, 8>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::unpack_kernel<4, 2>(const uint32_t*, half*, int, int64_t, int64_t);
-template __global__ void tf_exl3_linear::linear_kernel<5, 2, 2>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::linear_kernel<5, 2, 4>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::linear_kernel<5, 2, 8>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::unpack_kernel<5, 2>(const uint32_t*, half*, int, int64_t, int64_t);
-template __global__ void tf_exl3_linear::linear_kernel<6, 2, 2>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::linear_kernel<6, 2, 4>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::linear_kernel<6, 2, 8>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::unpack_kernel<6, 2>(const uint32_t*, half*, int, int64_t, int64_t);
-template __global__ void tf_exl3_linear::linear_kernel<7, 2, 2>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::linear_kernel<7, 2, 4>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::linear_kernel<7, 2, 8>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::unpack_kernel<7, 2>(const uint32_t*, half*, int, int64_t, int64_t);
-template __global__ void tf_exl3_linear::linear_kernel<8, 2, 2>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::linear_kernel<8, 2, 4>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::linear_kernel<8, 2, 8>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::unpack_kernel<8, 2>(const uint32_t*, half*, int, int64_t, int64_t);
-template __global__ void tf_exl3_linear::linear_kernel<10, 2, 2>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::linear_kernel<10, 2, 4>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::linear_kernel<10, 2, 8>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::unpack_kernel<10, 2>(const uint32_t*, half*, int, int64_t, int64_t);
-template __global__ void tf_exl3_linear::linear_kernel<12, 2, 2>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::linear_kernel<12, 2, 4>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::linear_kernel<12, 2, 8>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::unpack_kernel<12, 2>(const uint32_t*, half*, int, int64_t, int64_t);
-template __global__ void tf_exl3_linear::linear_kernel<14, 2, 2>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::linear_kernel<14, 2, 4>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::linear_kernel<14, 2, 8>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::unpack_kernel<14, 2>(const uint32_t*, half*, int, int64_t, int64_t);
-template __global__ void tf_exl3_linear::linear_kernel<16, 2, 2>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::linear_kernel<16, 2, 4>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::linear_kernel<16, 2, 8>(const half*, const uint32_t*, long long, long long, \
-    const half*, const half*, void*, int, float*, int*, int, int, int, int, const int*);
-template __global__ void tf_exl3_linear::unpack_kernel<16, 2>(const uint32_t*, half*, int, int64_t, int64_t);
+template __global__ void tf_exl3_linear::linear_kernel<2, 2, 2>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::linear_kernel<2, 2, 4>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::linear_kernel<2, 2, 8>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::unpack_kernel<2, 2>(const uint32_t* __restrict__, half* __restrict__, int,
+    int64_t, int64_t);
+template __global__ void tf_exl3_linear::linear_kernel<3, 2, 2>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::linear_kernel<3, 2, 4>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::linear_kernel<3, 2, 8>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::unpack_kernel<3, 2>(const uint32_t* __restrict__, half* __restrict__, int,
+    int64_t, int64_t);
+template __global__ void tf_exl3_linear::linear_kernel<4, 2, 2>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::linear_kernel<4, 2, 4>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::linear_kernel<4, 2, 8>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::unpack_kernel<4, 2>(const uint32_t* __restrict__, half* __restrict__, int,
+    int64_t, int64_t);
+template __global__ void tf_exl3_linear::linear_kernel<5, 2, 2>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::linear_kernel<5, 2, 4>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::linear_kernel<5, 2, 8>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::unpack_kernel<5, 2>(const uint32_t* __restrict__, half* __restrict__, int,
+    int64_t, int64_t);
+template __global__ void tf_exl3_linear::linear_kernel<6, 2, 2>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::linear_kernel<6, 2, 4>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::linear_kernel<6, 2, 8>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::unpack_kernel<6, 2>(const uint32_t* __restrict__, half* __restrict__, int,
+    int64_t, int64_t);
+template __global__ void tf_exl3_linear::linear_kernel<7, 2, 2>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::linear_kernel<7, 2, 4>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::linear_kernel<7, 2, 8>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::unpack_kernel<7, 2>(const uint32_t* __restrict__, half* __restrict__, int,
+    int64_t, int64_t);
+template __global__ void tf_exl3_linear::linear_kernel<8, 2, 2>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::linear_kernel<8, 2, 4>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::linear_kernel<8, 2, 8>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::unpack_kernel<8, 2>(const uint32_t* __restrict__, half* __restrict__, int,
+    int64_t, int64_t);
+template __global__ void tf_exl3_linear::linear_kernel<10, 2, 2>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::linear_kernel<10, 2, 4>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::linear_kernel<10, 2, 8>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::unpack_kernel<10, 2>(const uint32_t* __restrict__, half* __restrict__, int,
+    int64_t, int64_t);
+template __global__ void tf_exl3_linear::linear_kernel<12, 2, 2>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::linear_kernel<12, 2, 4>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::linear_kernel<12, 2, 8>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::unpack_kernel<12, 2>(const uint32_t* __restrict__, half* __restrict__, int,
+    int64_t, int64_t);
+template __global__ void tf_exl3_linear::linear_kernel<14, 2, 2>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::linear_kernel<14, 2, 4>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::linear_kernel<14, 2, 8>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::unpack_kernel<14, 2>(const uint32_t* __restrict__, half* __restrict__, int,
+    int64_t, int64_t);
+template __global__ void tf_exl3_linear::linear_kernel<16, 2, 2>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::linear_kernel<16, 2, 4>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::linear_kernel<16, 2, 8>(const half* __restrict__, \
+    const uint32_t* __restrict__, long long, long long, const half* __restrict__, const half* __restrict__, \
+    void* __restrict__, int, float* __restrict__, int* __restrict__, int, int, int, int, const int* __restrict__);
+template __global__ void tf_exl3_linear::unpack_kernel<16, 2>(const uint32_t* __restrict__, half* __restrict__, int,
+    int64_t, int64_t);
diff --git a/zig/kernels/cuda/exl3_block/pfdense_k10.cu b/zig/kernels/cuda/exl3_block/pfdense_k10.cu
index f343cde298..2c3ccf43ab 100644
--- a/zig/kernels/cuda/exl3_block/pfdense_k10.cu
+++ b/zig/kernels/cuda/exl3_block/pfdense_k10.cu
@@ -8,10 +8,12 @@
 #define DSV41_PFD_K2 10
 #define DSV41_PFD_LN false
 DSV41_PFD_CFGS(DSV41_PFD_INST)
-template __global__ void dsv41_pfd::pfd_dequant_kernel<10, false>(const uint32_t*, long long, long long, half*, int);
+template __global__ void dsv41_pfd::pfd_dequant_kernel<10, false>(const uint32_t* __restrict__, long long, long long,
+    half* __restrict__, int);
 #undef DSV41_PFD_LN
 #define DSV41_PFD_LN true
 DSV41_PFD_CFGS(DSV41_PFD_INST)
-template __global__ void dsv41_pfd::pfd_dequant_kernel<10, true>(const uint32_t*, long long, long long, half*, int);
+template __global__ void dsv41_pfd::pfd_dequant_kernel<10, true>(const uint32_t* __restrict__, long long, long long,
+    half* __restrict__, int);
 
 DSV41_PFD_WIDTH(10)
diff --git a/zig/kernels/cuda/exl3_block/pfdense_k12.cu b/zig/kernels/cuda/exl3_block/pfdense_k12.cu
index f429844f73..e77be85c6d 100644
--- a/zig/kernels/cuda/exl3_block/pfdense_k12.cu
+++ b/zig/kernels/cuda/exl3_block/pfdense_k12.cu
@@ -8,10 +8,12 @@
 #define DSV41_PFD_K2 12
 #define DSV41_PFD_LN false
 DSV41_PFD_CFGS(DSV41_PFD_INST)
-template __global__ void dsv41_pfd::pfd_dequant_kernel<12, false>(const uint32_t*, long long, long long, half*, int);
+template __global__ void dsv41_pfd::pfd_dequant_kernel<12, false>(const uint32_t* __restrict__, long long, long long,
+    half* __restrict__, int);
 #undef DSV41_PFD_LN
 #define DSV41_PFD_LN true
 DSV41_PFD_CFGS(DSV41_PFD_INST)
-template __global__ void dsv41_pfd::pfd_dequant_kernel<12, true>(const uint32_t*, long long, long long, half*, int);
+template __global__ void dsv41_pfd::pfd_dequant_kernel<12, true>(const uint32_t* __restrict__, long long, long long,
+    half* __restrict__, int);
 
 DSV41_PFD_WIDTH(12)
diff --git a/zig/kernels/cuda/exl3_block/pfdense_k16.cu b/zig/kernels/cuda/exl3_block/pfdense_k16.cu
index da8c9c588a..5613b0c562 100644
--- a/zig/kernels/cuda/exl3_block/pfdense_k16.cu
+++ b/zig/kernels/cuda/exl3_block/pfdense_k16.cu
@@ -8,7 +8,8 @@
 #define DSV41_PFD_K2 16
 #define DSV41_PFD_LN false
 DSV41_PFD_CFGS(DSV41_PFD_INST)
-template __global__ void dsv41_pfd::pfd_dequant_kernel<16, false>(const uint32_t*, long long, long long, half*, int);
+template __global__ void dsv41_pfd::pfd_dequant_kernel<16, false>(const uint32_t* __restrict__, long long, long long,
+    half* __restrict__, int);
 
 
 DSV41_PFD_WIDTH_STRIPS(16)
diff --git a/zig/kernels/cuda/exl3_block/pfdense_k8.cu b/zig/kernels/cuda/exl3_block/pfdense_k8.cu
index 1c8bbdbe82..37756c6607 100644
--- a/zig/kernels/cuda/exl3_block/pfdense_k8.cu
+++ b/zig/kernels/cuda/exl3_block/pfdense_k8.cu
@@ -8,10 +8,12 @@
 #define DSV41_PFD_K2 8
 #define DSV41_PFD_LN false
 DSV41_PFD_CFGS(DSV41_PFD_INST)
-template __global__ void dsv41_pfd::pfd_dequant_kernel<8, false>(const uint32_t*, long long, long long, half*, int);
+template __global__ void dsv41_pfd::pfd_dequant_kernel<8, false>(const uint32_t* __restrict__, long long, long long,
+    half* __restrict__, int);
 #undef DSV41_PFD_LN
 #define DSV41_PFD_LN true
 DSV41_PFD_CFGS(DSV41_PFD_INST)
-template __global__ void dsv41_pfd::pfd_dequant_kernel<8, true>(const uint32_t*, long long, long long, half*, int);
+template __global__ void dsv41_pfd::pfd_dequant_kernel<8, true>(const uint32_t* __restrict__, long long, long long,
+    half* __restrict__, int);
 
 DSV41_PFD_WIDTH(8)
diff --git a/zig/kernels/cuda/exl3_block/x3gm.cu b/zig/kernels/cuda/exl3_block/x3gm.cu
index b03465d641..84fb19d58f 100644
--- a/zig/kernels/cuda/exl3_block/x3gm.cu
+++ b/zig/kernels/cuda/exl3_block/x3gm.cu
@@ -461,17 +461,18 @@ DSV41_X3GM_DN(10, GM_DN1)
 DSV41_X3GM_DN(10, GM_DN2)
 DSV41_X3GM_DN(10, GM_DN3)
 DSV41_X3GM_DN(10, GM_DN4)
-#define DSV41_X3GM_ROT(T, M) template __global__ void dsv41_x3gm::rot_kernel<T, M>(const T*, int, const int*, const half*, \
-    const half*, half*, half*, int, int, int);
+#define DSV41_X3GM_ROT(T, M) template __global__ void dsv41_x3gm::rot_kernel<T, M>(const T* __restrict__, int, \
+    const int* __restrict__, const half* __restrict__, const half* __restrict__, half* __restrict__, \
+    half* __restrict__, int, int, int);
 DSV41_X3GM_ROT(__nv_bfloat16, 1)
 DSV41_X3GM_ROT(__nv_bfloat16, 2)
 DSV41_X3GM_ROT(half, 1)
 DSV41_X3GM_ROT(half, 2)
-template __global__ void dsv41_x3gm::dequant_kernel<3>(const uint32_t*, half*, int);
-template __global__ void dsv41_x3gm::dequant_kernel<4>(const uint32_t*, half*, int);
-template __global__ void dsv41_x3gm::dequant_kernel<5>(const uint32_t*, half*, int);
-template __global__ void dsv41_x3gm::dequant_kernel<6>(const uint32_t*, half*, int);
-template __global__ void dsv41_x3gm::dequant_kernel<7>(const uint32_t*, half*, int);
-template __global__ void dsv41_x3gm::dequant_kernel<8>(const uint32_t*, half*, int);
-template __global__ void dsv41_x3gm::dequant_kernel<10>(const uint32_t*, half*, int);
+template __global__ void dsv41_x3gm::dequant_kernel<3>(const uint32_t* __restrict__, half* __restrict__, int);
+template __global__ void dsv41_x3gm::dequant_kernel<4>(const uint32_t* __restrict__, half* __restrict__, int);
+template __global__ void dsv41_x3gm::dequant_kernel<5>(const uint32_t* __restrict__, half* __restrict__, int);
+template __global__ void dsv41_x3gm::dequant_kernel<6>(const uint32_t* __restrict__, half* __restrict__, int);
+template __global__ void dsv41_x3gm::dequant_kernel<7>(const uint32_t* __restrict__, half* __restrict__, int);
+template __global__ void dsv41_x3gm::dequant_kernel<8>(const uint32_t* __restrict__, half* __restrict__, int);
+template __global__ void dsv41_x3gm::dequant_kernel<10>(const uint32_t* __restrict__, half* __restrict__, int);
 #endif  // DSV41_X3GM_NO_INSTANCES
diff --git a/zig/kernels/cuda/exl3_block/x3ld.cu b/zig/kernels/cuda/exl3_block/x3ld.cu
index 47b1fdf948..f84159dc1e 100644
--- a/zig/kernels/cuda/exl3_block/x3ld.cu
+++ b/zig/kernels/cuda/exl3_block/x3ld.cu
@@ -249,8 +249,9 @@ __global__ void __launch_bounds__(W * 32, 3) ld_kernel<2, 8, 1, 8, 8, 0>(
 
 // The instantiations dsv41_x3ld_grouped_cuda dispatches: mul1, (nt, pd) x probe x the (lo, hi) ranges.
 #define DSV41_X3LD_INST(NT, PD, LO, HI, PROBE) template __global__ void dsv41_x3ld::ld_kernel<2, NT, PD, LO, HI, PROBE>( \
-    const half*, const half*, const int64_t*, const int64_t*, const int*, const int*, const int*, const int*, \
-    const int*, float*, int, int, int, int, int, int);
+    const half* __restrict__, const half* __restrict__, const int64_t* __restrict__, const int64_t* __restrict__, \
+    const int* __restrict__, const int* __restrict__, const int* __restrict__, const int* __restrict__, \
+    const int* __restrict__, float* __restrict__, int, int, int, int, int, int);
 DSV41_X3LD_INST(8, 1, 2, 10, 0)
 DSV41_X3LD_INST(8, 1, 2, 12, 0)
 DSV41_X3LD_INST(8, 2, 8, 8, 0)
diff --git a/zig/kernels/cuda/exl3_block/x3pf.cu b/zig/kernels/cuda/exl3_block/x3pf.cu
index 56fa065305..2f8f011c46 100644
--- a/zig/kernels/cuda/exl3_block/x3pf.cu
+++ b/zig/kernels/cuda/exl3_block/x3pf.cu
@@ -187,8 +187,9 @@ __global__ void __launch_bounds__(W * 32) pf_kernel(
 
 // The instantiations dsv41_x3pf_grouped_cuda dispatches: mul1, (nt, mtl) x the (lo, hi) ranges.
 #define DSV41_X3PF_INST(NT, MTL, LO, HI) template __global__ void dsv41_x3pf::pf_kernel<2, NT, MTL, LO, HI>( \
-    const half*, const half*, const int64_t*, const int64_t*, const int*, const int*, const int*, const int*, \
-    const int*, float*, int, int, int, int, int, int);
+    const half* __restrict__, const half* __restrict__, const int64_t* __restrict__, const int64_t* __restrict__, \
+    const int* __restrict__, const int* __restrict__, const int* __restrict__, const int* __restrict__, \
+    const int* __restrict__, float* __restrict__, int, int, int, int, int, int);
 DSV41_X3PF_INST(4, 4, 8, 8)
 DSV41_X3PF_INST(4, 4, 2, 10)
 DSV41_X3PF_INST(4, 4, 2, 12)
diff --git a/zig/kernels/cuda/experts.cu b/zig/kernels/cuda/experts.cu
index ea6b1a3a73..8cb065f047 100644
--- a/zig/kernels/cuda/experts.cu
+++ b/zig/kernels/cuda/experts.cu
@@ -329,10 +329,12 @@ __global__ void plan_scatter_rows(const int* __restrict__ picks, int P, int E, c
 TF_EXPERT(1)
 TF_EXPERT(0)
 // Groups of 32: SwiGLU gate and up together (epilogue 2), and the down projection to fp32.
-template __global__ void tf_experts::expert_kernel<32, 2, 2, 2, 4>(const __nv_bfloat16*, \
-    int, int, const uint4*, int, int, const int*, const int*, const int*, void*, int, float);
-template __global__ void tf_experts::expert_kernel<32, 1, 0, 2, 4>(const __nv_bfloat16*, \
-    int, int, const uint4*, int, int, const int*, const int*, const int*, void*, int, float);
+template __global__ void tf_experts::expert_kernel<32, 2, 2, 2, 4>(const __nv_bfloat16* __restrict__, int, int, \
+    const uint4* __restrict__, int, int, const int* __restrict__, const int* __restrict__, const int* __restrict__, \
+    void* __restrict__, int, float);
+template __global__ void tf_experts::expert_kernel<32, 1, 0, 2, 4>(const __nv_bfloat16* __restrict__, int, int, \
+    const uint4* __restrict__, int, int, const int* __restrict__, const int* __restrict__, const int* __restrict__, \
+    void* __restrict__, int, float);
 
 // Pair form: gate and up blocks interleaved per group (M = 2), SwiGLU in the epilogue (2) as fp8/nvfp4's up does.
 #define TF_EXPERT_PAIR(EPI) template __global__ void tf_experts::expert_kernel<64, 2, EPI, 2, 4>(const __nv_bfloat16* __restrict__,     int, int, const uint4* __restrict__, int, int, const int* __restrict__, const int* __restrict__, const int* __restrict__, void* __restrict__, int, float);
diff --git a/zig/kernels/cuda/experts_rows.cu b/zig/kernels/cuda/experts_rows.cu
index 8ec6d0d162..fff2ac0691 100644
--- a/zig/kernels/cuda/experts_rows.cu
+++ b/zig/kernels/cuda/experts_rows.cu
@@ -138,7 +138,9 @@ __global__ void __launch_bounds__(WARPS * 32)
 }  // namespace tf_experts_rows
 
 // Groups of 32, two row tiles (32 rows) a warp: SwiGLU gate and up together (epilogue 2), down to fp32 (epilogue 0).
-template __global__ void tf_experts_rows::expert_rows<32, 2, 2, 2, 4>(const __nv_bfloat16*, int, int, const uint4*,
-    int, int, const int*, const int*, const int*, void*, int, float);
-template __global__ void tf_experts_rows::expert_rows<32, 1, 0, 2, 4>(const __nv_bfloat16*, int, int, const uint4*,
-    int, int, const int*, const int*, const int*, void*, int, float);
+template __global__ void tf_experts_rows::expert_rows<32, 2, 2, 2, 4>(const __nv_bfloat16* __restrict__, int, int,
+    const uint4* __restrict__, int, int, const int* __restrict__, const int* __restrict__, const int* __restrict__,
+    void* __restrict__, int, float);
+template __global__ void tf_experts_rows::expert_rows<32, 1, 0, 2, 4>(const __nv_bfloat16* __restrict__, int, int,
+    const uint4* __restrict__, int, int, const int* __restrict__, const int* __restrict__, const int* __restrict__,
+    void* __restrict__, int, float);
diff --git a/zig/kernels/cuda/flashnext_attn.cu b/zig/kernels/cuda/flashnext_attn.cu
index 93b075f071..f39559d38b 100644
--- a/zig/kernels/cuda/flashnext_attn.cu
+++ b/zig/kernels/cuda/flashnext_attn.cu
@@ -491,4 +491,5 @@ extern "C" __global__ void __launch_bounds__(256) fn_attn_merge(const float* __r
 #include "flashnext_qsa_tile.cuh"
 
 // The indexer scores 8 rows a block, every head, the sums in fn_qsa_scores' order (its bits).
-template __global__ void qsa_tile::scores<8, 4>(const __nv_bfloat16*, const __nv_bfloat16*, const int*, float*, int, int, int, int);
+template __global__ void qsa_tile::scores<8, 4>(const __nv_bfloat16* __restrict__, const __nv_bfloat16* __restrict__,
+    const int* __restrict__, float* __restrict__, int, int, int, int);
diff --git a/zig/kernels/cuda/flashnext_fact.cu b/zig/kernels/cuda/flashnext_fact.cu
index bdb8b31268..50db0f35a2 100644
--- a/zig/kernels/cuda/flashnext_fact.cu
+++ b/zig/kernels/cuda/flashnext_fact.cu
@@ -256,8 +256,12 @@ __global__ void __launch_bounds__(WM * WN * 32, MINB) fact_kernel(
 }  // namespace fact
 
 // The prompt form the batched pass launches from 64 rows: 64 x 128 on 2 x 4 warps, bf16 out (one slice) or fp32 slices.
-template __global__ void fact::fact_kernel<64, 128, 2, 4, 3, 2, false, 2, false, true, false>(const __nv_bfloat16*,
-    const float*, const uint32_t*, const __nv_bfloat16*, const __nv_bfloat16*, void*, float*, int, int, int, int, int, int, int);
+template __global__ void fact::fact_kernel<64, 128, 2, 4, 3, 2, false, 2, false, true, false>(
+    const __nv_bfloat16* __restrict__, const float* __restrict__, const uint32_t* __restrict__,
+    const __nv_bfloat16* __restrict__, const __nv_bfloat16* __restrict__, void* __restrict__, float* __restrict__, int,
+    int, int, int, int, int, int);
 // The folded form from 64 rows: every K slice in the CTA, their totals added in slice order (one fp32 slice).
-template __global__ void fact::fact_kernel<64, 128, 2, 4, 3, 2, false, 2, true, true, true>(const __nv_bfloat16*,
-    const float*, const uint32_t*, const __nv_bfloat16*, const __nv_bfloat16*, void*, float*, int, int, int, int, int, int, int);
+template __global__ void fact::fact_kernel<64, 128, 2, 4, 3, 2, false, 2, true, true, true>(
+    const __nv_bfloat16* __restrict__, const float* __restrict__, const uint32_t* __restrict__,
+    const __nv_bfloat16* __restrict__, const __nv_bfloat16* __restrict__, void* __restrict__, float* __restrict__, int,
+    int, int, int, int, int, int);
diff --git a/zig/kernels/cuda/flashnext_fexp.cu b/zig/kernels/cuda/flashnext_fexp.cu
index 21cbf0e628..0451fa3bb3 100644
--- a/zig/kernels/cuda/flashnext_fexp.cu
+++ b/zig/kernels/cuda/flashnext_fexp.cu
@@ -198,7 +198,9 @@ __global__ void __launch_bounds__(WM * WN * 32)
 }  // namespace tf_fact_experts
 
 // gate and up with SwiGLU (bf16 out) and down to fp32, 64-pair tiles on 8 warps, 3 stages (2 CTAs an SM).
-template __global__ void tf_fact_experts::fact_kernel<2, 2, 1, 4, 2, 3>(const __nv_bfloat16*, int, int, const uint4*, int,
-    int, const int*, const int*, const int*, void*, int, float);
-template __global__ void tf_fact_experts::fact_kernel<1, 0, 2, 2, 4, 3>(const __nv_bfloat16*, int, int, const uint4*, int,
-    int, const int*, const int*, const int*, void*, int, float);
+template __global__ void tf_fact_experts::fact_kernel<2, 2, 1, 4, 2, 3>(const __nv_bfloat16* __restrict__, int, int,
+    const uint4* __restrict__, int, int, const int* __restrict__, const int* __restrict__, const int* __restrict__,
+    void* __restrict__, int, float);
+template __global__ void tf_fact_experts::fact_kernel<1, 0, 2, 2, 4, 3>(const __nv_bfloat16* __restrict__, int, int,
+    const uint4* __restrict__, int, int, const int* __restrict__, const int* __restrict__, const int* __restrict__,
+    void* __restrict__, int, float);
diff --git a/zig/kernels/cuda/flashnext_hcf.cu b/zig/kernels/cuda/flashnext_hcf.cu
index 2b5c4d7fbe..3657cc63f6 100644
--- a/zig/kernels/cuda/flashnext_hcf.cu
+++ b/zig/kernels/cuda/flashnext_hcf.cu
@@ -236,10 +236,12 @@ __global__ void __launch_bounds__(WM * WN * 32, MINB) hcf_kernel(
 }  // namespace hcf
 
 // HC down: 32 x 192 tiles, the streams normed in shared memory, the 32 K slices folded in slice order (fp32 out).
-template __global__ void hcf::hcf_kernel<32, 192, 2, 4, 4, 2, true, true, false, 2>(const __nv_bfloat16*, const float*,
-    const uint4*, const __nv_bfloat16*, const __nv_bfloat16*, void*, int, int, int, int, int, int, const float*, const float*,
-    float, int, int, int);
+template __global__ void hcf::hcf_kernel<32, 192, 2, 4, 4, 2, true, true, false, 2>(const __nv_bfloat16* __restrict__,
+    const float* __restrict__, const uint4* __restrict__, const __nv_bfloat16* __restrict__,
+    const __nv_bfloat16* __restrict__, void* __restrict__, int, int, int, int, int, int, const float* __restrict__,
+    const float* __restrict__, float, int, int, int);
 // HC up: 64 x 128 tiles, bf16 out (the up rows fn_hc_upmix mixes).
-template __global__ void hcf::hcf_kernel<64, 128, 2, 4, 3, 2, false, false, true, 2>(const __nv_bfloat16*, const float*,
-    const uint4*, const __nv_bfloat16*, const __nv_bfloat16*, void*, int, int, int, int, int, int, const float*, const float*,
-    float, int, int, int);
+template __global__ void hcf::hcf_kernel<64, 128, 2, 4, 3, 2, false, false, true, 2>(const __nv_bfloat16* __restrict__,
+    const float* __restrict__, const uint4* __restrict__, const __nv_bfloat16* __restrict__,
+    const __nv_bfloat16* __restrict__, void* __restrict__, int, int, int, int, int, int, const float* __restrict__,
+    const float* __restrict__, float, int, int, int);
diff --git a/zig/kernels/cuda/prefill_attention.cu b/zig/kernels/cuda/prefill_attention.cu
index 8ea83eecdf..6f44747980 100644
--- a/zig/kernels/cuda/prefill_attention.cu
+++ b/zig/kernels/cuda/prefill_attention.cu
@@ -248,8 +248,9 @@ vision_kernel(const __nv_bfloat16* __restrict__ qkv, __nv_bfloat16* __restrict__
 } // namespace tf_prefill_attention
 
 // Head dim 128, eight warps, eight query heads a block, eight staging slots (Nemotron's 16 heads a KV head).
-template __global__ void tf_prefill_attention::pattn_kernel<128, 8, 8, 8>(const __nv_bfloat16*,
-    const __nv_bfloat16*, const __nv_bfloat16*, __nv_bfloat16*, int, int, int, int, int, float, const int*, long long);
+template __global__ void tf_prefill_attention::pattn_kernel<128, 8, 8, 8>(const __nv_bfloat16* __restrict__,
+    const __nv_bfloat16* __restrict__, const __nv_bfloat16* __restrict__, __nv_bfloat16* __restrict__, int, int, int,
+    int, int, float, const int* __restrict__, long long);
 // Head dim 64, eight warps, four query heads a block, eight staging slots (LFM2's 4 heads a KV head, causal or trees).
 template __global__ void tf_prefill_attention::tree_kernel<64, 8, 4, 8>(const __nv_bfloat16* __restrict__,
     const __nv_bfloat16* __restrict__, const __nv_bfloat16* __restrict__, __nv_bfloat16* __restrict__, const int* __restrict__, int, int, int, int, int, float);
diff --git a/zig/kernels/cuda/qmm.cu b/zig/kernels/cuda/qmm.cu
index 4901978a3a..c41c99b154 100644
--- a/zig/kernels/cuda/qmm.cu
+++ b/zig/kernels/cuda/qmm.cu
@@ -226,5 +226,6 @@ __global__ void reduce_kernel(const float* __restrict__ part, void* __restrict__
 
 // Decode lane: groups of 32, 16 rows, 64 columns, one by four warps, four stages, bf16 out.
 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);
+    const __nv_bfloat16* __restrict__, const float* __restrict__, const uint32_t* __restrict__,
+    const __nv_bfloat16* __restrict__, const __nv_bfloat16* __restrict__, void* __restrict__, float* __restrict__, int,
+    int, int, int, int, int, int);
diff --git a/zig/kernels/cuda/window_attention.cu b/zig/kernels/cuda/window_attention.cu
index a39dd892f1..b4c749c105 100644
--- a/zig/kernels/cuda/window_attention.cu
+++ b/zig/kernels/cuda/window_attention.cu
@@ -531,7 +531,9 @@ extern "C" __global__ void __launch_bounds__(128) tf_window_decode_tc(
 } // namespace tf_window_prompt
 
 // Head dim 128, eight warps, four or eight query heads a block, eight staging slots.
-template __global__ void tf_window_prompt::wattn_kernel<128, 8, 4, 8>(const __nv_bfloat16*, const __nv_bfloat16*,
-    const __nv_bfloat16*, __nv_bfloat16*, int, int, int, int, int, float, int, int);
-template __global__ void tf_window_prompt::wattn_kernel<128, 8, 8, 8>(const __nv_bfloat16*, const __nv_bfloat16*,
-    const __nv_bfloat16*, __nv_bfloat16*, int, int, int, int, int, float, int, int);
+template __global__ void tf_window_prompt::wattn_kernel<128, 8, 4, 8>(const __nv_bfloat16* __restrict__,
+    const __nv_bfloat16* __restrict__, const __nv_bfloat16* __restrict__, __nv_bfloat16* __restrict__, int, int, int,
+    int, int, float, int, int);
+template __global__ void tf_window_prompt::wattn_kernel<128, 8, 8, 8>(const __nv_bfloat16* __restrict__,
+    const __nv_bfloat16* __restrict__, const __nv_bfloat16* __restrict__, __nv_bfloat16* __restrict__, int, int, int,
+    int, int, float, int, int);
diff --git a/zig/src/families/glm/cuda_native.zig b/zig/src/families/glm/cuda_native.zig
index c55edd63c0..eb9b8b0fa7 100644
--- a/zig/src/families/glm/cuda_native.zig
+++ b/zig/src/families/glm/cuda_native.zig
@@ -103,7 +103,9 @@ const Owned = struct {
     lanes: glm_lanes.Lanes,
 };
 
-extern "c" fn setenv(name: [*:0]const u8, value: [*:0]const u8, overwrite: c_int) c_int;
+const env = if (@import("builtin").os.tag == .windows) @import("sys").c else struct {
+    extern "c" fn setenv(name: [*:0]const u8, value: [*:0]const u8, overwrite: c_int) c_int;
+};
 
 /// Every peer's send/recv connected now, before the weights fill memory (connected on first use, ibv_reg_mr can fail).
 fn warmPeers(d: *const cuda.Driver, g: *const cuda.tp_group.Group) !void {
@@ -158,10 +160,10 @@ fn openRank(gpa: std.mem.Allocator, io: std.Io, ctx: *const cuda.Context, dir: [
     errdefer gpa.destroy(own);
     own.gpa = gpa;
     // every connection registered at init, while memory is free (lazily, a full GB10's ibv_reg_mr can refuse it)
-    if (std.c.getenv("NCCL_RUNTIME_CONNECT") == null) _ = setenv("NCCL_RUNTIME_CONNECT", "0", 0);
+    if (std.c.getenv("NCCL_RUNTIME_CONNECT") == null) _ = env.setenv("NCCL_RUNTIME_CONNECT", "0", 0);
     // NCCL's cuMem host buffers wore out GB10's driver pool under load (NV_ERR_NO_MEMORY; next ranks failed to start)
-    if (std.c.getenv("NCCL_CUMEM_HOST_ENABLE") == null) _ = setenv("NCCL_CUMEM_HOST_ENABLE", "0", 0);
-    if (std.c.getenv("NCCL_CUMEM_ENABLE") == null) _ = setenv("NCCL_CUMEM_ENABLE", "0", 0);
+    if (std.c.getenv("NCCL_CUMEM_HOST_ENABLE") == null) _ = env.setenv("NCCL_CUMEM_HOST_ENABLE", "0", 0);
+    if (std.c.getenv("NCCL_CUMEM_ENABLE") == null) _ = env.setenv("NCCL_CUMEM_ENABLE", "0", 0);
     var lib = try cuda.nccl.Library.open();
     errdefer lib.close();
     var id: cuda.nccl.UniqueId = undefined;