ling: Ling-3.1-flash on CUDA at TP=1 on #804's kernels (draft, stacked on #804)

docs/recipes/README.md

@@ -15,6 +15,7 @@ the native engine, or the Python engine frozen at 0.6.6.
 | Gemma 4 26B-A4B | [MLX, fused one-row decode](gemma-4.md) |
 | DeepSeek-V4-Flash | [MLX on a 256 GB Mac, DSpark and MTP drafts](deepseek-v4-flash.md) |
 | Qwen3.6-35B-A3B | [MLX with MTP drafts, one-GPU CUDA](qwen3.6-moe.md) |
+| Ling-3.1-flash | [Not served yet: CUDA weight loader, synthetic checkpoints](ling-3.1.md) |
 
 ## Capability floor
 

zig/build/cuda.zig

@@ -52,6 +52,7 @@ const kernels = [_]Kernel{
     .{ .name = "glm_glue", .flags = &.{"-O3"} }, // ours: the Python family's Triton glue by hand, its PTX op for op
     .{ .name = "glm_ops", .flags = &.{"-O3"} }, // ours: GLM's integer work the Python family did in torch (the DSA token lists)
     .{ .name = "glm_quant", .flags = torch_ops }, // ours: GLM's draft-only weight quantizers, torch at load in the Python family
+    .{ .name = "ling_glue", .flags = glue }, // ours: Ling-3.1's bf16 glue at the HF model's rounding points
     .{ .name = "torch_argmax", .src = "torch_ops/argmax", .flags = torch_ops },
     .{ .name = "torch_topk", .src = "torch_ops/topk", .flags = torch_ops },
     .{ .name = "torch_pointwise", .src = "torch_ops/pointwise", .flags = torch_ops },
@@ -154,6 +155,7 @@ pub fn targets(b: *std.Build, target: std.Build.ResolvedTarget, optimize: std.bu
     for (exlr8_build.denseModules(b, target, optimize, exlr8, exlr8_tests[0])) |m| runner.addImport(m.name, m.module);
     const glm_cuda = glmCuda(b, target, optimize, cuda, mods.core, mods.lanes, mods.heat);
     runner.addImport("glm_cuda", glm_cuda);
+    runner.addImport("ling_cuda", lingCuda(b, target, optimize, cuda, mods.core, mods.lanes, glm_cuda));
     const runner_exe = b.addExecutable(.{ .name = "tf-cuda-test", .root_module = runner });
     b.installArtifact(runner_exe);
     const controls = if (nvcc) |tool| fatbin(b, tool, version.?, .{ .name = "contract_controls", .flags = &.{"-O3"} }, sms) else null;
@@ -183,9 +185,17 @@ fn radioModule(b: *std.Build, target: std.Build.ResolvedTarget, optimize: std.bu
     return b.createModule(.{ .root_source_file = b.path("zig/src/families/radio/radio.zig"), .target = target, .optimize = optimize, .link_libc = true, .imports = &.{ .{ .name = "cuda", .module = cuda }, .{ .name = "core", .module = core }, .{ .name = "lanes", .module = lanes } } });
 }
 
+/// The Ling-3.1-flash CUDA family over the runtime, the core, the lane core and the GLM kernels it reuses.
+fn lingCuda(b: *std.Build, target: std.Build.ResolvedTarget, optimize: std.builtin.OptimizeMode, cuda: *std.Build.Module, core: *std.Build.Module, lanes: *std.Build.Module, glm: *std.Build.Module) *std.Build.Module {
+    const m = b.createModule(.{ .root_source_file = b.path("zig/src/families/ling/cuda.zig"), .target = target, .optimize = optimize, .link_libc = true });
+    for ([_]struct { []const u8, *std.Build.Module }{ .{ "glm_cuda", glm }, .{ "cuda", cuda }, .{ "core", core }, .{ "lanes", lanes } }) |x| m.addImport(x[0], x[1]);
+    return m;
+}
+
 /// The CUDA engines a native server opens (native/cuda.zig), over the given runtime and families.
 fn engines(b: *std.Build, target: std.Build.ResolvedTarget, optimize: std.builtin.OptimizeMode, cuda: *std.Build.Module, core: *std.Build.Module, lanes: *std.Build.Module, nemotron: *std.Build.Module, heat: *std.Build.Module) struct { api: *std.Build.Module, engines: *std.Build.Module } {
     const api = b.createModule(.{ .root_source_file = b.path("zig/src/core/engine_api.zig"), .target = target, .optimize = optimize, .link_libc = true, .imports = &.{.{ .name = "lanes", .module = lanes }} });
+    const glm = glmCuda(b, target, optimize, cuda, core, lanes, heat);
     const mod = b.createModule(.{
         .root_source_file = b.path("zig/src/native/cuda.zig"),
         .target = target,
@@ -198,7 +208,8 @@ fn engines(b: *std.Build, target: std.Build.ResolvedTarget, optimize: std.builti
             .{ .name = "nemotron", .module = nemotron },
             .{ .name = "lfm2", .module = lfm2Module(b, target, optimize, cuda, core) },
             .{ .name = "core", .module = core },
-            .{ .name = "glm_cuda", .module = glmCuda(b, target, optimize, cuda, core, lanes, heat) },
+            .{ .name = "glm_cuda", .module = glm },
+            .{ .name = "ling_cuda", .module = lingCuda(b, target, optimize, cuda, core, lanes, glm) },
         },
     });
     return .{ .api = api, .engines = mod };
@@ -272,6 +283,11 @@ pub fn hostTests(b: *std.Build, draft_ids: *std.Build.Module, build_options: *st
     // the GLM-5.3-Flash CUDA family's host side
     const glm = glmCuda(b, host, .debug, cuda, mods.core, mods.lanes, mods.heat);
     step.dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = glm })).step);
