Index: openmp/libomptarget/src/interface.cpp =================================================================== --- openmp/libomptarget/src/interface.cpp +++ openmp/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, Index: openmp/libomptarget/test/offloading/lone_target_exit_data.c =================================================================== --- /dev/null +++ openmp/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 + +int main() { + // CHECK: x = 98 + int x = 98; + #pragma omp target exit data map(from:x) + printf("x = %d\n", x); + return 0; +}