[llvm] [NVPTX] Update setmaxnreg intrinsic lowering (PR #125846)
Valery Chernov via llvm-commits
llvm-commits at lists.llvm.org
Thu Feb 6 01:00:31 PST 2025
https://github.com/vvchernov updated https://github.com/llvm/llvm-project/pull/125846
>From a460eb77e7c428d0d553ba3ac7b8a7dfc90a0e32 Mon Sep 17 00:00:00 2001
From: Valery Chernov <vchernov at nvidia.com>
Date: Tue, 4 Feb 2025 22:28:14 +0400
Subject: [PATCH] Update setmaxnreg intrinsic lowering
---
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 1 +
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 2 +-
llvm/test/CodeGen/NVPTX/setmaxnreg-sm100a.ll | 13 +++++++++++++
3 files changed, 15 insertions(+), 1 deletion(-)
create mode 100644 llvm/test/CodeGen/NVPTX/setmaxnreg-sm100a.ll
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index 74def43d825665b..bf1c9fb7dde1dcb 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -142,6 +142,7 @@ def hasLDU : Predicate<"Subtarget->hasLDU()">;
def hasPTXASUnreachableBug : Predicate<"Subtarget->hasPTXASUnreachableBug()">;
def noPTXASUnreachableBug : Predicate<"!Subtarget->hasPTXASUnreachableBug()">;
def hasOptEnabled : Predicate<"TM.getOptLevel() != CodeGenOptLevel::None">;
+def hasAcceleratedFeatures : Predicate<"Subtarget->hasAAFeatures()">;
def doF32FTZ : Predicate<"useF32FTZ()">;
def doNoF32FTZ : Predicate<"!useF32FTZ()">;
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index a0d00e4aac560a5..06c629c01d9ab09 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -7547,7 +7547,7 @@ multiclass SET_MAXNREG<string Action, Intrinsic Intr> {
def : NVPTXInst<(outs), (ins i32imm:$reg_count),
"setmaxnreg." # Action # ".sync.aligned.u32 $reg_count;",
[(Intr timm:$reg_count)]>,
- Requires<[hasSM90a, hasPTX<80>]>;
+ Requires<[hasAcceleratedFeatures, hasSM<90>, hasPTX<80>]>;
}
defm INT_SET_MAXNREG_INC : SET_MAXNREG<"inc", int_nvvm_setmaxnreg_inc_sync_aligned_u32>;
diff --git a/llvm/test/CodeGen/NVPTX/setmaxnreg-sm100a.ll b/llvm/test/CodeGen/NVPTX/setmaxnreg-sm100a.ll
new file mode 100644
index 000000000000000..fecc286c7a2fac7
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/setmaxnreg-sm100a.ll
@@ -0,0 +1,13 @@
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 | FileCheck --check-prefixes=CHECK %s
+; RUN: %if ptxas-12.8 %{ llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 | %ptxas-verify -arch=sm_100a %}
+
+; CHECK-LABEL: test_set_maxn_reg_sm100a
+define void @test_set_maxn_reg_sm100a() {
+ ; CHECK: setmaxnreg.inc.sync.aligned.u32 96;
+ call void @llvm.nvvm.setmaxnreg.inc.sync.aligned.u32(i32 96)
+
+ ; CHECK: setmaxnreg.dec.sync.aligned.u32 64;
+ call void @llvm.nvvm.setmaxnreg.dec.sync.aligned.u32(i32 64)
+
+ ret void
+}
More information about the llvm-commits
mailing list