Skip to content

[Metal] Derive pointer address spaces from TIR types - #2824

Closed
GY-Bai wants to merge 2 commits into
tile-ai:mainfrom
GY-Bai:fix/metal-address-space-codegen
Closed

[Metal] Derive pointer address spaces from TIR types#2824
GY-Bai wants to merge 2 commits into
tile-ai:mainfrom
GY-Bai:fix/metal-address-space-codegen

Conversation

@GY-Bai

@GY-Bai GY-Bai commented Jul 31, 2026

Copy link
Copy Markdown
Contributor

Problem

Metal requires every pointer type to carry an explicit address-space qualifier. On an Apple M2, a regular FP16 GEMM using dynamic shared-memory aliases named As and Bs failed during MSL compilation with errors such as:

error: pointer type must have explicit address space qualifier
*(half2*)(((half*)As) + ...)
simdgroup_load(..., (&(((half*)As)[...])), ...)

The existing source-rewrite passes handled a few hard-coded names and then scanned lines containing _shared. Therefore, correctness depended on variable names: A_shared could compile while an equivalent alias named As did not. The compiled and source-only Metal build entry points also exercised different rewrite logic.

This failure occurs in the Metal 3 simdgroup path and is independent of Metal 4 cooperative tensors. The invalid pointer types are rejected during MSL compilation, before GPU execution, so the issue can affect any Metal device that reaches this lowering path; the M2 is the hardware used to reproduce and validate it.

Root-cause trace

This was reproduced and traced from a clean main at commit e01c498b6499367e0a1cd64413e008d223a931b6.

  1. T.alloc_shared buffers are merged by MergeSharedMemoryAllocations::MakeAliasBindings() in src/transform/merge_shared_memory_allocations.cc.
  2. That pass creates tirx::Bind(buffer_var, handle_add_byte_offset(merged_buffer, offset)). The alias variable still has a typed PointerType, including its element dtype and shared.dyn storage scope.
  3. Generic CodeGenC::VisitStmt_(const BindNode*) did not register that type information in handle_data_type_ or alloc_storage_scope_.
  4. Metal codegen consequently emitted unqualified alias declarations and pointer casts. Five string post-processing passes in BuildTileLangMetal attempted to repair the generated MSL, but some depended on names such as A_shared, B_shared, or the _shared suffix.

The loss therefore happened at the TIR-to-MSL BindNode codegen boundary, not in the shared-memory merge pass and not in Metal runtime compilation.

Fix

  • Add a Metal-specific BindNode visitor that reads the alias PointerType, registers its pointee dtype and storage scope, and emits a typed, address-space-qualified declaration.
  • Map TIR storage scopes to MSL address spaces in one place:
    • global -> device
    • shared, shared.dyn, threadgroup -> threadgroup
    • local, local.var, and Metal register scopes -> thread
  • Derive handle_add_byte_offset address spaces from the pointer operand instead of hard-coding threadgroup.
  • Trace pointer provenance only through address_of(BufferLoad) and the base operand of handle_add_byte_offset; reject unknown scopes instead of silently defaulting to thread.
  • Remove all five name-based MSL source-rewrite passes. Both Metal build entry points now receive the same correct output directly from CodeGenTileLangMetal, while their existing callback behavior remains unchanged.

The failing aliases now generate:

threadgroup half* As = (threadgroup half*)...;
*(threadgroup half2*)(As + ...) = ...;

Tests

Run on a MacBook Air with Apple M2, macOS 15.6.1, and Metal 3:

ninja -C build -j4 tilelang
TILELANG_CACHE_DIR=... TILELANG_TMP_DIR=... python -m pytest testing/python/metal/test_metal_address_space.py -q -ra
# 3 passed (source-only + compiled codegen, tvm_ffi execution, torch execution)

TILELANG_CACHE_DIR=... TILELANG_TMP_DIR=... python -m pytest testing/python/metal -q -ra
# 36 passed, 3 skipped
# The skipped tests are guarded by TileLang's current Metal 4 detector,
# which requires macOS 26, SDK 26, and an M5-or-newer chip.

pre-commit run --all-files
# All hooks passed.

