ggml adapter: SM120/Qwen3.6-35B target — binding-driven LLM decode windows, MoE catalog family, packless artifacts

csrc/fp4_bindings.cpp

@@ -950,6 +950,86 @@ The downstream GEMM keeps its original K = N_il/2 weight.
         R"pbdoc(
 Skinny-M half-width fused GeGLU GEMM on the decoder tile (128x64x256);
 same contract as cutlass_fp4_gemm_geglu_il_hw.
+)pbdoc");
+
+  m.def("cutlass_fp4_gemm_geglu_il_hw_nod",
+        [](uintptr_t A_packed, uintptr_t SFA,
+           uintptr_t B_packed, uintptr_t SFB,
+           uintptr_t D_dummy, uintptr_t compact_packed, uintptr_t compact_sfa,
+           int M, int N_il, int K, uintptr_t stream) -> int {
+          const auto shape = fp4_kernel_shape({{"M", M}, {"N_il", N_il}, {"K", K}});
+          require_fp4_ptrs("cutlass_fp4_gemm_geglu_il_hw_nod",
+                           {{"A_packed", A_packed}, {"SFA", SFA},
+                            {"B_packed", B_packed}, {"SFB", SFB},
+                            {"D_dummy", D_dummy},
+                            {"compact_packed", compact_packed},
+                            {"compact_sfa", compact_sfa}}, shape);
+          require_fp4(M > 0 && N_il > 0 && K > 0 && (N_il % 32) == 0 &&
+                      (K % 16) == 0,
+                      "cutlass_fp4_gemm_geglu_il_hw_nod",
+                      "M must be positive, N_il a positive multiple of 32 "
+                      "and K a positive multiple of 16",
+                      shape);
+          return flash_rt::fp4::cutlass_fp4_gemm_geglu_il_hw_nod(
+              reinterpret_cast<void const*>(A_packed),
+              reinterpret_cast<void const*>(SFA),
+              reinterpret_cast<void const*>(B_packed),
+              reinterpret_cast<void const*>(SFB),
+              reinterpret_cast<void*>(D_dummy),
+              reinterpret_cast<void*>(compact_packed),
+              reinterpret_cast<void*>(compact_sfa),
+              M, N_il, K,
+              reinterpret_cast<cudaStream_t>(stream));
+        },
+        py::arg("A_packed"), py::arg("SFA"),
+        py::arg("B_packed"), py::arg("SFB"),
+        py::arg("D_dummy"), py::arg("compact_packed"), py::arg("compact_sfa"),
+        py::arg("M"), py::arg("N_il"), py::arg("K"),
+        py::arg("stream") = 0,
+        R"pbdoc(
+Half-width fused GeGLU GEMM with the collective's own D store elided:
+compact_packed/compact_sfa are the only outputs and D_dummy is never
+written (still validated; the host-side TMA descriptor needs a real
+pointer).  Same contract as cutlass_fp4_gemm_geglu_il_hw otherwise.
+)pbdoc");
+
+  m.def("cutlass_fp4_gemm_geglu_il_hw_nod_v10",
+        [](uintptr_t A_packed, uintptr_t SFA,
+           uintptr_t B_packed, uintptr_t SFB,
+           uintptr_t D_dummy, uintptr_t compact_packed, uintptr_t compact_sfa,
+           int M, int N_il, int K, uintptr_t stream) -> int {
+          const auto shape = fp4_kernel_shape({{"M", M}, {"N_il", N_il}, {"K", K}});
+          require_fp4_ptrs("cutlass_fp4_gemm_geglu_il_hw_nod_v10",
+                           {{"A_packed", A_packed}, {"SFA", SFA},
+                            {"B_packed", B_packed}, {"SFB", SFB},
+                            {"D_dummy", D_dummy},
+                            {"compact_packed", compact_packed},
+                            {"compact_sfa", compact_sfa}}, shape);
+          require_fp4(M > 0 && N_il > 0 && K > 0 && (N_il % 32) == 0 &&
+                      (K % 16) == 0,
+                      "cutlass_fp4_gemm_geglu_il_hw_nod_v10",
+                      "M must be positive, N_il a positive multiple of 32 "
+                      "and K a positive multiple of 16",
+                      shape);
+          return flash_rt::fp4::cutlass_fp4_gemm_geglu_il_hw_nod_v10(
+              reinterpret_cast<void const*>(A_packed),
+              reinterpret_cast<void const*>(SFA),
+              reinterpret_cast<void const*>(B_packed),
+              reinterpret_cast<void const*>(SFB),
+              reinterpret_cast<void*>(D_dummy),
+              reinterpret_cast<void*>(compact_packed),
+              reinterpret_cast<void*>(compact_sfa),
+              M, N_il, K,
+              reinterpret_cast<cudaStream_t>(stream));
+        },
+        py::arg("A_packed"), py::arg("SFA"),
+        py::arg("B_packed"), py::arg("SFB"),
+        py::arg("D_dummy"), py::arg("compact_packed"), py::arg("compact_sfa"),
+        py::arg("M"), py::arg("N_il"), py::arg("K"),
+        py::arg("stream") = 0,
+        R"pbdoc(
+Skinny-M no-D-store fused GeGLU GEMM on the decoder tile (128x64x256);
+same contract as cutlass_fp4_gemm_geglu_il_hw_nod.
 )pbdoc");
 
 #ifdef FLASHRT_HAVE_COSMOS3_EDGE

