Repository navigation
cuda: glm_mm takes common/qmm_frag.cuh, and its 117-line copy glm_qmm_frag.cuh goes, SASS unchanged - #1564
Merged
Conversation
…_frag.cuh goes, SASS unchanged Step 5 of ashhart#1409. glm_qmm_frag.cuh was a copy of qmm_frag.cuh's device code from before ashhart#1440 (its header names a gen_copies.py that is no longer in the tree): the same namespace qmm_frag, the same cp.async (cp16, cp16z, cp8, cp4, commit, wait), ldmatrix, mma, nibble and swizzle helpers, and fewer of them. common/qmm_frag.cuh is a superset (it adds ld_nc, ldmatrix4t and fp4pair, which glm_mm.cu does not call), so glm_mm.cu, its only user, includes that instead.
ashhart
force-pushed
the
prims-1409-step5
branch
from
October 29, 2026 23:40
e9e47a7 to
ba8f3df
Compare
ashhart
pushed a commit
to jschmied/TensorFold
that referenced
this pull request
Oct 29, 2026
ashhart#1409 step 5, the group fences. Eleven files kept their own cp.async.commit_group and cp.async.wait_group N under five names. common/cp_async.cuh now holds one tf_common::cp_commit and cp_wait<N>; the files take them with using declarations. The attention kernels' wait_group<N> and block_attention's wait_all() become cp_wait<N> and cp_wait<0>, exl3_prompt's cp_async_commit and cp_async_wait become cp_commit and cp_wait, and glm_index's cp_wait() becomes cp_wait<0>(). experts.cuh's gradient-serving pin takes the new bytes; prefill_attention's two pins already differ on main and are left. Not here: qmm_frag's commit and wait<N> and the copies that mirror them (h3_attention, glm_glue_attn, glm_glue_bmm, glm_qmm_frag) follow once ashhart#1564 lands. DeepSeek's copies carry a "memory" clobber and are not the same asm. cuobjdump -sass is identical for every function of the 27 images this reaches (2,245 functions, sm_80/86/89/120/121, nvcc 13.0). The tool's warp_sum control goes red; its fp4pair control has no user among these images. A planted control in cp_async.cuh (cp_wait<N> waiting for N - 1 groups) changes 1,190 functions.
ashhart
pushed a commit
to jschmied/TensorFold
that referenced
this pull request
Oct 29, 2026
ashhart#1409 step 6, the copies that are the same instruction: - prefill_attention.cu and qwen35_attn_tc.cuh kept their own ldmatrix4t, qmm_frag's text: both take qmm_frag::ldmatrix4t. - the exl3 kernels' ldsm4 (x3gm.cu, pfdense.cuh), ldsm2 (x3gm2.cuh) and ldmatrix_x4 (exl3_prompt.cuh) are qmm_frag's ldmatrix4 and ldmatrix2 with their own address cast; they take qmm_frag's, and their callers (x3gm, x3gm2, x3gm3, pfdense, exl3_prompt) the names. - experts_grouped.cuh's f16 mma16816 is exl3_decode.cuh's text: tf_exl3x takes tf_exl3's through a using declaration. Not here: mma_z (nvfp4_experts, flashnext_fexp) and experts.cuh's mma take scalar operands where qmm_frag's take arrays; pfdense's mma16816_bf16 is asm volatile where qmm_frag::mma is not; the glm glue copies follow ashhart#1564. cuobjdump -sass is identical for every function of the 14 images this reaches (1,935 functions, sm_80/86/89/120/121, nvcc 13.0), base ashhart#1588. The tool's warp_sum control goes red; a planted control in exl3_decode.cuh (the f16 mma as bf16) changes 2,238 of 2,803 functions.
ashhart
pushed a commit
that referenced
this pull request
Oct 29, 2026
#1409 step 5, the group fences. Eleven files kept their own cp.async.commit_group and cp.async.wait_group N under five names. common/cp_async.cuh now holds one tf_common::cp_commit and cp_wait<N>; the files take them with using declarations. The attention kernels' wait_group<N> and block_attention's wait_all() become cp_wait<N> and cp_wait<0>, exl3_prompt's cp_async_commit and cp_async_wait become cp_commit and cp_wait, and glm_index's cp_wait() becomes cp_wait<0>(). experts.cuh's gradient-serving pin takes the new bytes; prefill_attention's two pins already differ on main and are left. Not here: qmm_frag's commit and wait<N> and the copies that mirror them (h3_attention, glm_glue_attn, glm_glue_bmm, glm_qmm_frag) follow once #1564 lands. DeepSeek's copies carry a "memory" clobber and are not the same asm. cuobjdump -sass is identical for every function of the 27 images this reaches (2,245 functions, sm_80/86/89/120/121, nvcc 13.0). The tool's warp_sum control goes red; its fp4pair control has no user among these images. A planted control in cp_async.cuh (cp_wait<N> waiting for N - 1 groups) changes 1,190 functions.
ashhart
pushed a commit
that referenced
this pull request
Oct 29, 2026
#1409 step 6, the copies that are the same instruction: - prefill_attention.cu and qwen35_attn_tc.cuh kept their own ldmatrix4t, qmm_frag's text: both take qmm_frag::ldmatrix4t. - the exl3 kernels' ldsm4 (x3gm.cu, pfdense.cuh), ldsm2 (x3gm2.cuh) and ldmatrix_x4 (exl3_prompt.cuh) are qmm_frag's ldmatrix4 and ldmatrix2 with their own address cast; they take qmm_frag's, and their callers (x3gm, x3gm2, x3gm3, pfdense, exl3_prompt) the names. - experts_grouped.cuh's f16 mma16816 is exl3_decode.cuh's text: tf_exl3x takes tf_exl3's through a using declaration. Not here: mma_z (nvfp4_experts, flashnext_fexp) and experts.cuh's mma take scalar operands where qmm_frag's take arrays; pfdense's mma16816_bf16 is asm volatile where qmm_frag::mma is not; the glm glue copies follow #1564. cuobjdump -sass is identical for every function of the 14 images this reaches (1,935 functions, sm_80/86/89/120/121, nvcc 13.0), base #1588. The tool's warp_sum control goes red; a planted control in exl3_decode.cuh (the f16 mma as bf16) changes 2,238 of 2,803 functions.
Owner
|
Thank you @jschmied. This is in TensorFold 1.2.0, released today: https://github.com/ashhart/TensorFold/releases/tag/v1.2.0 It landed with your commits under your own name, so you're listed as a TensorFold contributor, and it's credited to you in the release notes and the changelog. |
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
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.
What
Step 5 of #1409. It is independent of #1550, #1555 and #1556, with no shared files. It is a pure refactor:
cuobjdump -sassofglm_mmis identical function by function.glm_qmm_frag.cuhwas a copy ofqmm_frag.cuh's device code from before #1440. Its first line still names agen_copies.pythat is no longer in the tree. It has the samenamespace qmm_fragand the samecp.asynchelpers (cp16,cp16z,cp8,cp4,commit,wait<N>), plus the sameldmatrix,mma, nibble, dependent-launch and swizzle helpers, only fewer of them.common/qmm_frag.cuhis a superset: it addsld_nc,ldmatrix4tandfp4pair, none of whichglm_mm.cucalls. Soglm_mm.cu, the copy's only user, now includes the common header, and the copy is deleted.2 files, +1 / −118.
Verification
dd0a9cc5e, heade9e47a744, nvcc 13.0, sm_80/86/89/120/121.glm_mmis the only image the change reaches. All 110 functions: 0 changed, 0 added, 0 removed, 0 unstable.glm_mmcalls neitherfp4pair(0 functions changed) norwarp_sum(the pattern is not in this image), so the tool exits 1.cp.async.cg ... 16incommon/qmm_frag.cuh(cp16) to.caand builtglm_mm.cuat sm_121 with the build's flags. 87 instructions change (LDGSTS.E.BYPASS138 → 51), so the comparison sees a change made in the included header.zig build test-gradient-serving: the same set of differing pins as main. Neither file is pinned.zig build test -Dnvcc=/usr/local/cuda/bin/nvcc: passes (RC 0).Part of #1409.