From a460eb77e7c428d0d553ba3ac7b8a7dfc90a0e32 Mon Sep 17 00:00:00 2001 From: Valery Chernov 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 74def43d82566..bf1c9fb7dde1d 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 a0d00e4aac560..06c629c01d9ab 100644 --- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td +++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td @@ -7547,7 +7547,7 @@ multiclass SET_MAXNREG { 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 0000000000000..fecc286c7a2fa --- /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 +}