[Metax][TLE] Metax TLE support local pointer.

[cjh9368]

METAX backend support for TLE GPU local pointers

This patch adds METAX/MACA support for TLE local pointers. TLE kernels can now create pointers to shared-memory allocations and access the referenced data through standard Triton load/store operations.

Supported APIs

• tle.gpu.alloc(..., scope=tle.gpu.smem)

  • Allocates a TLE shared-memory buffer.

• tle.gpu.local_ptr(buffer, indices=None)

  • Creates local pointers to a TLE shared-memory allocation.
  • Supports full-buffer views when indices is omitted.
  • Supports indexed 1-D and multi-dimensional views.
  • Supports scalar, tensor, static, and dynamic indices.

• tl.load(local_ptr, mask=None, other=None)

  • Loads data through TLE local pointers.
  • Supports masked and unmasked loads.

• tl.store(local_ptr, value, mask=None)

  • Stores data through TLE local pointers.
  • Supports masked and unmasked stores.

These APIs can be used in elementwise kernels, loops, and conditional control-flow regions. Local pointers can also be used together with tle.gpu.copy to stage global-memory data in TLE shared memory before accessing it through tl.load and tl.store.

Lowering path

Key points for the METAX implementation:

• Local-pointer loads and stores are lowered to METAX shared-memory load/store instructions.
• The lowering detects the address space of the generated LLVM pointers and selects the shared-memory path when necessary.
• Vectorized shared-memory loads and stores are supported when pointer alignment permits.
• Shared-memory side effects are modeled for Triton atomic operations.
• The TLE local-pointer barrier insertion pass remains enabled to preserve the required memory ordering.

Scope and limitations

This PR enables TLE local pointers for load/store-based kernels.

TLE kernels containing a tt.dot/tl.dot operation are not supported on the METAX backend yet. Such kernels are rejected during compilation. Dot-based local-pointer tests are therefore excluded from the METAX test suite for now.

[sunnycase]

Thanks for adding METAX support for tle.gpu.local_ptr! Have you observed any performance improvements from this implementation in specific operators or workloads? If so, could you please share the relevant benchmarks and comparisons?