[llvm] [LLVM][NVPTX] Add movmatrix intrinsic and PTX instruction support (PR #190109)
Varad Rahul Kamthe via llvm-commits
llvm-commits at lists.llvm.org
Tue Apr 14 06:37:47 PDT 2026
================
@@ -3936,6 +3936,36 @@ an event.
For more information on the pmevent instructions, refer to the `PTX ISA
<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#miscellaneous-instructions-pmevent>`__.
+Warp-level Matrix Transpose Intrinsics
+---------------------------------------
+
+'``llvm.nvvm.movmatrix.sync.aligned.m8n8.trans.b16``'
+^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+Syntax:
+"""""""
+
+.. code-block:: llvm
+
+ declare i32 @llvm.nvvm.movmatrix.sync.aligned.m8n8.trans.b16(i32 %src)
+
+Overview:
+"""""""""
+
+The '``@llvm.nvvm.movmatrix.sync.aligned.m8n8.trans.b16``' intrinsic
+transposes an 8x8 matrix of 16-bit elements distributed across all 32
----------------
varadk27 wrote:
Sticking with `i32 `as we wanted to keep the types similar to to ldmatrix which return `i32 `instead of `<2 x i16>`
https://github.com/llvm/llvm-project/pull/190109
More information about the llvm-commits
mailing list