[autoresearch] Goal: fold block-scaled FP8 into DeepGEMM's MegaMoE (the H200-now subset) - #6
Closed
wseaton wants to merge 127 commits into
Closed
Conversation
* Fix multicast bug and profile masked GEMM * Updates and lint --------- Co-authored-by: Kuai Yu <yukuai@deepseek.com> Co-authored-by: Chenggang Zhao <chenggangz@deepseek.com>
Signed-off-by: simon-mo <simon.mo@hey.com>
* build: Minor tweeks for wheel build Signed-off-by: oliver könig <okoenig@nvidia.com> * ci: Workflows for wheel build Signed-off-by: oliver könig <okoenig@nvidia.com> * fix Signed-off-by: oliver könig <okoenig@nvidia.com> * fix Signed-off-by: oliver könig <okoenig@nvidia.com> * build: Add CachedWheel Signed-off-by: oliver könig <okoenig@nvidia.com> * add version to init Signed-off-by: oliver könig <okoenig@nvidia.com> * revert Signed-off-by: oliver könig <okoenig@nvidia.com> * revert Signed-off-by: oliver könig <okoenig@nvidia.com> * revert Signed-off-by: oliver könig <okoenig@nvidia.com> * v2 Signed-off-by: oliver könig <okoenig@nvidia.com> * update Signed-off-by: oliver könig <okoenig@nvidia.com> * test Signed-off-by: oliver könig <okoenig@nvidia.com> * from packaging.version import parse Signed-off-by: oliver könig <okoenig@nvidia.com> * local version Signed-off-by: oliver könig <okoenig@nvidia.com> * remove file Signed-off-by: oliver könig <okoenig@nvidia.com> * revert Signed-off-by: oliver könig <okoenig@nvidia.com> * Updates and lint * revert missing cudaextension args Signed-off-by: oliver könig <okoenig@nvidia.com> * Add timeout * fix DG settings Signed-off-by: oliver könig <okoenig@nvidia.com> * DG_USE_LOCAL_VERSION Signed-off-by: oliver könig <okoenig@nvidia.com> * Update version * Detect local changes * Minor fix * Revert CUTLASS * Unify options --------- Signed-off-by: oliver könig <okoenig@nvidia.com> Co-authored-by: Chenggang Zhao <chenggangz@deepseek.com>
* build: Allow NGC builds Signed-off-by: oliver könig <okoenig@nvidia.com> * reduce grid Signed-off-by: oliver könig <okoenig@nvidia.com> * update grid Signed-off-by: oliver könig <okoenig@nvidia.com> * fix Signed-off-by: oliver könig <okoenig@nvidia.com> * upgrade cuda action Signed-off-by: oliver könig <okoenig@nvidia.com> * remove test Signed-off-by: oliver könig <okoenig@nvidia.com> * py3.8 Signed-off-by: oliver könig <okoenig@nvidia.com> * fix Signed-off-by: oliver könig <okoenig@nvidia.com> * exclude Signed-off-by: oliver könig <okoenig@nvidia.com> * fix Signed-off-by: oliver könig <okoenig@nvidia.com> * torch-version Signed-off-by: oliver könig <okoenig@nvidia.com> * py3.8/torch2.1/cuda12.3 Signed-off-by: oliver könig <okoenig@nvidia.com> * Update publish.yml * fix grid Signed-off-by: oliver könig <okoenig@nvidia.com> * fix Signed-off-by: oliver könig <okoenig@nvidia.com> * cuda11.8 Signed-off-by: oliver könig <okoenig@nvidia.com> * no hopper for 118 Signed-off-by: oliver könig <okoenig@nvidia.com> * fix Signed-off-by: oliver könig <okoenig@nvidia.com> * fix Signed-off-by: oliver könig <okoenig@nvidia.com> --------- Signed-off-by: oliver könig <okoenig@nvidia.com>
* py3.8 Signed-off-by: oliver könig <okoenig@nvidia.com> * chore: Rename from `deep_geem` to `deepgemm` Signed-off-by: oliver könig <okoenig@nvidia.com> --------- Signed-off-by: oliver könig <okoenig@nvidia.com>
Signed-off-by: oliver könig <okoenig@nvidia.com>
Signed-off-by: oliver könig <okoenig@nvidia.com>
…me-change Ko3n1g/chore/revert name change
…bf16_bias support bf16 bias in deepgemm2
Remove tests/sm120/ directory (development-only tests, not for upstream). Fix paged MQA test enumeration that accidentally restricted SM90's block_kv and next_n options — SM90 retains (32,64) and (1,2,4).
Enable split-K parallelism for AB-swap dense GEMM (M<=16), improving small-M decode latency by ~2x (e.g. M=1 N=2112 K=7168: 37.7->19.8 us). Three bugs fixed: - kSFTileKBlocks was hardcoded to 4 in the split-K heuristic, but the kernel's actual SF tile spans (4 * gran_k / block_k) k-blocks, which is 8 when gran_k=128 and BK=64. Misaligned split-K partitions read wrong SF bytes, causing ~2% cosine distance errors. - kSFKAlignment in the K-grouped scheduler used kGranKA only, ignoring kGranKB. Asymmetric gran_k would produce wrong SF offsets for B. - K-grouped assert required ks[i] % 128 == 0, but BK=64 only needs ks[i] % 64 == 0. Changes: - Add max_gran_k field to GemmDesc for split-K SF alignment calculation - Reduce kernel supports strided output (stride_cd_m/n) for swap path - Split-K epilogue uses per-element N bounds check (shape_n=1 fix) - Guard split-K against KernelNoSF (BF16) and kSFTileKBlocks==0
Remove the sizeof(cd_dtype_t) <= 2 guard that prevented TMA store for FP32 output. SM90/SM100 already use TMA store for FP32 unconditionally. The sub-tile epilogue (store_block_m=64) is automatically selected by the heuristic to keep 2 pipeline stages within the 99KB SMEM budget. Fix the SMEM D scatter to use pair_store_t (uint64_t for FP32, uint32_t for BF16) so both elements of a pair are written correctly. Wgrad (TN FP32 acc=1) improvement: +9% to +30% across shapes.
Signed-off-by: Jin Li <59594262+liji-nv@users.noreply.github.com>
Fix IMA guard in paged MQA logits scheduler
The FP32 TMA store epilogue (commit 553fe80) exposed a latent bug in the batched accumulation path of fp8_einsum 'bhd,bhr->hdr': 1. Double-count: the TMA-store epilogue read gmem_c into registers AND used SM90_TMA_REDUCE_ADD_3D, adding C (z_0) twice. Skip the register read for batched accumulation since REDUCE_ADD already accumulates into global C. 2. Swap exclusion: the einsum AB-swap path didn't exclude accumulation, so small head dims (d<=32) routed batched+accumulation through swap. The swapped/strided output is mishandled by both epilogue paths (wrong result for block_n=16, illegal memory access for block_n=32). Mirror the dense GEMM swap exclusion by requiring !c.has_value(). Verified: bhd,bhr->hdr passes for d in {16,32,64,1024} (diff ~0.0007); full einsum suite passes.
…rnels
The cuBLAS comparison kernels were only matched by ('nvjet', 'reduce'), so
shapes where cuBLAS picks a different kernel showed 0.00x (and crashed
test_cublaslt_gemm with ZeroDivisionError):
- FP8 large/small shapes fall back to sm89_xmma kernels on SM120.
- BF16 uses cutlass *gemm* kernels, and M=1 uses a dot_kernel +
reduce_1Block_kernel GEMV path.
Match all observed kernel name substrings and guard against zero time.
FP4/mixed shapes still show 0.00x by design (cuBLAS has no FP4 GEMM).
When num_q_blocks < num_sms (small query count), the persistent q-block grid leaves most SMs idle (S=128 → 64 q-blocks → ~34% of 188 SMs). Since logits[q,kv] are independent across kv (no reduction over KV), split each q-block's KV range across gridDim.y cooperating blocks that write disjoint kv-subranges — no combine, no atomics, no workspace. Implemented via a 2D grid: blockIdx.y/gridDim.y carry the KV split, so no kernel signature change is needed and SM90/SM100 (launched with gridDim.y=1) are byte-identical. The host picks kv_splits by minimizing waves/kv_splits (wave-aware, avoids overshooting a wave boundary), only when q-blocks underfill the SMs. Measured (FP8 dense, H=64, D=128): S=128 SKV=8192 217→397 TFLOPS (+83%), S=128 SKV=4096 196→335 (+71%). S>=256 resolve to kv_splits=1 (identical kernel) → no regression. FP4 dense uses a separate launcher, unaffected. Adds S=128/512 to the dense MQA test for split-path coverage.
The SM120 FP8/FP4 paged_mqa_logits kernels carried an incomplete port of
SM100's odd-next_n (kPadOddN) handling, causing a CUDA illegal access
("Invalid __shared__ read of 16 bytes") at non-varlen odd next_n >= 3
(e.g. DeepSeek-V4 MTP next_n=3). Bring SM120 to full SM100 parity:
- kNumNextNAtoms: drop a spurious "+1" in the kPadOddN branch; use
math::constexpr_ceil_div(kNextN, kNextNAtom). The consumer must decode the
same stride the host metadata kernel builds the schedule with
(num_next_n_atoms = ceil_div(next_n, 2)); the "+1" desynced the
block_table/context_lens decode.
- q_idx "no previous task" sentinel: batch_size -> batch_size * kNumNextNAtoms
in both the TMA and math warps. Valid atom indices span
[0, batch_size*kNumNextNAtoms); once kNumNextNAtoms >= 2 the old batch_size
sentinel can equal a real atom index, so the math warp's first task skips
the Q-pipeline wait and reads smem_q with an uninitialized q_stage_idx --
the reported invalid shared read. Dormant at next_n in {1,2} (kNumNextNAtoms == 1).
- Add the kPadOddN store arm: the odd leftover atom (last of each next_n
group) holds one real token; storing kNextNAtom rows wrote one logits row
past the tensor. Use compute_and_store(Int<1>) for that atom.
The odd atom's Q is still loaded as a full kNextNAtom box: the out-of-bounds
2nd token is zero-filled by TMA (CU_TENSOR_MAP_FLOAT_OOB_FILL_NONE) and is
never read once the store arm processes only token 0 -- identical to SM100,
so no special one-token descriptor is needed.
Host (csrc/apis/attention.hpp): remove the SM120-only "next_n <= 2" gates.
SMEM is sized per-atom (<= 2 tokens) and is independent of next_n, so SM120
has no next_n cap (full SM100 parity). Test: arch_major in (10, 12) now share
the (1,2,3,4,5,6) next_n arm.
Validated on sm_120a (RTX PRO 6000): test_attention.py::test_paged_mqa_logits
108/108 pass (12 cases each next_n 1-6, FP8+FP4); compute-sanitizer memcheck
reports 0 errors for odd next_n {3,5}.
* Optimize FP8/FP4 Mega MoE dispatch, scheduling, and shared memory layout * Add BF16 accumulation/output support updates for GEMM paths * Improve paged MQA scheduler and attention coverage * Minor fixes and test updates
…refactorings) into nv_dev Brings deepseek-ai PR deepseek-ai#347 (88965b0, also pulls 714dd1a 'Update test_mega_moe.py') onto nv_dev (currently at ac1f285 = deepseek-ai#316 merge deepseek-ai#328 + PR deepseek-ai#342 IMA guard). Merge base is 891d57b. Conflicts resolved (5 files): * scheduler/paged_mqa_logits.cuh: take deepseek-ai#347's refresh_num_kv_and_advance + reversed metadata allocation; drop nv_dev's PR deepseek-ai#342 exist_q_atom_idx guard and get_atom_advance (both subsumed — reversed alloc keeps current_q_atom_idx in-bounds, refresh_num_kv_and_advance reproduces the varlen 1-or-2 advance). NV's host-passed num_next_n_atoms survives in the auto-merged metadata kernel (verified: single param, no inline recompute, coherent with reversed alloc). * csrc/apis/gemm.hpp: drop redundant C/D dtype asserts (deepseek-ai#347 cleanup; the pre-existing d-dtype check + d==c check already cover them). * sm100_fp8_fp4_gemm_1d1d.cuh: take deepseek-ai#347's unconditional BF16/FP32 C/D assert (supersedes nv_dev's accumulation-only assert). * tests/generators.py: take deepseek-ai#347's broadened BF16-accumulation enumeration (out_dtype==bf16 and (dtype==bf16 or arch==10)) — strict superset of nv_dev's fp8+SM100 case. * tests/test_attention.py: merge both — keep NV's SM90 next_n=4 multicast + num_clusters metadata sizing; add deepseek-ai#347's pool-size limits, batch_size=4096, and block_table/context_lens sanity asserts. Auto-merged metadata kernel + smxx_fp8_fp4_paged_mqa_logits.hpp + attention.hpp verified coherent across all three layers (kernel signature / launch / host call all carry num_next_n_atoms + is_varlen + indices). Validated on B-card (SM100) and H200 (SM90) incl next_n=4 multicast. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
…_dev_347 Sync nv_dev with upstream deepseek-ai#347
Introduce sm100_fp8_mqa_logits_f16_weights, an FP16-weights version of the FP8 MQA logits kernel, and wire it up through the attention API. The kernel is selected based on the weights dtype. Also update test_mqa_logits to avoid f16 overflow when exercising the new path.
When all context_lens are zero (e.g. warmup decode with kv_lens=1 and indexer_compress_ratio=4 giving gen_seq_lens=0), the binary search pushes lo to batch_size. The subsequent prefix_sum[q_idx] access is OOB when q_idx == batch_size == kAlignedBatchSize, since prefix_sum only has indices [0, batch_size). This causes a deterministic IMA during DSv4-Pro disagg GEN warmup at batch_size=32 (warmup step: 32 tokens, 32 gen). Fix: clamp q_idx to batch_size - 1. When the binary search lands past all queries (total work = 0), num_segs_q correctly evaluates to 0, and the existing guards (num_segs_q > 0 ? ...) produce the sentinel values (atom_idx=0, kv_split_idx=0). Signed-off-by: Barry Kang <43644113+Barry-Delaney@users.noreply.github.com>
…etadata-oob Fix OOB in paged MQA logits metadata scheduler
Brings in 9 upstream nv_dev commits: paged-MQA metadata OOB fix (deepseek-ai#353), IMA guard in the paged-MQA scheduler (deepseek-ai#342), the FP16-weights FP8 MQA logits kernel, the mega-MoE refactor (deepseek-ai#347), and assorted updates. Conflicts resolved (3 files): - csrc/apis/attention.hpp (dense MQA dispatch): union of both sides — keep SM120 FP4 (sm120_fp4_mqa_logits) and SM120 FP8 (arch 12) arms, and add nv_dev's FP16-weights arm (sm100_fp8_mqa_logits_f16_weights, SM100-only, already guarded by the arch==10 host assert). - deep_gemm/include/deep_gemm/scheduler/paged_mqa_logits.cuh: take nv_dev's refactored scheduler. SM120's one-line OOB guard is subsumed by nv_dev's reversed-metadata allocation (the deepseek-ai#342/deepseek-ai#353 fix). The refactor replaces get_num_kv/get_atom_advance with refresh_num_kv_and_advance + get_last_advance. - tests/test_attention.py: keep SM120 small-S split-KV MQA shapes and add the upstream large-KV/32-head regime; next_n = (1..6) on SM120, (1,2,4,5,6) on SM100, (1,2,4) on SM90; take nv_dev's paged test body (memory caps + peak-mem reset). Semantic fixes required by the merge (no textual conflict, would not compile/run otherwise): - sm120_fp8/fp4_paged_mqa_logits.cuh: migrate the scheduler call get_atom_advance(next_q_idx, batch_size) -> get_last_advance() to match nv_dev's refactored scheduler API (same migration nv_dev applied to the SM100 paged kernel). - tests/test_attention.py: gate the FP16-weights path to arch 10. Upstream set FP16 weights unconditionally, which dispatches to the SM100-only kernel and trips the arch==10 host assert when test_attention runs on SM90/SM120. Validated on H200 (sm90): clean build; test_bf16, test_fp8_fp4, test_attention all pass. SM120 (arch 12) kernels are JIT and arch-gated, so they are neither compiled nor run on H200 — they require an RTX 6000 run before the final PR.
The FP8/FP4 1D1D epilogue has two store paths: a swizzled TMA-store path
(kUseTMAStoreEpilogue, writes via tensor_map_cd) and a strided global-store
path (honors stride_cd_n). The TMA path's tensor map assumes a contiguous-N
output and ignores stride_cd_n, so it cannot express the AB-swap path's
transposed output (stride_cd_n = original-N stride != 1).
kUseTMAStoreEpilogue depends only on swizzle_cd_mode and BLOCK_N, where
swizzle_cd_mode = 128 once BLOCK_N * sizeof(cd) >= 128. The dense GEMM swap
caps M at 16 (BLOCK_N = 16 -> swizzle 0 -> strided store, correct). The BMM
swap (fp8_einsum) caps M at 32: with M in 17..32 and an FP32 output,
BLOCK_N = 32 -> swizzle_cd_mode = 128 -> the TMA-store path runs for a
transposed output and writes out of bounds.
Repro (RTX PRO 6000): fp8_einsum('bhr,hdr->bhd') with batch b = 32 and an FP32
output raises CUDA illegal memory access; b <= 16, or BF16 output (swizzle
stays 0), are unaffected. Not covered by the test suite, which uses BF16
output for the swapping expressions.
Fix: make the invariant explicit. Add GemmDesc::cd_n_contiguous; the SM120
swizzle_cd_mode heuristic requires it, and the dense + BMM swap launchers set
it false, so the swap always uses the strided-store epilogue. The swap path is
small-M and latency-bound, so the TMA-store epilogue is not a loss there.
Validated on RTX PRO 6000 (arch 12): the repro now passes (b=32 FP32
diff 7e-4); test_split_k_swap, test_einsum, test_fp8_fp4 still pass.
…r changes The SM120 MQA-logits work modified the shared launchers in two ways that broke SM100 (and were latent on SM90), surfacing as test_attention failures on B300 (baseline=origin/nv_dev passes, sm120 HEAD fails). No SM100/SM90 device kernel (.cuh) changed, so this is purely a host launch-config regression. 1. Dense MQA SMEM under-allocation (smxx_fp8_fp4_mqa_logits.hpp). The SMEM size dropped the `(num_math_threads / 128) * 2` mbarrier pairs that the SM90/SM100 kernels still allocate (the SM120 kernel does not). On SM100 this under-sized the dynamic SMEM; compute-sanitizer reports an "Invalid __shared__ write" in sm100_fp8_mqa_logits, faulting (CUDA_ERROR_ILLEGAL_ADDRESS) at large configs such as seq_len_kv=130560. Restore the term for arch != 12. 2. Paged-MQA metadata capacity assert (smxx_fp8_fp4_paged_mqa_logits.hpp). An unconditional `smem_size <= SM120ArchSpec::smem_capacity` (99 KB) was added to the metadata launcher, which runs on every arch. On SM90/SM100 (228 KB) a large (varlen) batch's metadata SMEM legitimately exceeds 99 KB and falsely tripped the assert. Gate the capacity check to the running arch. Verified on B300 (sm100): compute-sanitizer memcheck clean on the previously out-of-bounds config (seq_len=510, seq_len_kv=130560); full test_attention (dense + paged, NextN 1-6, FP8/FP4) passes. Fix (1) also corrects a latent SMEM under-allocation on SM90.
Strip changes that should not ship upstream and align comment style with the existing codebase (terse, no decorative banners). No change to production GEMM/attention behavior — only removes a dev-only bench path and trims comments. - .gitignore: revert to upstream; personal ignores (docs_internal/, benchmarks/, ncu_reports/, internal/) moved to local .git/info/exclude. - Remove dev-only sm120_fp8_gemm_bench API and its override_layout tile-override param (gemm.hpp, __init__.py, sm120_fp8_fp4_gemm_1d1d.hpp); config selection collapses to the standard get_best_config path. - Trim changelog-/rationale-style block comments to single lines (shared rationale lives in commit history): MQA-logits SMEM, paged-MQA capacity assert, cd_n_contiguous, split-K SF alignment, latency model, einsum MN-major notes. - De-duplicate the AB-swap "transposed output" rationale (kept once in config.hpp) and the UE8M0 SF-tile note. - Remove decorative // ==== banner comments from sm120 .cuh (no other kernel uses them). - Remove an unused num_waves local in the SM120 latency model (clears a -Wunused-variable warning). - Remove tests/test_split_k_swap.py (standalone dev harness, not pytest-style).
The SM120 fp8/fp4 1D1D kernel iterated every ceil(M_sum/BLOCK_M) m-block and computed invalid (m_indices == -1) padding tiles against expert 0. SM90 skips these via Scheduler::is_computation_valid, but that helper is marked SM90-only and the SM120 producer issued the A/B/SF TMA copies unconditionally. At MoE decode the contiguous worst-case M_sum reserves a block per local expert (min(M*topk, local_experts)) while only a few are routed, so most m-blocks are empty padding. Computing them wastes a full-width GEMM tile — the dominant cost of EP decode. Measured on RTX PRO 6000, EP gate_up (N=4096, K=4096, 64 local experts, ~20 routed): worst-case 20-valid + 45-empty of 65 blocks = 304us vs the 20 real blocks alone = 170us; the empty tiles cost 134us (44%). Add the m_indices < 0 skip-continue to both the producer and consumer while-loops. The check is identical in both, so no barrier ops are issued for skipped blocks and the warp-specialized pipeline stays in sync. Empty experts now cost nothing: EP decode MoE GEMM drops ~1.8x (304 -> 169us), matching the TP-sharded layout. The all-valid path (prefill/dense) is unchanged. Validated: m-grouped contiguous correctness for fp8xfp8 / fp4xfp4 / fp8xfp4 (diff < 2%), worst-case-padded result == tight, all-65-valid timing unchanged.
Adds the reverse of the existing fp8-A x fp4-B mixed path: A is fp4 (e2m1), B is fp8 (e4m3), via mxf8f6f4 .e2m1.e4m3 (m16n8k32). Lets the expert weight sit on the well-filled M axis and the (few) decode tokens on a small N axis. - mma/sm120.cuh: fp4_fp8_mixed_mma_block_scaled (.e2m1.e4m3). - common/sm120_utils.cuh: ldmatrix_m8n16_x4_b4x16_p64 + load_a_fragment_b4x16 (4-reg b4x16 unpack for the fp4 A operand; same addressing as the fp8 A load). - impls/sm120_fp8_fp4_gemm_1d1d.cuh: kAIsFP4 template flag; A loaded as fp4-unpacked (b4x16 + <<2), TMA_A_BYTES = SMEM_A/2, MMA dispatch, perNTileX4 disabled for kAIsFP4. - jit_kernels/impls/sm120_fp8_fp4_gemm_1d1d.hpp: detect A=fp4 & B=fp8 (a_is_fp4); is_fp4 redefined to symmetric fp4xfp4 (value-identical for existing cases); emit kAIsFP4. - heuristics/sm120.hpp: a_padded_fp4 -> the fp4 A operand uses unpacked b4x16 SMEM (row = block_k, swizzle 128), mirroring the fp4 B handling. - tests/generators.py: add the FP4_A x FP8_B QuantConfig (32,128,True,False) to the SM120 sweep; the FP8 operand of a mixed config is 1D-scaled (per-token, recipe (1, gran)), not per-block. Validated on RTX PRO 6000: full fp8_fp4 sweep (normal / m-grouped contiguous / m-grouped masked, all dtype configs incl. the new FP4_A x FP8_B) passes (diff < max_diff, ~0.007 for fp4xfp8). Existing paths unchanged (a_is_fp4 defaults false). SM100 untouched. Note: for MoE decode this orientation is bandwidth-saturated (~92% DRAM) and only ~3-6% faster than the standard layout; the empty-tile skip is the real EP-decode lever. The K-grouped launcher A-stride is not yet wired for fp4-A.
Behavior-preserving cleanups from upstream review. gemm.hpp: pull the SM120 arm of fp8_fp4_gemm_nt into a dedicated fp8_fp4_gemm_nt_sm120() helper and dispatch in natural arch order (9/10 share the SF-transform path, then 12). Add a shared sm120_to_k_major() helper for the "repack packed-FP4 / contiguous to force K-major" idiom, reused by the FP8/FP4 NT GEMM, the m-grouped contiguous arm, and the BF16 GEMM. runtime.hpp: replace the coupled arch==12 ternaries in get_theoretical_mk_alignment_for_contiguous_layout() with a per-arch ContiguousMKAlignment policy plus a single shrink loop. SM90/SM100/SM120 results are unchanged; verified on SM120 with test_fp8_fp4.py and test_bf16.py.
The BF16 einsum tests compared DeepGEMM's bf16 output against cuBLAS's at 1e-10 (near bit-exact). On SM120 the two differ only in FP32 accumulation order (cuBLAS uses a multi-accumulator K-reduction at small M), so ~0.2% of low-magnitude outputs differ by 1 ULP and the comparison lands ~8e-9 -- not a DeepGEMM accuracy issue (both match the FP64 truth equally, ~1.4e-6). Keep the strict 1e-10-vs-cuBLAS bound for SM90/SM100, loosen to 1e-7 only on SM120, and add an architecture-independent FP32-reference correctness check on SM120 (calc_diff < 1e-5; DeepGEMM measures ~1.4e-6).
feat: add sm120 support for DeepGEMM
Signed-off-by: JaredforReal <w13431838023@gmail.com>
…eads_16 support indexer_n_heads = 16
Author
|
Duplicate of #5: the loop pod restart-looped after finishing (exit 1 + restartPolicy OnFailure) and republished each lap. |
wseaton
deleted the
autoresearch/20260804T041307Z-goal-fold-block-scaled-fp8-into-deepgemm/0
branch
August 4, 2026 04:50
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
19 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Autoresearch run
20260804T041307Z-goal-fold-block-scaled-fp8-into-deepgemm.Goal: Goal: fold block-scaled FP8 into DeepGEMM's MegaMoE (the H200-now subset)
Gate: correctness
Kept candidates:
Model
claude-opus-4-6· cost $3.13 · 2132s.🤖 opened by crucible autoresearch (draft — review and steer).