csrc/gemm/fp4/cutlass_fp4_gemm_bias_f32b_f16out_sm100.cu

@@ -0,0 +1,153 @@
+// ============================================================================
+//  FlashRT — NVFP4 GEMM with fp32 per-column bias and fp16 output
+//  (SM100/SM110). See header for the contract.
+// ============================================================================
+
+#include "gemm/fp4/cutlass_fp4_gemm_bias_f32b_f16out_sm100.cuh"
+
+#include "cutlass/cutlass.h"
+#include "cutlass/epilogue/dispatch_policy.hpp"
+#include "cutlass/epilogue/fusion/operations.hpp"
+#include "cutlass/gemm/dispatch_policy.hpp"
+#include "cutlass/gemm/collective/collective_builder.hpp"
+#include "cutlass/epilogue/collective/collective_builder.hpp"
+#include "cutlass/gemm/device/gemm_universal_adapter.h"
+#include "cutlass/gemm/kernel/gemm_universal.hpp"
+#include "cutlass/util/packed_stride.hpp"
+#include "cutlass/detail/sm100_blockscaled_layout.hpp"
+#include "cute/tensor.hpp"
+
+#include <mutex>
+#include <unordered_map>
+
+namespace flash_rt {
+namespace fp4 {
+
+namespace bias_f16out {
+
+using namespace cute;
+
+using ElementA = cutlass::nv_float4_t<cutlass::float_e2m1_t>;
+using LayoutATag = cutlass::layout::RowMajor;
+constexpr int AlignmentA = 32;
+
+using ElementB = cutlass::nv_float4_t<cutlass::float_e2m1_t>;
+using LayoutBTag = cutlass::layout::ColumnMajor;
+constexpr int AlignmentB = 32;
+
+using ElementAccumulator = float;
+using ElementCompute = float;
+using ArchTag = cutlass::arch::Sm100;
+using OperatorClass = cutlass::arch::OpClassBlockScaledTensorOp;
+
+using ElementD = cutlass::half_t;
+using ElementC = cutlass::half_t;
+constexpr int AlignmentCD = 8;
+
+using MmaTileShape = Shape<_128, _128, _256>;
+using ClusterShape = Shape<_1, _1, _1>;
+
+// per-shape CUTLASS workspace cache (capture-safe: growth happens during
+// the uncaptured warmup evaluation)
+struct ws_key {
+    int M, N, K;
+    bool operator==(const ws_key & o) const { return M == o.M && N == o.N && K == o.K; }
+};
+struct ws_key_hash {
+    size_t operator()(const ws_key & k) const noexcept {
+        return (size_t) k.M * 1315423911u ^ (size_t) k.N * 2654435761u ^ (size_t) k.K;
+    }
+};
+inline void * get_ws(int M, int N, int K, size_t needed) {
+    static std::unordered_map<ws_key, std::pair<void *, size_t>, ws_key_hash> cache;
+    static std::mutex mu;
+    std::lock_guard<std::mutex> lk(mu);
+    auto & e = cache[ws_key{M, N, K}];
+    if (e.second < needed) {
+        if (e.first) { cudaFree(e.first); }
+        cudaMalloc(&e.first, needed);
+        e.second = needed;
+    }
+    return e.first;
+}
+
+using FusionOperation = cutlass::epilogue::fusion::LinCombPerColBias<
+    ElementD, ElementCompute, float, ElementC, ElementCompute>;
+
+using CollectiveEpilogue =
+    typename cutlass::epilogue::collective::CollectiveBuilder<
+        ArchTag, OperatorClass, MmaTileShape, ClusterShape,
+        cutlass::epilogue::collective::EpilogueTileAuto,
+        ElementAccumulator, ElementAccumulator,
+        ElementC, cutlass::layout::RowMajor, AlignmentCD,
+        ElementD, cutlass::layout::RowMajor, AlignmentCD,
+        cutlass::epilogue::collective::EpilogueScheduleAuto,
+        FusionOperation>::CollectiveOp;
+
+using CollectiveMainloop =
+    typename cutlass::gemm::collective::CollectiveBuilder<
+        ArchTag, OperatorClass,
+        ElementA, LayoutATag, AlignmentA,
+        ElementB, LayoutBTag, AlignmentB,
+        ElementAccumulator, MmaTileShape, ClusterShape,
+        cutlass::gemm::collective::StageCountAutoCarveout<static_cast<int>(
+            sizeof(typename CollectiveEpilogue::SharedStorage))>,
+        cutlass::gemm::collective::KernelScheduleAuto>::CollectiveOp;
+
+using GemmKernel = cutlass::gemm::kernel::GemmUniversal<
+    Shape<int, int, int, int>, CollectiveMainloop, CollectiveEpilogue, void>;
+using Gemm = cutlass::gemm::device::GemmUniversalAdapter<GemmKernel>;
+
+}  // namespace bias_f16out
+
+int gemm_bias_f16out(
+    const void * A_packed, const void * SFA,
+    const void * B_packed, const void * SFB,
+    const void * bias_f32, void * D_f16,
+    int M, int N, int K,
+    cudaStream_t stream) {
+  using namespace bias_f16out;
+
+  auto stride_A = cutlass::make_cute_packed_stride(
+      typename Gemm::GemmKernel::StrideA{}, {M, K, 1});
+  auto stride_B = cutlass::make_cute_packed_stride(
+      typename Gemm::GemmKernel::StrideB{}, {N, K, 1});
+  auto stride_C = cutlass::make_cute_packed_stride(
+      typename Gemm::GemmKernel::StrideC{}, {M, N, 1});
+  auto stride_D = cutlass::make_cute_packed_stride(
+      typename Gemm::GemmKernel::StrideD{}, {M, N, 1});
+  using Cfg =
+      typename Gemm::GemmKernel::CollectiveMainloop::Sm1xxBlkScaledConfig;
+  auto layout_SFA = Cfg::tile_atom_to_shape_SFA(make_shape(M, N, K, 1));
+  auto layout_SFB = Cfg::tile_atom_to_shape_SFB(make_shape(M, N, K, 1));
+
+  using EA = typename ElementA::DataType;
+  using SA = typename ElementA::ScaleFactorType;
+
+  typename Gemm::Arguments args{
+      cutlass::gemm::GemmUniversalMode::kGemm, {M, N, K, 1},
+      {reinterpret_cast<EA const*>(A_packed), stride_A,
+       reinterpret_cast<EA const*>(B_packed), stride_B,
+       reinterpret_cast<SA const*>(SFA), layout_SFA,
+       reinterpret_cast<SA const*>(SFB), layout_SFB},
+      {{},
+       reinterpret_cast<ElementC const*>(D_f16), stride_C,
+       reinterpret_cast<ElementD*>(D_f16), stride_D}};
+  args.epilogue.thread.alpha = 1.0f;
+  args.epilogue.thread.beta = 0.0f;
+  args.epilogue.thread.bias_ptr = reinterpret_cast<float const*>(bias_f32);
+
+  Gemm gemm;
+  auto st = gemm.can_implement(args);
+  if (st != cutlass::Status::kSuccess) return static_cast<int>(st) | 0x10000;
+  size_t ws_sz = Gemm::get_workspace_size(args);
+  void* ws = ws_sz > 0 ? get_ws(M, N, K, ws_sz) : nullptr;
+  st = gemm.initialize(args, ws, stream);
+  if (st != cutlass::Status::kSuccess) return static_cast<int>(st) | 0x20000;
+  st = gemm.run(stream);
+  return (st == cutlass::Status::kSuccess) ? 0
+                                           : (static_cast<int>(st) | 0x30000);
+}
+
+}  // namespace fp4
+}  // namespace flash_rt

