[llvm] [LLVM][NVPTX]: Add intrinsic for setmaxnreg (PR #77289)

Artem Belevich via llvm-commits llvm-commits at lists.llvm.org
Mon Jan 8 11:52:26 PST 2024


================
@@ -0,0 +1,15 @@
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_90a -mattr=+ptx80| FileCheck --check-prefixes=CHECK %s
+; RUN: %if ptxas-12.0 %{ llc < %s -march=nvptx64 -mcpu=sm_90a -mattr=+ptx80| %ptxas-verify -arch=sm_90a %}
+
+declare void @llvm.nvvm.setmaxnreg.sync.aligned.u32(i32 %flags, i32 %reg_count)
+
+; CHECK-LABEL: test_set_maxn_reg
+define void @test_set_maxn_reg() {
+  ; CHECK: setmaxnreg.inc.sync.aligned.u32 96;
+  call void @llvm.nvvm.setmaxnreg.sync.aligned.u32(i32 0, i32 96)
+
+  ; CHECK: setmaxnreg.dec.sync.aligned.u32 64;
+  call void @llvm.nvvm.setmaxnreg.sync.aligned.u32(i32 1, i32 64)
----------------
Artem-B wrote:

I think separate intrinsics make more sense here. The `inc/dec` flags translate into a PTX instruction *name*, so matching intrinsic name would be a better fit, and has a bonus of simplifying the implementation.


https://github.com/llvm/llvm-project/pull/77289


More information about the llvm-commits mailing list