Skip to content

feat(sm120): vendor DeepGEMM-sm120 v0.1.0 device layer - #10

Merged
zyongye merged 5 commits into
vllm-project:devfrom
lucifer1004:vendor/sm120-deepgemm-sm120-v0.1.0
Sep 22, 2026
Merged

zyongye merged 5 commits into
vllm-project:devfrom
lucifer1004:vendor/sm120-deepgemm-sm120-v0.1.0

Conversation

@lucifer1004

@lucifer1004 lucifer1004 commented Sep 16, 2026 •

Copy link
Copy Markdown

Summary

Replace this fork's in-tree SM120 device headers with byte-identical vendored copies of lucifer1004/DeepGEMM-sm120 v0.1.4 — a standalone single-source-of-truth repository for the SM120 device layer, shared with the nv_dev-lineage SM120 port (deepseek-ai#447). Fixes and new SM120 kernels will land there once and propagate to each fork mechanically, instead of diverging per fork.

  • Nothing vendoring-specific enters this fork — no manifest, no tooling, no header markers; the vendored files are plain DeepGEMM-style headers. Provenance and drift control live on the canonical repo's side: per-tag sha256 manifest, tools/check_vendor.py --fork <path> (both directions), and a scheduled watcher that monitors registered downstream branches daily and files an issue when a tree matches no release tag.
  • 13 existing headers replaced, 3 added (impls/sm120_fp8_fp4_sparse_mqa_logits.cuh, impls/sm120_padding.cuh, scheduler/sm120_gemm.cuh); the vendored files are a strict superset of this fork's versions (both lineages' fixes merged).
  • Host launchers adapted to the superset kernel signatures (no behavior change for configurations this fork exercises).

Host adaptation details

  • sm120_bf16_gemm / sm120_fp8_fp4_gemm_1d1d: pass the new runtime args with neutral values (EpilogueIdentity{}, stride_c_m = 0, shape_cd_m = m); template additions are trailing-defaulted.
  • Masked m-grouped launchers now force swizzle_cd_mode = 0 when m % BLOCK_M != 0. This is required with the vendored kernels: they dropped the old boundary-tile scalar fallback (superseded by the nv_dev-lineage mechanism), so a boundary tile's full-tile TMA store would spill into the next group's rows. The scalar store path carries the masked row bounds (row_is_valid). Mirrors the proven nv_dev-lineage glue.
  • sm120_split_k_reduce: the removed kWithAccumulation template flag maps to the runtime C operand (gmem_c = D when accumulating, else nullptr).
  • Paged MQA logits: BLOCK_KV → PAGE_KV rename is semantics-preserving here (the fork already passes the real page size); the vendored kernels add within-page TMA offsets enabling PAGE_KV 128/256.

New in the vendored layer (available, inert until host glue opts in)

  • CUDA ≥ 13 compile-time guard for SM120 device passes (a pre-13 ptxas was reported to silently drop block_scale).
  • tensor_map_replace_global_dim_in_smem relocated into common/sm120_utils.cuh (main-lineage ptx/tma.cuh lacks it).
  • Sparse MQA logits kernels, padding kernels, PAGE_KV 128/256 support, HeadSplits epilogue, K-grouped constant-stride/PSUM layouts.

