[Openmp-commits] [openmp] [OpenMP][libomptarget] Fixes possible no-return warning (PR #70808)

[Jan Patrick Lehr via Openmp-commits]

https://github.com/jplehr updated https://github.com/llvm/llvm-project/pull/70808

>From c80f81c0c455d4b91d11d5265ab0af0561e5435b Mon Sep 17 00:00:00 2001
From: JP Lehr <JanPatrick.Lehr at amd.com>
Date: Tue, 31 Oct 2023 09:55:44 -0400
Subject: [PATCH 1/2] [OpenMP][libomptarget] Fixes possible no-return warning

The UNREACHABLE macro resolves to message + trap, which may still give
you no-return warnings. This silences these warnings.
---
 openmp/libomptarget/DeviceRTL/src/Mapping.cpp | 8 ++++++++
 1 file changed, 8 insertions(+)

diff --git a/openmp/libomptarget/DeviceRTL/src/Mapping.cpp b/openmp/libomptarget/DeviceRTL/src/Mapping.cpp
index 822b8dc2dd5e671..c64856c237f04e0 100644
--- a/openmp/libomptarget/DeviceRTL/src/Mapping.cpp
+++ b/openmp/libomptarget/DeviceRTL/src/Mapping.cpp
@@ -57,6 +57,7 @@ uint32_t getNumberOfThreadsInBlock(int32_t Dim) {
     return __builtin_amdgcn_workgroup_size_z();
   };
   UNREACHABLE("Dim outside range!");
+  return 0;
 }
 
 LaneMaskTy activemask() { return __builtin_amdgcn_read_exec(); }
@@ -91,6 +92,7 @@ uint32_t getThreadIdInBlock(int32_t Dim) {
     return __builtin_amdgcn_workitem_id_z();
   };
   UNREACHABLE("Dim outside range!");
+  return 0;
 }
 
 uint32_t getNumberOfThreadsInKernel() {
@@ -108,6 +110,7 @@ uint32_t getBlockIdInKernel(int32_t Dim) {
     return __builtin_amdgcn_workgroup_id_z();
   };
   UNREACHABLE("Dim outside range!");
+  return 0;
 }
 
 uint32_t getNumberOfBlocksInKernel(int32_t Dim) {
@@ -120,6 +123,7 @@ uint32_t getNumberOfBlocksInKernel(int32_t Dim) {
     return __builtin_amdgcn_grid_size_z() / __builtin_amdgcn_workgroup_size_z();
   };
   UNREACHABLE("Dim outside range!");
+  return 0;
 }
 
 uint32_t getWarpIdInBlock() {
@@ -150,6 +154,7 @@ uint32_t getNumberOfThreadsInBlock(int32_t Dim) {
     return __nvvm_read_ptx_sreg_ntid_z();
   };
   UNREACHABLE("Dim outside range!");
+  return 0;
 }
 
 const llvm::omp::GV &getGridValue() { return llvm::omp::NVPTXGridValues; }
@@ -182,6 +187,7 @@ uint32_t getThreadIdInBlock(int32_t Dim) {
     return __nvvm_read_ptx_sreg_tid_z();
   };
   UNREACHABLE("Dim outside range!");
