[MetaxGPU][testing] correct FP64 MMA store layout to fix the test pre…

[xxxzh99]

The corresponding fp64 types have been added to the macro files, and the appropriate fp64 store calculation logic has been added to the utils file.

Summary by CodeRabbit

• New Features

  • Extended float64 data type support across tensor core matrix multiply-accumulate operations and sparse tensor operations.

• Bug Fixes

  • Improved 64-bit accumulation result storage indexing for accurate data placement in MMA operations.

[coderabbitai]

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info

⚙️ Run configuration

Configuration used: defaults

Review profile: CHILL

Plan: Pro

Run ID: 3f057168-88bf-4a61-8ba6-deba839bcbd8

📥 Commits

Reviewing files that changed from the base of the PR and between b746f0b and 1ac49e8.

📒 Files selected for processing (4)

• src/maca/codegen/codegen_maca.cc
• tilelang/maca/intrinsics/layout/utils.py
• tilelang/maca/intrinsics/macro/mma_macro_generator.py
• tilelang/maca/intrinsics/macro/mma_sp_macro_generator.py

🚧 Files skipped from review as they are similar to previous changes (4)

• src/maca/codegen/codegen_maca.cc
• tilelang/maca/intrinsics/layout/utils.py
• tilelang/maca/intrinsics/macro/mma_macro_generator.py
• tilelang/maca/intrinsics/macro/mma_sp_macro_generator.py

---

📝 Walkthrough

Walkthrough

This PR adds float64 support to MACA tensor core MMA infrastructure by introducing a float64-aware store index mapping utility, updating both dense and sparse MMA macro generators to handle 64-bit types through dtype abbreviations and k_dim computation, and extending C++ codegen to rewrite float64x4 intrinsics.

Changes

Float64 Support in MACA Tensor Cores

Layer / File(s)	Summary
Float64 store index mapping foundation tilelang/maca/intrinsics/layout/utils.py	mma_store_index_map now accepts an optional is_float64 parameter and computes store indices directly from thread/local IDs for float64 instead of using the shared-access layout mapping.
Dense MMA macro generator float64 support tilelang/maca/intrinsics/macro/mma_macro_generator.py	TensorCoreIntrinEmitter adds float64 to dtype_abbrv (fp64), sets k_dim=4 for 64-bit input types, maps float64 to f64 in MMA prefix generation, and passes is_float64 computed from accumulator dtype through store index mapping in both shared and global store paths.
Sparse MMA macro generator float64 support tilelang/maca/intrinsics/macro/mma_sp_macro_generator.py	SparseTensorCoreIntrinEmitter mirrors dense MMA changes: adds float64 dtype abbreviation, updates k_dim and MMA prefix suffix generation, and wires is_float64 flag through store index mapping for shared and global fragment-to-buffer transfers.
C++ codegen float64 dtype mapping src/maca/codegen/codegen_maca.cc	MFMA code generation now includes float64x4 in the dtype_map lookup table to enable rewriting of 64-bit float intrinsics.

Estimated code review effort

🎯 3 (Moderate) | ⏱️ ~20 minutes

🐰 Float64 dreams now take flight,
Dense and sparse macros shine so bright,
Store indices dance with memory's grace,
Sixty-four bits now find their place! ✨

🚥 Pre-merge checks | ✅ 4 | ❌ 1

❌ Failed checks (1 warning)

Check name	Status	Explanation	Resolution
Docstring Coverage	⚠️ Warning	Docstring coverage is 3.03% which is insufficient. The required threshold is 80.00%.	Write docstrings for the functions missing them to satisfy the coverage threshold.

✅ Passed checks (4 passed)

Check name	Status	Explanation
Description Check	✅ Passed	Check skipped - CodeRabbit’s high-level summary is enabled.
Title check	✅ Passed	The title '[MetaxGPU][testing] correct FP64 MMA store layout to fix the test pre…' directly addresses the main change: correcting FP64 (float64) MMA store layout across multiple files to fix test precision issues.
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.

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

✨ 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

• X
• Mastodon
• Reddit
• LinkedIn

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

[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! 🚀