+    // Ling-3.1-flash's config loader and weight map, and its routing and limits as the GLM kernels take them
+    const ling = lingCuda(b, host, .debug, cuda, mods.core, mods.lanes, glm);
+    const ling_run = b.addRunArtifact(b.addTest(.{ .root_module = ling }));
+    step.dependOn(&ling_run.step);
+    b.step("test-ling", "Ling-3.1-flash's config loader and weight map (no GPU work)").dependOn(&ling_run.step);
     // the shares a GLM-5.3-Flash rank takes (cluster/canon.zig against the Python family's)
     const shares = b.createModule(.{ .root_source_file = b.path("zig/src/cluster/canon_glm_test.zig"), .target = host, .optimize = .debug });
     step.dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = shares })).step);

zig/kernels/cuda/experts.cuh

@@ -62,6 +62,17 @@ __device__ __forceinline__ float swiglu(float g, float u, float limit) {
   return bf(gv / (1.f + expf(-gv))) * uv;
 }
 
+// Ling's order (EPI 4): the limit on silu's output, min(silu(g), L) * clip(u, -L, L).
+__device__ __forceinline__ float swiglu_after(float g, float u, float limit) {
+  const float gv = bf(g);
+  float uv = bf(u), s = bf(gv / (1.f + expf(-gv)));
+  if (limit > 0.f) {
+    s = fminf(s, limit);
+    uv = fminf(fmaxf(uv, -limit), limit);
+  }
+  return s * uv;
+}
+
 __device__ __forceinline__ float relu2(float a) {
   const float u = fmaxf(bf(a), 0.f);
   return u * u;
@@ -88,6 +99,9 @@ __device__ __forceinline__ void epilogue(const float (&acc)[M][RT][NTW][4], int
         } else if constexpr (EPI == 2) {
           o0 = swiglu(a0, acc[M - 1][r][j][2 * h], limit);
           o1 = swiglu(a1, acc[M - 1][r][j][2 * h + 1], limit);
+        } else if constexpr (EPI == 4) {
+          o0 = swiglu_after(a0, acc[M - 1][r][j][2 * h], limit);
+          o1 = swiglu_after(a1, acc[M - 1][r][j][2 * h + 1], limit);
         } else {
           o0 = a0;
           o1 = a1;

zig/kernels/cuda/fp8_experts.cu

@@ -267,11 +267,13 @@ __global__ void __launch_bounds__(PWARPS * 32, PMIN)
   }
 }
 