csrc/gemm/fp4/cutlass_fp4_gemm_bias_f32b_f16out_sm100.cuh

@@ -0,0 +1,30 @@
+// ============================================================================
+//  FlashRT — NVFP4 GEMM with fp32 per-column bias and fp16 output
+//  (SM100/SM110).
+//
+//  D_f16[M, N] = A @ B^T + bias[N]. For hosts that keep biases in fp32 but
+//  consume the projection in fp16 (e.g. attention inputs cast for flash
+//  attention): the fp16 conversion happens once in the epilogue from the
+//  fp32 accumulator, matching an fp32-output GEMM followed by an fp16 cast
+//  bit for bit.
+// ============================================================================
+#pragma once
+
+#include <cuda_runtime.h>
+
+namespace flash_rt {
+namespace fp4 {
+
+// A: [M, K] NVFP4 packed row-major + SFA (tile-interleaved).
+// B: [N, K] NVFP4 packed column-major + SFB.
+// bias_f32: [N] fp32, broadcast over rows. D_f16: [M, N] fp16 row-major.
+// Returns 0 on success; CUTLASS status | stage flag otherwise.
+int gemm_bias_f16out(
+    const void * A_packed, const void * SFA,
+    const void * B_packed, const void * SFB,
+    const void * bias_f32, void * D_f16,
+    int M, int N, int K,
+    cudaStream_t stream);
+
+}  // namespace fp4
+}  // namespace flash_rt