Follow-up folded in (k-grouped guard)

  • K-grouped GEMM (bf16 + fp8/fp4) with m % BLOCK_M != 0 could spill a boundary tile's TMA store into the next group's output slab (reproduced 5/5: next group's head rows zeroed). Both k-grouped launchers now force the scalar store epilogue for unaligned m, mirroring the nv_dev-lineage glue; regression tests test_sm120_kgroup_unaligned_m_tail_tile_isolation added to the bf16 and fp8/fp4 suites. This exposure predates the vendor migration (the old kernels had no k-grouped boundary fallback either).

Validation (sm_120a)

  • tools/check_vendor.py: PASS (tree byte-identical to v0.1.4).
  • SM120 suite: 48/48 passed (46 baseline + 2 k-grouped guard) in plain mode at the current head. Memcheck (0 errors) and racecheck (0 hazards) evidence is from the 50-node three-mode run at the previous head: its sanitizer executions covered exactly these tests, and the dropped nodes are test-only (the removed host checker was inert unless env-enabled). The labels-contract checker and its regression tests are deferred to upstream feat(gemm): opt-in labels-contract checker for m-grouped contiguous GEMM deepseek-ai/DeepGEMM#450 and will arrive via sync rather than as a fork-local copy.
  • The vendored tree additionally passes the nv_dev-lineage SM120 suite (741 tests) on the sibling fork; the standalone repo's CI cross-compiles every header plus one representative instantiation per kernel for sm_120a against both base lineages and asserts on SASS/PTX opcodes (HMMA.16816, QMMA.*.SF.*, block_scale).

Merge-base friction

This branch merges into deepseek-ai/DeepGEMM main with zero conflicts (dry-run); all SM120-exclusive content stays in dedicated files, and VENDOR-sm120.json / tools/check_vendor.py are additive.

Follow-up folded in (v0.1.4 re-vendor)

Follow-up folded in (v0.1.5 re-vendor)

Re-vendored the device layer at DeepGEMM-sm120 v0.1.5 (3f4c73e): ports deepseek-ai#453 (thanks @Sunt-ing) and extends it to every site with the same hazard — 14 fence_view_async_shared() insertions across 8 SM120 kernel headers. The SM120 kernels read TMA-written smem through ldmatrix/ld_shared (generic proxy), so the empty-barrier arrive handing a stage back to the TMA producer must be ordered after those reads; otherwise the next TMA can overlap the consumer's reads (generic/async proxy WAR hazard). Sites: bf16_gemm (the deepseek-ai#453 report), bmk_bnk_mn, tf32_hc_prenorm_gemm, fp8_fp4_gemm_1d1d (×3), fp8/fp4 dense MQA logits (empty_kv + empty_q each), fp8/fp4 paged MQA logits (empty_q stage release + empty_kv each). Sparse MQA is exempt (legacy cp.async producer). Perf A/B on RTX PRO 6000 (bf16 GEMM, fp8 dense/paged MQA, 3 runs each): all deltas within run-to-run noise. Re-verified on sm_120a from a fresh JIT cache at this head: test_sm120_mqa.py + test_sm120_fp8_fp4.py 23/23 passed; tools/check_vendor.py PASS (tree byte-identical to v0.1.5).

Replace the in-tree SM120 device headers with byte-identical vendored
copies from lucifer1004/DeepGEMM-sm120 v0.1.3 (merged superset of both
lineages). Marker-free vendoring: no manifests, tooling, or header
comments enter this fork; provenance and drift control live on the
canonical repo's side (per-tag manifest, check_vendor.py --fork,
scheduled downstream watcher).

Signed-off-by: Zihua Wu <13583761+lucifer1004@users.noreply.github.com>
The vendored DeepGEMM-sm120 v0.1.0 device layer widens several kernel
ABIs; adapt the SM120 host glue with neutral mappings:

- sm120_bf16_gemm: pass the epilogue operator as a runtime argument
  (EpilogueArgs, same marshalling as SM100) and split the old
  int64_t stride_cd_m/stride_cd_batch into u32 stride_d_m/stride_c_m/
  stride_d_batch; stride_c_m=0 keeps C sharing D's row stride.
- sm120_fp8_fp4_gemm_1d1d: pass shape_cd_m (= shape_m), the runtime
  epilogue argument, and stride_c_m (= 0) after stride_cd_batch.
- sm120_split_k_reduce: the kWithAccumulation template flag is gone;
  instantiate the 3-arg template and map accumulation onto D in place
  to the new runtime C operand (gmem_c = D with D's own strides,
  with_alpha = false).
- Masked m-grouped GEMM (bf16 + fp8/fp4): disable the TMA-store
  epilogue (swizzle_cd_mode = 0) when m % BLOCK_M != 0. The vendored
  kernels dropped the old masked-boundary scalar fallback, so a
  boundary tile's full-tile TMA store would spill into the next
  group's rows; the scalar store path carries the masked row bounds.
  This mirrors the proven nv_dev-lineage launchers.

