[llvm] [NVPTX] Support intrinsics for reserved shared memory special registers (PR #182354)
via llvm-commits
llvm-commits at lists.llvm.org
Fri Feb 20 14:55:03 PST 2026
https://github.com/yasmincs updated https://github.com/llvm/llvm-project/pull/182354
>From 3a85099ba627eb6b28674d096792b7e7ff44b665 Mon Sep 17 00:00:00 2001
From: ysarita <ysarita at nvidia.com>
Date: Thu, 19 Feb 2026 19:48:09 +0000
Subject: [PATCH 1/3] Add NVPTX support for shared memory special register
intrinsics
---
llvm/include/llvm/IR/IntrinsicsNVVM.td | 12 +++++++
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 10 ++++++
llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp | 4 +++
llvm/lib/Target/NVPTX/NVPTXRegisterInfo.td | 6 ++++
.../CodeGen/NVPTX/reserved-smem-offset.ll | 35 +++++++++++++++++++
5 files changed, 67 insertions(+)
create mode 100644 llvm/test/CodeGen/NVPTX/reserved-smem-offset.ll
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index f2e1bcb5517c8..c0c80d92017b7 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -2450,6 +2450,18 @@ defm int_nvvm_read_ptx_sreg_cluster_nctaid : PTXReadSRegIntrinsicNB_v4i32<MAX_GR
def int_nvvm_read_ptx_sreg_cluster_ctarank : PTXReadSRegIntrinsicNB_r32;
def int_nvvm_read_ptx_sreg_cluster_nctarank : PTXReadSRegIntrinsicNB_r32;
+//
+// Reserved Shared Memory Intrinsics
+//
+def int_nvvm_read_ptx_sreg_reserved_smem_offset_begin
+ : PTXReadSRegIntrinsicNB_r32<[], "llvm.nvvm.read.ptx.sreg.reserved_smem_offset_begin">;
+def int_nvvm_read_ptx_sreg_reserved_smem_offset_end
+ : PTXReadSRegIntrinsicNB_r32<[], "llvm.nvvm.read.ptx.sreg.reserved_smem_offset_end">;
+def int_nvvm_read_ptx_sreg_reserved_smem_offset_cap
+ : PTXReadSRegIntrinsicNB_r32<[], "llvm.nvvm.read.ptx.sreg.reserved_smem_offset_cap">;
+def int_nvvm_read_ptx_sreg_reserved_smem_offset_0
+ : PTXReadSRegIntrinsicNB_r32<[], "llvm.nvvm.read.ptx.sreg.reserved_smem_offset_0">;
+
def int_nvvm_read_ptx_sreg_total_smem_size :
PTXReadSRegIntrinsicNB_r32<name = "llvm.nvvm.read.ptx.sreg.total_smem_size">;
def int_nvvm_read_ptx_sreg_aggr_smem_size :
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 7b7b11d14ecc8..6ce77651ac164 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -2986,6 +2986,16 @@ def : Pat<(int_nvvm_read_ptx_sreg_envreg29), (MOV_SPECIAL ENVREG29)>;
def : Pat<(int_nvvm_read_ptx_sreg_envreg30), (MOV_SPECIAL ENVREG30)>;
def : Pat<(int_nvvm_read_ptx_sreg_envreg31), (MOV_SPECIAL ENVREG31)>;
+// Reserved Shared Memory special register reads
+def : Pat<(int_nvvm_read_ptx_sreg_reserved_smem_offset_begin),
+ (MOV_SPECIAL RESERVED_SMEM_OFFSET_BEGIN)>, Requires<[hasPTX<76>, hasSM<80>]>;
+def : Pat<(int_nvvm_read_ptx_sreg_reserved_smem_offset_end),
+ (MOV_SPECIAL RESERVED_SMEM_OFFSET_END)>, Requires<[hasPTX<76>, hasSM<80>]>;
+def : Pat<(int_nvvm_read_ptx_sreg_reserved_smem_offset_cap),
+ (MOV_SPECIAL RESERVED_SMEM_OFFSET_CAP)>, Requires<[hasPTX<76>, hasSM<80>]>;
+def : Pat<(int_nvvm_read_ptx_sreg_reserved_smem_offset_0),
+ (MOV_SPECIAL RESERVED_SMEM_OFFSET_0)>, Requires<[hasPTX<76>, hasSM<80>]>;
+
//-----------------------------------
// Texture Intrinsics
//-----------------------------------
diff --git a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
index 646b554878c70..c7751fb50791a 100644
--- a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
@@ -98,6 +98,10 @@ BitVector NVPTXRegisterInfo::getReservedRegs(const MachineFunction &MF) const {
markSuperRegs(Reserved, NVPTX::VRFrame64);
markSuperRegs(Reserved, NVPTX::VRFrameLocal64);
markSuperRegs(Reserved, NVPTX::VRDepot);
+ markSuperRegs(Reserved, NVPTX::RESERVED_SMEM_OFFSET_BEGIN);
+ markSuperRegs(Reserved, NVPTX::RESERVED_SMEM_OFFSET_END);
+ markSuperRegs(Reserved, NVPTX::RESERVED_SMEM_OFFSET_CAP);
+ markSuperRegs(Reserved, NVPTX::RESERVED_SMEM_OFFSET_0);
return Reserved;
}
diff --git a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.td b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.td
index 913487b64617a..7cf6b91af5a3e 100644
--- a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.td
@@ -44,6 +44,12 @@ foreach i = 0...31 in {
def ENVREG#i : NVPTXReg<"%envreg"#i>;
}
+// Reserved Shared Memory special registers
+def RESERVED_SMEM_OFFSET_BEGIN : NVPTXReg<"%reserved_smem_offset_begin">;
+def RESERVED_SMEM_OFFSET_END : NVPTXReg<"%reserved_smem_offset_end">;
+def RESERVED_SMEM_OFFSET_CAP : NVPTXReg<"%reserved_smem_offset_cap">;
+def RESERVED_SMEM_OFFSET_0 : NVPTXReg<"%reserved_smem_offset_0">;
+
//===----------------------------------------------------------------------===//
// Register classes.
// NOTE: if you add new vector types for a register, you must update
diff --git a/llvm/test/CodeGen/NVPTX/reserved-smem-offset.ll b/llvm/test/CodeGen/NVPTX/reserved-smem-offset.ll
new file mode 100644
index 0000000000000..a72033d302c19
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/reserved-smem-offset.ll
@@ -0,0 +1,35 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_80 | FileCheck %s
+; RUN: %if ptxas-sm_80 %{ llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_80 | %ptxas-verify -arch=sm_80 %}
+
+declare i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_begin()
+declare i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_end()
+declare i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_cap()
+declare i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_0()
+
+define i32 @test() {
+; CHECK-LABEL: test(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<8>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: mov.b32 %r1, %reserved_smem_offset_begin;
+; CHECK-NEXT: mov.b32 %r2, %reserved_smem_offset_end;
+; CHECK-NEXT: mov.b32 %r3, %reserved_smem_offset_cap;
+; CHECK-NEXT: mov.b32 %r4, %reserved_smem_offset_0;
+; CHECK-NEXT: add.s32 %r5, %r1, %r2;
+; CHECK-NEXT: add.s32 %r6, %r5, %r3;
+; CHECK-NEXT: add.s32 %r7, %r6, %r4;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r7;
+; CHECK-NEXT: ret;
+ %begin = call i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_begin()
+ %end = call i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_end()
+ %cap = call i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_cap()
+ %offset0 = call i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_0()
+
+ %ret0 = add i32 %begin, %end
+ %ret1 = add i32 %ret0, %cap
+ %ret2 = add i32 %ret1, %offset0
+
+ ret i32 %ret2
+}
>From dfab4bc8d84bcbf91fd2dd3c2962a59a55b110e2 Mon Sep 17 00:00:00 2001
From: ysarita <ysarita at nvidia.com>
Date: Thu, 19 Feb 2026 22:16:31 +0000
Subject: [PATCH 2/3] changed smem offset intrinsics to be be like other sreg
intrinsics
---
llvm/include/llvm/IR/IntrinsicsNVVM.td | 19 ++++++-----
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 32 +++++++++++++------
llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp | 4 ---
llvm/lib/Target/NVPTX/NVPTXRegisterInfo.td | 5 ---
.../CodeGen/NVPTX/reserved-smem-offset.ll | 29 ++++++++++-------
5 files changed, 50 insertions(+), 39 deletions(-)
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index c0c80d92017b7..c1018c6f257c0 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -2350,7 +2350,8 @@ foreach vec = [TV_I8, TV_I16, TV_I32,
// Accessing special registers.
//
class PTXReadSRegIntrinsicNB_r32<list<IntrinsicProperty> properties = [], string name = "">
- : NVVMPureIntrinsic<[llvm_i32_ty], [], [NoUndef<RetIndex>] # properties, name>;
+ : NVVMPureIntrinsic<[llvm_i32_ty], [], [NoUndef<RetIndex>] # properties,
+ !if(!eq(name, ""), "", !strconcat("llvm.nvvm.read.ptx.sreg.", name))>;
class PTXReadSRegIntrinsic_r32<list<IntrinsicProperty> properties = []>
: PTXReadSRegIntrinsicNB_r32<properties>, NVVMBuiltin;
@@ -2454,20 +2455,22 @@ def int_nvvm_read_ptx_sreg_cluster_nctarank : PTXReadSRegIntrinsicNB_r32;
// Reserved Shared Memory Intrinsics
//
def int_nvvm_read_ptx_sreg_reserved_smem_offset_begin
- : PTXReadSRegIntrinsicNB_r32<[], "llvm.nvvm.read.ptx.sreg.reserved_smem_offset_begin">;
+ : PTXReadSRegIntrinsicNB_r32<name = "reserved_smem_offset_begin">;
def int_nvvm_read_ptx_sreg_reserved_smem_offset_end
- : PTXReadSRegIntrinsicNB_r32<[], "llvm.nvvm.read.ptx.sreg.reserved_smem_offset_end">;
+ : PTXReadSRegIntrinsicNB_r32<name = "reserved_smem_offset_end">;
def int_nvvm_read_ptx_sreg_reserved_smem_offset_cap
- : PTXReadSRegIntrinsicNB_r32<[], "llvm.nvvm.read.ptx.sreg.reserved_smem_offset_cap">;
+ : PTXReadSRegIntrinsicNB_r32<name = "reserved_smem_offset_cap">;
def int_nvvm_read_ptx_sreg_reserved_smem_offset_0
- : PTXReadSRegIntrinsicNB_r32<[], "llvm.nvvm.read.ptx.sreg.reserved_smem_offset_0">;
+ : PTXReadSRegIntrinsicNB_r32<name = "reserved_smem_offset_0">;
+def int_nvvm_read_ptx_sreg_reserved_smem_offset_1
+ : PTXReadSRegIntrinsicNB_r32<name = "reserved_smem_offset_1">;
def int_nvvm_read_ptx_sreg_total_smem_size :
- PTXReadSRegIntrinsicNB_r32<name = "llvm.nvvm.read.ptx.sreg.total_smem_size">;
+ PTXReadSRegIntrinsicNB_r32<name = "total_smem_size">;
def int_nvvm_read_ptx_sreg_aggr_smem_size :
- PTXReadSRegIntrinsicNB_r32<name = "llvm.nvvm.read.ptx.sreg.aggr_smem_size">;
+ PTXReadSRegIntrinsicNB_r32<name = "aggr_smem_size">;
def int_nvvm_read_ptx_sreg_dynamic_smem_size :
- PTXReadSRegIntrinsicNB_r32<name = "llvm.nvvm.read.ptx.sreg.dynamic_smem_size">;
+ PTXReadSRegIntrinsicNB_r32<name = "dynamic_smem_size">;
//
// SHUFFLE
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 6ce77651ac164..2b4dce2dabae2 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -2986,16 +2986,6 @@ def : Pat<(int_nvvm_read_ptx_sreg_envreg29), (MOV_SPECIAL ENVREG29)>;
def : Pat<(int_nvvm_read_ptx_sreg_envreg30), (MOV_SPECIAL ENVREG30)>;
def : Pat<(int_nvvm_read_ptx_sreg_envreg31), (MOV_SPECIAL ENVREG31)>;
-// Reserved Shared Memory special register reads
-def : Pat<(int_nvvm_read_ptx_sreg_reserved_smem_offset_begin),
- (MOV_SPECIAL RESERVED_SMEM_OFFSET_BEGIN)>, Requires<[hasPTX<76>, hasSM<80>]>;
-def : Pat<(int_nvvm_read_ptx_sreg_reserved_smem_offset_end),
- (MOV_SPECIAL RESERVED_SMEM_OFFSET_END)>, Requires<[hasPTX<76>, hasSM<80>]>;
-def : Pat<(int_nvvm_read_ptx_sreg_reserved_smem_offset_cap),
- (MOV_SPECIAL RESERVED_SMEM_OFFSET_CAP)>, Requires<[hasPTX<76>, hasSM<80>]>;
-def : Pat<(int_nvvm_read_ptx_sreg_reserved_smem_offset_0),
- (MOV_SPECIAL RESERVED_SMEM_OFFSET_0)>, Requires<[hasPTX<76>, hasSM<80>]>;
-
//-----------------------------------
// Texture Intrinsics
//-----------------------------------
@@ -4936,6 +4926,28 @@ def INT_PTX_SREG_PM1 : PTX_READ_SREG_R32<"pm1", int_nvvm_read_ptx_sreg_pm1>;
def INT_PTX_SREG_PM2 : PTX_READ_SREG_R32<"pm2", int_nvvm_read_ptx_sreg_pm2>;
def INT_PTX_SREG_PM3 : PTX_READ_SREG_R32<"pm3", int_nvvm_read_ptx_sreg_pm3>;
+// Reserved shared memory special register reads
+def INT_PTX_SREG_RESERVED_SMEM_OFFSET_BEGIN :
+ PTX_READ_SREG_R32<"reserved_smem_offset_begin",
+ int_nvvm_read_ptx_sreg_reserved_smem_offset_begin,
+ [hasPTX<76>, hasSM<80>]>;
+def INT_PTX_SREG_RESERVED_SMEM_OFFSET_END :
+ PTX_READ_SREG_R32<"reserved_smem_offset_end",
+ int_nvvm_read_ptx_sreg_reserved_smem_offset_end,
+ [hasPTX<76>, hasSM<80>]>;
+def INT_PTX_SREG_RESERVED_SMEM_OFFSET_CAP :
+ PTX_READ_SREG_R32<"reserved_smem_offset_cap",
+ int_nvvm_read_ptx_sreg_reserved_smem_offset_cap,
+ [hasPTX<76>, hasSM<80>]>;
+def INT_PTX_SREG_RESERVED_SMEM_OFFSET_0 :
+ PTX_READ_SREG_R32<"reserved_smem_offset_0",
+ int_nvvm_read_ptx_sreg_reserved_smem_offset_0,
+ [hasPTX<76>, hasSM<80>]>;
+def INT_PTX_SREG_RESERVED_SMEM_OFFSET_1 :
+ PTX_READ_SREG_R32<"reserved_smem_offset_1",
+ int_nvvm_read_ptx_sreg_reserved_smem_offset_1,
+ [hasPTX<76>, hasSM<80>]>;
+
// TODO: It would be nice to use PTX_READ_SREG here, but it doesn't
// handle the constant.
def INT_PTX_SREG_WARPSIZE :
diff --git a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
index c7751fb50791a..646b554878c70 100644
--- a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
@@ -98,10 +98,6 @@ BitVector NVPTXRegisterInfo::getReservedRegs(const MachineFunction &MF) const {
markSuperRegs(Reserved, NVPTX::VRFrame64);
markSuperRegs(Reserved, NVPTX::VRFrameLocal64);
markSuperRegs(Reserved, NVPTX::VRDepot);
- markSuperRegs(Reserved, NVPTX::RESERVED_SMEM_OFFSET_BEGIN);
- markSuperRegs(Reserved, NVPTX::RESERVED_SMEM_OFFSET_END);
- markSuperRegs(Reserved, NVPTX::RESERVED_SMEM_OFFSET_CAP);
- markSuperRegs(Reserved, NVPTX::RESERVED_SMEM_OFFSET_0);
return Reserved;
}
diff --git a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.td b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.td
index 7cf6b91af5a3e..ca9b641047e6c 100644
--- a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.td
@@ -44,11 +44,6 @@ foreach i = 0...31 in {
def ENVREG#i : NVPTXReg<"%envreg"#i>;
}
-// Reserved Shared Memory special registers
-def RESERVED_SMEM_OFFSET_BEGIN : NVPTXReg<"%reserved_smem_offset_begin">;
-def RESERVED_SMEM_OFFSET_END : NVPTXReg<"%reserved_smem_offset_end">;
-def RESERVED_SMEM_OFFSET_CAP : NVPTXReg<"%reserved_smem_offset_cap">;
-def RESERVED_SMEM_OFFSET_0 : NVPTXReg<"%reserved_smem_offset_0">;
//===----------------------------------------------------------------------===//
// Register classes.
diff --git a/llvm/test/CodeGen/NVPTX/reserved-smem-offset.ll b/llvm/test/CodeGen/NVPTX/reserved-smem-offset.ll
index a72033d302c19..46c010fb614a1 100644
--- a/llvm/test/CodeGen/NVPTX/reserved-smem-offset.ll
+++ b/llvm/test/CodeGen/NVPTX/reserved-smem-offset.ll
@@ -1,35 +1,40 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
-; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_80 | FileCheck %s
-; RUN: %if ptxas-sm_80 %{ llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_80 | %ptxas-verify -arch=sm_80 %}
+; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_80 -mattr=+ptx76 | FileCheck %s
+; RUN: %if ptxas-sm_80 && ptxas-isa-7.6 %{ llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_80 -mattr=+ptx76 | %ptxas-verify -arch=sm_80 %}
declare i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_begin()
declare i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_end()
declare i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_cap()
declare i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_0()
+declare i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_1()
define i32 @test() {
; CHECK-LABEL: test(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<8>;
+; CHECK-NEXT: .reg .b32 %r<10>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: mov.b32 %r1, %reserved_smem_offset_begin;
-; CHECK-NEXT: mov.b32 %r2, %reserved_smem_offset_end;
-; CHECK-NEXT: mov.b32 %r3, %reserved_smem_offset_cap;
-; CHECK-NEXT: mov.b32 %r4, %reserved_smem_offset_0;
-; CHECK-NEXT: add.s32 %r5, %r1, %r2;
-; CHECK-NEXT: add.s32 %r6, %r5, %r3;
-; CHECK-NEXT: add.s32 %r7, %r6, %r4;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r7;
+; CHECK-NEXT: mov.u32 %r1, %reserved_smem_offset_begin;
+; CHECK-NEXT: mov.u32 %r2, %reserved_smem_offset_end;
+; CHECK-NEXT: mov.u32 %r3, %reserved_smem_offset_cap;
+; CHECK-NEXT: mov.u32 %r4, %reserved_smem_offset_0;
+; CHECK-NEXT: mov.u32 %r5, %reserved_smem_offset_1;
+; CHECK-NEXT: add.s32 %r6, %r1, %r2;
+; CHECK-NEXT: add.s32 %r7, %r6, %r3;
+; CHECK-NEXT: add.s32 %r8, %r7, %r4;
+; CHECK-NEXT: add.s32 %r9, %r8, %r5;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r9;
; CHECK-NEXT: ret;
%begin = call i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_begin()
%end = call i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_end()
%cap = call i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_cap()
%offset0 = call i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_0()
+ %offset1 = call i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_1()
%ret0 = add i32 %begin, %end
%ret1 = add i32 %ret0, %cap
%ret2 = add i32 %ret1, %offset0
+ %ret3 = add i32 %ret2, %offset1
- ret i32 %ret2
+ ret i32 %ret3
}
>From 24c024b0a8642fb89ad13b5aa7b60590a486c0d7 Mon Sep 17 00:00:00 2001
From: ysarita <ysarita at nvidia.com>
Date: Fri, 20 Feb 2026 22:54:33 +0000
Subject: [PATCH 3/3] addressed reviewer comments
---
llvm/docs/NVPTXUsage.rst | 36 ++++++++++++++++++++++
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 26 ++++------------
llvm/lib/Target/NVPTX/NVPTXRegisterInfo.td | 1 -
3 files changed, 42 insertions(+), 21 deletions(-)
diff --git a/llvm/docs/NVPTXUsage.rst b/llvm/docs/NVPTXUsage.rst
index 0e7e21ad46b8d..f4dd87ea82260 100644
--- a/llvm/docs/NVPTXUsage.rst
+++ b/llvm/docs/NVPTXUsage.rst
@@ -300,6 +300,42 @@ including the portion reserved for system software use.
The '``dynamic_smem_size``' variant returns the amount of dynamic shared
memory allocated per CTA for the kernel at launch time.
+'``llvm.nvvm.read.ptx.sreg.reserved_smem_offset_*``'
+^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+Syntax:
+"""""""
+
+.. code-block:: llvm
+
+ declare i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_begin()
+ declare i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_end()
+ declare i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_cap()
+ declare i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_0()
+ declare i32 @llvm.nvvm.read.ptx.sreg.reserved_smem_offset_1()
+
+Overview:
+"""""""""
+
+The '``@llvm.nvvm.read.ptx.sreg.reserved_smem_offset_*``' intrinsics provide
+access to PTX special registers that hold information about reserved shared
+memory offsets.
+
+The '``reserved_smem_offset_begin``' intrinsic reads the start offset of
+the reserved shared memory region.
+
+The '``reserved_smem_offset_end``' intrinsic reads the end offset of the
+reserved shared memory region.
+
+The '``reserved_smem_offset_cap``' intrinsic reads the capacity limit of
+the reserved shared memory region.
+
+The '``reserved_smem_offset_0``' and '``reserved_smem_offset_1``' intrinsics
+read additional offsets in the reserved shared memory region.
+
+For more information, refer `PTX ISA
+<https://docs.nvidia.com/cuda/parallel-thread-execution/#special-registers-reserved-smem>`__.
+
Barriers
--------
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 2b4dce2dabae2..dd5c7fe680440 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -4927,26 +4927,12 @@ def INT_PTX_SREG_PM2 : PTX_READ_SREG_R32<"pm2", int_nvvm_read_ptx_sreg_pm2>;
def INT_PTX_SREG_PM3 : PTX_READ_SREG_R32<"pm3", int_nvvm_read_ptx_sreg_pm3>;
// Reserved shared memory special register reads
-def INT_PTX_SREG_RESERVED_SMEM_OFFSET_BEGIN :
- PTX_READ_SREG_R32<"reserved_smem_offset_begin",
- int_nvvm_read_ptx_sreg_reserved_smem_offset_begin,
- [hasPTX<76>, hasSM<80>]>;
-def INT_PTX_SREG_RESERVED_SMEM_OFFSET_END :
- PTX_READ_SREG_R32<"reserved_smem_offset_end",
- int_nvvm_read_ptx_sreg_reserved_smem_offset_end,
- [hasPTX<76>, hasSM<80>]>;
-def INT_PTX_SREG_RESERVED_SMEM_OFFSET_CAP :
- PTX_READ_SREG_R32<"reserved_smem_offset_cap",
- int_nvvm_read_ptx_sreg_reserved_smem_offset_cap,
- [hasPTX<76>, hasSM<80>]>;
-def INT_PTX_SREG_RESERVED_SMEM_OFFSET_0 :
- PTX_READ_SREG_R32<"reserved_smem_offset_0",
- int_nvvm_read_ptx_sreg_reserved_smem_offset_0,
- [hasPTX<76>, hasSM<80>]>;
-def INT_PTX_SREG_RESERVED_SMEM_OFFSET_1 :
- PTX_READ_SREG_R32<"reserved_smem_offset_1",
- int_nvvm_read_ptx_sreg_reserved_smem_offset_1,
- [hasPTX<76>, hasSM<80>]>;
+foreach suffix = ["begin", "end", "cap", "0", "1"] in {
+ defvar regname = "reserved_smem_offset_" # suffix;
+ defvar intr = !cast<Intrinsic>("int_nvvm_read_ptx_sreg_" # regname);
+ def "INT_PTX_SREG_RESERVED_SMEM_OFFSET_" # !toupper(suffix) :
+ PTX_READ_SREG_R32<regname, intr, [hasPTX<76>, hasSM<80>]>;
+}
// TODO: It would be nice to use PTX_READ_SREG here, but it doesn't
// handle the constant.
diff --git a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.td b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.td
index ca9b641047e6c..913487b64617a 100644
--- a/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.td
@@ -44,7 +44,6 @@ foreach i = 0...31 in {
def ENVREG#i : NVPTXReg<"%envreg"#i>;
}
-
//===----------------------------------------------------------------------===//
// Register classes.
// NOTE: if you add new vector types for a register, you must update
More information about the llvm-commits
mailing list