[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:05:17 PDT 2026
https://github.com/varadk27 updated https://github.com/llvm/llvm-project/pull/190109
>From c3fe471d72f5db1d256bf4da3eceeaf2092bbc97 Mon Sep 17 00:00:00 2001
From: Varad Rahul Kamthe <vkamthe at nvidia.com>
Date: Mon, 23 Mar 2026 13:05:54 +0000
Subject: [PATCH 1/4] Add movmatrix intrinsic and PTX instruction support
---
llvm/include/llvm/IR/IntrinsicsNVVM.td | 5 +++++
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 11 +++++++++++
llvm/test/CodeGen/NVPTX/movmatrix.ll | 20 ++++++++++++++++++++
3 files changed, 36 insertions(+)
create mode 100644 llvm/test/CodeGen/NVPTX/movmatrix.ll
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index b3e0033d005a9..cf6dea651c9d5 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -2802,6 +2802,11 @@ foreach transposed = [0, 1] in {
}
}
+// movmatrix
+def int_nvvm_movmatrix_sync_aligned_m8n8_trans_b16
+ : DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty],
+ [IntrNoMem, IntrConvergent, IntrNoCallback]>;
+
// 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..4e4acc25f6305 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
+ : NVPTXInst<(outs B32:$dst),
+ (ins B32:$src),
+ "movmatrix.sync.aligned.m8n8.trans.b16 $dst, $src;",
+ [(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..0d38e88dacf26
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/movmatrix.ll
@@ -0,0 +1,20 @@
+; 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
+}
>From 33584c62d5400b3e4dfdfceeab765bf9a5705db9 Mon Sep 17 00:00:00 2001
From: Varad Rahul Kamthe <vkamthe at nvidia.com>
Date: Thu, 2 Apr 2026 06:19:28 +0000
Subject: [PATCH 2/4] [NVPTX] Add documentation for movmatrix intrinsic
---
llvm/docs/NVPTXUsage.rst | 28 ++++++++++++++++++++++++++++
1 file changed, 28 insertions(+)
diff --git a/llvm/docs/NVPTXUsage.rst b/llvm/docs/NVPTXUsage.rst
index 70341e87c4770..9ebcc7c6708a0 100644
--- a/llvm/docs/NVPTXUsage.rst
+++ b/llvm/docs/NVPTXUsage.rst
@@ -3936,6 +3936,34 @@ 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>`__.
+movmatrix 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 generates
+the ``movmatrix.sync.aligned.m8n8.trans.b16`` PTX instruction, which performs
+a warp-synchronous register shuffle to transpose an 8x8 matrix of 16-bit
+values held in warp registers.
+
+The 32 threads of a warp collectively hold the elements of an 8x8 matrix of
+``.b16`` values, with two elements packed into each thread's 32-bit register.
+The transposed result is returned in the same packed layout.
+
+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
----------------
>From 000d27c9af9e66c833d3f514e210b440b62cc4af Mon Sep 17 00:00:00 2001
From: Varad Rahul Kamthe <vkamthe at nvidia.com>
Date: Wed, 8 Apr 2026 13:12:48 +0000
Subject: [PATCH 3/4] [NVPTX] Address review comments for movmatrix intrinsic
---
llvm/docs/NVPTXUsage.rst | 20 +++++++++++---------
llvm/include/llvm/IR/IntrinsicsNVVM.td | 2 +-
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 4 ++--
llvm/test/CodeGen/NVPTX/movmatrix.ll | 20 ++++++++++++++++++++
4 files changed, 34 insertions(+), 12 deletions(-)
diff --git a/llvm/docs/NVPTXUsage.rst b/llvm/docs/NVPTXUsage.rst
index 9ebcc7c6708a0..d94432d237c89 100644
--- a/llvm/docs/NVPTXUsage.rst
+++ b/llvm/docs/NVPTXUsage.rst
@@ -3936,8 +3936,8 @@ 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>`__.
-movmatrix Intrinsics
---------------------
+Warp-level Matrix Transpose Intrinsics
+---------------------------------------
'``llvm.nvvm.movmatrix.sync.aligned.m8n8.trans.b16``'
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -3952,18 +3952,20 @@ Syntax:
Overview:
"""""""""
-The '``@llvm.nvvm.movmatrix.sync.aligned.m8n8.trans.b16``' intrinsic generates
-the ``movmatrix.sync.aligned.m8n8.trans.b16`` PTX instruction, which performs
-a warp-synchronous register shuffle to transpose an 8x8 matrix of 16-bit
-values held in warp registers.
+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 32 threads of a warp collectively hold the elements of an 8x8 matrix of
-``.b16`` values, with two elements packed into each thread's 32-bit register.
-The transposed result is returned in the same packed layout.
+This is a warp-synchronous operation — all threads in the warp must
+execute the intrinsic together, and the result in each thread depends on
+values provided by other threads.
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 cf6dea651c9d5..8b81d35794658 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -2805,7 +2805,7 @@ foreach transposed = [0, 1] in {
// movmatrix
def int_nvvm_movmatrix_sync_aligned_m8n8_trans_b16
: DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty],
- [IntrNoMem, IntrConvergent, IntrNoCallback]>;
+ [IntrInaccessibleMemOnly, IntrConvergent]>;
// MAPA
let IntrProperties = [IntrNoMem, IntrSpeculatable, NoCapture<ArgIndex<0>>] in {
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 4e4acc25f6305..8979276bc5afb 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -5634,9 +5634,9 @@ defset list<WMMA_INSTR> STMATRIXs = {
// movmatrix
let isConvergent = true in {
def MOVMATRIX_SYNC_ALIGNED_M8N8_TRANS_B16
- : NVPTXInst<(outs B32:$dst),
+ : BasicNVPTXInst<(outs B32:$dst),
(ins B32:$src),
- "movmatrix.sync.aligned.m8n8.trans.b16 $dst, $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>]>;
diff --git a/llvm/test/CodeGen/NVPTX/movmatrix.ll b/llvm/test/CodeGen/NVPTX/movmatrix.ll
index 0d38e88dacf26..baf403843d516 100644
--- a/llvm/test/CodeGen/NVPTX/movmatrix.ll
+++ b/llvm/test/CodeGen/NVPTX/movmatrix.ll
@@ -18,3 +18,23 @@ define i32 @test_movmatrix(i32 %a) {
%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
+}
>From 7458a433803742e1871d0da82e4a8e1f9c98d4e1 Mon Sep 17 00:00:00 2001
From: Varad Rahul Kamthe <vkamthe at nvidia.com>
Date: Tue, 14 Apr 2026 13:03:27 +0000
Subject: [PATCH 4/4] [NVPTX] Address latest review comments for movmatrix
intrinsic
---
llvm/docs/NVPTXUsage.rst | 12 +++++++++---
llvm/include/llvm/IR/IntrinsicsNVVM.td | 5 +++--
2 files changed, 12 insertions(+), 5 deletions(-)
diff --git a/llvm/docs/NVPTXUsage.rst b/llvm/docs/NVPTXUsage.rst
index d94432d237c89..2b32d35510a16 100644
--- a/llvm/docs/NVPTXUsage.rst
+++ b/llvm/docs/NVPTXUsage.rst
@@ -3958,9 +3958,15 @@ 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.
-This is a warp-synchronous operation — all threads in the warp must
-execute the intrinsic together, and the result in each thread depends on
-values provided by other threads.
+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>`__.
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 8b81d35794658..8ed062d767477 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -2804,8 +2804,9 @@ foreach transposed = [0, 1] in {
// movmatrix
def int_nvvm_movmatrix_sync_aligned_m8n8_trans_b16
- : DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty],
- [IntrInaccessibleMemOnly, IntrConvergent]>;
+ : Intrinsic<[llvm_i32_ty], [llvm_i32_ty],
+ [IntrInaccessibleMemOnly, IntrConvergent,
+ IntrNoCallback, IntrNoFree]>;
// MAPA
let IntrProperties = [IntrNoMem, IntrSpeculatable, NoCapture<ArgIndex<0>>] in {
More information about the llvm-commits
mailing list