Skip to content

Detect tcgen05 load/store API by signature - #119

Merged
icavan merged 2 commits into
mainfrom
codex/cutedsl-tcgen05-signature-detection
Aug 5, 2026
Merged

Detect tcgen05 load/store API by signature#119
icavan merged 2 commits into
mainfrom
codex/cutedsl-tcgen05-signature-detection

Conversation

@icavan

@icavan icavan commented Aug 5, 2026

Copy link
Copy Markdown
Collaborator

Summary

  • replace the CuTeDSL version gate with capability detection based on the generated tcgen05_ld and tcgen05_st function signatures
  • detect the optional num argument independently for load and store, and select the supported store value keyword (r or val)
  • fail fast with an explicit error when an unknown or ambiguous generated binding is encountered
  • add unit coverage for the legacy, inferred-shape, mixed, and unsupported signatures
  • migrate the SM100 intracard CP tests from the removed private _no_cp guard to the public use_intracard_cp=False control

Motivation

The low-level generated NVVM binding changed between CuTeDSL 4.5.2 and 4.5.3: num was removed and the store value keyword changed from r to val. A package-version comparison is brittle because this is an internal API and distributions may backport or ship the binding change independently. Inspecting the callable interface directly makes the compatibility decision match the API that is actually installed.

The intracard CP tests also retained two calls to the private _no_cp keyword after production dispatch was migrated to the shared CPMode interface. Using use_intracard_cp=False restores the intended serial baseline and fallback behavior without reintroducing a second, conflicting dispatch control.

Validation

Validated on an SM100 GPU in isolated environments:

  • CuTeDSL 4.5.2 detected as ld_has_num=True, st_has_num=True, st_value_keyword="r"
  • CuTeDSL 4.5.3 detected as ld_has_num=False, st_has_num=False, st_value_keyword="val"
  • tests/test_cutedsl_compat.py: 4 passed
  • tests/test_ptx_umma_ws.py::test_ws_ss_tf32: passed on both 4.5.2 and 4.5.3
  • KDA forward/backward smoke test, including TMEM load and store: passed on both 4.5.2 and 4.5.3
  • KDA fast group: 18 passed, 38 skipped, 136 deselected
  • pre-commit run --all-files: passed
  • Ruff, formatting, bytecode compilation, and git diff --check: passed

Related to #118.

icavan added 2 commits August 5, 2026 14:23
Signed-off-by: shuyan.ycf <shuyan.ycf@antgroup.com>
Signed-off-by: shuyan.ycf <shuyan.ycf@antgroup.com>
@icavan
icavan marked this pull request as ready for review August 5, 2026 09:58
@icavan
icavan requested review from cherhh and zheyang0825 August 5, 2026 09:58

@cherhh cherhh left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM.

@zheyang0825 zheyang0825 left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

lgtm

@icavan
icavan merged commit 5370b14 into main Aug 5, 2026
2 checks passed
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.

3 participants