The skip condition above reflects TileLang's current capability detector, not Apple's published hardware boundary. Apple's May 2026 Metal Feature Set Tables list the Metal 4 programming model and tensors starting with Apple7 (the M1 generation). Metal 4 capability detection is separate from this Metal 3 codegen fix and is intentionally unchanged in this PR.

Potentially affected hardware

This defect is in TIR-to-MSL source generation rather than GPU execution. Within TileLang's current macOS/arm64 deployment scope, any Apple silicon Mac generation from M1 through M5 can potentially hit it when a kernel uses merged dynamic shared-memory aliases and reaches the affected pointer-cast path. This does not mean every Metal kernel fails: the previous name-based rewrites masked many aliases containing _shared.

Only the M2/Metal 3 configuration has been validated on hardware in this PR. The M1, M3, M4, and M5 impact is inferred from the generation-independent MSL address-space rule and the shared TileLang codegen path, not from direct testing on each chip. iOS and iPadOS targets are not claimed here because TileLang's current availability detector only enables the backend on arm64 macOS.

Summary

  • Derive Metal pointer address spaces from TIR types and storage scopes.
  • Add BindNode handling for typed pointer bindings.
  • Support threadgroup, metal.simdgroup, and metal.cooperative_tensor scopes.
  • Infer address spaces for pointer expressions and handle_add_byte_offset.
  • Reject unknown or unresolved storage scopes.
  • Remove five name-based MSL rewrite passes.
  • Add regression tests for dynamic shared-memory aliases through TVM FFI and Torch backends.
  • Verify identical compiled and source-only Metal output.
  • Tests passed on Apple M2: 36 Metal tests passed, with 3 skips for unsupported Metal 4 cooperative tensors.
  • All pre-commit hooks passed.

C++ style / lint notes

  • The PR changes C++ code but does not change documented rules in docs/developer_guide/cpp_style.md.
  • The C++ API Style Audit remains warning-only. No correctness or build issue is reported.
  • No new public API, FFI risk, or maintainability risk requires blocking merge for advisory TLCPP003/TLCPP004 findings.

@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 Jul 31, 2026

Copy link
Copy Markdown
Contributor

Review Change Stack

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Pro Plus

Run ID: 4184b56f-b115-4ba2-bbc3-9715356c775d

📥 Commits

Reviewing files that changed from the base of the PR and between 559bd7e and bea3ede.

📒 Files selected for processing (1)
  • testing/python/metal/test_metal_address_space.py
🚧 Files skipped from review as they are similar to previous changes (1)
  • testing/python/metal/test_metal_address_space.py

📝 Walkthrough

Walkthrough

Metal codegen now uses typed, storage-scope-aware pointer inference for declarations and byte-offset casts. It supports additional Metal scopes, removes string-based source rewriting, and adds regression tests for generated source and MPS execution.

Changes

Metal address-space code generation

Layer / File(s) Summary
Storage-scope mapping and typed bindings
src/metal/codegen/codegen_metal.cc, src/metal/codegen/codegen_metal.h
Metal storage scopes map to qualified address spaces. BindNode bindings validate typed pointers and record scope and pointee metadata.
Pointer inference and byte offsets
src/metal/codegen/codegen_metal.cc
Pointer metadata is inferred from allocations, annotations, address-of expressions, and byte-offset bases. Byte-offset casts use the inferred address space. String-based postprocessing is removed.
Address-space regression coverage
testing/python/metal/test_metal_address_space.py
Tests compare compiled and source-only Metal output, check threadgroup casts, reject generic half* aliases, and execute GEMM tests on MPS through TVM FFI and Torch backends.

Estimated code review effort: 4 (Complex) | ~45 minutes

Sequence Diagram(s)

sequenceDiagram
  participant BindNode
  participant MetalCodeGen
  participant GeneratedMetal
  participant MPS
  BindNode->>MetalCodeGen: provide typed pointer binding
  MetalCodeGen->>MetalCodeGen: infer storage scope and pointee type
  MetalCodeGen->>GeneratedMetal: emit qualified declarations and casts
  GeneratedMetal->>MPS: execute generated GEMM kernel
  MPS-->>GeneratedMetal: return computed result
Loading

Possibly related PRs

  • tile-ai/tilelang#2766: This PR refines related Metal shared-memory pointer address-space handling.
  • tile-ai/tilelang#2767: This PR replaces related string-based threadgroup pointer qualification with typed inference.
  • tile-ai/tilelang#2808: Both PRs modify Metal pointer address-space qualification in codegen_metal.cc.

Suggested reviewers: leiwang1999, siriusneo

🚥 Pre-merge checks | ✅ 5
✅ 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 summarizes the main change: deriving Metal pointer address spaces from TIR types.
Docstring Coverage ✅ Passed No functions found in the changed files to evaluate docstring coverage. Skipping docstring coverage check.
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 PR with unit tests

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.

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

Actionable comments posted: 2

🧹 Nitpick comments (1)
testing/python/metal/test_metal_address_space.py (1)

63-68: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low value

Make the assertions independent of the buffer name hints.

The assertions require the identifiers As and Bs in the generated Metal source. Those identifiers come from TIR var name hints through AllocVarID, so any renaming pass breaks this test for a reason unrelated to address spaces. Match the qualifier and the pointee type with a regular expression instead.

♻️ Proposed name-independent assertions
-    assert "threadgroup half* As = (threadgroup half*)" in compiled_source
-    assert "threadgroup half* Bs = (threadgroup half*)" in compiled_source
-    assert "*(threadgroup half2*)(As +" in compiled_source
-    assert "*(threadgroup half2*)(Bs +" in compiled_source
-    assert "(half*)As" not in compiled_source
-    assert "(half*)Bs" not in compiled_source
+    decls = re.findall(r"threadgroup half\* (\w+) = \(threadgroup half\*\)", compiled_source)
+    assert len(decls) >= 2, compiled_source
+    for name in decls:
+        assert f"*(threadgroup half2*)({name} +" in compiled_source
+        assert f"(half*){name}" not in compiled_source

Add import re at the top of the file.

🤖 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 `@testing/python/metal/test_metal_address_space.py` around lines 63 - 68,
Update the assertions in the Metal address-space test to avoid depending on the
generated buffer names As and Bs. Add the re module import and use
regular-expression matching to verify the threadgroup qualifier and half/half2
pointee types, while retaining the checks that reject invalid half-pointer
casts.
🤖 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.

Inline comments:
In `@src/metal/codegen/codegen_metal.cc`:
- Around line 926-947: Update GetAddrSpaceOf and its pointer-expression handling
to unwrap opaque wrappers such as builtin::tvm_access_ptr and
builtin::reinterpret before resolving the address space. Reuse the underlying
buffer/data expression’s storage scope where appropriate, while preserving
existing handling for Var, address_of(BufferLoad), and handle_add_byte_offset;
avoid allowing unsupported wrapped operands to reach the unconditional
GetStorageScopeOf failure.

In `@testing/python/metal/test_metal_address_space.py`:
- Around line 57-59: Update test_shared_alias_address_spaces_are_type_driven to
call lower_to_metal with enable_device_compile=False, since this source-only
test should use the Metal without-compile lowering path. Remove the ineffective
TVM_COMPILE_FORCE_FALLBACK environment setup unless it is needed elsewhere.

---

Nitpick comments:
In `@testing/python/metal/test_metal_address_space.py`:
- Around line 63-68: Update the assertions in the Metal address-space test to
avoid depending on the generated buffer names As and Bs. Add the re module
import and use regular-expression matching to verify the threadgroup qualifier
and half/half2 pointee types, while retaining the checks that reject invalid
half-pointer casts.
🪄 Autofix (Beta)

Fix all unresolved CodeRabbit comments on this PR:

  • Push a commit to this branch (recommended)
  • Create a new PR with the fixes

ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Pro Plus

Run ID: e99d3463-5a81-4302-921c-5ff5b8e2a6fe

📥 Commits

Reviewing files that changed from the base of the PR and between e01c498 and 559bd7e.

📒 Files selected for processing (3)
  • src/metal/codegen/codegen_metal.cc
  • src/metal/codegen/codegen_metal.h
  • testing/python/metal/test_metal_address_space.py

Comment thread src/metal/codegen/codegen_metal.cc
Comment thread testing/python/metal/test_metal_address_space.py Outdated
@GY-Bai

GY-Bai commented Aug 9, 2026

Copy link
Copy Markdown
Contributor Author

Closing this in favor of a smaller TileLang-side implementation aligned with Apache TVM #20101, which has now been merged. I’ll send the updated version in a new PR.

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.

1 participant