[SM12x] Restore pure-fp8 1d1d path (removed in 8b1392b); fix kernel TU preamble by maci0 · Pull Request #419 · deepseek-ai/DeepGEMM

maci0 · 2026-08-26T02:56:36Z

Summary

8b1392b deleted the pure-fp8 1d1d kernels (sm100_fp8_gemm_1d1d.{hpp,cuh}, −983 lines) and rewrote the SM12x fp8 dispatch. On arch 12, fp8xfp8 inputs then run the combined sm120_fp8_fp4_gemm_1d1d kernel, which misreads fp8 weights as fp4 — silent numerical corruption of FP8 linear layers.

Measured on 2x DGX Spark (GB10 / sm_121a), DeepSeek-V4-Flash-0731, vLLM --linear-backend deep_gemm:

8b1392b pin: France output degenerates (' Septy Septy…'), DSpark draft acceptance collapses, decode ~25.8 → ~4.4 tok/s.
a6b593d pin (last commit with the pure-fp8 kernel): coherent output.

Tracked in #417; vLLM-side interim pin: vllm-project/vllm#53680.

What this restores

csrc/jit_kernels/impls/sm100_fp8_gemm_1d1d.hpp (JIT kernel runtime)
deep_gemm/include/deep_gemm/dg25_fp8/ — the pure-fp8 kernel + its common/ dependencies as a self-contained module
gemm.hpp: fp8xfp8 on arch 10/12 routes to sm100_fp8_gemm_1d1d; fp4 keeps the combined kernel (AB-swap paths handled).

Header fixes the `a6b593d` TU needs to compile at all

The original 1d1d TU cannot be compiled as generated:

common/math.cuh was never included by the 1d1d include chain → align/swap in common/scheduler.cuh:192 / common/sm100_utils.cuh:127 were undefined (identifier "align" is undefined). The TU now includes it and pulls in deep_gemm::math (include before the impl header, so the using-directive precedes the dependent code).
math.cuh's cast_into_bf16_and_pack needs #include <cuda_bf16.h> (__float22bfloat162_rn).
std::is_same_v → cute::is_same_v for NVRTC (libcu++ does not expose variable templates in std under NVRTC).

Validation status (honest)

NVRTC frontend (CUDA 13.3): with the header fixes the TU compiles to PTX for compute_121a (~231 KB) — the preamble fix is verified.
GB10 JIT toolchain gap (same as [Regression] 8b1392b removes SM12x pure-fp8 1d1d kernels and aliases fp8_gemm_nt to the fp4 dispatcher (silent corruption on GB10) #417): stock CUDA 13.0/13.3 ptxas rejects tcgen05.mma/.cta_group::1/.block32 for every sm_12x target; the NVRTC cubin is rejected by the driver (CUDA_ERROR_INVALID_IMAGE); the driver-JIT PTX route is rejected too (CUDA_ERROR_INVALID_PTX). So the kernel needs either a toolchain that can assemble tcgen05 for sm_121a, or build-time compilation (the golden/2.5.0 stack precompiles it).
The fp8_einsum o_proj path runs correctly on SM121a with the restored stack (T=10/96/8192, finite output).

Happy to adjust the module naming (e.g. rename dg25_fp8 → sm100_fp8 or fold into the existing tree) to match maintainer preference.

8b1392b deleted the pure-fp8 1d1d kernels (sm100_fp8_gemm_1d1d.{hpp,cuh}) and rewrote the SM12x fp8 dispatch. On arch 12, fp8xfp8 inputs then run the combined fp8xfp4 kernel, which misreads fp8 weights as fp4 - silent numerical corruption of FP8 linear layers (measured on GB10: France output degenerates, decode drops ~25.8 -> ~4.4 tok/s with vLLM --linear-backend deep_gemm; a6b593d serves coherent output). This restores the pure-fp8 path as a self-contained module (deep_gemm/dg25_fp8/) with the header fixes the a6b593d TU needs to compile: - common/math.cuh was never included by the original 1d1d impl chain, so align/swap in common/scheduler.cuh / sm100_utils.cuh were undefined; the TU now includes it and pulls in deep_gemm::math. - math.cuh's cast_into_bf16_and_pack needs cuda_bf16.h. - std::is_same_v -> cute::is_same_v for NVRTC (libcu++ does not expose it in std under NVRTC). Dispatch: fp8xfp8 on arch 10/12 routes to sm100_fp8_gemm_1d1d; fp4 keeps the combined kernel. Validation note: on GB10 with stock CUDA 13.0/13.3, none of the JIT routes (ptxas, NVRTC cubin, driver-JIT PTX) can assemble the tcgen05 instructions in this kernel for sm_121a (see deepseek-ai#417); the NVRTC frontend does compile the TU to PTX once the header fixes above are in.

ds-review-bot · 2026-08-26T03:16:49Z

+            sm100_fp8_gemm_1d1d(b_data, sfa, a_data, sfb, std::nullopt, d,
+                                eff_m, eff_n, k, gran_k_a, gran_k_b,
+                                k_major, k_major, compiled_dims);


🔴 critical: 为纯 FP8 的 AB 交换实现转置写回: 当 SM12 上 m <= 16、输出连续且无累加时，swap_ab 为真；这里交换输入和 M/N 后调用的旧内核仍按普通连续 [eff_m, eff_n] 输出，而不像下方 SM120 内核通过 swap_ab=true 使用转置步长。对实际形状 [m,n] 的 d，这会按 [n,m] 写入，造成越界写入和错误结果；应禁用该分支的交换或为恢复的内核增加转置写回。

_{🤖 v6}

ds-review-bot · 2026-08-26T03:16:49Z

+        .tc_util = device_runtime->get_tc_util(),
+        .compiled_dims = compiled_dims
+    };
+    const auto config = get_best_config<SM100ArchSpec>(desc);


