mirror of
https://github.com/deepseek-ai/DeepGEMM.git
synced 2026-11-05 00:47:26 +00:00
Labels
No labels
pull-request
No milestone
No project
No assignees
1 participant
Notifications
Due date
No due date set.
Dependencies
No dependencies set.
Reference
deepseek-ai/DeepGEMM#431
Loading…
Add table
Add a link
Reference in a new issue
No description provided.
Delete branch "%!s()"
Deleting a branch is permanent. Although the deleted branch may continue to exist for a short time before it actually gets removed, it CANNOT be undone in most cases. Continue?
📋 Pull Request Information
Original PR: https://github.com/deepseek-ai/DeepGEMM/pull/419
Author: @maci0
Created: 8/26/2026
Status: 🔄 Open
Base:
nv_dev← Head:fix/sm12x-pure-fp8-1d1d-restore📝 Commits (4)
405f341[SM12x] Restore pure-fp8 1d1d path (removed in8b1392b)44d9d2e[SM12x] Address review: no AB-swap for pure fp8, filter swap layouts, fix arch-10 routing14d3548ci: re-trigger ds-review-bot after review fixes (44d9d2e)54d5a3e[SM12x] Size pipeline stages from the device's actual SMEM capacity📊 Changes
20 files changed (+2611 additions, -4 deletions)
View changed files
📝
csrc/apis/gemm.hpp(+22 -3)📝
csrc/jit_kernels/heuristics/config.hpp(+4 -0)📝
csrc/jit_kernels/heuristics/sm100.hpp(+11 -1)➕
csrc/jit_kernels/impls/sm100_fp8_gemm_1d1d.hpp(+172 -0)➕
deep_gemm/include/deep_gemm/dg25_fp8/common/compile.cuh(+18 -0)➕
deep_gemm/include/deep_gemm/dg25_fp8/common/cute_tie.cuh(+50 -0)➕
deep_gemm/include/deep_gemm/dg25_fp8/common/epilogue_utils.cuh(+27 -0)➕
deep_gemm/include/deep_gemm/dg25_fp8/common/exception.cuh(+43 -0)➕
deep_gemm/include/deep_gemm/dg25_fp8/common/math.cuh(+150 -0)➕
deep_gemm/include/deep_gemm/dg25_fp8/common/reduction.cuh(+44 -0)➕
deep_gemm/include/deep_gemm/dg25_fp8/common/scheduler.cuh(+288 -0)➕
deep_gemm/include/deep_gemm/dg25_fp8/common/sm100_utils.cuh(+266 -0)➕
deep_gemm/include/deep_gemm/dg25_fp8/common/sm120_utils.cuh(+257 -0)➕
deep_gemm/include/deep_gemm/dg25_fp8/common/sm90_utils.cuh(+332 -0)➕
deep_gemm/include/deep_gemm/dg25_fp8/common/tma_copy.cuh(+92 -0)➕
deep_gemm/include/deep_gemm/dg25_fp8/common/tma_utils.cuh(+116 -0)➕
deep_gemm/include/deep_gemm/dg25_fp8/common/types.cuh(+43 -0)➕
deep_gemm/include/deep_gemm/dg25_fp8/common/types.hpp(+41 -0)➕
deep_gemm/include/deep_gemm/dg25_fp8/common/utils.cuh(+68 -0)➕
deep_gemm/include/deep_gemm/dg25_fp8/impls/sm100_fp8_gemm_1d1d.cuh(+567 -0)📄 Description
Summary
8b1392bdeleted 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 combinedsm120_fp8_fp4_gemm_1d1dkernel, 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:8b1392bpin: France output degenerates (' Septy Septy…'), DSpark draft acceptance collapses, decode ~25.8 → ~4.4 tok/s.a6b593dpin (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 modulegemm.hpp: fp8xfp8 on arch 10/12 routes tosm100_fp8_gemm_1d1d; fp4 keeps the combined kernel (AB-swap paths handled).Header fixes the
a6b593dTU needs to compile at allThe original 1d1d TU cannot be compiled as generated:
common/math.cuhwas never included by the 1d1d include chain →align/swapincommon/scheduler.cuh:192/common/sm100_utils.cuh:127were undefined (identifier "align" is undefined). The TU now includes it and pulls indeep_gemm::math(include before the impl header, so the using-directive precedes the dependent code).math.cuh'scast_into_bf16_and_packneeds#include <cuda_bf16.h>(__float22bfloat162_rn).std::is_same_v→cute::is_same_vfor NVRTC (libcu++ does not expose variable templates instdunder NVRTC).Validation status (honest)
compute_121a(~231 KB) — the preamble fix is verified.tcgen05.mma/.cta_group::1/.block32for 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).fp8_einsumo_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_fp8or fold into the existing tree) to match maintainer preference.🔄 This issue represents a GitHub Pull Request. It cannot be merged through Gitea due to API limitations.