[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

92–100 ↗

(On Diff #238062)

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

92–100 ↗

(On Diff #238062)

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

92–100 ↗

(On Diff #238062)

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 ↗

(On Diff #238160)

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 ↗

(On Diff #238160)

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 ↗

(On Diff #238160)

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 ↗

(On Diff #238160)

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/rtl.cpp

 }
 
 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;
-
-    // Scan the RTLs that have associated images until we find one that supports
-    // the current image.
+  // Load all plugins available in the system and initialize them
+  std::call_once(initFlag, [this]() {
+    // 1. Load plugins
+    this->LoadRTLs();
+    // 2. Initialize them if any one has not been initialized yet
     for (auto &R : RTLs.AllRTLs) {
-      if (!R.is_valid_binary(img)) {
-        DP("Image " DPxMOD " is NOT compatible with RTL %s!\n",
-            DPxPTR(img->ImageStart), R.RTLName.c_str());
-        continue;
-      }
-
-      DP("Image " DPxMOD " is compatible with RTL %s!\n",
-          DPxPTR(img->ImageStart), R.RTLName.c_str());
-
       // If this RTL is not already in use, initialize it.
       if (!R.isUsed) {
         // Initialize the device information for the RTL we are about to use.

         size_t start = Devices.size();
         Devices.resize(start + R.NumberOfDevices, device);
         for (int32_t device_id = 0; device_id < R.NumberOfDevices;
              device_id++) {
           // global device ID
           Devices[start + device_id].DeviceID = start + device_id;
           // RTL local device ID

                     ? 0
                     : RTLs.UsedRTLs.back()->Idx +
                           RTLs.UsedRTLs.back()->NumberOfDevices;
-        assert((size_t) R.Idx == start &&
+        assert((size_t)R.Idx == start &&
                "RTL index should equal the number of devices used so far.");
         R.isUsed = true;
         RTLs.UsedRTLs.push_back(&R);
 
         DP("RTL " DPxMOD " has index %d!\n", DPxPTR(R.LibraryHandler), R.Idx);
       }
+    }
+  });
+
+  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;
+
+    // Scan the RTLs that have associated images until we find one that supports
+    // the current image.
+    for (auto &R : RTLs.AllRTLs) {
+      if (!R.is_valid_binary(img)) {
+        DP("Image " DPxMOD " is NOT compatible with RTL %s!\n",
+           DPxPTR(img->ImageStart), R.RTLName.c_str());
+        continue;
+      }
+
+      DP("Image " DPxMOD " is compatible with RTL %s!\n",
+         DPxPTR(img->ImageStart), R.RTLName.c_str());
+
+      assert(R.isUsed && "RTL should have already been initialized!");
 
       // Initialize (if necessary) translation table for this library.
       TrlTblMtx.lock();
-      if(!HostEntriesBeginToTransTable.count(desc->HostEntriesBegin)){
+      if (!HostEntriesBeginToTransTable.count(desc->HostEntriesBegin)) {
         TranslationTable &tt =
             HostEntriesBeginToTransTable[desc->HostEntriesBegin];
         tt.HostTable.EntriesBegin = desc->HostEntriesBegin;