Index: openmp/libomptarget/src/interface.cpp =================================================================== --- openmp/libomptarget/src/interface.cpp +++ openmp/libomptarget/src/interface.cpp @@ -164,6 +164,12 @@ DP("Use default device id %" PRId64 "\n", device_id); } + 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 +252,12 @@ device_id = omp_get_default_device(); } + 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 +343,12 @@ device_id = omp_get_default_device(); } + 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 +417,13 @@ device_id = omp_get_default_device(); } + 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 +509,13 @@ device_id = omp_get_default_device(); } + 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 +595,12 @@ device_id = omp_get_default_device(); } + 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); Index: openmp/libomptarget/test/offloading/host_as_target.c =================================================================== --- /dev/null +++ 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; +}