diff --git a/openmp/libomptarget/src/interface.cpp b/openmp/libomptarget/src/interface.cpp --- a/openmp/libomptarget/src/interface.cpp +++ b/openmp/libomptarget/src/interface.cpp @@ -164,6 +164,19 @@ DP("Use default device id %" PRId64 "\n", device_id); } + // Proposed behavior for OpenMP 5.2 in OpenMP spec github issue 2669. + if (omp_get_num_devices() == 0) { + DP("omp_get_num_devices() == 0 but offload is manadatory\n"); + HandleTargetOutcome(false, loc); + return; + } + + if (device_id == omp_get_initial_device()) { + DP("Device is host (%" PRId64 "), returning as if offload is disabled\n", + device_id); + return; + } + if (CheckDeviceAndCtors(device_id) != OFFLOAD_SUCCESS) { DP("Failed to get device %" PRId64 " ready\n", device_id); HandleTargetOutcome(false, loc); @@ -246,6 +259,19 @@ device_id = omp_get_default_device(); } + // Proposed behavior for OpenMP 5.2 in OpenMP spec github issue 2669. + if (omp_get_num_devices() == 0) { + DP("omp_get_num_devices() == 0 but offload is manadatory\n"); + HandleTargetOutcome(false, loc); + return; + } + + if (device_id == omp_get_initial_device()) { + DP("Device is host (%" PRId64 "), returning as if offload is disabled\n", + device_id); + return; + } + PM->RTLsMtx.lock(); size_t DevicesSize = PM->Devices.size(); PM->RTLsMtx.unlock(); @@ -331,6 +357,19 @@ device_id = omp_get_default_device(); } + // Proposed behavior for OpenMP 5.2 in OpenMP spec github issue 2669. + if (omp_get_num_devices() == 0) { + DP("omp_get_num_devices() == 0 but offload is manadatory\n"); + HandleTargetOutcome(false, loc); + return; + } + + if (device_id == omp_get_initial_device()) { + DP("Device is host (%" PRId64 "), returning as if offload is disabled\n", + device_id); + return; + } + if (CheckDeviceAndCtors(device_id) != OFFLOAD_SUCCESS) { DP("Failed to get device %" PRId64 " ready\n", device_id); HandleTargetOutcome(false, loc); @@ -399,6 +438,20 @@ device_id = omp_get_default_device(); } + // Proposed behavior for OpenMP 5.2 in OpenMP spec github issue 2669. + if (omp_get_num_devices() == 0) { + DP("omp_get_num_devices() == 0 but offload is manadatory\n"); + HandleTargetOutcome(false, loc); + return OFFLOAD_FAIL; + } + + if (device_id == omp_get_initial_device()) { + DP("Device is host (%" PRId64 "), returning OFFLOAD_FAIL as if offload is " + "disabled\n", + device_id); + return OFFLOAD_FAIL; + } + if (CheckDeviceAndCtors(device_id) != OFFLOAD_SUCCESS) { REPORT("Failed to get device %" PRId64 " ready\n", device_id); HandleTargetOutcome(false, loc); @@ -484,6 +537,20 @@ device_id = omp_get_default_device(); } + // Proposed behavior for OpenMP 5.2 in OpenMP spec github issue 2669. + if (omp_get_num_devices() == 0) { + DP("omp_get_num_devices() == 0 but offload is manadatory\n"); + HandleTargetOutcome(false, loc); + return OFFLOAD_FAIL; + } + + if (device_id == omp_get_initial_device()) { + DP("Device is host (%" PRId64 "), returning OFFLOAD_FAIL as if offload is " + "disabled\n", + device_id); + return OFFLOAD_FAIL; + } + if (CheckDeviceAndCtors(device_id) != OFFLOAD_SUCCESS) { REPORT("Failed to get device %" PRId64 " ready\n", device_id); HandleTargetOutcome(false, loc); @@ -563,6 +630,19 @@ device_id = omp_get_default_device(); } + // Proposed behavior for OpenMP 5.2 in OpenMP spec github issue 2669. + if (omp_get_num_devices() == 0) { + DP("omp_get_num_devices() == 0 but offload is manadatory\n"); + HandleTargetOutcome(false, loc); + return; + } + + if (device_id == omp_get_initial_device()) { + DP("Device is host (%" PRId64 "), returning as if offload is disabled\n", + device_id); + return; + } + if (CheckDeviceAndCtors(device_id) != OFFLOAD_SUCCESS) { DP("Failed to get device %" PRId64 " ready\n", device_id); HandleTargetOutcome(false, loc); diff --git a/openmp/libomptarget/test/offloading/host_as_target.c b/openmp/libomptarget/test/offloading/host_as_target.c new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/offloading/host_as_target.c @@ -0,0 +1,153 @@ +// Check that specifying device as omp_get_initial_device(): +// - Doesn't cause the runtime to fail. +// - Offloads code to the host. +// - Doesn't transfer data. In this case, just check that neither host data nor +// default device data are affected by the specified transfers. +// - Works whether it's specified directly or as the default device. + +// 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 +#include + +static void check(char *X, int Dev) { + printf(" host X = %c\n", *X); + #pragma omp target device(Dev) + printf("device X = %c\n", *X); +} + +#define CHECK_DATA() check(&X, DevDefault) + +int main(void) { + int DevDefault = omp_get_default_device(); + int DevInit = omp_get_initial_device(); + + //-------------------------------------------------- + // Initialize data on the host and default device. + //-------------------------------------------------- + + // CHECK: host X = h + // CHECK-NEXT: device X = d + char X = 'd'; + #pragma omp target enter data map(to:X) + X = 'h'; + CHECK_DATA(); + + //-------------------------------------------------- + // Check behavior when specifying host directly. + //-------------------------------------------------- + + // CHECK-NEXT: omp_is_initial_device() = 1 + // CHECK-NEXT: host X = h + // CHECK-NEXT: device X = d + #pragma omp target device(DevInit) map(always,tofrom:X) + printf("omp_is_initial_device() = %d\n", omp_is_initial_device()); + CHECK_DATA(); + + // CHECK-NEXT: omp_is_initial_device() = 1 + // CHECK-NEXT: host X = h + // CHECK-NEXT: device X = d + #pragma omp target teams device(DevInit) num_teams(1) map(always,tofrom:X) + printf("omp_is_initial_device() = %d\n", omp_is_initial_device()); + CHECK_DATA(); + + // Check that __kmpc_push_target_tripcount doesn't fail. I'm not sure how to + // check that it actually pushes to the initial device. + #pragma omp target teams device(DevInit) num_teams(1) + #pragma omp distribute + for (int i = 0; i < 2; ++i) + ; + + // CHECK-NEXT: host X = h + // CHECK-NEXT: device X = d + #pragma omp target data device(DevInit) map(always,tofrom:X) + ; + CHECK_DATA(); + + // CHECK-NEXT: host X = h + // CHECK-NEXT: device X = d + #pragma omp target enter data device(DevInit) map(always,to:X) + ; + CHECK_DATA(); + + // CHECK-NEXT: host X = h + // CHECK-NEXT: device X = d + #pragma omp target exit data device(DevInit) map(always,from:X) + ; + CHECK_DATA(); + + // CHECK-NEXT: host X = h + // CHECK-NEXT: device X = d + #pragma omp target update device(DevInit) to(X) + ; + CHECK_DATA(); + + // CHECK-NEXT: host X = h + // CHECK-NEXT: device X = d + #pragma omp target update device(DevInit) from(X) + ; + CHECK_DATA(); + + //-------------------------------------------------- + // Check behavior when device defaults to host. + //-------------------------------------------------- + + omp_set_default_device(DevInit); + + // CHECK-NEXT: omp_is_initial_device() = 1 + // CHECK-NEXT: host X = h + // CHECK-NEXT: device X = d + #pragma omp target map(always,tofrom:X) + printf("omp_is_initial_device() = %d\n", omp_is_initial_device()); + CHECK_DATA(); + + // CHECK-NEXT: omp_is_initial_device() = 1 + // CHECK-NEXT: host X = h + // CHECK-NEXT: device X = d + #pragma omp target teams num_teams(1) map(always,tofrom:X) + printf("omp_is_initial_device() = %d\n", omp_is_initial_device()); + CHECK_DATA(); + + // Check that __kmpc_push_target_tripcount doesn't fail. I'm not sure how to + // check that it actually pushes to the initial device. + #pragma omp target teams num_teams(1) + #pragma omp distribute + for (int i = 0; i < 2; ++i) + ; + + // CHECK-NEXT: host X = h + // CHECK-NEXT: device X = d + #pragma omp target data map(always,tofrom:X) + ; + CHECK_DATA(); + + // CHECK-NEXT: host X = h + // CHECK-NEXT: device X = d + #pragma omp target enter data map(always,to:X) + ; + CHECK_DATA(); + + // CHECK-NEXT: host X = h + // CHECK-NEXT: device X = d + #pragma omp target exit data map(always,from:X) + ; + CHECK_DATA(); + + // CHECK-NEXT: host X = h + // CHECK-NEXT: device X = d + #pragma omp target update to(X) + ; + CHECK_DATA(); + + // CHECK-NEXT: host X = h + // CHECK-NEXT: device X = d + #pragma omp target update from(X) + ; + CHECK_DATA(); + + return 0; +} diff --git a/openmp/libomptarget/test/offloading/mandatory_but_no_devices.c b/openmp/libomptarget/test/offloading/mandatory_but_no_devices.c new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/offloading/mandatory_but_no_devices.c @@ -0,0 +1,54 @@ +// Check that mandatory offloading causes various offloading directives to fail +// when omp_get_num_devices() == 0 even if the requested device is the initial +// device. This behavior is proposed for OpenMP 5.2 in OpenMP spec github +// issue 2669. + +// RUN: %libomptarget-compile-nvptx64-nvidia-cuda -DDIR=target +// RUN: env OMP_TARGET_OFFLOAD=mandatory CUDA_VISIBLE_DEVICES= \ +// RUN: %libomptarget-run-fail-nvptx64-nvidia-cuda 2>&1 | \ +// RUN: %fcheck-nvptx64-nvidia-cuda + +// RUN: %libomptarget-compile-nvptx64-nvidia-cuda -DDIR='target teams' +// RUN: env OMP_TARGET_OFFLOAD=mandatory CUDA_VISIBLE_DEVICES= \ +// RUN: %libomptarget-run-fail-nvptx64-nvidia-cuda 2>&1 | \ +// RUN: %fcheck-nvptx64-nvidia-cuda + +// RUN: %libomptarget-compile-nvptx64-nvidia-cuda -DDIR='target data map(X)' +// RUN: env OMP_TARGET_OFFLOAD=mandatory CUDA_VISIBLE_DEVICES= \ +// RUN: %libomptarget-run-fail-nvptx64-nvidia-cuda 2>&1 | \ +// RUN: %fcheck-nvptx64-nvidia-cuda + +// RUN: %libomptarget-compile-nvptx64-nvidia-cuda \ +// RUN: -DDIR='target enter data map(to:X)' +// RUN: env OMP_TARGET_OFFLOAD=mandatory CUDA_VISIBLE_DEVICES= \ +// RUN: %libomptarget-run-fail-nvptx64-nvidia-cuda 2>&1 | \ +// RUN: %fcheck-nvptx64-nvidia-cuda + +// RUN: %libomptarget-compile-nvptx64-nvidia-cuda \ +// RUN: -DDIR='target exit data map(from:X)' +// RUN: env OMP_TARGET_OFFLOAD=mandatory CUDA_VISIBLE_DEVICES= \ +// RUN: %libomptarget-run-fail-nvptx64-nvidia-cuda 2>&1 | \ +// RUN: %fcheck-nvptx64-nvidia-cuda + +// RUN: %libomptarget-compile-nvptx64-nvidia-cuda \ +// RUN: -DDIR='target update to(X)' +// RUN: env OMP_TARGET_OFFLOAD=mandatory CUDA_VISIBLE_DEVICES= \ +// RUN: %libomptarget-run-fail-nvptx64-nvidia-cuda 2>&1 | \ +// RUN: %fcheck-nvptx64-nvidia-cuda + +// RUN: %libomptarget-compile-nvptx64-nvidia-cuda \ +// RUN: -DDIR='target update from(X)' +// RUN: env OMP_TARGET_OFFLOAD=mandatory CUDA_VISIBLE_DEVICES= \ +// RUN: %libomptarget-run-fail-nvptx64-nvidia-cuda 2>&1 | \ +// RUN: %fcheck-nvptx64-nvidia-cuda + +#include +#include + +// CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory +int main(void) { + int X; + #pragma omp DIR device(omp_get_initial_device()) + ; + return 0; +}