[OpenMP] Fix support for device as host

[jdenny]

Summary

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.

[grokos]

This is a correct observation. The host device is not a target device, so there is no offloading available.

However, in all __tgt_target_* functions we call CheckDeviceAndCtors() which in turn calls device_is_ready(). The latter checks whether the device ID is greater than or equal to the number of available devices. If it is, then it returns false to CheckDeviceAndCtors() which ultimately returns OFFLOAD_FAIL to the caller __tgt_target_* function. Why doesn't the existing implementation work? What extra checks is this patch adding?

[jdenny]

In every case where this patch inserts a check, there is code immediately following it that calls HandleTargetOutcome(false, loc);, usually because CheckDeviceAndCtors(device_id) != OFFLOAD_SUCCESS. The inserted check prevents that from code from executing. Is that what you're asking?

[grokos]

OK, clear. You are trying to avoid the call to HandleTargetOutcome(false,...) which erroneously says that offloading has failed. LGTM.

[jdenny]

Thanks!

[RaviNarayanaswamy]

If OMP_TARGET_OFFLOAD=mandatory it will not fail but just return

[jdenny]

In D97616#2595403, @RaviNarayanaswamy wrote:

If OMP_TARGET_OFFLOAD=mandatory it will not fail but just return

Are you suggesting that's incorrect behavior when offloading to host?

OpenMP 5.1 sec. 6.17 "OMP_TARGET_OFFLOAD" says:

The mandatory value specifies that program execution is terminated if a device construct or
device memory routine is encountered and the device is not available or is not supported by the
implementation.

I assumed that host is always considered available and supported.

[RaviNarayanaswamy]

This is where we are heading towards in OpenMP spec
The mandatory value specifies that program execution is terminated if a device construct or
device memory routine is encountered and the device is not available or is not supported by the
implementation. or if omp_get_num_devices() would return 0.

[jdenny]

Thanks for pointing that out (OpenMP spec issue 2669). I wasn't aware it was changing. I'll wait until that settles.

[jdenny]

In anticipation of the spec change, I've updated the patch to check the omp_get_num_devices() == 0 case. A few questions:

• Does it match your understanding of what's being proposed? I realize it's possible that will change, so I'm not suggesting we push this yet.
• Is it compatible with OpenMP 5.1 or 5.0? Should we skip the check in those cases?
• The modified functions all have a common prologue. In a separate patch, I think that most of that should be moved to CheckDeviceAndCtors. Agreed?

[jdenny]

One more question: How do I test the new check? For cuda offload, I can set CUDA_VISIBLE_DEVICES to an empty string, but I haven't found a general way.

[grokos]

In D97616#2597237, @jdenny wrote:

• The modified functions all have a common prologue. In a separate patch, I think that most of that should be moved to CheckDeviceAndCtors. Agreed?

Sure!

In D97616#2597240, @jdenny wrote:

One more question: How do I test the new check? For cuda offload, I can set CUDA_VISIBLE_DEVICES to an empty string, but I haven't found a general way.

I don't think there is a general way to do this. An ugly hack (just for testing purposes) would be to move all plugin libraries (i.e. libomptarget.rtl.x86_64.so, libomptarget.rtl.cuda.so etc) to a directory where libomptarget cannot find them, then it will appear as if no supported devices are available in the system.

[jdenny]

• Rebased.
• Added test for mandatory offload with omp_get_num_devices() == 0.
• Added comments pointing to OpenMP spec issue 2669.

[jdenny]

In D97616#2600586, @grokos wrote:

In D97616#2597237, @jdenny wrote:

• The modified functions all have a common prologue. In a separate patch, I think that most of that should be moved to CheckDeviceAndCtors. Agreed?

Sure!

Will do after this one.

In D97616#2597240, @jdenny wrote:

One more question: How do I test the new check? For cuda offload, I can set CUDA_VISIBLE_DEVICES to an empty string, but I haven't found a general way.

I don't think there is a general way to do this. An ugly hack (just for testing purposes) would be to move all plugin libraries (i.e. libomptarget.rtl.x86_64.so, libomptarget.rtl.cuda.so etc) to a directory where libomptarget cannot find them, then it will appear as if no supported devices are available in the system.

Good to know for manual testing. For the test suite, I've added testing just for cuda.

Thanks.

openmp/libomptarget/src/interface.cpp

DP("Entering data begin region for device %" PRId64 " with %d mappings\n",
      device_id, arg_num);
 
   // No devices available?
   if (device_id == OFFLOAD_DEVICE_DEFAULT) {
     device_id = omp_get_default_device();
     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);
     return;
   }
 
   DeviceTy &Device = PM->Devices[device_id];
 
if (IsOffloadDisabled())
     return;
   DP("Entering data end region with %d mappings\n", arg_num);
 
   // No devices available?
   if (device_id == OFFLOAD_DEVICE_DEFAULT) {
     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();
   if (DevicesSize <= (size_t)device_id) {
     DP("Device ID  %" PRId64 " does not have a matching RTL.\n", device_id);
     HandleTargetOutcome(false, loc);
     return;
   }
if (IsOffloadDisabled())
     return;
   DP("Entering data update with %d mappings\n", arg_num);
 
   // No devices available?
   if (device_id == OFFLOAD_DEVICE_DEFAULT) {
     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);
     return;
   }
 
   if (getInfoLevel() & OMP_INFOTYPE_KERNEL_ARGS)
     printKernelArguments(loc, device_id, arg_num, arg_sizes, arg_types,
EXTERN int __tgt_target_mapper(ident_t *loc, int64_t device_id, void *host_ptr,
   DP("Entering target region with entry point " DPxMOD " and device Id %" PRId64
      "\n",
      DPxPTR(host_ptr), device_id);
 
   if (device_id == OFFLOAD_DEVICE_DEFAULT) {
     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);
     return OFFLOAD_FAIL;
   }
 
   if (getInfoLevel() & OMP_INFOTYPE_KERNEL_ARGS)
     printKernelArguments(loc, device_id, arg_num, arg_sizes, arg_types,
EXTERN int __tgt_target_teams_mapper(ident_t *loc, int64_t device_id,
   DP("Entering target region with entry point " DPxMOD " and device Id %" PRId64
      "\n",
      DPxPTR(host_ptr), device_id);
 
   if (device_id == OFFLOAD_DEVICE_DEFAULT) {
     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);
     return OFFLOAD_FAIL;
   }
 
   if (getInfoLevel() & OMP_INFOTYPE_KERNEL_ARGS)
     printKernelArguments(loc, device_id, arg_num, arg_sizes, arg_types,
EXTERN void __kmpc_push_target_tripcount(ident_t *loc, int64_t device_id,
   TIMESCOPE_WITH_IDENT(loc);
   if (IsOffloadDisabled())
     return;
 
   if (device_id == OFFLOAD_DEVICE_DEFAULT) {
     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);
     return;
   }
 
   DP("__kmpc_push_target_tripcount(%" PRId64 ", %" PRIu64 ")\n", device_id,
      loop_tripcount);
   PM->TblMapMtx.lock();
   PM->Devices[device_id].LoopTripCnt.emplace(__kmpc_global_thread_num(NULL),
                                              loop_tripcount);
   PM->TblMapMtx.unlock();
 }

openmp/libomptarget/test/offloading/host_as_target.c

+// 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;
+}

openmp/libomptarget/test/offloading/mandatory_but_no_devices.c

+// 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;
+}