-// The instances fp8_experts_cuda and fp8_experts_prompt_cuda launch: SwiGLU gate-up, fp32 and bf16 down.
+// The instances: SwiGLU gate-up (EPI 4 in Ling's clamp order), fp32 and bf16 down; decode and prompt forms.
 template __global__ void fp8_expert_kernel<2, 2, 4>(const __nv_bfloat16* __restrict__, int, int, const uint4* __restrict__, const float* __restrict__, int, int, const int* __restrict__, const int* __restrict__, const int* __restrict__, void* __restrict__, int, float, int);
+template __global__ void fp8_expert_kernel<2, 4, 4>(const __nv_bfloat16* __restrict__, int, int, const uint4* __restrict__, const float* __restrict__, int, int, const int* __restrict__, const int* __restrict__, const int* __restrict__, void* __restrict__, int, float, int);
 template __global__ void fp8_expert_kernel<1, 0, 4>(const __nv_bfloat16* __restrict__, int, int, const uint4* __restrict__, const float* __restrict__, int, int, const int* __restrict__, const int* __restrict__, const int* __restrict__, void* __restrict__, int, float, int);
 template __global__ void fp8_expert_kernel<1, 3, 4>(const __nv_bfloat16* __restrict__, int, int, const uint4* __restrict__, const float* __restrict__, int, int, const int* __restrict__, const int* __restrict__, const int* __restrict__, void* __restrict__, int, float, int);
 template __global__ void fp8_expert_prompt_kernel<2, 2, 2, 2>(const __nv_bfloat16* __restrict__, int, int, const uint4* __restrict__, const float* __restrict__, int, int, const int* __restrict__, const int* __restrict__, const int* __restrict__, void* __restrict__, int, float, int);
+template __global__ void fp8_expert_prompt_kernel<2, 4, 2, 2>(const __nv_bfloat16* __restrict__, int, int, const uint4* __restrict__, const float* __restrict__, int, int, const int* __restrict__, const int* __restrict__, const int* __restrict__, void* __restrict__, int, float, int);
 template __global__ void fp8_expert_prompt_kernel<1, 0, 4, 4>(const __nv_bfloat16* __restrict__, int, int, const uint4* __restrict__, const float* __restrict__, int, int, const int* __restrict__, const int* __restrict__, const int* __restrict__, void* __restrict__, int, float, int);
 template __global__ void fp8_expert_prompt_kernel<1, 3, 4, 4>(const __nv_bfloat16* __restrict__, int, int, const uint4* __restrict__, const float* __restrict__, int, int, const int* __restrict__, const int* __restrict__, const int* __restrict__, void* __restrict__, int, float, int);
 }

zig/kernels/cuda/glm_glue_attn.cuh

@@ -39,26 +39,28 @@ __device__ __forceinline__ void ldm4(uint32_t* r, const void* p) {
                  : "r"(a));
 }
 
-// A 512-wide bf16 row's 16-byte unit u in shared memory, swizzled by the row so ldmatrix's eight rows hit eight banks.
-__device__ __forceinline__ int wide(int row, int u) { return row * LW + ((u ^ (row & 7)) << 3); }
+// A KW-wide bf16 row's 16-byte unit u in shared memory, swizzled by the row so ldmatrix's eight rows hit eight banks.
+template <int KW = LW>
+__device__ __forceinline__ int wide(int row, int u) { return row * KW + ((u ^ (row & 7)) << 3); }
 // p's [32][32] rows (64 bytes): unit u of row r.
 __device__ __forceinline__ int narrow(int row, int u) { return row * KT + ((u ^ ((row >> 1) & 3)) << 3); }
 
 // latent._dense_chunks/_sparse_chunks (sparse 0/1), block (row, HB heads, chunk), 8 warps: online softmax by 32 keys.
