[jschmied]

#1409, step 5: the cp.async group fences. This is independent of #1550, #1555, #1556 and #1564, and based on current main.

What

Eleven files each kept their own cp.async.commit_group and cp.async.wait_group N, under five names (cp_commit/cp_wait, cp_async_commit/cp_async_wait, wait_group, wait_all). The new common/cp_async.cuh holds one tf_common::cp_commit() and cp_wait<N>(), and the files take them with using declarations, as in #1440:

• exl3_block/pfdense.cuh, exl3_block/x3gm.cu, experts.cuh: using only, no caller changes.
• exl3_prompt.cuh: cp_async_commit/cp_async_wait are renamed to cp_commit/cp_wait.
• glm_index_scores.cuh: cp_wait() (a hard-coded wait_group 0) becomes cp_wait<0>().
• prefill_attention.cu and prefill_attention_rows.cuh, plus tree_attention.cu, window_attention.cu, window_rows.cuh: wait_group<N> becomes cp_wait<N>.
• block_attention.cu: wait_all() becomes cp_wait<0>().

experts.cuh's gradient-serving pin takes the new bytes. The two prefill_attention pins already differ on main and are left as they are.

Not here:

• qmm_frag's commit/wait<N> and the copies that mirror them (h3_attention.cuh, glm_glue_attn.cuh, glm_glue_bmm.cuh, glm_qmm_frag.cuh). These follow once cuda: glm_mm takes common/qmm_frag.cuh, and its 117-line copy glm_qmm_frag.cuh goes, SASS unchanged #1564 lands, which removes glm_qmm_frag.cuh.
• DeepSeek's copies: they carry a "memory" clobber, so they aren't the same asm.

Checks

• tools/zig/cuda_sass_identity.py origin/main HEAD: identical cuobjdump -sass for every function of the 27 images this change reaches. That is 2,245 functions over sm_80/86/89/120/121 with nvcc 13.0.
  • The tool's warp_sum control goes red.
  • Its fp4pair control changes nothing here, because none of these images uses fp4pair; that is why the tool exits 1.
• A planted control in cp_async.cuh (cp_wait<N> waiting for N − 1 groups) changes 1,190 functions, so the comparison can see this header.
• Lean checker: no new problems.
• zig build test: passes.

[ashhart]

Thank you, this is a nice cleanup, and the SASS receipt with the planted control is just what we need.

There is one thing to fix before we can land it. tools/glm_index_emulate/emulate.cpp builds glm_index_scores.cuh as host C++ with GLM_INDEX_HOST defined, and it supplies its own stub inline void cp_wait() {}. Now that the kernel calls cp_wait<0>(), the host build fails:

glm_index_scores.cuh:95:5: error: use of undeclared identifier 'cp_wait'

You can reproduce it with zig build test-glm-index-emulate. That step is not part of zig build test, so it is easy to miss. Making the stub a template fixes it:

template <int N> inline void cp_wait() {}

With that change the emulator builds and passes (0 failures) on our side. Could you add it to this PR? Everything else looks good to land.

[jschmied]

Thanks @ashhart. Done in 6c92b2685: the stub is now template <int N> inline void cp_wait() {}. zig build test-glm-index-emulate builds and passes here, and zig build test plus the CI set (test-kernel-contract test-golden test-server-security test-pull-security fuzz-cortheon), git diff --check and lean (0 new) pass on the new head.

[ashhart]

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.