[OpenMP] Implementing omp_get_device_num()

[josemonsalve2]

Summary

This patch implements omp_get_device_num() in the host and the device.

It uses the already existing getDeviceNum in the device config for the device.
And in the host it uses the omp_get_num_devices().

Two simple tests added

[tianshilei1992]

openmp/libomptarget/src/api.cpp

35

The spec says:

When called on the host device, it will return the same value as the omp_get_initial_device routine.

Therefore I suggest we call omp_get_initial_device here.

[josemonsalve2]

Ah! I missed that. Thanks

[josemonsalve2]

Changing omp_get_num_devices() for omp_get_initial_device()

[tianshilei1992]

LGTM with one code style issue.

openmp/libomptarget/src/api.cpp

39

HostDevice should be the LLVM style variable name.

[josemonsalve2]

Fixing var name to coding standard

[josemonsalve2]

I'm not sure what to do about the buildbot errors I am getting. I'm accessing them but there's no info that tells me what's going on.

[josemonsalve2]

Rebasing to master

openmp/libomptarget/DeviceRTL/include/Configuration.h

enum DebugKind : uint32_t {
   FunctionTracing = 1U << 1,
   CommonIssues = 1U << 2,
 };
 
 /// Return the number of devices in the system, same number as returned on the
 /// host by omp_get_num_devices.
 uint32_t getNumDevices();
 
-/// Return the number of devices in the system, same number as returned on the
-/// host by omp_get_num_devices.
+/// Return the device number in the system for omp_get_device_num.
 uint32_t getDeviceNum();
 
 /// Return the user choosen debug level.
 uint32_t getDebugKind();
 
 /// Return the amount of dynamic shared memory that was allocated at launch.
 uint64_t getDynamicMemorySize();
 

openmp/libomptarget/DeviceRTL/include/Interface.h

 int omp_get_cancellation(void);
 
 void omp_set_default_device(int deviceId);
 
 int omp_get_default_device(void);
 
 int omp_get_num_devices(void);
 
+int omp_get_device_num(void);
+
 int omp_get_num_teams(void);
 
 int omp_get_team_num();
 
 int omp_get_initial_device(void);
 
 void *llvm_omp_target_dynamic_shared_alloc();
 

openmp/libomptarget/DeviceRTL/src/State.cpp

 int omp_get_cancellation(void) { return 0; }
 
 void omp_set_default_device(int) {}
 
 int omp_get_default_device(void) { return -1; }
 
 int omp_get_num_devices(void) { return config::getNumDevices(); }
 
+int omp_get_device_num(void) { return config::getDeviceNum(); }
+
 int omp_get_num_teams(void) { return mapping::getNumberOfBlocks(); }
 
 int omp_get_team_num() { return mapping::getBlockId(); }
 
 int omp_get_initial_device(void) { return -1; }
 }
 
 extern "C" {

openmp/libomptarget/include/omptarget.h

struct __tgt_device_info {
   void *Device = nullptr;
 };
 
 #ifdef __cplusplus
 extern "C" {
 #endif
 
 int omp_get_num_devices(void);
+int omp_get_device_num(void);
 int omp_get_initial_device(void);
 void *omp_target_alloc(size_t size, int device_num);
 void omp_target_free(void *device_ptr, int device_num);
 int omp_target_is_present(const void *ptr, int device_num);
 int omp_target_memcpy(void *dst, const void *src, size_t length,
                       size_t dst_offset, size_t src_offset, int dst_device,
                       int src_device);
 int omp_target_memcpy_rect(void *dst, const void *src, size_t element_size,

openmp/libomptarget/src/api.cpp

EXTERN int omp_get_num_devices(void) {
   size_t DevicesSize = PM->Devices.size();
   PM->RTLsMtx.unlock();
 
   DP("Call to omp_get_num_devices returning %zd\n", DevicesSize);
 
   return DevicesSize;
 }
 
+EXTERN int omp_get_device_num(void) {
+  TIMESCOPE();
+  int HostDevice = omp_get_initial_device();
+
+  DP("Call to omp_get_device_num returning %d\n", HostDevice);
+
+  return HostDevice;
+}
+
 EXTERN int omp_get_initial_device(void) {
   TIMESCOPE();
   int hostDevice = omp_get_num_devices();
   DP("Call to omp_get_initial_device returning %d\n", hostDevice);
   return hostDevice;
 }
 
 EXTERN void *omp_target_alloc(size_t size, int device_num) {

openmp/libomptarget/src/exports

global:
     __tgt_target_data_update_nowait_mapper;
     __tgt_target_nowait_mapper;
     __tgt_target_teams_nowait_mapper;
     __tgt_mapper_num_components;
     __tgt_push_mapper_component;
     __kmpc_push_target_tripcount;
     __kmpc_push_target_tripcount_mapper;
     omp_get_num_devices;
+    omp_get_device_num;
     omp_get_initial_device;
     omp_target_alloc;
     omp_target_free;
     omp_target_is_present;
     omp_target_memcpy;
     omp_target_memcpy_rect;
     omp_target_associate_ptr;
     omp_target_disassociate_ptr;

openmp/libomptarget/test/api/omp_get_device_num.c

+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <stdio.h>
+#include <omp.h>
+
+int test_omp_get_device_num()
+{
+  /* checks that omp_get_device_num() == omp_get_num_devices() in the host */
+  int device_num = omp_get_device_num();
+  printf("device_num = %d\n", device_num);
+
+  #pragma omp target
+  {}
+
+  return (device_num == omp_get_num_devices());
+}
+
+int main()
+{
+  int i;
+  int failed=0;
+
+  if (!test_omp_get_device_num()) {
+    failed++;
+  }
+  if (failed)
+    printf("FAIL\n");
+  else
+    printf("PASS\n");
+  return failed;
+}
+
+// CHECK: PASS

openmp/runtime/test/api/omp_get_device_num.c

+// RUN: %libomp-compile-and-run
+// Linking fails for icc 18
+// UNSUPPORTED: icc-18
+
+#include <stdio.h>
+#include "omp_testsuite.h"
+
+int test_omp_get_device_num()
+{
+  /* checks that omp_get_device_num */
+  int device_num = omp_get_device_num();
+
+  return (device_num == omp_get_num_devices());
+}
+
+int main()
+{
+  int i;
+  int num_failed=0;
+
+  for(i = 0; i < REPETITIONS; i++) {
+    if(!test_omp_get_device_num()) {
+      num_failed++;
+    }
+  }
+  return num_failed;
+}