Files
tilelang/testing/python/language
Lei Wang bdf436d4c3 [Feature] Support TMA store in T.tma_copy() (#1981)
* [Feature] Support TMA store in T.tma_copy() with user-managed synchronization

T.tma_copy() previously only supported TMA loads (global -> shared) with
user-managed barrier synchronization. This extends it to also support TMA
stores (shared -> global), where:

- The barrier parameter is now optional (not needed for stores)
- For stores, only tma_store + tma_store_arrive is emitted (no wait)
- Users call T.tma_store_wait() explicitly to synchronize

* Refine tma_store_wait to accept count parameter and clean up builtin signatures

- Wire tma_store_wait count parameter through the full stack (Python -> C++
  builtin registration -> CUDA/CuTeDSL codegen)
- Replace vague *args with proper signatures for fence_proxy_async() and
  tma_store_arrive() (both take no arguments)
- Improve docstrings for internal TMA APIs (create_tma_descriptor, tma_load,
  tma_load_2sm)
2026-03-27 15:29:48 +08:00
..