[llvm] 882be8f - [LLVM][NVPTX] Add movmatrix intrinsic and PTX instruction support (#190109)
via llvm-commits
llvm-commits at lists.llvm.org
Thu Apr 16 00:33:53 PDT 2026
Author: Varad Rahul Kamthe
Date: 2026-04-16T13:03:48+05:30
New Revision: 882be8f173d49fbbb58660fa556d782947bba33f
URL: https://github.com/llvm/llvm-project/commit/882be8f173d49fbbb58660fa556d782947bba33f
DIFF: https://github.com/llvm/llvm-project/commit/882be8f173d49fbbb58660fa556d782947bba33f.diff
LOG: [LLVM][NVPTX] Add movmatrix intrinsic and PTX instruction support (#190109)
This change adds NVVM intrinsic and NVPTX backend support for the movmatrix PTX instruction. Tests are added in `movmatrix.ll` and verified through `ptxas-13.1`.
PTX Spec Reference: https://docs.nvidia.com/cuda/parallel-thread-execution/#warp-level-matrix-instructions-movmatrix
Assisted-by: Cursor
Added:
llvm/test/CodeGen/NVPTX/movmatrix.ll
Modified:
llvm/docs/NVPTXUsage.rst
llvm/include/llvm/IR/IntrinsicsNVVM.td
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
Removed:
################################################################################
diff --git a/llvm/docs/NVPTXUsage.rst b/llvm/docs/NVPTXUsage.rst
index 70341e87c4770..2b32d35510a16 100644
--- a/llvm/docs/NVPTXUsage.rst
+++ b/llvm/docs/NVPTXUsage.rst
@@ -3936,6 +3936,42 @@ 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
+threads of a warp. Each thread provides a 32-bit register containing two
+packed ``.b16`` elements, and receives back two packed ``.b16`` elements
+from the transposed matrix in the same format.
+
+The mandatory ``.sync`` qualifier indicates that ``movmatrix`` causes the
+executing thread to wait until all threads in the warp execute the same
+``movmatrix`` intrinsic before resuming execution.
+
+The mandatory ``.aligned`` qualifier indicates that all threads in the warp
+must execute the same ``movmatrix`` intrinsic. In conditionally executed
+code, a ``movmatrix`` intrinsic should only be used if it is known that
+all threads in the warp evaluate the condition identically, otherwise the
+behavior is undefined.
+
+For more information, refer to the `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-movmatrix>`__.
+
+
Other Intrinsics
----------------
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 3fa219ee89325..b81d347eb3e32 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -2811,6 +2811,12 @@ foreach transposed = [0, 1] in {
}
}
+// movmatrix
+def int_nvvm_movmatrix_sync_aligned_m8n8_trans_b16
+ : Intrinsic<[llvm_i32_ty], [llvm_i32_ty],
+ [IntrInaccessibleMemOnly, IntrConvergent,
+ IntrNoCallback, IntrNoFree]>;
+
// MAPA
let IntrProperties = [IntrNoMem, IntrSpeculatable, NoCapture<ArgIndex<0>>] in {
def int_nvvm_mapa
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index cde6b2c8760f4..8979276bc5afb 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -5631,6 +5631,17 @@ defset list<WMMA_INSTR> STMATRIXs = {
} // transposed
} // defset
+// movmatrix
+let isConvergent = true in {
+def MOVMATRIX_SYNC_ALIGNED_M8N8_TRANS_B16
+ : BasicNVPTXInst<(outs B32:$dst),
+ (ins B32:$src),
+ "movmatrix.sync.aligned.m8n8.trans.b16",
+ [(set B32:$dst,
+ (int_nvvm_movmatrix_sync_aligned_m8n8_trans_b16 B32:$src))]>,
+ Requires<[hasSM<75>, hasPTX<78>]>;
+} // isConvergent = true
+
// Constructing non-flat DAGs is still a pain. I can't !subst a dag node with a
// dag, so the ptx.version must be appended *after* foreach replaces 'ins' with
// the instruction record.
diff --git a/llvm/test/CodeGen/NVPTX/movmatrix.ll b/llvm/test/CodeGen/NVPTX/movmatrix.ll
new file mode 100644
index 0000000000000..baf403843d516
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/movmatrix.ll
@@ -0,0 +1,40 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_75 -mattr=+ptx78 | FileCheck %s
+; RUN: %if ptxas-sm_75 && ptxas-isa-7.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_75 -mattr=+ptx78 | %ptxas-verify -arch=sm_75 %}
+
+declare i32 @llvm.nvvm.movmatrix.sync.aligned.m8n8.trans.b16(i32)
+
+; CHECK-LABEL: test_movmatrix
+define i32 @test_movmatrix(i32 %a) {
+; CHECK-LABEL: test_movmatrix(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [test_movmatrix_param_0];
+; CHECK-NEXT: movmatrix.sync.aligned.m8n8.trans.b16 %r2, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %d = call i32 @llvm.nvvm.movmatrix.sync.aligned.m8n8.trans.b16(i32 %a)
+ ret i32 %d
+}
+
+; Test that LLVM does not CSE two movmatrix calls with the same input,
+; as the result depends on values from other threads in the warp.
+define i32 @test_movmatrix_cse(i32 %a) {
+; CHECK-LABEL: test_movmatrix_cse(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [test_movmatrix_cse_param_0];
+; CHECK-NEXT: movmatrix.sync.aligned.m8n8.trans.b16 %r2, %r1;
+; CHECK-NEXT: movmatrix.sync.aligned.m8n8.trans.b16 %r3, %r1;
+; CHECK-NEXT: add.s32 %r4, %r2, %r3;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
+; CHECK-NEXT: ret;
+ %d1 = call i32 @llvm.nvvm.movmatrix.sync.aligned.m8n8.trans.b16(i32 %a)
+ %d2 = call i32 @llvm.nvvm.movmatrix.sync.aligned.m8n8.trans.b16(i32 %a)
+ %sum = add i32 %d1, %d2
+ ret i32 %sum
+}
More information about the llvm-commits
mailing list