Repository navigation
cuda: ldmatrix and the f16 mma from their one home each, SASS unchanged (#1409 step 6, stacked on #1588) - #1591
Merged
Merged
Conversation
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#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
force-pushed
the
prims-1409-step6
branch
from
October 29, 2026 23:41
1bbbe8a to
20f05de
Compare
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.
Stacked on #1588: only the top commit is new here. #1409, step 6: helper copies that compile to the same instruction.
What
ldmatrix4t:prefill_attention.cuandqwen35_attn_tc.cuheach kept a copy ofqmm_frag's text. Both now takeqmm_frag::ldmatrix4t.ldmatrixhelpers:ldsm4(x3gm.cu,pfdense.cuh),ldsm2(x3gm2.cuh) andldmatrix_x4(exl3_prompt.cuh) areqmm_frag'sldmatrix4andldmatrix2with their own address cast. They now takeqmm_frag's, and their callers (x3gm,x3gm2,x3gm3,pfdense,exl3_prompt) use those names.mma16816:experts_grouped.cuh's copy isexl3_decode.cuh's text, sotf_exl3xtakestf_exl3::mma16816through ausingdeclaration.Not here:
mma_z(nvfp4_experts,flashnext_fexp) andexperts.cuh'smmatake scalar operands whereqmm_frag's take arrays.pfdense'smma16816_bf16isasm volatilewhereqmm_frag::mmais not.ldm4,ldm4s,ldm2s) follow cuda: glm_mm takes common/qmm_frag.cuh, and its 117-line copy glm_qmm_frag.cuh goes, SASS unchanged #1564.Checks
cuda_sass_identity.py(base cuda: cp.async commit and wait from common/cp_async.cuh, SASS unchanged (#1409 step 5) #1588): identicalcuobjdump -sassfor every function of the 14 images this reaches. That is 1,935 functions over sm_80/86/89/120/121 with nvcc 13.0.warp_sumcontrol goes red.exl3_decode.cuh(the f16 mma as bf16) changes 2,238 of 2,803 functions.prefill_attention.cu's gradient-serving pin already differs on main; it is left as it is.zig build test: passes.