[clang] [llvm] [NVPTX] Add builtin for 'exit' handling (PR #79777)
via cfe-commits
cfe-commits at lists.llvm.org
Sun Jan 28 19:29:56 PST 2024
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-llvm-ir
@llvm/pr-subscribers-clang
Author: Joseph Huber (jhuber6)
<details>
<summary>Changes</summary>
Summary:
The PTX ISA has always supported the 'exit' instruction to terminate
individual threads. This patch adds a builtin to handle it. See the PTX
documentation for further details.
https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#control-flow-instructions-exit
---
Full diff: https://github.com/llvm/llvm-project/pull/79777.diff
5 Files Affected:
- (modified) clang/include/clang/Basic/BuiltinsNVPTX.def (+1)
- (modified) clang/test/CodeGen/builtins-nvptx.c (+8)
- (modified) llvm/include/llvm/IR/IntrinsicsNVVM.td (+4)
- (modified) llvm/lib/Target/NVPTX/NVPTXIntrinsics.td (+3)
- (modified) llvm/test/CodeGen/NVPTX/intrinsics.ll (+8)
``````````diff
diff --git a/clang/include/clang/Basic/BuiltinsNVPTX.def b/clang/include/clang/Basic/BuiltinsNVPTX.def
index 0f2e8260143be78..1ae23a32c2adcf0 100644
--- a/clang/include/clang/Basic/BuiltinsNVPTX.def
+++ b/clang/include/clang/Basic/BuiltinsNVPTX.def
@@ -155,6 +155,7 @@ BUILTIN(__nvvm_read_ptx_sreg_pm3, "i", "n")
// MISC
BUILTIN(__nvvm_prmt, "UiUiUiUi", "")
+BUILTIN(__nvvm_exit, "v", "r")
// Min Max
diff --git a/clang/test/CodeGen/builtins-nvptx.c b/clang/test/CodeGen/builtins-nvptx.c
index 353f3ebb608c2b1..0a19e40a01aedb1 100644
--- a/clang/test/CodeGen/builtins-nvptx.c
+++ b/clang/test/CodeGen/builtins-nvptx.c
@@ -165,6 +165,14 @@ __device__ void sync() {
}
+__device__ void exit() {
+
+// CHECK: call void @llvm.nvvm.exit()
+
+ __nvvm_exit();
+
+}
+
// NVVM intrinsics
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 5a5ba2592e1467e..b751ffa27e0203d 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -4801,4 +4801,8 @@ def int_nvvm_setmaxnreg_dec_sync_aligned_u32
[IntrConvergent, IntrNoMem, IntrHasSideEffects, ImmArg<ArgIndex<0>>],
"llvm.nvvm.setmaxnreg.dec.sync.aligned.u32">;
+// Exit
+def int_nvvm_exit : ClangBuiltin<"__nvvm_exit">,
+ Intrinsic<[], [], [IntrConvergent, IntrInaccessibleMemOnly, IntrNoReturn]>;
+
} // let TargetPrefix = "nvvm"
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 33f1e4a43e072af..0db351a33f2a6d5 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -6832,4 +6832,7 @@ multiclass SET_MAXNREG<string Action, Intrinsic Intr> {
defm INT_SET_MAXNREG_INC : SET_MAXNREG<"inc", int_nvvm_setmaxnreg_inc_sync_aligned_u32>;
defm INT_SET_MAXNREG_DEC : SET_MAXNREG<"dec", int_nvvm_setmaxnreg_dec_sync_aligned_u32>;
+
} // isConvergent
+
+def INT_EXIT : NVPTXInst<(outs), (ins), "exit;", [(int_nvvm_exit)]>;
diff --git a/llvm/test/CodeGen/NVPTX/intrinsics.ll b/llvm/test/CodeGen/NVPTX/intrinsics.ll
index c09c7a72fd10181..4b7d5c8f2390769 100644
--- a/llvm/test/CodeGen/NVPTX/intrinsics.ll
+++ b/llvm/test/CodeGen/NVPTX/intrinsics.ll
@@ -133,6 +133,13 @@ define i64 @test_clock64() {
ret i64 %ret
}
+; CHECK-LABEL: test_exit
+define void @test_exit() {
+; CHECK: exit;
+ call void @llvm.nvvm.exit()
+ ret void
+}
+
declare float @llvm.fabs.f32(float)
declare double @llvm.fabs.f64(double)
declare float @llvm.nvvm.sqrt.f(float)
@@ -146,3 +153,4 @@ declare i64 @llvm.ctpop.i64(i64)
declare i32 @llvm.nvvm.read.ptx.sreg.tid.x()
declare i32 @llvm.nvvm.read.ptx.sreg.clock()
declare i64 @llvm.nvvm.read.ptx.sreg.clock64()
+declare void @llvm.nvvm.exit()
``````````
</details>
https://github.com/llvm/llvm-project/pull/79777
More information about the cfe-commits
mailing list