part of https://github.com/tile-ai/tilelang/pull/1869
## Summary
Add TIR builtins and storage scope for Metal cooperative_tensor
operations (MetalPerformancePrimitives / Metal 4).
## Motivation
Apple Metal 4 introduces MetalPerformancePrimitives (MPP) with
`matmul2d` using `cooperative_tensor` operands. On M5, this routes to
NAX tensor cores; on M1-M4, it falls back to simdgroup matrix
instructions. These TIR primitives enable backend codegen to emit MPP
calls.
## Changes
### New TIR builtins
- `cooperative_tensor_fill(d, index, value, rows, cols)`
- `cooperative_tensor_load(d, index, ptr, stride, rows, cols,
transpose)`
- `cooperative_tensor_store(d, index, ptr, stride, rows, cols,
transpose)`
- `cooperative_tensor_multiply_accumulate(d, di, a, ai, b, bi, c, ci, M,
N, K, trans_a, trans_b)`
### New storage scope
- `metal.cooperative_tensor` (`StorageRank::kMetalCooperativeTensor`)
### Files changed
- `include/tvm/tirx/builtin.h` — Op declarations
- `src/tirx/op/builtin.cc` — Op registrations
- `python/tvm/tirx/op.py` — Python wrappers
- `python/tvm/script/ir_builder/tirx/ir.py` — Script parser exports
- `src/runtime/thread_storage_scope.h` — StorageRank enum + scope
parsing
These builtins mirror the existing `simdgroup_*` builtins for the older
Metal simdgroup matrix API, extended with M/N/K dimension parameters for
the matmul2d descriptor.
Yichen Yan
c309e4ea5f