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

Durgadoss R via llvm-commits llvm-commits at lists.llvm.org
Mon Jan 8 11:37:43 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)
----------------
durga4github wrote:

This is what I initially thought, but this would leave us with no option to extend the intrinsic in future (may be with different action settings).

If you think, we should add two explicit intrinsics (since this is just two and not a lot more now), I am fine with that too. Let me know what you think,

Also, I will add verifier logic for the %8 and range checks.

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


More information about the llvm-commits mailing list