[OpenMP][Offloading] Fix the issue that omp_get_num_devices returns wrong number of devices

[tianshilei1992]

Summary

This patch is to fix issue in the following simple case:

#include <omp.h>
#include <stdio.h>

int main(int argc, char *argv[]) {
  int num = omp_get_num_devices();
  printf("%d\n", num);

  return 0;
}

Currently it returns 0 even devices exist. Since this file doesn't contain any
target region, the host entry is empty so further actions like initialization
will not be proceeded, leading to wrong device number returned by runtime
function call.

[ABataev]

"Revert" is wrong here since one of the tests will definitely fail after this. Prepare a patch with the proper fix.

[tianshilei1992]

In D72576#1816111, @ABataev wrote:

"Revert" is wrong here since one of the tests will definitely fail after this. Prepare a patch with the proper fix.

Yep. I got your point. Didn't notice what you tried to fix before. Will update correspondingly.

[ABataev]

In D72576#1816159, @tianshilei1992 wrote:

In D72576#1816111, @ABataev wrote:

"Revert" is wrong here since one of the tests will definitely fail after this. Prepare a patch with the proper fix.

Yep. I got your point. Didn't notice what you tried to fix before. Will update correspondingly.

Ok, thanks.

[tianshilei1992]

This comment was removed by tianshilei1992.

[protze.joachim]

From OpenMP spec perspective I see no issue in returning 0 for the provided code even if there would be devices available, as long as this is consistent with the remaining execution of the application.

As I understand the modification, the function would not return 0 if afterwards a target region will be offloaded to a device?

[tianshilei1992]

In D72576#1818720, @protze.joachim wrote:

From OpenMP spec perspective I see no issue in returning 0 for the provided code even if there would be devices available, as long as this is consistent with the remaining execution of the application.

As I understand the modification, the function would not return 0 if afterwards a target region will be offloaded to a device?

Actually, no because of empty PendingCtorsDtors and empty target entries. And what about the following code, what value do you expect to be returned?

#include <omp.h>
#include <stdint.h>
#include <stdio.h>
#include <stdlib.h>

int main() {
  const size_t N = 1UL << 30;
  const int ndev = omp_get_num_devices();
  const int host = omp_get_initial_device();

  printf("Found %d devices, host is %d\n", ndev, host);

  uint8_t *v = (uint8_t *)malloc(N);
  uint8_t *u[ndev];

  for (int i = 0; i < ndev; i++) {
    u[i] = (uint8_t *)omp_target_alloc(N, i);
  }

  for (size_t i = 0; i < N; i++) {
    v[i] = 0;
  }

  for (int i = 0; i < ndev; i++) {
    uint8_t *ui = u[i];

#pragma omp target data device(i) map(tofrom : v [0:N])
    {
#pragma omp target data device(i) use_device_ptr(v)
      { printf("Device %d: v = %p, u = %p\n", i, v, ui); }
    }
  }

  for (int i = 0; i < ndev; i++) {
    omp_target_free(u[i], i);
  }
  free(v);
}

[ABataev]

Could you provide full diff of your modificaions?

[tianshilei1992]

In D72576#1820230, @ABataev wrote:

Could you provide full diff of your modificaions?

This is exactly the full diff.

[ABataev]

In D72576#1820258, @tianshilei1992 wrote:

In D72576#1820230, @ABataev wrote:

Could you provide full diff of your modificaions?

This is exactly the full diff.