csrc/gemm/fp4/cutlass_fp4_gemm_geglu_il_sm100.cu

@@ -24,6 +24,7 @@
 #include "cute/tensor.hpp"
 
 #include "gemm/fp4/sm100_gelu_mul_blockscale_visitor.hpp"
+#include "gemm/fp4/sm100_epilogue_nod.hpp"
 
 namespace flash_rt {
 namespace fp4 {
@@ -173,6 +174,48 @@ using GemmKernelHwV10 = cutlass::gemm::kernel::GemmUniversal<
 
 using GemmHwV10 = cutlass::gemm::device::GemmUniversalAdapter<GemmKernelHwV10>;
 
+// ── No-D-store instantiations ──
+// The builder's collective still stages the unread D tile through smem and
+// TMA-stores it (row-aliased, but the smem->L2 traffic and store instructions
+// remain).  Rebind the built epilogue onto CollectiveEpilogueNoD
+// (sm100_epilogue_nod.hpp), which elides that store entirely; the compact
+// store node is the only writer.  SharedStorage is unchanged, so the
+// mainloop carveout from the builder output stays valid.
+template <class BuiltEpilogue>
+struct MakeNoD;
+
+template <int StagesC, int StagesD, int FragmentSize,
+          bool ReuseSmemC, bool DelayTmaStore, class... Rest>
+struct MakeNoD<cutlass::epilogue::collective::CollectiveEpilogue<
+    cutlass::epilogue::Sm100TmaWarpSpecialized<
+        StagesC, StagesD, FragmentSize, ReuseSmemC, DelayTmaStore>,
+    Rest...>> {
+  using type = cutlass::epilogue::collective::CollectiveEpilogueNoD<
+      StagesC, StagesD, FragmentSize, ReuseSmemC, DelayTmaStore, Rest...>;
+};
+
+using CollectiveEpilogueHwNoD = typename MakeNoD<CollectiveEpilogueHw>::type;
+static_assert(sizeof(typename CollectiveEpilogueHwNoD::SharedStorage) ==
+              sizeof(typename CollectiveEpilogueHw::SharedStorage),
+              "NoD epilogue must keep the builder's smem footprint");
+
+using GemmKernelHwNoD = cutlass::gemm::kernel::GemmUniversal<
+    Shape<int, int, int, int>,
+    CollectiveMainloopHw, CollectiveEpilogueHwNoD, void>;
+
+using GemmHwNoD = cutlass::gemm::device::GemmUniversalAdapter<GemmKernelHwNoD>;
+
+using CollectiveEpilogueHwNoDV10 = typename MakeNoD<CollectiveEpilogueHwV10>::type;
+static_assert(sizeof(typename CollectiveEpilogueHwNoDV10::SharedStorage) ==
+              sizeof(typename CollectiveEpilogueHwV10::SharedStorage),
+              "NoD epilogue must keep the builder's smem footprint");
+
+using GemmKernelHwNoDV10 = cutlass::gemm::kernel::GemmUniversal<
+    Shape<int, int, int, int>,
+    CollectiveMainloopHwV10, CollectiveEpilogueHwNoDV10, void>;
+
+using GemmHwNoDV10 = cutlass::gemm::device::GemmUniversalAdapter<GemmKernelHwNoDV10>;
+
 }  // namespace geglu_il
 
 int cutlass_fp4_gemm_geglu_il(
@@ -253,7 +296,12 @@ static int run_geglu_il_hw(
   auto stride_A = cutlass::make_cute_packed_stride(StrideAT{}, {M, K, 1});
   auto stride_B = cutlass::make_cute_packed_stride(StrideBT{}, {N_il, K, 1});
   auto stride_C = cutlass::make_cute_packed_stride(StrideCT{}, {M, N_il, 1});
+  // The compact store node writes the real output itself; the collective's
+  // own D store lands in a buffer nothing reads. Aliasing every row onto
+  // row 0 shrinks that write from M*N_il to a single row, which at encoder
+  // shape is 12.8 MB of pure waste per call.
   auto stride_D = cutlass::make_cute_packed_stride(StrideDT{}, {M, N_il, 1});
+  cute::get<0>(stride_D) = 0;
   auto layout_SFA = CfgT::tile_atom_to_shape_SFA(make_shape(M, N_il, K, 1));
   auto layout_SFB = CfgT::tile_atom_to_shape_SFB(make_shape(M, N_il, K, 1));
 
@@ -319,5 +367,31 @@ int cutlass_fp4_gemm_geglu_il_hw_v10(
       M, N_il, K, stream);
 }
 
+int cutlass_fp4_gemm_geglu_il_hw_nod(
+    void const* A_packed, void const* SFA,
+    void const* B_packed, void const* SFB,
+    void*       D_dummy,
+    void*       compact_packed,
+    void*       compact_sfa,
+    int M, int N_il, int K,
+    cudaStream_t stream) {
+  return geglu_il::run_geglu_il_hw<geglu_il::GemmHwNoD>(
+      A_packed, SFA, B_packed, SFB, D_dummy, compact_packed, compact_sfa,
+      M, N_il, K, stream);
+}
+
+int cutlass_fp4_gemm_geglu_il_hw_nod_v10(
+    void const* A_packed, void const* SFA,
+    void const* B_packed, void const* SFB,
+    void*       D_dummy,
+    void*       compact_packed,
+    void*       compact_sfa,
+    int M, int N_il, int K,
+    cudaStream_t stream) {
+  return geglu_il::run_geglu_il_hw<geglu_il::GemmHwNoDV10>(
+      A_packed, SFA, B_packed, SFB, D_dummy, compact_packed, compact_sfa,
+      M, N_il, K, stream);
+}
+
 }  // namespace fp4
 }  // namespace flash_rt