[OpenMP] Fix lone target exit data
Without this patch, an `omp target exit data` before the runtime is
initialized produces a runtime error. This patch fixes that by
changing `__tgt_target_data_end_mapper` to call `CheckDeviceAndCtors`
like many other runtime routines.
Discussed at
<https://lists.llvm.org/pipermail/openmp-dev/2021-March/003920.html>.
Reviewed By: grokos
Differential Revision: https://reviews.llvm.org/D97907
GitOrigin-RevId: bfe5452b93a7d859f612bf7e8b46aa626c8c8e83
diff --git a/libomptarget/src/interface.cpp b/libomptarget/src/interface.cpp
index 233ef23..2470600 100644
--- a/libomptarget/src/interface.cpp
+++ b/libomptarget/src/interface.cpp
@@ -272,21 +272,13 @@
return;
}
- PM->RTLsMtx.lock();
- size_t DevicesSize = PM->Devices.size();
- PM->RTLsMtx.unlock();
- if (DevicesSize <= (size_t)device_id) {
- DP("Device ID %" PRId64 " does not have a matching RTL.\n", device_id);
+ if (CheckDeviceAndCtors(device_id) != OFFLOAD_SUCCESS) {
+ DP("Failed to get device %" PRId64 " ready\n", device_id);
HandleTargetOutcome(false, loc);
return;
}
DeviceTy &Device = PM->Devices[device_id];
- if (!Device.IsInit) {
- DP("Uninit device: ignore");
- HandleTargetOutcome(false, loc);
- return;
- }
if (getInfoLevel() & OMP_INFOTYPE_KERNEL_ARGS)
printKernelArguments(loc, device_id, arg_num, arg_sizes, arg_types,
diff --git a/libomptarget/test/offloading/lone_target_exit_data.c b/libomptarget/test/offloading/lone_target_exit_data.c
new file mode 100644
index 0000000..0bbbb9e
--- /dev/null
+++ b/libomptarget/test/offloading/lone_target_exit_data.c
@@ -0,0 +1,18 @@
+// Check that a target exit data directive behaves correctly when the runtime
+// has not yet been initialized.
+
+// RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-nvptx64-nvidia-cuda
+
+#include <stdio.h>
+
+int main() {
+ // CHECK: x = 98
+ int x = 98;
+ #pragma omp target exit data map(from:x)
+ printf("x = %d\n", x);
+ return 0;
+}