[Openmp-commits] [openmp] 10c18c6 - [OpenMP] Fix support for device as host
Joel E. Denny via Openmp-commits
openmp-commits at lists.llvm.org
Thu Mar 4 09:12:57 PST 2021
Author: Joel E. Denny
Date: 2021-03-04T12:03:42-05:00
New Revision: 10c18c69f2a8372c6a85e4654e00d11b479119aa
URL: https://github.com/llvm/llvm-project/commit/10c18c69f2a8372c6a85e4654e00d11b479119aa
DIFF: https://github.com/llvm/llvm-project/commit/10c18c69f2a8372c6a85e4654e00d11b479119aa.diff
LOG: [OpenMP] Fix support for device as host
Without this patch, when the offload device is set to
`omp_get_initial_device()`, the runtime fails with an error diagnostic
when entering target regions or target data regions.
However, OpenMP 5.1, sec. 2.14.5 "target Construct", "Restrictions",
p. 203, L3-5 states:
> The device clause expression must evaluate to a non-negative integer
> value that is less than or equal to the value of
> omp_get_num_devices().
Sec. 3.7.7 "omp_get_initial_device", p. 412, L2-3 states:
> The value of the device number is the value returned by the
> omp_get_num_devices routine.
Similarly, OpenMP 5.0, sec. 2.12.5 "target Construct", "Restrictions",
p. 174 L30-32 states:
> The device clause expression must evaluate to a non-negative integer
> value less than the value of omp_get_num_devices() or to the value
> of omp_get_initial_device().
This patch fixes this behavior by changing the runtime to behave as if
offloading is disabled whenever it finds the offload device (either
from a `device` clause or the default device) is set to the host
device. In the case of mandatory offloading when
`omp_get_num_devices() == 0`, it incorporates the behavior proposed
for OpenMP 5.2 in OpenMP spec github issue 2669.
Reviewed By: grokos, RaviNarayanaswamy
Differential Revision: https://reviews.llvm.org/D97616
Added:
openmp/libomptarget/test/offloading/host_as_target.c
openmp/libomptarget/test/offloading/mandatory_but_no_devices.c
Modified:
openmp/libomptarget/src/interface.cpp
Removed:
################################################################################
diff --git a/openmp/libomptarget/src/interface.cpp b/openmp/libomptarget/src/interface.cpp
index 095c4d31d1a8..233ef23083a3 100644
--- a/openmp/libomptarget/src/interface.cpp
+++ b/openmp/libomptarget/src/interface.cpp
@@ -164,6 +164,19 @@ EXTERN void __tgt_target_data_begin_mapper(ident_t *loc, int64_t device_id,
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 @@ EXTERN void __tgt_target_data_end_mapper(ident_t *loc, int64_t device_id,
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 @@ EXTERN void __tgt_target_data_update_mapper(ident_t *loc, int64_t device_id,
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 @@ EXTERN int __tgt_target_mapper(ident_t *loc, int64_t device_id, void *host_ptr,
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 @@ EXTERN int __tgt_target_teams_mapper(ident_t *loc, int64_t device_id,
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 @@ EXTERN void __kmpc_push_target_tripcount(ident_t *loc, int64_t device_id,
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
index 000000000000..56772a312d5a
--- /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 <stdio.h>
+#include <omp.h>
+
+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
index 000000000000..911d98b93b50
--- /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 <omp.h>
+#include <stdio.h>
+
+// 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;
+}
More information about the Openmp-commits
mailing list