[OpenMP] Add function for setting LIBOMPTARGET_INFO at runtime

[jhuber6]

Summary

This patch adds a new runtime function __tgt_set_info_flag that allows the
user to set the information level at runtime without using the environment
variable. Using this will require an extern function, but will eventually be
added into an auxilliary library for OpenMP support functions.

This patch required moving the current InfoLevel to a global variable which must
be instantiated by each plugin. Communication between the host and device is
done through the __tgt_rtl_set_info_flag which should be defined by any plugin
that want's to use the information level.

[jdoerfert]

LG

[jhuber6]

Changing the interface, previously each plugin / file had its own version of the static variable, now to change it we need a global atomic variable to represent it. Whenever the value is updated it also needs to be sent to the currently loaded plugins if they support it. I added the __tgt_rtl_set_info_flag function to the plugin library for this purpose.

[JonChesterfield]

Surprising amount of code here, probably because it changes from one function to some code * number plugins.

Was the idea to allow directly writing to static uint32_t InfoLevel = 0;, which is inconvenient when it's a static function? If so, that is still achievable in header only fashion without adding global variables to program scope, e.g. by returning a reference from getInfoLevel and adding a setter function, or going with a detail function that is called by the getter and setter.

[jhuber6]

In D100774#2709580, @JonChesterfield wrote:

Surprising amount of code here, probably because it changes from one function to some code * number plugins.

Was the idea to allow directly writing to static uint32_t InfoLevel = 0;, which is inconvenient when it's a static function? If so, that is still achievable in header only fashion without adding global variables to program scope, e.g. by returning a reference from getInfoLevel and adding a setter function, or going with a detail function that is called by the getter and setter.

The problem I had was that in libomptarget there wasn't a single instance of #include "Debug.h". This meant that there were about three different versions of InfoLevel. This meant if the variable was static it wouldn't set everyone's version of InfoLevel when called from the function. It might've been a better solution to just require that Debug.h is included only once per library, then use a method like you described. It would certainly be nicer to not require setting the extern variables for each plugin, which is what I originally removed.

[JonChesterfield]

In D100774#2709610, @jhuber6 wrote:

The problem I had was that in libomptarget there wasn't a single instance of #include "Debug.h". This meant that there were about three different versions of InfoLevel. This meant if the variable was static it wouldn't set everyone's version of InfoLevel when called from the function. It might've been a better solution to just require that Debug.h is included only once per library, then use a method like you described. It would certainly be nicer to not require setting the extern variables for each plugin, which is what I originally removed.

I think that was collateral damage from getInfoLevel being static. Changing that to just inline will mean a single copy per shared library, which seems to be the right semantics. Probably still want the local variable to be atomic qualified though.

[jhuber6]

In D100774#2709722, @JonChesterfield wrote:

In D100774#2709610, @jhuber6 wrote:

The problem I had was that in libomptarget there wasn't a single instance of #include "Debug.h". This meant that there were about three different versions of InfoLevel. This meant if the variable was static it wouldn't set everyone's version of InfoLevel when called from the function. It might've been a better solution to just require that Debug.h is included only once per library, then use a method like you described. It would certainly be nicer to not require setting the extern variables for each plugin, which is what I originally removed.

I think that was collateral damage from getInfoLevel being static. Changing that to just inline will mean a single copy per shared library, which seems to be the right semantics. Probably still want the local variable to be atomic qualified though.

Seems reasonable. I'd probably have another function that returns a reference to the internal atomic variable, then getInfoLevel() just returns the value inside the atomic. Should this be a another review? Or a revert + fix.

[JonChesterfield]

In D100774#2709750, @jhuber6 wrote:

Seems reasonable. I'd probably have another function that returns a reference to the internal atomic variable, then getInfoLevel() just returns the value inside the atomic.

Sounds good.

In D100774#2709750, @jhuber6 wrote:

Should this be a another review? Or a revert + fix.

I don't think anything is broken here so would suggest another patch + review on top, leaving this in place. Plus that way if the header only version doesn't work out so well we still have this fix.

