[PR] [TIR] Support more intrinsics and `get_mma_intrin_group` utility [tvm]

["Ubospica (via GitHub)" <gi...@apache.org>]

Ubospica opened a new pull request, #16073:
URL: https://github.com/apache/tvm/pull/16073

   This PR focuses on adding supporting various intrinsics for matmul scheduling. To be specific, this PR:
   - Adds support for transposed A in ldmatrix and mma_sync
   - Changes all `T.launch_thread(tx, 32)` annotations to for loops `for tx in T.thread_binding(0, WARP_SIZE, "threadIdx.x")`. This is for the convenience of later transformation.
   - Refactor some logic and add a utility get_mma_intrin_group to get a group of intrinsics:
   ```
   def get_mma_intrin_group(
       load_scope: Literal["shared", "shared.dyn"],
       store_scope: Literal["global", "shared", "shared.dyn"],
       in_dtype: Literal["float16", "int8"],
       out_dtype: Literal["float16", "float32", "int32"],
       trans_a: bool,
       trans_b: bool,
       not_use_mma_store_intrinic: bool = True,
       store_to_smem_dtype: Optional[Literal["float16", "float32", "int32"]] = None,
   ) -> Dict[str, str]
   ```
   - Avoid use the current mma_store intrinsic. Instread, use BufferStore statements. 
     - This is because if we use mma_store intrinsic, during swizzling shared memory visits, our rearrangement scheme will involve areas accessed by different mma_store calls. This makes swizzling quite complex. But BufferStore will not face this problem
   
   


-- 
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.

To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org

For queries about this service, please contact Infrastructure at:
users@infra.apache.org

Re: [PR] [TIR] Support more mma intrinsics and `get_mma_intrin_group` utility [tvm]

["vinx13 (via GitHub)" <gi...@apache.org>]

vinx13 merged PR #16073:
URL: https://github.com/apache/tvm/pull/16073


-- 
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.

To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org

For queries about this service, please contact Infrastructure at:
users@infra.apache.org