mirror of
https://github.com/tile-ai/tilelang.git
synced 2026-10-21 06:40:23 +08:00
Lei WangandClaude Opus 4.6
853e805b2a
* [Codegen] Add lexical_alloc_scope for scoped local variable lifetime
Introduce a `lexical_alloc_scope` AttrStmt that generates `{ ... }` in
C/CUDA codegen, giving the underlying compiler accurate variable lifetime
information for better register allocation.
- Define `tl::attr::kLexicalAllocScope` constant
- LowerOpaqueBlock wraps block-local allocations in the new AttrStmt
- StorageRewrite treats it as a scope boundary with proper thread_scope_
save/restore so allocations are not hoisted past the boundary
- CUDA and HIP codegen emit scoped `{ }` blocks
- Add tests for IR insertion, StorageRewrite preservation, and codegen output
* lint fix
* [Codegen] Refine lexical_alloc_scope: skip top-level blocks and decouple from thread_scope
Two improvements to the lexical_alloc_scope mechanism:
1. LowerOpaqueBlock: only insert lexical_alloc_scope for blocks inside
loops (inside_loop_ > 0). Top-level blocks already have function-body
lifetime, so the extra `{ }` in codegen is pointless.
2. StorageRewrite: introduce a separate `lexical_scope_` / `effective_scope()`
instead of overriding `thread_scope_`. This avoids breaking
PlanNewScope's toggle protocol when lexical_alloc_scope is nested
inside thread_extent, fixing a `ICHECK(thread_scope_ == op)` crash.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* [Transform] Fix alloc-scope placement regressions
* [Infra] Remove clang-tidy integration
* Remove unused local descriptor allocation pass and related tests; update example to disable main execution and print kernel source for debugging.
* Preserve lexical alloc scopes for nested register buffers
* Limit lexical alloc scopes to local storage
* refactor
* Unify lexical alloc scope annotations
* Clean up dead code and fix unsafe cast
- Remove unused block_nesting_ member from OpaqueBlockLower
- Use static_cast instead of reinterpret_cast in ResolveAllocationSite
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
---------
Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com>