openmp/docs/design/Runtimes.rst

 host pointers to the newly created device data. The data mappings in the OpenMP
 device kernel show the default mappings being used for all the variables used
 implicitly on the device. Because ``X`` and ``Y`` are already mapped in the
 device's table, no new entries are created. Additionally, the default mapping
 shows that ``D`` will be copied back from the device once the OpenMP device
 kernel region ends even though it isn't written to. Finally, at the end of the
 OpenMP data region the entries for ``X`` and ``Y`` are removed from the table.
 
+The information level can be controlled at runtime using an internal
+libomptarget library call ``__tgt_set_info_flag``. This allows for different
+levels of information to be enabled or disabled for certain regions of code.
+Using this requires declaring the function signature as an external function so
+it can be linked with the runtime library.
+
+.. code-block:: c++
+
+    extern "C" void __tgt_set_info_flag(uint32_t);
+
+    extern foo();
+
+    int main() {
+      __tgt_set_info_flag(0x10);
+    #pragma omp target
+      foo();
+    }
+
 .. _libopenmptarget_errors:
 
 Errors:
 ^^^^^^^
 
 ``libomptarget`` provides error messages when the program fails inside the
 OpenMP target region. Common causes of failure could be an invalid pointer
 access, running out of device memory, or trying to offload when the device is

openmp/libomptarget/include/Debug.h

 //                               // 8 digits for 32bit systems
 //                               // 16 digits for 64bit
 //   (uintptr_t) ptr);
 //
 //===----------------------------------------------------------------------===//
 #ifndef _OMPTARGET_DEBUG_H
 #define _OMPTARGET_DEBUG_H
 
+#include <atomic>
 #include <mutex>
 
 /// 32-Bit field data attributes controlling information presented to the user.
 enum OpenMPInfoType : uint32_t {
   // Print data arguments and attributes upon entering an OpenMP device kernel.
   OMP_INFOTYPE_KERNEL_ARGS = 0x0001,
   // Indicate when an address already exists in the device mapping table.
   OMP_INFOTYPE_MAPPING_EXISTS = 0x0002,
#define GCC_VERSION                                                            \
   (__GNUC__ * 10000 + __GNUC_MINOR__ * 100 + __GNUC_PATCHLEVEL__)
 
 #if !defined(__clang__) && defined(__GNUC__) && GCC_VERSION < 70100
 #define USED __attribute__((used))
 #else
 #define USED
 #endif
 
+// Interface to the InfoLevel variable defined by each library.
+extern std::atomic<uint32_t> InfoLevel;
+
 // Add __attribute__((used)) to work around a bug in gcc 5/6.
 USED static inline uint32_t getInfoLevel() {
-  static uint32_t InfoLevel = 0;
   static std::once_flag Flag{};
   std::call_once(Flag, []() {
     if (char *EnvStr = getenv("LIBOMPTARGET_INFO"))
-      InfoLevel = std::stoi(EnvStr);
+      InfoLevel.store(std::stoi(EnvStr));
   });
 
-  return InfoLevel;
+  return InfoLevel.load();
 }
 
 // Add __attribute__((used)) to work around a bug in gcc 5/6.
 USED static inline uint32_t getDebugLevel() {
   static uint32_t DebugLevel = 0;
   static std::once_flag Flag{};
   std::call_once(Flag, []() {
     if (char *EnvStr = getenv("LIBOMPTARGET_DEBUG"))

openmp/libomptarget/include/omptarget.h

int __tgt_target_teams_nowait_mapper(
     int32_t thread_limit, int32_t depNum, void *depList, int32_t noAliasDepNum,
     void *noAliasDepList);
 
 void __kmpc_push_target_tripcount(int64_t device_id, uint64_t loop_tripcount);
 
 void __kmpc_push_target_tripcount_mapper(ident_t *loc, int64_t device_id,
                                          uint64_t loop_tripcount);
 
+void __tgt_set_info_flag(uint32_t);
+
 #ifdef __cplusplus
 }
 #endif
 
 #ifdef __cplusplus
 #define EXTERN extern "C"
 #else
 #define EXTERN extern
 #endif
 
 #endif // _OMPTARGET_H_

openmp/libomptarget/include/omptargetplugin.h

int32_t __tgt_rtl_run_target_team_region_async(
     int32_t ID, void *Entry, void **Args, ptrdiff_t *Offsets, int32_t NumArgs,
     int32_t NumTeams, int32_t ThreadLimit, uint64_t loop_tripcount,
     __tgt_async_info *AsyncInfo);
 
 // Device synchronization. In case of success, return zero. Otherwise, return an
 // error code.
 int32_t __tgt_rtl_synchronize(int32_t ID, __tgt_async_info *AsyncInfo);
 
+// Set plugin's internal information flag externally.
+void __tgt_rtl_set_info_flag(uint32_t);
+
 #ifdef __cplusplus
 }
 #endif
 
 #endif // _OMPTARGETPLUGIN_H_

openmp/libomptarget/plugins/amdgpu/src/rtl.cpp

int32_t __tgt_rtl_synchronize(int32_t device_id, __tgt_async_info *AsyncInfo) {
   // Cuda asserts that AsyncInfo->Queue is non-null, but this invariant
   // is not ensured by devices.cpp for amdgcn
   // assert(AsyncInfo->Queue && "AsyncInfo->Queue is nullptr");
   if (AsyncInfo->Queue) {
     finiAsyncInfo(AsyncInfo);
   }
   return OFFLOAD_SUCCESS;
 }
+
+// AMDGPU plugin's internal InfoLevel.
+std::atomic<uint32_t> InfoLevel;

openmp/libomptarget/plugins/cuda/src/rtl.cpp

int32_t __tgt_rtl_synchronize(int32_t device_id,
                               __tgt_async_info *async_info_ptr) {
   assert(DeviceRTL.isValidDeviceId(device_id) && "device_id is invalid");
   assert(async_info_ptr && "async_info_ptr is nullptr");
   assert(async_info_ptr->Queue && "async_info_ptr->Queue is nullptr");
 
   return DeviceRTL.synchronize(device_id, async_info_ptr);
 }
 
+void __tgt_rtl_set_info_flag(uint32_t NewInfoLevel) {
+  InfoLevel.store(NewInfoLevel);
+}
+
 #ifdef __cplusplus
 }
 #endif
+
+// Cuda plugin's internal InfoLevel.
+std::atomic<uint32_t> InfoLevel;

openmp/libomptarget/plugins/exports

global:
     __tgt_rtl_run_target_team_region;
     __tgt_rtl_run_target_team_region_async;
     __tgt_rtl_run_target_region;
     __tgt_rtl_run_target_region_async;
     __tgt_rtl_synchronize;
     __tgt_rtl_register_lib;
     __tgt_rtl_unregister_lib;
     __tgt_rtl_supports_empty_images;
+    __tgt_rtl_set_info_flag;
   local:
     *;
 };