Signed-off-by: Zihua Wu <13583761+lucifer1004@users.noreply.github.com>
The vendored v0.1.0 GEMM kernels have no boundary-tile scalar fallback
in the TMA-store epilogue, and the k-grouped D descriptor is one flat
2D map over all groups, so TMA can only clamp at the outermost M dim:
with m % BLOCK_M != 0, a group's partial tail M tile TMA-store spills
into the next group's slab. Mirror the masked-launcher guard (and the
nv_dev-lineage reference glue) in both k-grouped launchers: force
swizzle_cd_mode = 0 (group-bounded scalar store path) for unaligned m.

Repro evidence (num_sms=2, groups=2, m=48, n=128, ks=[8192, 128]; the
asymmetric K makes group 0's tile store last, so the spill overwrites
group 1's head rows deterministically):
- before: bf16/bf16-out 5/5 iterations contaminated (d[1][0:16] = 0.0,
  expected 768); fp8/fp32-out 5/5 contaminated (d[1][0:16] = 0.0,
  expected 771)
- bf16/fp32-out was already clean (0/5): the bf16 kernel gates its
  TMA-store epilogue on sizeof(cd_dtype_t) <= 2, so fp32 outputs take
  the scalar path; the fp8 kernel has no such dtype gate
- after: all three configurations 0/5 contaminated

Adds test_sm120_kgroup_unaligned_m_tail_tile_isolation to
tests/test_sm120_bf16.py (fp32 and bf16 outputs) and
tests/test_sm120_fp8_fp4.py (NT and TN layouts), asserting per-group
correctness including slab head rows plus flat-storage guard regions.
Full SM120 suite: 48 passed (46 baseline nodes + 2 new). Memcheck on
the two affected files: 20 passed, 0 sanitizer errors.

Signed-off-by: Zihua Wu <13583761+lucifer1004@users.noreply.github.com>
@lucifer1004
lucifer1004 force-pushed the vendor/sm120-deepgemm-sm120-v0.1.0 branch 3 times, most recently from 2cf7301 to 0aecffa Compare September 16, 2026 23:15
@lucifer1004

Copy link
Copy Markdown
Author

Force-pushed (2cf7301 → 0aecffa): dropped the fork-local labels-contract checker (DG_CHECK_CONTIGUOUS_LABELS) and its two regression nodes. Both are deferred to upstream deepseek-ai#450 and will arrive via sync, keeping this fork free of a divergent copy. Re-verified at the current head: SM120 suite 48/48 passed on sm_120a (plain mode, fresh JIT cache); the test_filter collection errors are the pre-existing helper-discovery quirk (deepseek-ai#446), not failures.

FP8 paged MQA logits gains PAGE_KV=32 (BLOCK_KV derived as
min(PAGE_KV, 64), mirroring the FP4 sibling). Device-only update: this
fork's host launcher still restricts paged FP8 to page 64, so page32
stays inert until host glue opts in (vllm-project#14).

Validated on sm_120a: test_sm120_mqa.py + test_sm120_fp8_fp4.py 23/23
passed from a fresh JIT cache (only the pre-existing test_filter
collection quirk remains, deepseek-ai#446).

Signed-off-by: Zihua Wu <13583761+lucifer1004@users.noreply.github.com>
Vendor lucifer1004/DeepGEMM-sm120 v0.1.5 (b31a688): 14
fence_view_async_shared() sites across 8 SM120 kernel headers, ordering
generic-proxy smem reads (ldmatrix/ld_shared) before the empty-barrier
arrive that hands each stage back to the TMA producer. Ports
deepseek-ai#453 (bf16 GEMM) and fixes the same hazard found by
audit in bmk_bnk_mn, tf32_hc_prenorm, fp8_fp4_gemm_1d1d, and the
fp8/fp4 dense+paged MQA logits kernels. Sparse MQA is exempt (legacy
cp.async producer).

Perf A/B on RTX PRO 6000: no regression beyond run-to-run noise.

Signed-off-by: Zihua Wu <13583761+lucifer1004@users.noreply.github.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants