New model family: MiniMax-M3.1-Flash VL (minimax_m3_vl) EXL3, CUDA, TP2-served

zig/build/cuda.zig

@@ -117,6 +117,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("minimax_m3_vl_cuda", minimaxM3Native(b, target, optimize, cuda, mods.core, mods.lanes));
     runner.addImport("kolibri1", kolibri1Module(b, target, optimize, cuda, mods.core, mods.lanes));
     const runner_exe = b.addExecutable(.{ .name = "tf-cuda-test", .root_module = runner });
     b.installArtifact(runner_exe);
@@ -158,6 +159,15 @@ fn glmCuda(b: *std.Build, target: std.Build.ResolvedTarget, optimize: std.builti
     return m;
 }
 
+/// MiniMax-M3.1-Flash's family module: the openWith API, the pack's admission and the MSA contracts.
+fn minimaxM3Native(b: *std.Build, target: std.Build.ResolvedTarget, optimize: std.builtin.OptimizeMode, cuda: *std.Build.Module, core: *std.Build.Module, lanes: *std.Build.Module) *std.Build.Module {
+    const m = b.createModule(.{ .root_source_file = b.path("zig/src/families/minimax_m3_vl/cuda_native.zig"), .target = target, .optimize = optimize, .link_libc = true });
+    m.addImport("cuda", cuda);
+    m.addImport("core", core);
+    m.addImport("lanes", lanes);
+    return m;
+}
+
 /// The Kolibri-1 CUDA family over the runtime, the core and the lane core.
 const kolibri1Module = @import("kolibri.zig").module;
 
@@ -205,6 +215,7 @@ fn engines(b: *std.Build, target: std.Build.ResolvedTarget, optimize: std.builti
             .{ .name = "core", .module = core },
             .{ .name = "flashnext", .module = flashnextModule(b, target, optimize, cuda, core, lanes, qwen_vision) },
             .{ .name = "glm_cuda", .module = glmCuda(b, target, optimize, cuda, core, lanes, heat) },
+            .{ .name = "minimax_m3_vl_cuda", .module = minimaxM3Native(b, target, optimize, cuda, core, lanes) },
             .{ .name = "qwen35", .module = qwen35 },
             .{ .name = "kolibri1", .module = kolibri1Module(b, target, optimize, cuda, core, lanes) },
         },
@@ -483,7 +494,7 @@ pub fn hostTests(b: *std.Build, draft_ids: *std.Build.Module, build_options: *st
         @import("cortheon_fuzz.zig").add(b, host, native_modules.api, mods.tokenizer, template);
     }
     // the tokenizer's own tests (memo, steps, regex, unicode) run here: no other test binary reaches them
-    for ([_]*std.Build.Module{ cuda, mods.core, mods.lanes, mods.nemotron, mods.qwen35, mods.qwen_vision, flashnextModule(b, host, .debug, cuda, mods.core, mods.lanes, mods.qwen_vision), lfm2Module(b, host, .debug, cuda, mods.core), mods.radio, mods.heat, mods.tokenizer, native_modules.api, native, stagger(b, host, .debug) }) |m| step.dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = m })).step);
+    for ([_]*std.Build.Module{ cuda, mods.core, mods.lanes, mods.nemotron, mods.qwen35, mods.qwen_vision, flashnextModule(b, host, .debug, cuda, mods.core, mods.lanes, mods.qwen_vision), lfm2Module(b, host, .debug, cuda, mods.core), mods.radio, mods.heat, mods.tokenizer, native_modules.api, native, minimaxM3Native(b, host, .debug, cuda, mods.core, mods.lanes), stagger(b, host, .debug) }) |m| step.dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = m })).step);
     const selected = b.option([]const u8, "modelopt-host-filter", "One portable ModelOpt regression") orelse "modelopt host:";
     const run = b.addRunArtifact(b.addTest(.{ .root_module = mods.nemotron, .filters = &.{selected} }));
     run.has_side_effects = true;

zig/build/cuda_kernels.zig

@@ -90,6 +90,7 @@ pub const kernels = [_]Kernel{
     .{ .name = "modelopt_pack", .flags = &.{"-O3"} }, // ours: ModelOpt tensors into the lane, prompt and expert layouts
     .{ .name = "glm_kda", .flags = &.{ "-O3", "--fmad=false" } }, // GLM: the Python family's glm5_next/cuda/kda.cu
     .{ .name = "glm_glue", .flags = &.{"-O3"} }, // ours: the Python family's Triton glue by hand, its PTX op for op
+    .{ .name = "minimax_m3", .flags = glue }, // ours: MiniMax M3.1's MSA indexer, the fork m3 contracts (glue flags keep the fp16 dots' products exact)
     .{ .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 = "exl3_linear", .flags = &.{ "-O3", "--expt-relaxed-constexpr" } }, // exl3/linear.cu's device code, tensorfold_exl3_linear_v5

zig/kernels/cuda/exl3_experts.cu

@@ -247,6 +247,11 @@ extern "C" __global__ void tf_exl3_experts_rot_in(const RotInArgs a) {
         const __nv_bfloat16* xv = reinterpret_cast<const __nv_bfloat16*>(xr);
 #pragma unroll
         for (int j = 0; j < 4; ++j) v[j] = __bfloat162float(xv[j]) * __half2float(suh[j]);
+    } else if (a.x_dtype == 2) {
+        // f32 input (the minimax engine's x32): element indexing must be 32-bit, not the 16-bit xr stride
+        const float* xv = reinterpret_cast<const float*>(a.x) + (size_t)row * a.x_stride + blk * 128 + 4 * lane;
+#pragma unroll
+        for (int j = 0; j < 4; ++j) v[j] = xv[j] * __half2float(suh[j]);
     } else {
         const half* xv = reinterpret_cast<const half*>(xr);
 #pragma unroll
@@ -304,6 +309,16 @@ extern "C" __global__ void tf_exl3_gateup_epilogue(const GateUpArgs a) {
             float uu = bf16r(uv[j] * HAD_SCALE * __half2float(a.svh_u[(size_t)e * a.n + n + j]));
             if (a.limit > 0) { gg = upTo(gg, a.limit); uu = within(uu, a.limit); }
             act = bf16r(bf16r(gg / (1.f + expf(-gg))) * uu);
+        } else if (a.act_mode == 2) {
+            // the minimax M3.1 OAI swiglu contract at the fork's quantization points — the dense
+            // path's shape: the gate/up rows are fp16 (the fork's f16 GEMM outputs), the sigmoid
+            // form is the tf_m3_swigluoai contract (gate clipped from above only, up clamped both
+            // ways and +1, sigmoid alpha = 1.702), and the activation is an fp16 row BEFORE this
+            // epilogue's down-side rotation (the fork's down kernel rotates an fp16 z in fp32)
+            float gg = __half2float(__float2half_rn(gv[j] * HAD_SCALE * __half2float(a.svh_g[(size_t)e * a.n + n + j])));
+            float uu = __half2float(__float2half_rn(uv[j] * HAD_SCALE * __half2float(a.svh_u[(size_t)e * a.n + n + j])));
+            if (a.limit > 0) { gg = upTo(gg, a.limit); uu = within(uu, a.limit); }
+            act = __half2float(__float2half_rn((uu + 1.0f) * (gg / (1.0f + __expf(-1.702f * gg)))));
         } else {
             float gg = gv[j] * HAD_SCALE * __half2float(a.svh_g[(size_t)e * a.n + n + j]);
             float uu = uv[j] * HAD_SCALE * __half2float(a.svh_u[(size_t)e * a.n + n + j]);

zig/kernels/cuda/exl3_prompt.cuh

@@ -179,10 +179,20 @@ __device__ __forceinline__ void tgateup(const half* __restrict__ Xg, const half*
                 float uu = uv[j] * HAD_SCALE * __half2float(svh_u[(size_t)e * N + n + j]);
                 if (act_mode == 0) { gg = bf16r(gg); uu = bf16r(uu); }
                 if (limit > 0) { gg = upTo(gg, limit); uu = within(uu, limit); }
-                float act = gg / (1.f + expf(-gg));
-                if (act_mode == 0) act = bf16r(act);
-                act *= uu;
-                if (act_mode == 0) act = bf16r(act);
+                float act;
+                if (act_mode == 2) {
+                    // the minimax M3.1 OAI swiglu contract at the fork's quantization points — the
+                    // same shape as the decode epilogue's mode 2: fp16 gate/up rows, then a fp16
+                    // activation row before the down-side rotation
+                    gg = __half2float(__float2half_rn(gg));
+                    uu = __half2float(__float2half_rn(uu));
+                    act = __half2float(__float2half_rn((uu + 1.0f) * (gg / (1.0f + __expf(-1.702f * gg)))));
+                } else {
+                    act = gg / (1.f + expf(-gg));
+                    if (act_mode == 0) act = bf16r(act);
+                    act *= uu;
+                    if (act_mode == 0) act = bf16r(act);
+                }
                 v[j] = act * __half2float(suh_d[(size_t)e * N + n + j]);
             }
             fwht128(v, lane);