openmp/libomptarget/plugins/generic-elf-64bit/src/rtl.cpp

int32_t __tgt_rtl_run_target_region(int32_t device_id, void *tgt_entry_ptr,
   // use one team and one thread.
   return __tgt_rtl_run_target_team_region(device_id, tgt_entry_ptr, tgt_args,
                                           tgt_offsets, arg_num, 1, 1, 0);
 }
 
 #ifdef __cplusplus
 }
 #endif
+
+// Elf-64 plugin's internal InfoLevel.
+std::atomic<uint32_t> InfoLevel;

openmp/libomptarget/plugins/remote/src/rtl.cpp

return Manager->runTargetTeamRegionAsync(
       DeviceId, TgtEntryPtr, TgtArgs, TgtOffsets, ArgNum, TeamNum, ThreadLimit,
       LoopTripCount, AsyncInfo);
 }
 
 // Exposed library API function
 #ifdef __cplusplus
 }
 #endif
+
+// Remote Offloading interal InfoLevel.
+std::atomic<uint32_t> InfoLevel;

openmp/libomptarget/plugins/ve/src/rtl.cpp

 // Otherwise, return an error code.
 int32_t __tgt_rtl_run_target_region(int32_t ID, void *Entry, void **Args,
                                     ptrdiff_t *Offsets, int32_t NumArgs) {
   return __tgt_rtl_run_target_team_region(ID, Entry, Args, Offsets, NumArgs, 1,
                                           1, 0);
 }
 
 int32_t __tgt_rtl_supports_empty_images() { return 1; }
+
+// VEC plugin's internal InfoLevel.
+std::atomic<uint32_t> InfoLevel;

openmp/libomptarget/src/exports

global:
     omp_target_is_present;
     omp_target_memcpy;
     omp_target_memcpy_rect;
     omp_target_associate_ptr;
     omp_target_disassociate_ptr;
     llvm_omp_target_alloc_host;
     llvm_omp_target_alloc_shared;
     llvm_omp_target_alloc_device;
+    __tgt_set_info_flag;
   local:
     *;
 };
 

openmp/libomptarget/src/interface.cpp

EXTERN void __kmpc_push_target_tripcount_mapper(ident_t *loc, int64_t device_id,
 
   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();
 }
+
+EXTERN void __tgt_set_info_flag(uint32_t NewInfoLevel) {
+  InfoLevel.store(NewInfoLevel);
+  for (auto &R : PM->RTLs.AllRTLs) {
+    if (R.set_info_flag)
+      R.set_info_flag(NewInfoLevel);
+  }
+}
+
+// Libomptarget's InfoLevel storage.
+std::atomic<uint32_t> InfoLevel;

openmp/libomptarget/src/rtl.h

struct RTLInfoTy {
   typedef int32_t(run_team_region_async_ty)(int32_t, void *, void **,
                                             ptrdiff_t *, int32_t, int32_t,
                                             int32_t, uint64_t,
                                             __tgt_async_info *);
   typedef int64_t(init_requires_ty)(int64_t);
   typedef int64_t(synchronize_ty)(int32_t, __tgt_async_info *);
   typedef int32_t (*register_lib_ty)(__tgt_bin_desc *);
   typedef int32_t(supports_empty_images_ty)();
+  typedef void(set_info_flag_ty)(uint32_t);
 
   int32_t Idx = -1;             // RTL index, index is the number of devices
                                 // of other RTLs that were registered before,
                                 // i.e. the OpenMP index of the first device
                                 // to be registered with this RTL.
   int32_t NumberOfDevices = -1; // Number of devices this RTL deals with.
 
   void *LibraryHandler = nullptr;
#endif
   run_region_async_ty *run_region_async = nullptr;
   run_team_region_ty *run_team_region = nullptr;
   run_team_region_async_ty *run_team_region_async = nullptr;
   init_requires_ty *init_requires = nullptr;
   synchronize_ty *synchronize = nullptr;
   register_lib_ty register_lib = nullptr;
   register_lib_ty unregister_lib = nullptr;
   supports_empty_images_ty *supports_empty_images = nullptr;
+  set_info_flag_ty *set_info_flag = nullptr;
 
   // Are there images associated with this RTL.
   bool isUsed = false;
 
   // Mutex for thread-safety when calling RTL interface functions.
   // It is easier to enforce thread-safety at the libomptarget level,
   // so that developers of new RTLs do not have to worry about it.
   std::mutex Mtx;

openmp/libomptarget/src/rtl.cpp

#endif
     *((void **)&R.is_data_exchangable) =
         dlsym(dynlib_handle, "__tgt_rtl_is_data_exchangable");
     *((void **)&R.register_lib) =
         dlsym(dynlib_handle, "__tgt_rtl_register_lib");
     *((void **)&R.unregister_lib) =
         dlsym(dynlib_handle, "__tgt_rtl_unregister_lib");
     *((void **)&R.supports_empty_images) =
         dlsym(dynlib_handle, "__tgt_rtl_supports_empty_images");
+    *((void **)&R.set_info_flag) =
+        dlsym(dynlib_handle, "__tgt_rtl_set_info_flag");
   }
 
   DP("RTLs loaded!\n");
 
   return;
 }
 
 ////////////////////////////////////////////////////////////////////////////////

openmp/libomptarget/test/offloading/info.c

 // RUN: %libomptarget-compile-nvptx64-nvidia-cuda -gline-tables-only && env LIBOMPTARGET_INFO=31 %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | %fcheck-nvptx64-nvidia-cuda -allow-empty -check-prefix=INFO
 
 #include <stdio.h>
 #include <omp.h>
 
 #define N 64
 