+  return 0;
 }
 
 uint32_t getThreadIdInWarp() {
@@ -199,6 +205,7 @@ uint32_t getBlockIdInKernel(int32_t Dim) {
     return __nvvm_read_ptx_sreg_ctaid_z();
   };
   UNREACHABLE("Dim outside range!");
+  return 0;
 }
 
 uint32_t getNumberOfBlocksInKernel(int32_t Dim) {
@@ -211,6 +218,7 @@ uint32_t getNumberOfBlocksInKernel(int32_t Dim) {
     return __nvvm_read_ptx_sreg_nctaid_z();
   };
   UNREACHABLE("Dim outside range!");
+  return 0;
 }
 
 uint32_t getNumberOfThreadsInKernel() {

>From be4ead7afa6805f7e2e018f79929f0981f5f5ed3 Mon Sep 17 00:00:00 2001
From: JP Lehr <JanPatrick.Lehr at amd.com>
Date: Tue, 31 Oct 2023 11:14:16 -0400
Subject: [PATCH 2/2] fixup! [OpenMP][libomptarget] Fixes possible no-return
 warning

---
 openmp/libomptarget/DeviceRTL/include/Debug.h | 2 +-
 openmp/libomptarget/DeviceRTL/src/Mapping.cpp | 8 --------
 2 files changed, 1 insertion(+), 9 deletions(-)

diff --git a/openmp/libomptarget/DeviceRTL/include/Debug.h b/openmp/libomptarget/DeviceRTL/include/Debug.h
index bd4d40e8f24fd5b..477ff06d19836c2 100644
--- a/openmp/libomptarget/DeviceRTL/include/Debug.h
+++ b/openmp/libomptarget/DeviceRTL/include/Debug.h
@@ -36,7 +36,7 @@ void __assert_fail_internal(const char *expr, const char *msg, const char *file,
   }
 #define UNREACHABLE(msg)                                                       \
   PRINT(msg);                                                                  \
-  __builtin_trap();
+  __builtin_unreachable();
 
 ///}
 
diff --git a/openmp/libomptarget/DeviceRTL/src/Mapping.cpp b/openmp/libomptarget/DeviceRTL/src/Mapping.cpp
index c64856c237f04e0..822b8dc2dd5e671 100644
--- a/openmp/libomptarget/DeviceRTL/src/Mapping.cpp
+++ b/openmp/libomptarget/DeviceRTL/src/Mapping.cpp
@@ -57,7 +57,6 @@ uint32_t getNumberOfThreadsInBlock(int32_t Dim) {
     return __builtin_amdgcn_workgroup_size_z();
   };
   UNREACHABLE("Dim outside range!");
-  return 0;
 }
 
 LaneMaskTy activemask() { return __builtin_amdgcn_read_exec(); }
@@ -92,7 +91,6 @@ uint32_t getThreadIdInBlock(int32_t Dim) {
     return __builtin_amdgcn_workitem_id_z();
   };
   UNREACHABLE("Dim outside range!");
-  return 0;
 }
 
 uint32_t getNumberOfThreadsInKernel() {
@@ -110,7 +108,6 @@ uint32_t getBlockIdInKernel(int32_t Dim) {
     return __builtin_amdgcn_workgroup_id_z();
   };
   UNREACHABLE("Dim outside range!");
-  return 0;
 }
 
 uint32_t getNumberOfBlocksInKernel(int32_t Dim) {
@@ -123,7 +120,6 @@ uint32_t getNumberOfBlocksInKernel(int32_t Dim) {
     return __builtin_amdgcn_grid_size_z() / __builtin_amdgcn_workgroup_size_z();
   };
   UNREACHABLE("Dim outside range!");
-  return 0;
 }
 
 uint32_t getWarpIdInBlock() {
@@ -154,7 +150,6 @@ uint32_t getNumberOfThreadsInBlock(int32_t Dim) {
     return __nvvm_read_ptx_sreg_ntid_z();
   };
   UNREACHABLE("Dim outside range!");
-  return 0;
 }
 
 const llvm::omp::GV &getGridValue() { return llvm::omp::NVPTXGridValues; }
@@ -187,7 +182,6 @@ uint32_t getThreadIdInBlock(int32_t Dim) {
     return __nvvm_read_ptx_sreg_tid_z();
   };
   UNREACHABLE("Dim outside range!");
-  return 0;
 }
 
 uint32_t getThreadIdInWarp() {
@@ -205,7 +199,6 @@ uint32_t getBlockIdInKernel(int32_t Dim) {
     return __nvvm_read_ptx_sreg_ctaid_z();
   };
   UNREACHABLE("Dim outside range!");
-  return 0;
 }
 
 uint32_t getNumberOfBlocksInKernel(int32_t Dim) {
@@ -218,7 +211,6 @@ uint32_t getNumberOfBlocksInKernel(int32_t Dim) {
     return __nvvm_read_ptx_sreg_nctaid_z();
   };
   UNREACHABLE("Dim outside range!");
-  return 0;
 }
 
 uint32_t getNumberOfThreadsInKernel() {