🔴 critical: 限制旧 FP8 内核只选择非 swap-AB 配置: get_best_config<SM100ArchSpec> 仍会选择 layout.swap_ab=true，但恢复的旧内核没有 swap 模板参数，并明确要求 LOAD_BLOCK_M == BLOCK_M。例如偶数 SM 设备上的纯 FP8 m=1,n=7168 会优先选 cluster_n=2 的 swap 配置，导致 JIT 触发静态断言而无法运行；应过滤 swap 候选或移植新版的 swap-aware 内核配置。

_{🤖 v6}

ds-review-bot · 2026-08-26T03:16:49Z

+    // Pure fp8xfp8 must not run the fp8xfp4 kernel (fp8 weights are
+    // misread as fp4: silent corruption). Route to the fp8 1d1d kernel.
+    if (a_data.scalar_type() != kPackedFP4 and b_data.scalar_type() != kPackedFP4) {
+        if (swap_ab) {


🔴 critical: SM120 swap_ab 分支调用 sm100_fp8_gemm_1d1d(b_data, sfa, a_data, sfb, ..., eff_m=n, eff_n=m, ...) 会产生转置/错位输出。sm120_fp8_fp4_gemm_1d1d 通过 swap_ab=true + strided-store epilogue（stride_cd_m=d.stride(-1), stride_cd_n=d.stride(-2)，cd_n_contiguous=!swap_ab）正确回写；而恢复的 sm100 kernel 只会经 tensor_map_cd 按 (m_idx,n_idx) 正常 TMA store（dg25_fp8/impls/sm100_fp8_gemm_1d1d.cuh:540-545），host 侧 make_tma_cd_desc(d, eff_m=n, d.size(-1)=n, ..., d.stride(-2)) 会把 D^T 当作 n×m 行主序写入 m×n 的 d。这正是 kSwapAbMMax=16 覆盖的 decode 主路径。建议把纯 fp8 判断提前到 swap_ab 决策之前，纯 fp8 时强制不做 host 侧 AB-swap（sm100 heuristics 对小 M 本身有 block_m<=32 候选）。

_{🤖 v5}

ds-review-bot · 2026-08-26T03:16:49Z

+        .compiled_dims = compiled_dims
+    };
+    const auto config = get_best_config<SM100ArchSpec>(desc);
+


🔴 critical: get_best_config(desc) 对 Normal GEMM 会枚举 swap_ab∈{0,1}（heuristics/sm100.hpp:47），小 M 场景 compare() 很可能选中 swap_ab==1 布局，但本 runtime 完全忽略 config.layout.swap_ab：dg25 kernel 模板没有 swap 参数。后果：swap 布局下 store_block_m=umma_step_n(16)、smem_cd 公式不同（sm100.hpp:157/180），与 kernel 内 STORE_BLOCK_M/STORE_BLOCK_N 假设不符 → 静默数值错误；swap 且 cluster_n>1 时映射为 kIsMulticastOnA=true、kNumMulticast=2，触发 kernel 内 DG_STATIC_ASSERT(not kIsMulticastOnA or kNumMulticast == 1)（.cuh:80）导致 JIT 编译失败。建议在此处过滤/禁用 swap 候选（例如给 GemmDesc 加开关让 get_layout_candidates 跳过 swap_ab==1，或至少 DG_HOST_ASSERT(config.layout.swap_ab == 0)）。

_{🤖 v5}

ds-review-bot · 2026-08-26T03:16:49Z

        } else if (arch_major == 10 and sfa.scalar_type() == torch::kInt) {
-            sm100_fp8_fp4_gemm_1d1d(a.first, sfa, b.first, sfb, c, d, m, n, k, gran_k_a, gran_k_b,
-                                    major_a, major_b, compiled_dims);
+            if (b.first.scalar_type() != kPackedFP4) {


🔴 critical: arch 10 路由条件只检查 b.first.scalar_type() != kPackedFP4：fp4_A x fp8_B（sm120 分支中明确存在的 a_is_fp4 组合）会被误路由到纯 fp8 kernel，把 fp4 A 当 fp8 读，造成静默数值错误。应与 sm120 分支一致，同时检查两个操作数：a.first.scalar_type() != kPackedFP4 and b.first.scalar_type() != kPackedFP4。

_{🤖 v5}

ds-review-bot · 2026-08-26T03:16:49Z

+/// Casting
+template <typename old_t>
+CUTLASS_DEVICE int cast_into_bf16_and_pack(old_t& x, old_t& y) {
+    auto bf16x2 = __float22bfloat162_rn({*reinterpret_cast<float*>(&x), *reinterpret_cast<float*>(&y)});


🔵 suggestion: cast_into_bf16_and_pack 使用 __float22bfloat162_rn，但 math.cuh 本身没有 #include <cuda_bf16.h>，仅靠 TU preamble（sm100_fp8_gemm_1d1d.hpp:41）先行包含。JIT 路径可编译，但头文件不自包含，且与 MR 描述所述的修复位置不符；建议直接在本头文件加 include。

_{🤖 v5}

ds-review-bot · 2026-08-26T03:16:49Z

+        .tc_util = device_runtime->get_tc_util(),
+        .compiled_dims = compiled_dims
+    };
+    const auto config = get_best_config<SM100ArchSpec>(desc);


🟡 warning: arch 12 路径复用 SM100ArchSpec（smem_capacity=232448）计算 num_stages/smem_size。若 GB10/sm_121 的 per-SM 动态 shared memory 上限低于该值，launch 会失败。鉴于 MR 描述承认该 kernel 在 GB10 上因 toolchain 限制未实际跑通，建议对 sm_12x 加容量校验或注释说明。

_{🤖 v5}

ds-review-bot · 2026-08-26T03:16:49Z

+        .tc_util = device_runtime->get_tc_util(),
+        .compiled_dims = compiled_dims
+    };
+    const auto config = get_best_config<SM100ArchSpec>(desc);


🔴 critical: get_best_config(desc) 会枚举 layout.swap_ab=true（以及 cluster_n=2 的 multicast-on-A）布局，但本 kernel 模板没有 kSwapAB 参数，且 static_assert 只支持 multicast-on-B（kIsMulticastOnA=true 且 kNumMulticast>1 会直接编译失败）。SM100ArchSpec 对 Normal fp8 GEMM 在 M 小、N 较大时会优先选 swap_ab 布局（例如 M=10、N=8192、num_sms 为偶数时，swap_ab+cluster_n=2 的单波布局会因 cluster_size 更大而胜出），此时生成的内核会编译失败；若选中 cluster_n=1 的 swap_ab 布局，则 storage/pipeline 按 swap_ab 配置（store_block_m=16、smem_cd 小），而内核按非 swap 布局使用 SMEM（STORE_BLOCK_M=min(BLOCK_M,128)），动态共享内存分配小于内核实际使用，导致越界。需要在此禁止 swap_ab 布局（或为纯 fp8 提供只枚举非 swap 布局的启发式）再传给本 kernel。

_{🤖 v4p}

ds-review-bot · 2026-08-26T03:16:49Z

+    // misread as fp4: silent corruption). Route to the fp8 1d1d kernel.
+    if (a_data.scalar_type() != kPackedFP4 and b_data.scalar_type() != kPackedFP4) {
+        if (swap_ab) {
+            sm100_fp8_gemm_1d1d(b_data, sfa, a_data, sfb, std::nullopt, d,


🔴 critical: SM120 纯 fp8 的 swap_ab 分支把 b_data/a_data 交换后调用 sm100_fp8_gemm_1d1d，但该 SM100 kernel 没有 AB-swap/转置 epilogue（其 make_tma_cd_desc 使用传入的 m 和 d.size(-1) 构造普通行优先 TMA descriptor）。该分支在 m<=16、无 c 且 d 连续时生效，kernel 实际计算的是 D^T（形状 eff_m x eff_n = n x m），却按非转置方式写入 D[m,n]；m_idx 会覆盖 0..n-1，而 D 只有 m 行，n>m 时直接越界写，n<=m 时结果也是转置错误。纯 fp8 不应走 swap_ab 路径，建议用原始 a_data/b_data 和原始 m/n 调用 sm100_fp8_gemm_1d1d，并相应重算 SF。

_{🤖 v4p}

ds-review-bot · 2026-08-26T03:16:49Z

        } else if (arch_major == 10 and sfa.scalar_type() == torch::kInt) {
-            sm100_fp8_fp4_gemm_1d1d(a.first, sfa, b.first, sfb, c, d, m, n, k, gran_k_a, gran_k_b,
-                                    major_a, major_b, compiled_dims);
+            if (b.first.scalar_type() != kPackedFP4) {


🟡 warning: arch10 的分派只判断 b.first.scalar_type() != kPackedFP4，会把 A=FP4、B=FP8 的混合情况也送进纯 fp8 kernel sm100_fp8_gemm_1d1d，绕过 sm100_fp8_fp4_gemm_1d1d 的 AB-swap 支持（MR 描述也说明 FP4 应保留给 combined kernel）。该 kernel 没有 kSwapAB，不能按 A=FP4/B=FP8 的交换布局处理。建议改为 a.first.scalar_type() != kPackedFP4 and b.first.scalar_type() != kPackedFP4，只让纯 fp8xfp8 走新 kernel。

_{🤖 v4p}

ds-review-bot · 2026-08-26T03:16:53Z

🤖 ds-review-bot Code Review

v6

恢复的内核未适配现有的两种 AB-swap 机制：SM12 小 M 路径会错误写回输出，而现代 SM100 启发式生成的 swap 配置可直接导致 JIT 编译失败。

v5

本 MR 恢复了 8b1392b 删除的纯 fp8 1d1d 路径（sm100_fp8_gemm_1d1d.hpp + dg25_fp8 自包含内核模块），并修正了 TU preamble（math.cuh 包含顺序、cuda_bf16.h、cute::is_same_v）。方向正确：模板实参与 kernel 26 个模板参数逐一对齐（含 cluster_size/cluster_n>1 → kNumMulticast/kIsMulticastOnA 在 swap_ab==0 布局下的正确映射）；dg25 模块自包含，与主树头文件的宏均有 #ifndef 守卫，无冲突。但存在 3 个正确性问题：(1) SM120 swap_ab 分支调用无转置存储支持的 sm100_fp8_gemm_1d1d，会在 m<=16 decode 场景写出转置/错位的 D；(2) 新版 SM100 heuristics 会为 Normal GEMM 枚举并可能选中 layout.swap_ab==1，而恢复的 runtime 完全忽略该字段（dg25 kernel 无 swap 支持），导致 store/smem 配置错配（静默错误）或 kIsMulticastOnA 静态断言编译失败；(3) arch 10 分支只检查 B 的 dtype，fp4_A x fp8_B 混合组合会被误路由到纯 fp8 kernel。次要问题：epilogue_type 参数被硬编码 EpilogueIdentity 忽略；math.cuh 本身仍未包含 cuda_bf16.h（仅 TU preamble）；SM100ArchSpec::smem_capacity=232448 复用于 sm_12x 设备的 stage 计算存在超限风险。跨仓一致性（vllm-project/vllm#53680）：无接口冲突，vLLM 仅回钉 a6b593d，fp8_gemm_nt 等 pybind 别名与签名两侧一致，本 MR 修复上述问题合入后可解除 pin。建议合入前修复问题 1-3。

v4p

本 MR 旨在恢复被 8b1392b 删除的 pure-fp8 1d1d kernel（csrc JIT 运行时 + deep_gemm/dg25_fp8 自包含模块），并修正 SM12x fp8 分派，使 fp8xfp8 不再落入会误读 FP4 的 sm120_fp8_fp4_gemm_1d1d。整体方向正确，NVRTC preamble 修复（cuda_bf16/cstdio 前置、cute::is_same_v）也有针对性；但当前分派仍有两个关键路径会把纯 fp8 送入不支持 swap 的 SM100 kernel，造成编译失败或越界写，另有 arch10 混合 FP4xFP8 误路由问题，建议修复后再合入。

Files reviewed: 18
配对 MR checkout: 1
Issues found: 🔴 7 critical | 🟡 2 warning | 🔵 2 suggestion
Inline comments posted: 11

… fix arch-10 routing Review feedback (ds-review-bot on deepseek-ai#419): - SM120 pure-fp8 no longer participates in the host-side AB-swap: the restored sm100_fp8_gemm_1d1d has no transposed-store support, so the swapped path wrote D^T as row-major (OOB/transposed for m<=16 decode). swap_ab now excludes pure_fp8 and the pure-fp8 branch always calls with the original m/n (SF transformed unswapped). - get_best_config<SM100ArchSpec> may select swap_ab==1 layouts for Normal GEMMs; the kernel has no swap template param. Added GemmDesc::allow_swap_ab (default true, preserving all existing callers); get_layout_candidates skips swap_ab==1 when false; the restored runtime sets it false. - arch-10 routing now checks both operands (fp4_A x fp8_B stays on the combined kernel), matching the SM120 branch. - Minor: epilogue_type wired back through get_default_epilogue_type; dg25_fp8/common/math.cuh now includes cuda_bf16.h itself; noted the SM100ArchSpec smem_capacity reuse for sm_12x.

maci0 · 2026-08-26T06:12:04Z

Thanks for the detailed review — all three criticals are addressed in 44d9d2e:

SM120 swap_ab (transposed D): pure fp8 is now excluded from the host-side AB-swap decision (swap_ab requires !pure_fp8). The pure-fp8 branch always calls sm100_fp8_gemm_1d1d with the original a/b and m/n, with SF transformed unswapped — the restored kernel has no transposed-store support, and the SM100 heuristics cover small M with block_m<=32 candidates.
SM100 swap layouts from heuristics: added GemmDesc::allow_swap_ab (default true, so all existing kernels keep their candidate set); SM100ArchSpec::get_layout_candidates skips the swap_ab==1 iteration when it is false; the restored runtime sets it false. This prevents both the silent store/smem mismatch and the kIsMulticastOnA static-assert JIT failure.
arch-10 mixed FP4×FP8 routing: the pure-fp8 branch now requires a and b to be non-FP4 (a.first.scalar_type() != kPackedFP4 and b.first.scalar_type() != kPackedFP4), matching the SM120 branch — fp4_A x fp8_B stays on the combined kernel.

Also folded in the minors: epilogue_type is wired through get_default_epilogue_type(args.epilogue_type) instead of the hardcoded literal; dg25_fp8/common/math.cuh now includes cuda_bf16.h itself (self-contained header); noted the SM100ArchSpec::smem_capacity reuse for sm_12x (GB10 matches SM100's 227 KB dynamic smem, so stage sizing is safe there).

Happy to re-run anything or adjust the allow_swap_ab approach if you'd prefer a different mechanism for filtering swap candidates.

All critical/suggestion comments from the v5/v6 reviews are addressed in 44d9d2e (AB-swap disabled for pure fp8, swap layout candidates filtered, arch-10 routing checks both operands, epilogue type wired back, math.cuh self-contained). Requesting re-review. Signed-off-by: maci0 <maci0@users.noreply.github.com>

maci0 · 2026-08-27T08:26:08Z

The re-review comments duplicate the earlier v4p/v5/v6 findings — all of them were addressed in 44d9d2e (the current head, plus the ci: re-trigger commit on top). Point-by-point:

AB-swap for pure fp8 (gemm.hpp, fp8_fp4_gemm_nt_sm120) — resolved: const bool pure_fp8 = (a_data.scalar_type() != kPackedFP4) and (b_data.scalar_type() != kPackedFP4) and the swap condition now includes !pure_fp8, so the restored kernel only ever runs non-swapped.
get_best_config<SM100ArchSpec> may select swap_ab=true (sm100_fp8_gemm_1d1d.hpp / heuristics) — resolved: GemmDesc.allow_swap_ab=false (config.hpp) is set in the restored kernel runtime, and get_layout_candidates (sm100.hpp:52) does if (swap_ab == 1 and not desc.allow_swap_ab) continue; — swapped layouts are never enumerated. The multicast-on-A hazard from the swapped cluster_n=2 path can no longer be reached; with swap_ab=0, cluster_n>1 is also filtered (sm100.hpp:93), so kIsMulticastOnA stays false and the kNumMulticast == 1 static assert is not tripped.
arch-10 routing checks only B's dtype (gemm.hpp:218) — resolved: the pure-fp8 route condition now checks both operands (a_data.scalar_type() != kPackedFP4 and b_data.scalar_type() != kPackedFP4), so fp4-A × fp8-B stays on the combined kernel.
sm_12x smem capacity — comment added in the kernel runtime: GB10 per-SM dynamic smem matches SM100 (227 KB), so SM100ArchSpec stage sizing is safe on sm_12x. (Runtime verification on GB10 is pending a toolchain-capable wheel; the kernel JIT path itself compiles.)
EpilogueIdentity hardcoded — resolved: get_default_epilogue_type(args.epilogue_type) is wired back in the template args.
math.cuh not self-contained — resolved: #include <cuda_bf16.h> added at the top of dg25_fp8/common/math.cuh.

Happy to adjust anything that still looks off on re-review.

lucifer1004 · 2026-08-27T11:06:41Z

SM12x SMEM capacity is 100 KiB only.

maci0 · 2026-08-27T14:36:50Z

Related upstream scale-path work in the same area, for cross-reference:

transform_sf_into_required_layout: no SM120/121 (arch_major=12) branch for (gran_mn=1, gran_k=32) NVFP4 expert scales -- blocks DeepSeek-V4 weight loading #372 / Fix SM120 scale-factor layout transformation #403 — transform_sf_into_required_layout has no arch-12 branch for (gran_mn=1, gran_k=32) NVFP4 expert scales (blocks DSV4 weight loading on SM120/121); Fix SM120 scale-factor layout transformation #403 allows explicit SM120 recipes to use the SM100 UE8M0 packing/validation.
Fix ue8m0 packing: mask mantissa bits when extracting fp32 exponents #337 — transpose_and_pack_fp32_into_ue8m0 leaks mantissa bits into adjacent exponent bytes when packing non-power-of-2 FP32 scales.

This restore (pure-fp8 1d1d) and those two are complementary: kernel path restore + packer/transform correctness on SM12x. Worth reviewing together.

SM100ArchSpec::smem_capacity (232448 = 227 KiB) is hardcoded; SM12x (GB10) only has 100 KiB of opt-in per-block SMEM, so stage sizing from the SM100 value over-allocates and the restored 1d1d kernel would not fit. get_pipeline_config now reads sharedMemPerBlockOptin from the device runtime, which is correct on both archs. Signed-off-by: maci0 <maci0@users.noreply.github.com>

maci0 · 2026-08-28T07:18:42Z

Thanks for the catch — fixed in 54d5a3e.

SM100ArchSpec::smem_capacity (232448 = 227 KiB) was hardcoded and reused for sm_12x; with GB10 at only 100 KiB of opt-in per-block SMEM, the restored 1d1d kernel would over-allocate stages. get_pipeline_config now sizes stages from the device runtime's sharedMemPerBlockOptin, which is correct on both SM100 and SM12x, and the misleading "GB10 matches SM100" comment in the kernel runtime is corrected.

maci0 · 2026-08-28T13:59:02Z

Reproducible recipe: https://github.com/maci0/vllm-spark-0731 — the SM12x findings live in docs/knowledge/09-golden-deepgemm.md and docs/UPSTREAM.md (DeepGEMM section); deepgemm-fp8-1d1d-port.diff is our local pre-merge port of this restore.

maci0 mentioned this pull request Aug 26, 2026

[Regression] 8b1392b removes SM12x pure-fp8 1d1d kernels and aliases fp8_gemm_nt to the fp4 dispatcher (silent corruption on GB10) #417

Open

ds-review-bot reviewed Aug 26, 2026

View reviewed changes

maci0 mentioned this pull request Aug 26, 2026

[Bugfix][DeepGEMM] Pin DeepGEMM to a6b593d to restore pure-fp8 1d1d kernels until DeepGEMM#419 vllm-project/vllm#53680

Closed

Conversation

maci0 commented Aug 26, 2026

Summary

What this restores

Header fixes the a6b593d TU needs to compile at all

Validation status (honest)

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

ds-review-bot commented Aug 26, 2026

🤖 ds-review-bot Code Review

v6

v5

v4p

Uh oh!

maci0 commented Aug 26, 2026

Uh oh!

maci0 commented Aug 27, 2026

Uh oh!

lucifer1004 commented Aug 27, 2026

Uh oh!

maci0 commented Aug 27, 2026

Uh oh!

maci0 commented Aug 28, 2026

Uh oh!

maci0 commented Aug 28, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Header fixes the `a6b593d` TU needs to compile at all