[Openmp-commits] [openmp] [OpenMP][libomptarget] Fixes possible no-return warning (PR #70808)
Jan Patrick Lehr via Openmp-commits
openmp-commits at lists.llvm.org
Mon Nov 6 01:10:13 PST 2023
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/5] [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/5] 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() {
>From 7316eb643ff1f66eb2aca5dd5670295268b9bc2d Mon Sep 17 00:00:00 2001
From: JP Lehr <JanPatrick.Lehr at amd.com>
Date: Wed, 1 Nov 2023 06:47:37 -0400
Subject: [PATCH 3/5] fixup! fixup! [OpenMP][libomptarget] Fixes possible
no-return warning
---
openmp/libomptarget/DeviceRTL/include/Debug.h | 4 ++++
openmp/libomptarget/DeviceRTL/src/Mapping.cpp | 16 ++++++++--------
2 files changed, 12 insertions(+), 8 deletions(-)
diff --git a/openmp/libomptarget/DeviceRTL/include/Debug.h b/openmp/libomptarget/DeviceRTL/include/Debug.h
index 477ff06d19836c2..5aa0e5639774a05 100644
--- a/openmp/libomptarget/DeviceRTL/include/Debug.h
+++ b/openmp/libomptarget/DeviceRTL/include/Debug.h
@@ -38,6 +38,10 @@ void __assert_fail_internal(const char *expr, const char *msg, const char *file,
PRINT(msg); \
__builtin_unreachable();
+#define MSG_AND_TRAP(msg) \
+ PRINT(msg); \
+ __builtin_trap();
+
///}
#define PRINTF(fmt, ...) (void)printf(fmt, ##__VA_ARGS__);
diff --git a/openmp/libomptarget/DeviceRTL/src/Mapping.cpp b/openmp/libomptarget/DeviceRTL/src/Mapping.cpp
index 822b8dc2dd5e671..fa812656dac353a 100644
--- a/openmp/libomptarget/DeviceRTL/src/Mapping.cpp
+++ b/openmp/libomptarget/DeviceRTL/src/Mapping.cpp
@@ -56,7 +56,7 @@ uint32_t getNumberOfThreadsInBlock(int32_t Dim) {
case 2:
return __builtin_amdgcn_workgroup_size_z();
};
- UNREACHABLE("Dim outside range!");
+ MSG_AND_TRAP("Dim outside range!");
}
LaneMaskTy activemask() { return __builtin_amdgcn_read_exec(); }
@@ -90,7 +90,7 @@ uint32_t getThreadIdInBlock(int32_t Dim) {
case 2:
return __builtin_amdgcn_workitem_id_z();
};
- UNREACHABLE("Dim outside range!");
+ MSG_AND_TRAP("Dim outside range!");
}
uint32_t getNumberOfThreadsInKernel() {
@@ -107,7 +107,7 @@ uint32_t getBlockIdInKernel(int32_t Dim) {
case 2:
return __builtin_amdgcn_workgroup_id_z();
};
- UNREACHABLE("Dim outside range!");
+ MSG_AND_TRAP("Dim outside range!");
}
uint32_t getNumberOfBlocksInKernel(int32_t Dim) {
@@ -119,7 +119,7 @@ uint32_t getNumberOfBlocksInKernel(int32_t Dim) {
case 2:
return __builtin_amdgcn_grid_size_z() / __builtin_amdgcn_workgroup_size_z();
};
- UNREACHABLE("Dim outside range!");
+ MSG_AND_TRAP("Dim outside range!");
}
uint32_t getWarpIdInBlock() {
@@ -149,7 +149,7 @@ uint32_t getNumberOfThreadsInBlock(int32_t Dim) {
case 2:
return __nvvm_read_ptx_sreg_ntid_z();
};
- UNREACHABLE("Dim outside range!");
+ MSG_AND_TRAP("Dim outside range!");
}
const llvm::omp::GV &getGridValue() { return llvm::omp::NVPTXGridValues; }
@@ -181,7 +181,7 @@ uint32_t getThreadIdInBlock(int32_t Dim) {
case 2:
return __nvvm_read_ptx_sreg_tid_z();
};
- UNREACHABLE("Dim outside range!");
+ MSG_AND_TRAP("Dim outside range!");
}
uint32_t getThreadIdInWarp() {
@@ -198,7 +198,7 @@ uint32_t getBlockIdInKernel(int32_t Dim) {
case 2:
return __nvvm_read_ptx_sreg_ctaid_z();
};
- UNREACHABLE("Dim outside range!");
+ MSG_AND_TRAP("Dim outside range!");
}
uint32_t getNumberOfBlocksInKernel(int32_t Dim) {
@@ -210,7 +210,7 @@ uint32_t getNumberOfBlocksInKernel(int32_t Dim) {
case 2:
return __nvvm_read_ptx_sreg_nctaid_z();
};
- UNREACHABLE("Dim outside range!");
+ MSG_AND_TRAP("Dim outside range!");
}
uint32_t getNumberOfThreadsInKernel() {
>From f16a4235ee6b7e0dacb0fde43acfe126b3b6a314 Mon Sep 17 00:00:00 2001
From: JP Lehr <JanPatrick.Lehr at amd.com>
Date: Wed, 1 Nov 2023 06:48:05 -0400
Subject: [PATCH 4/5] fixup! fixup! fixup! [OpenMP][libomptarget] Fixes
possible no-return warning
---
openmp/libomptarget/DeviceRTL/include/Debug.h | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/openmp/libomptarget/DeviceRTL/include/Debug.h b/openmp/libomptarget/DeviceRTL/include/Debug.h
index 5aa0e5639774a05..053984bb75fcbe3 100644
--- a/openmp/libomptarget/DeviceRTL/include/Debug.h
+++ b/openmp/libomptarget/DeviceRTL/include/Debug.h
@@ -38,8 +38,8 @@ void __assert_fail_internal(const char *expr, const char *msg, const char *file,
PRINT(msg); \
__builtin_unreachable();
-#define MSG_AND_TRAP(msg) \
- PRINT(msg); \
+#define MSG_AND_TRAP(msg) \
+ PRINT(msg); \
__builtin_trap();
///}
>From ec264a113fe7c341a2a22d4c651a4984c3cbf69f Mon Sep 17 00:00:00 2001
From: JP Lehr <JanPatrick.Lehr at amd.com>
Date: Mon, 6 Nov 2023 04:09:29 -0500
Subject: [PATCH 5/5] fixup! fixup! fixup! fixup! [OpenMP][libomptarget] Fixes
possible no-return warning
---
openmp/libomptarget/DeviceRTL/include/Debug.h | 5 +----
openmp/libomptarget/DeviceRTL/src/Mapping.cpp | 16 ++++++++--------
2 files changed, 9 insertions(+), 12 deletions(-)
diff --git a/openmp/libomptarget/DeviceRTL/include/Debug.h b/openmp/libomptarget/DeviceRTL/include/Debug.h
index 053984bb75fcbe3..22998f44a5bea5d 100644
--- a/openmp/libomptarget/DeviceRTL/include/Debug.h
+++ b/openmp/libomptarget/DeviceRTL/include/Debug.h
@@ -36,12 +36,9 @@ void __assert_fail_internal(const char *expr, const char *msg, const char *file,
}
#define UNREACHABLE(msg) \
PRINT(msg); \
+ __builtin_trap(); \
__builtin_unreachable();
-#define MSG_AND_TRAP(msg) \
- PRINT(msg); \
- __builtin_trap();
-
///}
#define PRINTF(fmt, ...) (void)printf(fmt, ##__VA_ARGS__);
diff --git a/openmp/libomptarget/DeviceRTL/src/Mapping.cpp b/openmp/libomptarget/DeviceRTL/src/Mapping.cpp
index fa812656dac353a..822b8dc2dd5e671 100644
--- a/openmp/libomptarget/DeviceRTL/src/Mapping.cpp
+++ b/openmp/libomptarget/DeviceRTL/src/Mapping.cpp
@@ -56,7 +56,7 @@ uint32_t getNumberOfThreadsInBlock(int32_t Dim) {
case 2:
return __builtin_amdgcn_workgroup_size_z();
};
- MSG_AND_TRAP("Dim outside range!");
+ UNREACHABLE("Dim outside range!");
}
LaneMaskTy activemask() { return __builtin_amdgcn_read_exec(); }
@@ -90,7 +90,7 @@ uint32_t getThreadIdInBlock(int32_t Dim) {
case 2:
return __builtin_amdgcn_workitem_id_z();
};
- MSG_AND_TRAP("Dim outside range!");
+ UNREACHABLE("Dim outside range!");
}
uint32_t getNumberOfThreadsInKernel() {
@@ -107,7 +107,7 @@ uint32_t getBlockIdInKernel(int32_t Dim) {
case 2:
return __builtin_amdgcn_workgroup_id_z();
};
- MSG_AND_TRAP("Dim outside range!");
+ UNREACHABLE("Dim outside range!");
}
uint32_t getNumberOfBlocksInKernel(int32_t Dim) {
@@ -119,7 +119,7 @@ uint32_t getNumberOfBlocksInKernel(int32_t Dim) {
case 2:
return __builtin_amdgcn_grid_size_z() / __builtin_amdgcn_workgroup_size_z();
};
- MSG_AND_TRAP("Dim outside range!");
+ UNREACHABLE("Dim outside range!");
}
uint32_t getWarpIdInBlock() {
@@ -149,7 +149,7 @@ uint32_t getNumberOfThreadsInBlock(int32_t Dim) {
case 2:
return __nvvm_read_ptx_sreg_ntid_z();
};
- MSG_AND_TRAP("Dim outside range!");
+ UNREACHABLE("Dim outside range!");
}
const llvm::omp::GV &getGridValue() { return llvm::omp::NVPTXGridValues; }
@@ -181,7 +181,7 @@ uint32_t getThreadIdInBlock(int32_t Dim) {
case 2:
return __nvvm_read_ptx_sreg_tid_z();
};
- MSG_AND_TRAP("Dim outside range!");
+ UNREACHABLE("Dim outside range!");
}
uint32_t getThreadIdInWarp() {
@@ -198,7 +198,7 @@ uint32_t getBlockIdInKernel(int32_t Dim) {
case 2:
return __nvvm_read_ptx_sreg_ctaid_z();
};
- MSG_AND_TRAP("Dim outside range!");
+ UNREACHABLE("Dim outside range!");
}
uint32_t getNumberOfBlocksInKernel(int32_t Dim) {
@@ -210,7 +210,7 @@ uint32_t getNumberOfBlocksInKernel(int32_t Dim) {
case 2:
return __nvvm_read_ptx_sreg_nctaid_z();
};
- MSG_AND_TRAP("Dim outside range!");
+ UNREACHABLE("Dim outside range!");
}
uint32_t getNumberOfThreadsInKernel() {
More information about the Openmp-commits
mailing list