[Refactor][CodeGen] Refactor CodeGen part for multi-backend decoupling#2121
Conversation
|
👋 Hi! Thank you for contributing to the TileLang project. Please remember to run We appreciate you taking this step! Our team will review your contribution, and we look forward to your awesome work! 🚀 |
📝 WalkthroughWalkthroughThis PR moves backend-specific codegen implementations and stub sources from ChangesBackend relocation & include refactor
Metal entry + public C host API
Sequence Diagram(s)sequenceDiagram
participant Caller as Client / FFI caller
participant Registry as TVM FFI registry
participant MetalEntrypoint as BuildTileLangMetal
participant CHost as tl::BuildTileLangCHost
Caller->>Registry: invoke "target.build.tilelang_metal"(IRModule, Target)
Registry->>MetalEntrypoint: call BuildTileLangMetal(IRModule, Target)
MetalEntrypoint->>CHost: delegate to BuildTileLangCHost(mod, target)
CHost->>CHost: generate C host module (Metal context respected)
CHost-->>Registry: return ffi::Module
Registry-->>Caller: return ffi::Module
Estimated code review effort🎯 3 (Moderate) | ⏱️ ~25 minutes Possibly related issues
Possibly related PRs
Suggested labels
Suggested reviewers
Poem
🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 passed)
✏️ Tip: You can configure your own custom pre-merge checks in the settings. ✨ Finishing Touches🧪 Generate unit tests (beta)
Tip 💬 Introducing Slack Agent: The best way for teams to turn conversations into code.Slack Agent is built on CodeRabbit's deep understanding of your code, so your team can collaborate across the entire SDLC without losing context.
Built for teams:
One agent for your entire SDLC. Right inside Slack. 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. Comment |
SiriusNEO
force-pushed
the
chaofan/refactor_cmake_0428
branch
from
f941a85 to
423b0bf
Compare
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Caution
Some comments are outside the diff and can’t be posted inline due to platform limitations.
⚠️ Outside diff range comments (2)
src/backend/cuda/codegen/codegen_cutedsl.cc (1)
2-3:⚠️ Potential issue | 🟡 Minor | ⚡ Quick winStale
\filepath in doc comment.The file has moved to
src/backend/cuda/codegen/codegen_cutedsl.ccbut the doxygen tag still references the old location.📝 Proposed fix
-/*! - * \file target/codegen_cutedsl.cc - */ +/*! + * \file backend/cuda/codegen/codegen_cutedsl.cc + */🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@src/backend/cuda/codegen/codegen_cutedsl.cc` around lines 2 - 3, Update the doxygen file tag at the top of this source: replace the stale "\file target/codegen_cutedsl.cc" comment with the current path "\file src/backend/cuda/codegen/codegen_cutedsl.cc" so the documentation reflects the file's new location; locate the doxygen header in codegen_cutedsl.cc and update the \file tag accordingly.src/backend/cuda/codegen/codegen_cuda.cc (1)
2-2:⚠️ Potential issue | 🟡 Minor | ⚡ Quick winUpdate stale
\filedoxygen path.The file has been relocated to
src/backend/cuda/codegen/codegen_cuda.cc, but the doc comment still references the old pathtarget/codegen.cc.📝 Proposed fix
-/*! - * \file target/codegen.cc - */ +/*! + * \file backend/cuda/codegen/codegen_cuda.cc + */🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@src/backend/cuda/codegen/codegen_cuda.cc` at line 2, Update the stale doxygen file tag: replace the existing "\file target/codegen.cc" doc comment with the correct path for this translation unit (e.g. "\file src/backend/cuda/codegen/codegen_cuda.cc") so the top-of-file doxygen \file tag in codegen_cuda.cc matches the file's new location.
🧹 Nitpick comments (2)
src/backend/metal/codegen/rt_mod_metal.cc (1)
1-27: 💤 Low valueMetal FFI proxy looks correct; cppcheck warning is a false positive.
BuildTileLangMetalcleanly delegates totl::BuildTileLangCHost, and the FFI registration viaTVM_FFI_STATIC_INIT_BLOCKis the standard pattern. The cppcheck "syntax error" reported on line 21 is the usual false positive that tools raise on TVM's FFI init macros and can be ignored.One small note: since
BuildTileLangMetalis only referenced via the FFI registry, consider giving it internal linkage (staticor anonymous namespace) to avoid leaking a symbol intotvm::codegen.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@src/backend/metal/codegen/rt_mod_metal.cc` around lines 1 - 27, The BuildTileLangMetal symbol is only used via the FFI registry and can be given internal linkage to avoid leaking a public symbol; change the free function BuildTileLangMetal to have internal linkage (e.g., mark it static or place it in an anonymous namespace) and keep the TVM_FFI_STATIC_INIT_BLOCK registration intact so the refl::GlobalDef still registers "target.build.tilelang_metal" with BuildTileLangMetal.CMakeLists.txt (1)
174-174: 💤 Low valueInclude precedence: prepending
src/at position 0 may shadow TVM headers.
list(INSERT TILE_LANG_INCLUDES 0 ...)puts TileLang'ssrc/ahead ofTVM_INCLUDES. TileLang'ssrc/target/andsrc/op/overlap with TVM's owntarget/andop/namespaces in path (e.g.,target/utils.h,op/builtin.h); if TVM ever exposes a same-named header through its include dirs, the TileLang one will silently win. Consider appending instead of inserting at 0, unless the override is intentional.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@CMakeLists.txt` at line 174, The current insertion "list(INSERT TILE_LANG_INCLUDES 0 \"${CMAKE_CURRENT_SOURCE_DIR}/src\")" prepends TileLang's src/ ahead of TVM headers and can unintentionally shadow TVM headers; change it to append or insert after TVM include entries so TILE_LANG_INCLUDES is added after TVM_INCLUDES (use list(APPEND TILE_LANG_INCLUDES ...) or insert at the end) to preserve TVM include precedence and avoid accidental header overrides.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Outside diff comments:
In `@src/backend/cuda/codegen/codegen_cuda.cc`:
- Line 2: Update the stale doxygen file tag: replace the existing "\file
target/codegen.cc" doc comment with the correct path for this translation unit
(e.g. "\file src/backend/cuda/codegen/codegen_cuda.cc") so the top-of-file
doxygen \file tag in codegen_cuda.cc matches the file's new location.
In `@src/backend/cuda/codegen/codegen_cutedsl.cc`:
- Around line 2-3: Update the doxygen file tag at the top of this source:
replace the stale "\file target/codegen_cutedsl.cc" comment with the current
path "\file src/backend/cuda/codegen/codegen_cutedsl.cc" so the documentation
reflects the file's new location; locate the doxygen header in
codegen_cutedsl.cc and update the \file tag accordingly.
---
Nitpick comments:
In `@CMakeLists.txt`:
- Line 174: The current insertion "list(INSERT TILE_LANG_INCLUDES 0
\"${CMAKE_CURRENT_SOURCE_DIR}/src\")" prepends TileLang's src/ ahead of TVM
headers and can unintentionally shadow TVM headers; change it to append or
insert after TVM include entries so TILE_LANG_INCLUDES is added after
TVM_INCLUDES (use list(APPEND TILE_LANG_INCLUDES ...) or insert at the end) to
preserve TVM include precedence and avoid accidental header overrides.
In `@src/backend/metal/codegen/rt_mod_metal.cc`:
- Around line 1-27: The BuildTileLangMetal symbol is only used via the FFI
registry and can be given internal linkage to avoid leaking a public symbol;
change the free function BuildTileLangMetal to have internal linkage (e.g., mark
it static or place it in an anonymous namespace) and keep the
TVM_FFI_STATIC_INIT_BLOCK registration intact so the refl::GlobalDef still
registers "target.build.tilelang_metal" with BuildTileLangMetal.
ℹ️ Review info
⚙️ Run configuration
Configuration used: defaults
Review profile: CHILL
Plan: Pro
Run ID: 6ff4ad41-17f5-4fbe-b67b-a6e8bd6b8e52
📒 Files selected for processing (33)
CMakeLists.txtsrc/backend/cuda/CMakeLists.txtsrc/backend/cuda/codegen/codegen_cuda.ccsrc/backend/cuda/codegen/codegen_cuda.hsrc/backend/cuda/codegen/codegen_cutedsl.ccsrc/backend/cuda/codegen/codegen_cutedsl.hsrc/backend/cuda/codegen/codegen_py.ccsrc/backend/cuda/codegen/codegen_py.hsrc/backend/cuda/codegen/intrin_rule_cuda.ccsrc/backend/cuda/codegen/ptx.ccsrc/backend/cuda/codegen/ptx.hsrc/backend/cuda/codegen/rt_mod_cuda.ccsrc/backend/cuda/codegen/rt_mod_cutedsl.ccsrc/backend/cuda/codegen/stubs/cuda.ccsrc/backend/cuda/codegen/stubs/cuda.hsrc/backend/cuda/codegen/stubs/cudart.ccsrc/backend/cuda/codegen/stubs/nvrtc.ccsrc/backend/cuda/codegen/stubs/vendor/cuda.hsrc/backend/metal/CMakeLists.txtsrc/backend/metal/codegen/rt_mod_metal.ccsrc/backend/rocm/CMakeLists.txtsrc/backend/rocm/codegen/codegen_hip.ccsrc/backend/rocm/codegen/codegen_hip.hsrc/backend/rocm/codegen/intrin_rule_hip.ccsrc/backend/rocm/codegen/rt_mod_hip.ccsrc/backend/rocm/codegen/stubs/hip.ccsrc/backend/rocm/codegen/stubs/hip.hsrc/backend/rocm/codegen/stubs/hiprtc.ccsrc/backend/rocm/codegen/stubs/vendor/hip_runtime.hsrc/op/builtin.ccsrc/op/utils.hsrc/runtime/runtime.ccsrc/target/codegen_c_host.h
🚧 Files skipped from review as they are similar to previous changes (2)
- src/op/utils.h
- src/target/codegen_c_host.h
1 participant
As titled. Currently we put the
candc_hostunder the original folder. Note that metal backend usesCodeGenCto build. So the wrapper underbackend/metal/codegenis just a proxy.Summary by CodeRabbit
New Features
Chores