Nope, it does not have context (see this https://www.llvm.org/docs/Phabricator.html). Also, add a test.

[tianshilei1992]

Upload full diff information

[ABataev]

openmp/libomptarget/src/omptarget.cpp

97–101

Maybe just check if TransTable->HostTable.EntriesBegin == TransTable->HostTable.EntriesEnd? Along with the check if entries have processed already. Plus, I don't think it is correct to exit from the loop here.

[tianshilei1992]

openmp/libomptarget/src/omptarget.cpp

97–101

Maybe we can make it in this way?

if (TransTable->TargetsTable[device_id] != 0) {
  // Library entries have already been processed
  continue;
}

if (TransTable->HostTable.EntriesBegin == TransTable->HostTable.EntriesEnd) {
  // No entries so no need to process
  continue;
}

[ABataev]

openmp/libomptarget/src/omptarget.cpp

97–101

Sure, this one is also fine, if it works correctly. Just, maybe, make this check the first.

[ABataev]

Need a test

[tianshilei1992]

Reused existing test case with just a little tiny modification to test this patch

[ABataev]

openmp/libomptarget/test/api/omp_get_num_devices.c

15

Hmm, better to add a new test, I think, rather than modify this one. We need to increase the coverage, not just change it in some way.

[protze.joachim]

In D72576#1820159, @tianshilei1992 wrote:

In D72576#1818720, @protze.joachim wrote:

From OpenMP spec perspective I see no issue in returning 0 for the provided code even if there would be devices available, as long as this is consistent with the remaining execution of the application.

As I understand the modification, the function would not return 0 if afterwards a target region will be offloaded to a device?

Actually, no because of empty PendingCtorsDtors and empty target entries.

Ok.

And what about the following code, what value do you expect to be returned?

I would expect as many devices as available for this execution.

[tianshilei1992]

I would expect as many devices as available for this execution.

Yeah, it should be. However, w/o this patch, it returns 0.

[tianshilei1992]

openmp/libomptarget/test/api/omp_get_num_devices.c

15

Would you recommend to add a real new file or add a new function to cover the case?

[ABataev]

openmp/libomptarget/test/api/omp_get_num_devices.c

15

It must be a new file, otherwise there won't be an empty image.

[tianshilei1992]

openmp/libomptarget/test/api/omp_get_num_devices.c

15

Oh! That's right. Shame on me. I should have thought of that. :-|

[tianshilei1992]

Add a new test case instead of modifying existing one to improve test coverage

[ABataev]

openmp/libomptarget/test/api/omp_get_num_devices_with_empty_target.c

25 ↗

(On Diff #238385)

No // CHECK: directives.

[tianshilei1992]

openmp/libomptarget/test/api/omp_get_num_devices_with_empty_target.c

25 ↗

(On Diff #238385)

It is at the last line.

[ABataev]

LG.

openmp/libomptarget/test/api/omp_get_num_devices_with_empty_target.c

25 ↗

(On Diff #238385)

Missed it.

[tianshilei1992]

Is anything else I need to do to make it merge?

[ABataev]

In D72576#1831644, @tianshilei1992 wrote:

Is anything else I need to do to make it merge?

You need to commit it.

[tianshilei1992]

In D72576#1831660, @ABataev wrote:

In D72576#1831644, @tianshilei1992 wrote:

Is anything else I need to do to make it merge?

You need to commit it.

Hmm, seems like I didn't have commit access?

[ABataev]

In D72576#1831715, @tianshilei1992 wrote:

In D72576#1831660, @ABataev wrote:

In D72576#1831644, @tianshilei1992 wrote:

Is anything else I need to do to make it merge?

You need to commit it.

Hmm, seems like I didn't have commit access?

Ok, will commit it for you.

[tianshilei1992]

In D72576#1831732, @ABataev wrote:

In D72576#1831715, @tianshilei1992 wrote:

In D72576#1831660, @ABataev wrote:

In D72576#1831644, @tianshilei1992 wrote:

Is anything else I need to do to make it merge?

You need to commit it.

Hmm, seems like I didn't have commit access?

Ok, will commit it for you.

Thank you!

openmp/libomptarget/src/omptarget.cpp

static int InitLibrary(DeviceTy& Device) {
   int rc = OFFLOAD_SUCCESS;
 
   Device.PendingGlobalsMtx.lock();
   TrlTblMtx.lock();
   for (HostEntriesBeginToTransTableTy::iterator
       ii = HostEntriesBeginToTransTable.begin();
       ii != HostEntriesBeginToTransTable.end(); ++ii) {
     TranslationTable *TransTable = &ii->second;
+    if (TransTable->HostTable.EntriesBegin ==
+        TransTable->HostTable.EntriesEnd) {
+      // No host entry so no need to proceed
+      continue;
+    }
     if (TransTable->TargetsTable[device_id] != 0) {
       // Library entries have already been processed
       continue;
     }
 
     // 1) get image.
     assert(TransTable->TargetsImages.size() > (size_t)device_id &&
            "Not expecting a device ID outside the table's bounds!");
     __tgt_device_image *img = TransTable->TargetsImages[device_id];
     if (!img) {
       DP("No image loaded for device id %d.\n", device_id);
       rc = OFFLOAD_FAIL;
       break;
     }
     // 2) load image into the target table.
     __tgt_target_table *TargetTable =
         TransTable->TargetsTable[device_id] = Device.load_binary(img);
     // Unable to get table for this image: invalidate image and fail.
     if (!TargetTable) {
       DP("Unable to generate entries table for device id %d.\n", device_id);
       TransTable->TargetsImages[device_id] = 0;
       rc = OFFLOAD_FAIL;
       break;
     }
 
     // Verify whether the two table sizes match.
     size_t hsize =
         TransTable->HostTable.EntriesEnd - TransTable->HostTable.EntriesBegin;
     size_t tsize = TargetTable->EntriesEnd - TargetTable->EntriesBegin;
 

openmp/libomptarget/src/rtl.cpp

void RTLsTy::RegisterRequires(int64_t flags) {
   DP("New requires flags %ld compatible with existing %ld!\n",
      flags, RequiresFlags);
 }
 
 void RTLsTy::RegisterLib(__tgt_bin_desc *desc) {
   // Attempt to load all plugins available in the system.
   std::call_once(initFlag, &RTLsTy::LoadRTLs, this);
 
-  if (desc->HostEntriesBegin == desc->HostEntriesEnd)
-    return;
   RTLsMtx.lock();
   // Register the images with the RTLs that understand them, if any.
   for (int32_t i = 0; i < desc->NumDeviceImages; ++i) {
     // Obtain the image.
     __tgt_device_image *img = &desc->DeviceImages[i];
 
     RTLInfoTy *FoundRTL = NULL;
 
void RTLsTy::RegisterLib(__tgt_bin_desc *desc) {
 
 
   DP("Done registering entries!\n");
 }
 
 void RTLsTy::UnregisterLib(__tgt_bin_desc *desc) {
   DP("Unloading target library!\n");
 
-  if (desc->HostEntriesBegin == desc->HostEntriesEnd)
-    return;
   RTLsMtx.lock();
   // Find which RTL understands each image, if any.
   for (int32_t i = 0; i < desc->NumDeviceImages; ++i) {
     // Obtain the image.
     __tgt_device_image *img = &desc->DeviceImages[i];
 
     RTLInfoTy *FoundRTL = NULL;
 

openmp/libomptarget/test/api/omp_get_num_devices.c

 // 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
 
 #include <stdio.h>
 #include <omp.h>
 
 int test_omp_get_num_devices()
 {
   /* checks that omp_get_num_devices() > 0 */
   int num_devices = omp_get_num_devices();
   printf("num_devices = %d\n", num_devices);
 
-  #pragma omp target
-  {}
-
   return (num_devices > 0);
 }
 
 int main()
 {
   int i;
   int failed=0;