Skip to content

[BugFix] Only record validated TMA descriptor names - #3484

Merged
LeiWang1999 merged 4 commits into
tile-ai:mainfrom
adenzhou1350:codex/fix-tma-descriptor-name-map
Oct 28, 2026
Merged

LeiWang1999 merged 4 commits into
tile-ai:mainfrom
adenzhou1350:codex/fix-tma-descriptor-name-map

Conversation

@adenzhou1350

@adenzhou1350 adenzhou1350 commented Oct 12, 2026 •

Copy link
Copy Markdown
Contributor

maybe_desc_name updates desc_name_map before checking the neighboring CUtensorMap declaration. A normal pointer parameter whose name starts with another buffer's descriptor name therefore leaves a false mapping behind.

For example, TMA lowering with input A and ordinary output A_desc produces:

mixed_kernel(__grid_constant__ const CUtensorMap A_desc,
             half_t* __restrict__ A_desc_1)

The extra A_desc_1 mapping has no corresponding descriptor variable, so wrapper construction raises Handle name A_desc_1 not found in desc_name_var_map.

Record the mapping only after confirming the descriptor type. Remove testing/python/jit/test_tilelang_jit_adapter_wrapper.py as requested during review.

Validation:

  • Before removing the test file, focused checks loaded the complete adapter modules over an existing native build. The base and current main utility modules each produced two failures and one pass; the fixed adapter passed all three tests, including real TMA lowering. With CUDA hidden, the fixed adapter produced two passes and one skip.
  • The previous CUDA CI run also passed all three tests before the file was removed.
  • For the deletion commit, git diff --check and bash format.sh --files tilelang/jit/adapter/utils.py passed.

No local GPU kernel execution or fresh full native rebuild was performed.

Summary

  • Fix maybe_desc_name so it records a descriptor-name mapping only after confirming a neighboring declaration is a CUtensorMap. This prevents mappings for non-descriptor parameters.
  • Add regression coverage for descriptor-map mutation and CUDA TMA lowering.
  • Remove the adapter-wrapper test that expected an error when lowered modules were missing.

Validation

  • The reported tests passed after the fix, including real TMA lowering. With CUDA hidden, pytest reported 2 passed and 1 skipped.
  • git diff --check passed. Ruff was not run. No GPU kernel execution or fresh native rebuild was performed.

@github-actions

Copy link
Copy Markdown

👋 Hi! Thank you for contributing to the TileLang project.

Please remember to run pre-commit run --all-files in the root directory of the project to ensure your changes are properly linted and formatted. This will help ensure your contribution passes the format check.

We appreciate you taking this step! Our team will review your contribution, and we look forward to your awesome work! 🚀

@coderabbitai

coderabbitai Bot commented Oct 12, 2026 •

Copy link
Copy Markdown
Contributor

Review in Change Stack →

Caution

Review failed

The pull request is closed.

No actionable comments were generated in the recent review.

ℹ️ Recent review info
⚙️ Run configuration
  • Configuration used: Repository: tile-ai/tilelang/.coderabbit.yaml
  • Review profile: CHILL
  • Plan: Advanced
  • Run ID: 8ec55556-9bce-4b75-a622-dca5211b1434
📥 Commits

Reviewing files that changed from the base of the PR and between 4d39358 and 368574f.

📒 Files selected for processing: 1
  • testing/python/jit/test_tilelang_jit_adapter_wrapper.py
🚧 Files skipped from review as they are similar to previous changes: 1
  • testing/python/jit/test_tilelang_jit_adapter_wrapper.py

Included review availability: This review used your included allowance. Your plan provides up to 8 included reviews per hour; 7 remain after this review.

Walkthrough

maybe_desc_name now checks neighboring declarations before recording a descriptor-name mapping. The CPU wrapper test for missing lowered modules and its direct pytest runner were removed.

Changes: Descriptor-name detection

Layer / File(s) Summary
Descriptor detection
tilelang/jit/adapter/utils.py
maybe_desc_name records a descriptor-name mapping only after a neighboring declaration identifies a CUtensorMap descriptor.

CPU wrapper test removal

Layer / File(s) Summary
CPU wrapper test removal
testing/python/jit/test_tilelang_jit_adapter_wrapper.py
The test that expected a ValueError when device_mod and host_mod were missing, along with the file’s direct pytest runner, was removed.

Priority: Not assessed

Estimated code review effort: 2 (Simple) | ~10 minutes

Change: Bug fix

Possibly related PRs

  • tile-ai/tilelang#3634: Changes CPU source-wrapper coverage, which is also affected by the removed TLCPUSourceWrapper test.
  • tile-ai/tilelang#3638: Changes TMA descriptor and kernel-argument binding, which also involves descriptor argument parsing.
  • tile-ai/tilelang#3285: Changes maybe_desc_name and CUDA/NVRTC descriptor parsing.

Suggested reviewers: siriusneo

Merge Risk: Low · up to 267d6

The fix is in place, so this is mergeable with a bounded follow-up: restore the focused regression test to protect against the original wrapper failure returning.

Pre-merge checks: Passed 5
Passed Passed checks: 5 passed
Check name Status Explanation
Description Check Passed Check skipped - CodeRabbit’s high-level summary is enabled.
Title check Passed The title clearly and concisely describes the main change: recording TMA descriptor names only after validation.
Docstring Coverage Passed Docstring coverage is 100.00% which is sufficient. The required threshold is 80.00%. Docstring coverage is scoped to functions touched by this diff. Analyzed 6 functions across 2 files.
Linked Issues check Passed Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check Passed Check skipped because no linked issues were found for this pull request.
Finishing Touches
🧪 Generate unit tests (beta)
  • Create a new PR
  • Autofix · Keep fixing CodeRabbit findings and required CI, and resolving merge conflicts

Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out.

❤️ Share

Comment @coderabbitai help to get the list of available commands.

@adenzhou1350

Copy link
Copy Markdown
Contributor Author

Fixed the Ruff 0.16.10 formatting failure in 368574f. The hook required one blank line after the new test docstring; Ruff format/check and git diff --check now pass, and the Python AST is unchanged. The new push restarts CI.

@SiriusNEO SiriusNEO added bug Something isn't working python Pull requests that update python code labels Oct 15, 2026
@adenzhou1350

Copy link
Copy Markdown
Contributor Author

Could a maintainer check the self-hosted NVIDIA CI queue? #3484 is waiting on the CUDA-auto test job, #3477 on the CuTeDSL examples job (its ordinary CUDA test passed), and #3384 on the CUDA-auto test job plus one workflow approval. The completed checks are green. Keeping the runner/approval follow-up in one place rather than posting on all three PRs.

@LeiWang1999 LeiWang1999 left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

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

Recording the requested test-file removal; the descriptor-name mapping fix is retained.

@coderabbitai
coderabbitai Bot requested a review from SiriusNEO October 28, 2026 09:11
@LeiWang1999
LeiWang1999 merged commit 1086665 into tile-ai:main Oct 28, 2026
7 of 8 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

bug Something isn't working python Pull requests that update python code

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants