fix(intrinsics): add missing _legalize_to_buffer_region in SM70 emitter

[Coloured-glaze]

#1782

device1 simple test:

device2 simple test :

pytest testing/python/kernel/ :
https://github.com/tile-ai/tilelang/blob/main/CONTRIBUTING.md#test-locally

Summary by CodeRabbit

• Refactor
  • Enhanced internal buffer handling to normalize various buffer input types into a unified representation, improving code consistency and robustness.

[github-actions]

👋 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]

📝 Walkthrough

Walkthrough

A new static method _legalize_to_buffer_region is added to the TensorCoreIntrinEmitter class in mma_sm70_macro_generator.py. This method normalizes Buffer, BufferRegion, or BufferLoad inputs into a unified BufferRegion representation. Supporting imports are expanded to include tvm.tir, tvm.ir.Range, BufferLoad, and utility functions for buffer region extraction.

Changes

Cohort / File(s)	Summary
Buffer Normalization Utility tilelang/intrinsics/mma_sm70_macro_generator.py	Added _legalize_to_buffer_region static method to normalize diverse buffer-like types (Buffer, BufferRegion, BufferLoad) into a unified BufferRegion format; expanded imports to support tvm.tir, tvm.ir.Range, and buffer region extraction utilities.

Estimated code review effort

🎯 2 (Simple) | ⏱️ ~10 minutes

Possibly related issues

• [BUG] AttributeError: 'TensorCoreIntrinEmitter' object has no attribute '_legalize_to_buffer_region' #1782: Directly addresses the missing TensorCoreIntrinEmitter._legalize_to_buffer_region method that was causing an AttributeError in tensor core intrinsic emission workflows.

Poem

🐰 A method hops into place, so neat,
Buffers of all kinds now dance as one beat,
BufferLoad and Region in harmony blend,
Legalization ensures they comprehend! ✨

🚥 Pre-merge checks | ✅ 3

✅ Passed checks (3 passed)

Check name	Status	Explanation
Description Check	✅ Passed	Check skipped - CodeRabbit’s high-level summary is enabled.
Title check	✅ Passed	The PR title accurately describes the main change: adding the missing _legalize_to_buffer_region method to the SM70 emitter, which is the primary modification shown in the code summary.
Docstring Coverage	✅ Passed	Docstring coverage is 100.00% which is sufficient. The required threshold is 80.00%.

_{✏️ Tip: You can configure your own custom pre-merge checks in the settings.}

✨ Finishing touches

• 📝 Generate docstrings

🧪 Generate unit tests (beta)

• Create PR with unit tests
• Post copyable unit tests in a comment

---

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

• X
• Mastodon
• Reddit
• LinkedIn

_{Comment @coderabbitai help to get the list of available commands and usage tips.}

[Coloured-glaze]

There are too many questions about "pytest", and I'm not sure how to solve them.

[LeiWang1999]

@Coloured-glaze Thanks, that makes sense. Just a note: some tests haven't been run on sm70, as we are currently using sm90 for testing

[Copilot]

Pull request overview

Adds the missing _legalize_to_buffer_region helper to the SM70 Tensor Core intrinsic emitter so that Volta/SM70 GEMM lowering no longer fails with AttributeError and behaves consistently with the generic MMA and MFMA emitters.

Changes:

• Add the necessary imports (tvm.tir as tir, Range, BufferLoad, and get_buffer_region_from_load) to support buffer-region legalization.
• Implement TensorCoreIntrinEmitter._legalize_to_buffer_region in mma_sm70_macro_generator.py, matching the semantics used in mma_macro_generator.py and mfma_macro_generator.py and wiring it into ldmatrix_a/ldmatrix_b.

---

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

[coderabbitai]

Actionable comments posted: 1

🤖 Fix all issues with AI agents

In `@tilelang/intrinsics/mma_sm70_macro_generator.py`:
- Around line 512-523: The code builds BufferRegion ranges with hard-coded
tir.IntImm("int32", ...) causing mixed dtypes; update the mins/extents to use
tir.const(value, dtype) (or equivalent) using the actual dtype of each
shape/index element so types match TIR (e.g., when constructing mins from
obj.shape use each shape element's dtype, and when constructing ones for indices
use the dtype of obj.indices elements); update both the BufferRegion(obj,
ranges) path (where mins come from obj.shape) and the BufferRegion(obj.buffer,
ranges) fallback for BufferLoad so Range.from_min_extent receives matching-dtype
min and extent values (refer to BufferRegion, Range.from_min_extent, BufferLoad,
get_buffer_region_from_load, obj.shape and obj.indices).