+extern void __tgt_set_info_flag(unsigned);
+
 int main() {
   int A[N];
   int B[N];
   int C[N];
   int val = 1;
 
 // INFO: CUDA device 0 info: Device supports up to {{.*}} CUDA blocks and {{.*}} threads with a warp size of {{.*}}
-// INFO: Libomptarget device 0 info: Entering OpenMP data region at info.c:39:1 with 3 arguments:
+// INFO: Libomptarget device 0 info: Entering OpenMP data region at info.c:{{[0-9]+}}:1 with 3 arguments:
 // INFO: Libomptarget device 0 info: alloc(A[0:64])[256]
 // INFO: Libomptarget device 0 info: tofrom(B[0:64])[256]
 // INFO: Libomptarget device 0 info: to(C[0:64])[256]
 // INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=A[0:64]
 // INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=B[0:64]
 // INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=C[0:64]
-// INFO: Libomptarget device 0 info: OpenMP Host-Device pointer mappings after block at info.c:39:1:
+// INFO: Libomptarget device 0 info: OpenMP Host-Device pointer mappings after block at info.c:{{[0-9]+}}:1:
 // INFO: Libomptarget device 0 info: Host Ptr           Target Ptr         Size (B) RefCount Declaration
-// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256      1        C[0:64] at info.c:11:7
-// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256      1        B[0:64] at info.c:10:7
-// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256      1        A[0:64] at info.c:9:7
-// INFO: Libomptarget device 0 info: Entering OpenMP kernel at info.c:40:1 with 1 arguments:
+// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256      1        C[0:64] at info.c:{{[0-9]+}}:7
+// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256      1        B[0:64] at info.c:{{[0-9]+}}:7
+// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256      1        A[0:64] at info.c:{{[0-9]+}}:7
+// INFO: Libomptarget device 0 info: Entering OpenMP kernel at info.c:{{[0-9]+}}:1 with 1 arguments:
 // INFO: Libomptarget device 0 info: firstprivate(val)[4]
 // INFO: CUDA device 0 info: Launching kernel {{.*}} with {{.*}} and {{.*}} threads in {{.*}} mode
-// INFO: Libomptarget device 0 info: OpenMP Host-Device pointer mappings after block at info.c:40:1:
+// INFO: Libomptarget device 0 info: OpenMP Host-Device pointer mappings after block at info.c:{{[0-9]+}}:1:
 // INFO: Libomptarget device 0 info: Host Ptr           Target Ptr         Size (B) RefCount Declaration
-// INFO: Libomptarget device 0 info: 0x{{.*}} 0x{{.*}} 256      1        C[0:64] at info.c:11:7
-// INFO: Libomptarget device 0 info: 0x{{.*}} 0x{{.*}} 256      1        B[0:64] at info.c:10:7
-// INFO: Libomptarget device 0 info: 0x{{.*}} 0x{{.*}} 256      1        A[0:64] at info.c:9:7
-// INFO: Libomptarget device 0 info: Exiting OpenMP data region at info.c:39:1
+// INFO: Libomptarget device 0 info: 0x{{.*}} 0x{{.*}} 256      1        C[0:64] at info.c:{{[0-9]+}}:7
+// INFO: Libomptarget device 0 info: 0x{{.*}} 0x{{.*}} 256      1        B[0:64] at info.c:{{[0-9]+}}:7
+// INFO: Libomptarget device 0 info: 0x{{.*}} 0x{{.*}} 256      1        A[0:64] at info.c:{{[0-9]+}}:7
+// INFO: Libomptarget device 0 info: Exiting OpenMP data region at info.c:{{[0-9]+}}:1
 // INFO: Libomptarget device 0 info: Removing map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=C[0:64]
 // INFO: Libomptarget device 0 info: Removing map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=B[0:64]
 // INFO: Libomptarget device 0 info: Removing map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=A[0:64]
 #pragma omp target data map(alloc:A[0:N]) map(tofrom:B[0:N]) map(to:C[0:N])
 #pragma omp target firstprivate(val)
   { val = 1; }
 
+  __tgt_set_info_flag(0x0);
+// INFO-NOT: Libomptarget device 0 info: {{.*}}
+#pragma omp target
+  { }
+
   return 0;
 }