-template <int HB>
+template <int HB, int RW = 0>
 constexpr int wide_smem() {
-  return (HB * LW + 2 * KT * LW + HB * KT) * 2 + (HB * 4 + 2 * HB) * 4;
+  return (HB * (LW + RW) + 2 * KT * (LW + RW) + HB * KT) * 2 + (HB * 4 + 2 * HB) * 4;
 }
 
-template <int HB>
+// RW rope columns (Ling's 64) follow the latent in q and keys: scores over LW + RW, values from the latent alone.
+template <int HB, int RW = 0>
 __device__ void chunks_wide(const __nv_bfloat16* qa, const __nv_bfloat16* lc, const int* pos, const int* tok,
                             const int* cnt, float* po, float* pm, float* pl, int64_t rows, int64_t heads, int64_t w,
                             float scale, int sparse) {
-  constexpr int MT = HB / 16;
+  constexpr int MT = HB / 16, KW = LW + RW;
   extern __shared__ __align__(16) unsigned char smem[];
   __nv_bfloat16* qs = reinterpret_cast<__nv_bfloat16*>(smem);
-  __nv_bfloat16* kvs = qs + HB * LW;
-  __nv_bfloat16* ps = kvs + 2 * KT * LW;
+  __nv_bfloat16* kvs = qs + HB * KW;
+  __nv_bfloat16* ps = kvs + 2 * KT * KW;
   float* red = reinterpret_cast<float*>(ps + HB * KT);
   float* stat = red + HB * 4;
   const int t = threadIdx.x, lane = t & 31, warp = t >> 5, g = lane >> 2, q = lane & 3;
@@ -67,14 +69,14 @@ __device__ void chunks_wide(const __nv_bfloat16* qa, const __nv_bfloat16* lc, co
   const int64_t n = sparse ? cnt[r] : 0, limit = sparse ? 0 : (int64_t)pos[0] + r, start = c * CH;
   const bool live = sparse ? start < n : start <= limit;
   auto load_kv = [&](int tile) {
-    __nv_bfloat16* dst = kvs + (tile & 1) * KT * LW;
+    __nv_bfloat16* dst = kvs + (tile & 1) * KT * KW;
     const int64_t first = start + tile * KT;
-    for (int idx = t; idx < KT * (LW / 8); idx += blockDim.x) {
-      const int key = idx / (LW / 8), u = idx % (LW / 8);
+    for (int idx = t; idx < KT * (KW / 8); idx += blockDim.x) {
+      const int key = idx / (KW / 8), u = idx % (KW / 8);
       const int64_t ki = first + key;
       const bool ok = sparse ? ki < n : ki <= limit;
       const int64_t src = ok ? (sparse ? (int64_t)tok[r * w + ki] : ki) : 0;
-      cp16(dst + wide(key, u), lc + src * LW + u * 8, ok);
+      cp16(dst + wide<KW>(key, u), lc + src * KW + u * 8, ok);
     }
   };
   float m[MT][2], l[MT][2], o[MT][8][4];
@@ -85,10 +87,10 @@ __device__ void chunks_wide(const __nv_bfloat16* qa, const __nv_bfloat16* lc, co
       for (int e = 0; e < 4; ++e) o[i][j][e] = 0.0f;
   constexpr int tiles = CH / KT;
   if (live) {
-    for (int idx = t; idx < HB * (LW / 8); idx += blockDim.x) {
-      const int hd = idx / (LW / 8), u = idx % (LW / 8);
+    for (int idx = t; idx < HB * (KW / 8); idx += blockDim.x) {
+      const int hd = idx / (KW / 8), u = idx % (KW / 8);
       const bool ok = h0 + hd < heads;
-      cp16(qs + wide(hd, u), qa + (r * heads + (ok ? h0 + hd : 0)) * LW + u * 8, ok);
+      cp16(qs + wide<KW>(hd, u), qa + (r * heads + (ok ? h0 + hd : 0)) * KW + u * 8, ok);
     }
     load_kv(0);
     commit();
@@ -101,17 +103,17 @@ __device__ void chunks_wide(const __nv_bfloat16* qa, const __nv_bfloat16* lc, co
     else
       wait<0>();
     __syncthreads();
-    const __nv_bfloat16* kv = kvs + (tile & 1) * KT * LW;
+    const __nv_bfloat16* kv = kvs + (tile & 1) * KT * KW;
     const int64_t first = start + tile * KT;
     float sc[4];
     if (hh < MT) {
       float acc[4] = {0.0f, 0.0f, 0.0f, 0.0f};
 #pragma unroll 4
-      for (int kk = 0; kk < LW; kk += 32) {
+      for (int kk = 0; kk < KW; kk += 32) {
         uint32_t a0[4], a1[4], b[4];
-        ldm4<false>(a0, qs + wide(hh * 16 + (lane & 15), (kk >> 3) + (lane >> 4)));
-        ldm4<false>(a1, qs + wide(hh * 16 + (lane & 15), ((kk + 16) >> 3) + (lane >> 4)));
-        ldm4<false>(b, kv + wide(kg * 8 + (lane & 7), (kk >> 3) + (lane >> 3)));
+        ldm4<false>(a0, qs + wide<KW>(hh * 16 + (lane & 15), (kk >> 3) + (lane >> 4)));
+        ldm4<false>(a1, qs + wide<KW>(hh * 16 + (lane & 15), ((kk + 16) >> 3) + (lane >> 4)));
+        ldm4<false>(b, kv + wide<KW>(kg * 8 + (lane & 7), (kk >> 3) + (lane >> 3)));
         glm_router::mma16816(acc, a0, b);
         glm_router::mma16816(acc, a1, b + 2);
       }
@@ -172,7 +174,7 @@ __device__ void chunks_wide(const __nv_bfloat16* qa, const __nv_bfloat16* lc, co
 #pragma unroll
       for (int j = 0; j < 8; j += 2) {
         uint32_t b[4];
-        ldm4<true>(b, kv + wide(k16 + (lane & 15), ((warp * 64 + j * 8) >> 3) + (lane >> 4)));
+        ldm4<true>(b, kv + wide<KW>(k16 + (lane & 15), ((warp * 64 + j * 8) >> 3) + (lane >> 4)));
         for (int i = 0; i < MT; ++i) {
           glm_router::mma16816(o[i][j], a[i], b);
           glm_router::mma16816(o[i][j + 1], a[i], b + 2);
@@ -213,6 +215,13 @@ extern "C" __global__ void __launch_bounds__(256)
   glm_attn::chunks_wide<32>(qa, lc, pos, tok, cnt, po, pm, pl, rows, heads, w, scale, sparse);
 }
 
+// Ling's 16-head form: 576-wide q and keys (latent, then 64 roped columns), in wide_smem<16, 64>() bytes.
+extern "C" __global__ void __launch_bounds__(256)
+    tf_glm_chunks_r16(const __nv_bfloat16* qa, const __nv_bfloat16* lc, const int* pos, const int* tok, const int* cnt,
+                      float* po, float* pm, float* pl, int64_t rows, int64_t heads, int64_t w, float scale, int sparse) {
+  glm_attn::chunks_wide<16, 64>(qa, lc, pos, tok, cnt, po, pm, pl, rows, heads, w, scale, sparse);
+}
+
 // latent._merge, block (row, head), 128 threads: chunks merged in order, o / l by div.full; skips sparse count-0 rows.
 extern "C" __global__ void __launch_bounds__(128)
     tf_glm_merge(const float* po, const float* pm, const float* pl, __nv_bfloat16* out, const int* cnt, int64_t rows,