[OpenMP] Device Time Profile

[nrj5k]

Summary

Adding functionality to
— collect time information from the target devices
— sideloads the collected data TimeProfiler Entrys

Resulting in a Libomptarget Profile containing host + target information.

[jdoerfert]

Can you 1) upload the patches with full context (see https://llvm.org/docs/Contributing.html#how-to-submit-a-patch) and 2) split it into the refactor of the timer facility and the OpenMP patch.

[nrj5k]

• only contains OpenMP patch for Device Time Profile.
• patch contains full context

[jdoerfert]

I left some initial comments, I think this is not in a "working" state, it looks like parts are missing, beyond the LLVM changes.

openmp/libomptarget/DeviceRTL/src/Kernel.cpp

25

Leftover

45

Turn this around.

if (Idx >= OMP_ARRAY_SIZE)
  return OMP_ARRAY_SIZE
...

Where does OMP_ARRAY_SIZE and OMP_THREADS come from? Shouldn't those be read from the struct the runtime initializes?

50

Don't use parallel here. This is the device runtime.
Depending on where this is called, all threads will execute it anyway.
Just use mapping::getThreadIdInBlock() to get the thread id and mapping::getBlockId() to get the block id.
Same above.

120

What is this?

170

And this?

[dhruvachak]

Can you please describe at a high level what data "Device Time Profile" is supposed to collect and how you are collecting it? Is there a specification somewhere for this functionality?

[jdoerfert]

In D133463#3986224, @dhruvachak wrote:

Can you please describe at a high level what data "Device Time Profile" is supposed to collect and how you are collecting it? Is there a specification somewhere for this functionality?

The profile collects time spend between user (or compiler) placed calls to start/stop a timer. The patch is not complete and I am not sure if Neeraj monitors these emails.
The values are spliced into our LIBOMPTARGET_PROFILE output, so you view it the same way, e.g., via Chrome tracing.

llvm/test/Transforms/OpenMP/spmdization_constant_prop.ll

 @__omp_rtl_assume_teams_oversubscription = weak_odr hidden local_unnamed_addr addrspace(1) constant i32 0
 @__omp_rtl_assume_threads_oversubscription = weak_odr hidden local_unnamed_addr addrspace(1) constant i32 0
 @0 = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00", align 1
 @1 = private unnamed_addr addrspace(1) constant %struct.ident_t { i32 0, i32 2, i32 0, i32 22, ptr @0 }, align 8
 @__omp_offloading_20_11e3950_main_l12_exec_mode = weak addrspace(1) constant i8 1
 @__omp_rtl_debug_kind = weak_odr hidden local_unnamed_addr addrspace(1) constant i32 0
 @__omp_rtl_assume_no_thread_state = weak_odr hidden local_unnamed_addr addrspace(1) constant i32 0
 @omptarget_device_environment = weak protected addrspace(4) global %struct.DeviceEnvironmentTy undef, align 4
+@omptarget_device_time_profile = weak protected addrspace(4) global %struct.DeviceTimeProfileTy undef, align 4
 @IsSPMDMode = weak hidden addrspace(3) global i32 undef, align 4
 @.str.12 = private unnamed_addr addrspace(4) constant [47 x i8] c"ValueRAII initialization with wrong old value!\00", align 1
 @_ZN12_GLOBAL__N_122SharedMemorySmartStackE = internal addrspace(3) global %"struct.(anonymous namespace)::SharedMemorySmartStackTy" undef, align 16
 @_ZN12_GLOBAL__N_19TeamStateE = internal unnamed_addr addrspace(3) global %"struct.(anonymous namespace)::TeamStateTy" undef, align 8
 @_ZN12_GLOBAL__N_112ThreadStatesE = internal unnamed_addr addrspace(3) global [1024 x ptr] undef, align 16
 @.str.12.47 = private unnamed_addr addrspace(4) constant [49 x i8] c"Thread state modified while explicitly disabled!\00", align 1
 @_ZL29SharedMemVariableSharingSpace = internal unnamed_addr addrspace(3) global [64 x ptr] undef, align 16
 @G = global i32 undef
 @llvm.used = appending addrspace(1) global [2 x ptr] [ptr addrspacecast (ptr addrspace(3) @IsSPMDMode to ptr), ptr addrspacecast (ptr addrspace(4) @omptarget_device_environment to ptr)], section "llvm.metadata"
[e-mail] = appending addrspace(1) global [2 x ptr] [ptr addrspacecast (ptr addrspace(3) @IsSPMDMode to ptr), ptr addrspacecast (ptr addrspace(4) @omptarget_device_time_profile to ptr)], section "llvm.metadata"
 @llvm.compiler.used = appending addrspace(1) global [1 x ptr] [ptr addrspacecast (ptr addrspace(1) @__omp_offloading_20_11e3950_main_l12_exec_mode to ptr)], section "llvm.metadata"
 
 ; Function Attrs: alwaysinline convergent norecurse nounwind
 define weak_odr amdgpu_kernel void @__omp_offloading_20_11e3950_main_l12(i64 noundef %nxyz, i64 noundef %ng, ptr noundef nonnull align 8 dereferenceable(8) %aa) local_unnamed_addr #0 {
 entry:
   %ng1 = alloca i32, align 4
   %captured_vars_addrs = alloca [2 x ptr], align 8, addrspace(5)
   %0 = call i32 @__kmpc_target_init(ptr addrspacecast (ptr addrspace(1) @1 to ptr), i8 1, i1 true, i1 true)

openmp/libomptarget/DeviceRTL/include/Configuration.h

 //===--- Configuration.h - OpenMP device configuration interface -- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // API to query the global (constant) device environment.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef OMPTARGET_CONFIGURATION_H
 #define OMPTARGET_CONFIGURATION_H
 
+#include "DeviceTimeProfile.h"
 #include "Types.h"
+#include "targetprofile.h"
 
 namespace _OMP {
+
+namespace targetprof {
+
+int getNumRecords();
+_OMP::prof::Record *getRecords();
+
+} // namespace targetprof
+
 namespace config {
 
 enum DebugKind : uint32_t {
   Assertion = 1U << 0,
   FunctionTracing = 1U << 1,
   CommonIssues = 1U << 2,
 };
 

openmp/libomptarget/DeviceRTL/include/Interface.h

 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_num_timeprofile_records();
+
 int omp_get_initial_device(void);
 
 void *llvm_omp_target_dynamic_shared_alloc();
 
 /// Synchronization
 ///
 ///{
 void omp_init_lock(omp_lock_t *Lock);

openmp/libomptarget/DeviceRTL/src/Configuration.cpp

 //===- Configuration.cpp - OpenMP device configuration interface -- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file contains the data object of the constant device environment and the
 // query API.
 //
 //===----------------------------------------------------------------------===//
 
 #include "Configuration.h"
 #include "DeviceEnvironment.h"
+#include "DeviceTimeProfile.h"
 #include "State.h"
 #include "Types.h"
 
 using namespace _OMP;
 
 #pragma omp begin declare target device_type(nohost)
 
 // defined by CGOpenMPRuntimeGPU
 extern uint32_t __omp_rtl_debug_kind;
 extern uint32_t __omp_rtl_assume_no_thread_state;
 extern uint32_t __omp_rtl_assume_no_nested_parallelism;
 
 // TODO: We want to change the name as soon as the old runtime is gone.
 // This variable should be visibile to the plugin so we override the default
 // hidden visibility.
 DeviceEnvironmentTy CONSTANT(omptarget_device_environment)
     __attribute__((used, retain, weak, visibility("protected")));
 
+DeviceTimeProfileTy omptarget_device_time_profile
+    __attribute__((used, retain, weak, visibility("protected")));
+
+int targetprof::getNumRecords() { return omptarget_device_time_profile.StartC; }
+
+_OMP::prof::Record *targetprof::getRecords() {
+  return omptarget_device_time_profile.Records;
+}
+
 uint32_t config::getDebugKind() {
   return __omp_rtl_debug_kind & omptarget_device_environment.DebugKind;
 }
 
 uint32_t config::getNumDevices() {
   return omptarget_device_environment.NumDevices;
 }
 

openmp/libomptarget/DeviceRTL/src/Kernel.cpp

 //===--- Kernel.cpp - OpenMP device kernel interface -------------- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file contains the kernel entry points for the device.
 //
 //===----------------------------------------------------------------------===//
 
 #include "Debug.h"
+#include "DeviceTimeProfile.h"
 #include "Interface.h"
 #include "Mapping.h"
 #include "State.h"
 #include "Synchronization.h"
 #include "Types.h"
 
 using namespace _OMP;
 
 #pragma omp begin declare target device_type(nohost)
 
+// CLOCKINIT();
+extern DeviceTimeProfileTy omptarget_device_time_profile;
+
+namespace _OMP {
+namespace prof {
+
+__attribute__((used)) int clockIn(const char *FuncName, int Idx) {
+  if (Idx < OMP_ARRAY_SIZE) {
+#pragma omp parallel
+    {
+      int thread_ = omp_get_thread_num();
+      if (thread_ < OMP_THREADS) {
+        omptarget_device_time_profile.Records[Idx].Start[thread_] =
+            omp_get_wtime();
+        omptarget_device_time_profile.Records[Idx].Thread[thread_] = thread_;
+      }
+    }
+    return Idx;
+  } else {
+    return OMP_ARRAY_SIZE;
+  }
+}
+
+__attribute__((used)) void clockOut(const char *FuncName, int Idx) {
+  if (Idx < OMP_ARRAY_SIZE) {
+#pragma omp parallel
+    {
+      int thread_ = omp_get_thread_num();
+      if (thread_ < OMP_THREADS) {
+        omptarget_device_time_profile.Records[Idx].Stop[thread_] =
+            omp_get_wtime();
+      }
+    }
+  }
+}
+} // namespace prof
+} // namespace _OMP
+
 static void inititializeRuntime(bool IsSPMD) {
   // Order is important here.
   synchronize::init(IsSPMD);
   mapping::init(IsSPMD);
   state::init(IsSPMD);
 }
 
 /// Simple generic state machine for worker threads.
if (IsSPMD) {
     inititializeRuntime(/* IsSPMD */ true);
     synchronize::threadsAligned();
   } else {
     inititializeRuntime(/* IsSPMD */ false);
     // No need to wait since only the main threads will execute user
     // code and workers will run into a barrier right away.
   }
 
+  CLOCKIN();
   if (IsSPMD) {
     state::assumeInitialState(IsSPMD);
     return -1;
   }
 
   if (mapping::isInitialThreadInLevel0(IsSPMD))
     return -1;
 

 ///
 /// In non-SPMD, this function releases the workers trapped in a state machine
 /// and also any memory dynamically allocated by the runtime.
 ///
 /// \param Ident Source location identification, can be NULL.
 ///
 void __kmpc_target_deinit(IdentTy *Ident, int8_t Mode, bool) {
   FunctionTracingRAII();
+  CLOCKOUT();
   const bool IsSPMD = Mode & OMP_TGT_EXEC_MODE_SPMD;
   state::assumeInitialState(IsSPMD);
   if (IsSPMD)
     return;
 
   // Signal the workers to exit the state machine and exit the kernel.
   state::ParallelRegionFn = nullptr;
 }
 
 int8_t __kmpc_is_spmd_exec_mode() {
   FunctionTracingRAII();
   return mapping::isSPMDMode();
 }
 }
 
 #pragma omp end declare target

openmp/libomptarget/DeviceRTL/src/State.cpp

 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_num_target_timeprofile_records() {
+  return targetprof::getNumRecords();
+}
+
+_OMP::prof::Record *omp_get_target_timeprofile_records() {
+  return targetprof::getRecords();
+}
+
 int omp_get_initial_device(void) { return -1; }
 }
 
 extern "C" {
 __attribute__((noinline)) void *__kmpc_alloc_shared(uint64_t Bytes) {
   FunctionTracingRAII();
   return memory::allocShared(Bytes, "Frontend alloc shared");
 }

openmp/libomptarget/include/device.h

struct DeviceTy {
   DeviceTy(const DeviceTy &D) = delete;
   DeviceTy &operator=(const DeviceTy &D) = delete;
 
   ~DeviceTy();
 
   // Return true if data can be copied to DstDevice directly
   bool isDataExchangable(const DeviceTy &DstDevice);
 
+  int32_t getTimeProfilePtr(void **ptr);
+
   /// Lookup the mapping of \p HstPtrBegin in \p HDTTMap. The accessor ensures
   /// exclusive access to the HDTT map.
   LookupResult lookupMapping(HDTTMapAccessorTy &HDTTMap, void *HstPtrBegin,
                              int64_t Size);
 
   /// Get the target pointer based on host pointer begin and base. If the
   /// mapping already exists, the target pointer will be returned directly. In
   /// addition, if required, the memory region pointed by \p HstPtrBegin of size

openmp/libomptarget/include/omptargetplugin.h

 int32_t __tgt_rtl_sync_event(int32_t ID, void *Event);
 
 int32_t __tgt_rtl_destroy_event(int32_t ID, void *Event);
 // }
 
 int32_t __tgt_rtl_init_async_info(int32_t ID, __tgt_async_info **AsyncInfoPtr);
 int32_t __tgt_rtl_init_device_info(int32_t ID, __tgt_device_info *DeviceInfoPtr,
                                    const char **ErrStr);
+int32_t __tgt_rtl_get_timeprofile_data(int32_t DeviceId, void **DataPtr);
 
 #ifdef __cplusplus
 }
 #endif
 
 #endif // _OMPTARGETPLUGIN_H_

openmp/libomptarget/include/rtl.h

struct RTLInfoTy {
   typedef int32_t (*register_lib_ty)(__tgt_bin_desc *);
   typedef int32_t(supports_empty_images_ty)();
   typedef void(print_device_info_ty)(int32_t);
   typedef void(set_info_flag_ty)(uint32_t);
   typedef int32_t(create_event_ty)(int32_t, void **);
   typedef int32_t(record_event_ty)(int32_t, void *, __tgt_async_info *);
   typedef int32_t(wait_event_ty)(int32_t, void *, __tgt_async_info *);
   typedef int32_t(sync_event_ty)(int32_t, void *);
+  typedef int32_t(time_profile_ptr_ty)(int32_t, void **);
   typedef int32_t(destroy_event_ty)(int32_t, void *);
   typedef int32_t(release_async_info_ty)(int32_t, __tgt_async_info *);
   typedef int32_t(init_async_info_ty)(int32_t, __tgt_async_info **);
   typedef int64_t(init_device_into_ty)(int64_t, __tgt_device_info *,
                                        const char **);
 
   int32_t Idx = -1;             // RTL index, index is the number of devices
                                 // of other RTLs that were registered before,
#endif
   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;
   print_device_info_ty *print_device_info = nullptr;
+  time_profile_ptr_ty *time_profile_ptr = nullptr;
   create_event_ty *create_event = nullptr;
   record_event_ty *record_event = nullptr;
   wait_event_ty *wait_event = nullptr;
   sync_event_ty *sync_event = nullptr;
   destroy_event_ty *destroy_event = nullptr;
   init_async_info_ty *init_async_info = nullptr;
   init_device_into_ty *init_device_info = nullptr;
   release_async_info_ty *release_async_info = nullptr;

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

 #include "ELFSymbols.h"
 #include "impl_runtime.h"
 #include "interop_hsa.h"
 
 #include "internal.h"
 #include "rt.h"
 
 #include "DeviceEnvironment.h"
+#include "DeviceTimeProfile.h"
 #include "get_elf_mach_gfx_name.h"
 #include "omptargetplugin.h"
 #include "print_tracing.h"
 
 using namespace llvm;
 using namespace llvm::object;
 using namespace llvm::ELF;
 
if (Valid) {
         return DeviceInfo().freesignalpoolMemcpyH2D(StatePtr, &HostDeviceEnv,
                                                     StatePtrSize, DeviceId);
       }
     }
     return HSA_STATUS_SUCCESS;
   }
 };
 
+struct DeviceTimeProfile {
+  // initialise an DeviceTimeProfileTy in the deviceRTL
+  // patches around differences in the deviceRTL between trunk, aomp,
+  // rocmcc. Over time these differences will tend to zero and this class
+  // simplified.
+  // Symbol may be in .data or .bss, and may be missing fields, todo:
+  // review aomp/trunk/rocm and simplify the following
+
+  // The symbol may also have been deadstripped because the device side
+  // accessors were unused.
+
+  // If the symbol is in .data (aomp, rocm) it can be written directly.
+  // If it is in .bss, we must wait for it to be allocated space on the
+  // gpu (trunk) and initialize after loading.
+  const char *sym() { return "omptarget_device_time_profile"; }
+
+  DeviceTimeProfileTy HostDeviceTimeProfile;
+  SymbolInfo SI;
+  bool Valid = false;
+
+  __tgt_device_image *Image;
+  const size_t ImgSize;
+
+  DeviceTimeProfile(int DeviceId, _OMP::prof::Record records_[OMP_ARRAY_SIZE],
+                    int startc_, int idx_, __tgt_device_image *Image,
+                    const size_t ImgSize)
+      : Image(Image), ImgSize(ImgSize) {
+
+    HostDeviceTimeProfile.StartC = startc_;
+    HostDeviceTimeProfile.Idx = idx_;
+    HostDeviceTimeProfile.DeviceId = DeviceId;
+    for (int i = 0; i < startc_; i++) {
+      for (int j = 0; j < OMP_THREADS; j++) {
+        HostDeviceTimeProfile.Records[i].Thread[j] = records_[i].Thread[j];
+        // HostDeviceTimeProfile.Records[i].Event[j] = records_[i].Event[j];
+        HostDeviceTimeProfile.Records[i].Stop[j] = records_[i].Stop[j];
+        HostDeviceTimeProfile.Records[i].Start[j] = records_[i].Start[j];
+      }
+    }
+    /*
+        if (char *EnvStr = getenv("LIBOMPTARGET_DEVICE_RTL_DEBUG"))
+          HostDeviceTimeProfile.DebugKind = std::stoi(EnvStr);
+     */
+    int Rc = getSymbolInfoWithoutLoading((char *)Image->ImageStart, ImgSize,
+                                         sym(), &SI);
+    if (Rc != 0) {
+      DP("Finding global device time profile '%s' - symbol missing.\n", sym());
+      return;
+    }
+
+    if (SI.Size > sizeof(HostDeviceTimeProfile)) {
+      DP("Symbol '%s' has size %u, expected at most %zu.\n", sym(), SI.Size,
+         sizeof(HostDeviceTimeProfile));
+      return;
+    }
+
+    Valid = true;
+  }
+
+  bool inImage() { return SI.ShType != SHT_NOBITS; }
+
+  hsa_status_t beforeLoading(void *Data, size_t Size) {
+    if (Valid) {
+      if (inImage()) {
+        DP("Setting global device environment before load (%u bytes)\n",
+           SI.Size);
+        uint64_t Offset = (char *)SI.Addr - (char *)Image->ImageStart;
+        void *Pos = (char *)Data + Offset;
+        memcpy(Pos, &HostDeviceTimeProfile, SI.Size);
+      }
+    }
+    return HSA_STATUS_SUCCESS;
+  }
+
+  hsa_status_t afterLoading() {
+    if (Valid) {
+      if (!inImage()) {
+        DP("Setting global device environment after load (%u bytes)\n",
+           SI.Size);
+        int DeviceId = HostDeviceTimeProfile.DeviceId;
+        auto &SymbolInfo = DeviceInfo().SymbolInfoTable[DeviceId];
+        void *StatePtr;
+        uint32_t StatePtrSize;
+        hsa_status_t Err = interop_hsa_get_symbol_info(
+            SymbolInfo, DeviceId, sym(), &StatePtr, &StatePtrSize);
+        if (Err != HSA_STATUS_SUCCESS) {
+          DP("failed to find %s in loaded image\n", sym());
+          return Err;
+        }
+
+        if (StatePtrSize != SI.Size) {
+          DP("Symbol had size %u before loading, %u after\n", StatePtrSize,
+             SI.Size);
+          return HSA_STATUS_ERROR;
+        }
+
+        return DeviceInfo().freesignalpoolMemcpyH2D(
+            StatePtr, &HostDeviceTimeProfile, StatePtrSize, DeviceId);
+      }
+    }
+    return HSA_STATUS_SUCCESS;
+  }
+};
+
 hsa_status_t implCalloc(void **RetPtr, size_t Size, int DeviceId) {
   uint64_t Rounded = 4 * ((Size + 3) / 4);
   void *Ptr;
   hsa_amd_memory_pool_t MemoryPool = DeviceInfo().getDeviceMemoryPool(DeviceId);
   hsa_status_t Err = hsa_amd_memory_pool_allocate(MemoryPool, Rounded, 0, &Ptr);
   if (Err != HSA_STATUS_SUCCESS) {
     return Err;
   }

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

 #include <list>
 #include <memory>
 #include <mutex>
 #include <string>
 #include <vector>
 
 #include "Debug.h"
 #include "DeviceEnvironment.h"
+#include "DeviceTimeProfile.h"
 #include "omptarget.h"
 #include "omptargetplugin.h"
 
 #define TARGET_NAME CUDA
 #define DEBUG_PREFIX "Target " GETNAME(TARGET_NAME) " RTL"
 
 #include "MemoryManager.h"
 
 #include "llvm/Frontend/OpenMP/OMPConstants.h"
 
+DeviceTimeProfileTy DeviceTimeProf;
+CUdeviceptr DeviceTimeProfPtr;
+
 using namespace llvm;
 
 // Utility for retrieving and printing CUDA error string.
 #ifdef OMPTARGET_DEBUG
 #define CUDA_ERR_STRING(err)                                                   \
   do {                                                                         \
     if (getDebugLevel() > 0) {                                                 \
       const char *errStr = nullptr;                                            \
DP("Error when copying data from device to device. Pointers: src "
        DPxPTR(SrcPtr), DPxPTR(DstPtr), Size);
     CUDA_ERR_STRING(Err);
     return OFFLOAD_FAIL;
   }
 
   return OFFLOAD_SUCCESS;
 }
 
+int32_t getTimeProfilePtr(void **Ptr) {
+  *Ptr = (void *)DeviceTimeProfPtr;
+  return CUDA_SUCCESS;
+}
+
 int recordEvent(void *EventPtr, __tgt_async_info *AsyncInfo) {
   CUstream Stream = reinterpret_cast<CUstream>(AsyncInfo->Queue);
   CUevent Event = reinterpret_cast<CUevent>(EventPtr);
 
   CUresult Err = cuEventRecord(Event, Stream);
   if (Err != CUDA_SUCCESS) {
     DP("Error when recording event. stream = " DPxMOD ", event = " DPxMOD "\n",
        DPxPTR(Stream), DPxPTR(Event));
// send device environment data to the device
       } else {
         DP("Finding global device environment '%s' - symbol missing.\n",
            DeviceEnvName);
         DP("Continue, considering this is a device RTL which does not accept "
            "environment setting.\n");
       }
     }
 
+    {
+      // TODO: The device ID used here is not the real device ID used by OpenMP.
+
+      const char *DeviceTimeProfName = "omptarget_device_time_profile";
+      size_t CUSize;
+
+      Err = cuModuleGetGlobal(&DeviceTimeProfPtr, &CUSize, Module,
+                              DeviceTimeProfName);
+      if (Err == CUDA_SUCCESS) {
+        if (CUSize != sizeof(DeviceTimeProfileTy)) {
+          REPORT(
+              "Global device_time_profile '%s' - size mismatch (%zu != %zu)\n",
+              DeviceTimeProfName, CUSize, sizeof(DeviceTimeProfileTy));
+          CUDA_ERR_STRING(Err);
+          return nullptr;
+        }
+
+        DP("Sending global device time profile data %zu bytes\n", CUSize);
+      } else {
+        DP("Finding global device time profile '%s' - symbol missing.\n",
+           DeviceTimeProfName);
+        DP("Continue, considering this is a device RTL which does not accept "
+           "environment setting.\n");
+      }
+    }
+
     return getOffloadEntriesTable(DeviceId);
   }
 
   void *dataAlloc(const int DeviceId, const int64_t Size,
                   const TargetAllocTy Kind) {
     switch (Kind) {
     case TARGET_ALLOC_DEFAULT:
     case TARGET_ALLOC_DEVICE:
int32_t __tgt_rtl_init_device_info(int32_t DeviceId,
   assert(DeviceInfoPtr && "device_info_ptr is nullptr");
 
   if (DeviceRTL.setContext(DeviceId) != OFFLOAD_SUCCESS)
     return OFFLOAD_FAIL;
 
   return DeviceRTL.initDeviceInfo(DeviceId, DeviceInfoPtr, ErrStr);
 }
 
+int32_t __tgt_rtl_get_timeprofile_data(int32_t DeviceId, void **DataPtr) {
+  // assert(!data_ptr && "data_ptr is not nullptr");
+  return getTimeProfilePtr(DataPtr);
+}
+
 #ifdef __cplusplus
 }
 #endif

openmp/libomptarget/src/device.cpp

 //===--------- device.cpp - Target independent OpenMP target RTL ----------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // Functionality for managing devices that are handled by RTL plugins.
 //
 //===----------------------------------------------------------------------===//
 
 #include "device.h"
+#include "DeviceTimeProfile.h"
 #include "omptarget.h"
 #include "private.h"
 #include "rtl.h"
+#include "targetprofile.h"
 
+#include <llvm/Support/JSON.h>
+#include <llvm/Support/TimeProfiler.h>
+
+#include <bits/chrono.h>
 #include <cassert>
+#include <chrono>
 #include <climits>
+#include <cstddef>
 #include <cstdint>
 #include <cstdio>
+#include <ratio>
 #include <string>
 #include <thread>
 
+CLOCKDEC();
+
+void makeJsonRecordEntry(llvm::raw_fd_ostream &OS,
+                         _OMP::prof::Record recs[OMP_ARRAY_SIZE],
+                         int total_entries) {
+  // SecToUsFactor * time to convert seconds to microseconds as mentioned by the
+  // chrome trace format
+  const int32_t SecToUsFactor = 1000000;
+  llvm::json::OStream J(OS);
+  J.objectBegin();
+  J.attributeBegin("traceEvents");
+  J.arrayBegin();
+
+  auto writeEvent = [&](_OMP::prof::ThreadRecord rec) {
+    J.object([&] {
+      J.attribute("pid", -1);
+      J.attribute("tid", int64_t(rec.Thread));
+      J.attribute("ph", "X");
+      J.attribute("ts", rec.Start);
+      J.attribute("dur", rec.Stop - rec.Start);
+      J.attribute("name", rec.Event);
+    });
+  };
+  auto writeMetadataEvent = [&](const char *Name, uint64_t Tid,
+                                llvm::StringRef arg) {
+    J.object([&] {
+      J.attribute("cat", "Target Event");
+      J.attribute("pid", -1);
+      J.attribute("tid", int64_t(Tid));
+      J.attribute("ts", 0);
+      J.attribute("ph", "M");
+      J.attribute("name", Name);
+      J.attributeObject("args", [&] { J.attribute("name", arg); });
+    });
+  };
+
+  for (int i = 0; i < total_entries; i++) {
+    for (int j = 0; j < OMP_THREADS; j++) {
+      _OMP::prof::ThreadRecord t;
+      t.Stop = recs[i].Stop[j] * SecToUsFactor;
+      t.Start = recs[i].Start[j] * SecToUsFactor;
+      t.Thread = recs[i].Thread[j];
+      // sprintf(t.Event, "%d-Target Thread#%d", i, t.Thread);
+      // sprintf(t.Event, "Target Event#%d", t.Thread);
+      sprintf(t.Event, "Target Event");
+      writeEvent(t);
+      writeMetadataEvent("thread_name", t.Thread, t.Event);
+    }
+  }
+
+  writeMetadataEvent("process_name", -1,
+                     "clang-libomptarget-offloading-process");
+  writeMetadataEvent("process_name", -1,
+                     "clang-libomptarget-offloading-thread");
+
+  J.arrayEnd();
+  J.attributeEnd();
+
+  double BeginningOfTime = (recs[0].Start[0] - 100) * SecToUsFactor;
+  J.attribute("beginningOfTime", BeginningOfTime);
+
+  J.objectEnd();
+}
+
+void sideloadTimeProfileEntries(_OMP::prof::Record Recs[OMP_ARRAY_SIZE],
+                                int total_entries) {
+  auto makeRecEntry = [](const _OMP::prof::Record R, size_t Idx) {
+    char TargetDetail[64];
+    sprintf(TargetDetail, "Target Thread#%d", R.Thread[Idx]);
+    const long long int Multiplier = 1000000000;
+    long long int Start_ = R.Start[Idx] * Multiplier;
+    std::chrono::duration<long long int, std::ratio<1, Multiplier>> Start_dur(
+        Start_);
+    long long int Stop_ = R.Stop[Idx] * Multiplier;
+    std::chrono::duration<long long int, std::ratio<1, Multiplier>> Stop_dur(
+        Stop_);
+
+    llvm::Entry Temp(
+        std::chrono::time_point<std::chrono::steady_clock>(Start_dur),
+        std::chrono::time_point<std::chrono::steady_clock>(Stop_dur),
+        std::string("Target Op"), std::string(TargetDetail));
+    return Temp;
+  };
+
+  bool TPINotInit = false;
+  llvm::TimeTraceProfiler *TPI = llvm::getTimeTraceProfilerInstance();
+  if (TPI == nullptr) {
+    llvm::TimeTraceProfiler TPV =
+        llvm::TimeTraceProfiler(0, "Target Operations");
+    TPI = &TPV;
+    TPINotInit = true;
+  }
+  for (size_t i = 0; i < total_entries; i++) {
+    for (size_t j = 0; j < OMP_THREADS; j++) {
+      TPI->Entries.emplace_back(makeRecEntry(Recs[i], j));
+    }
+  }
+  if (TPINotInit)
+    llvm::Instances.List.emplace_back(TPI);
+}
+
+void clockFinalize(DeviceTy *Device) {
+  int32_t Idx = omptarget_device_time_profile.Idx;
+  assert(Idx >= 0 && Idx < OMP_ARRAY_SIZE && "Out of beyond bounds");
+  // TODO: use env var
+  if (Idx > 0) {
+    llvm::StringRef out("Target-profile-only.json");
+    std::error_code EC;
+    llvm::raw_fd_ostream OS(out, EC);
+    llvm::json::OStream J(OS);
+    sideloadTimeProfileEntries(omptarget_device_time_profile.Records, Idx);
+    makeJsonRecordEntry(OS, omptarget_device_time_profile.Records, Idx);
+    OS.close();
+  } else
+    printf("No Target\n");
+}
+
 int HostDataToTargetTy::addEventIfNecessary(DeviceTy &Device,
                                             AsyncInfoTy &AsyncInfo) const {
   // First, check if the user disabled atomic map transfer/malloc/dealloc.
   if (!PM->UseEventsForAtomicTransfers)
     return OFFLOAD_SUCCESS;
 
   void *Event = getEvent();
   bool NeedNewEvent = Event == nullptr;

 }
 
 DeviceTy::DeviceTy(RTLInfoTy *RTL)
     : DeviceID(-1), RTL(RTL), RTLDeviceID(-1), IsInit(false), InitFlag(),
       HasPendingGlobals(false), PendingCtorsDtors(), ShadowPtrMap(),
       PendingGlobalsMtx(), ShadowMtx() {}
 
 DeviceTy::~DeviceTy() {
+  CLOCKFINALIZE();
   if (DeviceID == -1 || !(getInfoLevel() & OMP_INFOTYPE_DUMP_TABLE))
     return;
 
   ident_t Loc = {0, 0, 0, 0, ";libomptarget;libomptarget;0;0;;"};
   dumpTargetPointerMappings(&Loc, *this);
 }
 
 int DeviceTy::associatePtr(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size) {
if (Ret != OFFLOAD_SUCCESS)
     return;
 
   IsInit = true;
 }
 
 /// Thread-safe method to initialize the device only once.
 int32_t DeviceTy::initOnce() {
   std::call_once(InitFlag, &DeviceTy::init, this);
+  CLOCKINIT();
 
   // At this point, if IsInit is true, then either this thread or some other
   // thread in the past successfully initialized the device, so we can return
   // OFFLOAD_SUCCESS. If this thread executed init() via call_once() and it
   // failed, return OFFLOAD_FAIL. If call_once did not invoke init(), it means
   // that some other thread already attempted to execute init() and if IsInit
   // is still false, return OFFLOAD_FAIL.
   if (IsInit)
int32_t DeviceTy::submitData(void *TgtPtrBegin, void *HstPtrBegin, int64_t Size,
   }
 
   if (!AsyncInfo || !RTL->data_submit_async || !RTL->synchronize)
     return RTL->data_submit(RTLDeviceID, TgtPtrBegin, HstPtrBegin, Size);
   return RTL->data_submit_async(RTLDeviceID, TgtPtrBegin, HstPtrBegin, Size,
                                 AsyncInfo);
 }
 
+int32_t DeviceTy::getTimeProfilePtr(void **ptr) {
+  if (RTL->time_profile_ptr)
+    return RTL->time_profile_ptr(RTLDeviceID, ptr);
+  return OFFLOAD_SUCCESS;
+}
+
 // Retrieve data from device
 int32_t DeviceTy::retrieveData(void *HstPtrBegin, void *TgtPtrBegin,
                                int64_t Size, AsyncInfoTy &AsyncInfo) {
   if (getInfoLevel() & OMP_INFOTYPE_DATA_TRANSFER) {
     HDTTMapAccessorTy HDTTMap = HostDataToTargetMap.getExclusiveAccessor();
     LookupResult LR = lookupMapping(HDTTMap, HstPtrBegin, Size);
     auto *HT = &*LR.Entry;
     INFO(OMP_INFOTYPE_DATA_TRANSFER, DeviceID,

openmp/libomptarget/src/interface.cpp

 //===-------- interface.cpp - Target independent OpenMP target RTL --------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // Implementation of the interface to be used by Clang during the codegen of a
 // target region.
 //
 //===----------------------------------------------------------------------===//
 
+#include "DeviceTimeProfile.h"
 #include "device.h"
 #include "omptarget.h"
 #include "private.h"
 #include "rtl.h"
 
 #include <cassert>
 #include <cstdio>
 #include <cstdlib>
 #include <mutex>
 
+extern DeviceTimeProfileTy omptarget_device_time_profile;
+
 ////////////////////////////////////////////////////////////////////////////////
 /// adds requires flags
 EXTERN void __tgt_register_requires(int64_t Flags) {
   TIMESCOPE();
   PM->RTLs.registerRequires(Flags);
 }
 
 ////////////////////////////////////////////////////////////////////////////////
EXTERN void __tgt_target_data_begin_nowait_mapper(
     void **ArgMappers, int32_t DepNum, void *DepList, int32_t NoAliasDepNum,
     void *NoAliasDepList) {
   TIMESCOPE_WITH_IDENT(Loc);
 
   __tgt_target_data_begin_mapper(Loc, DeviceId, ArgNum, ArgsBase, Args,
                                  ArgSizes, ArgTypes, ArgNames, ArgMappers);
 }
 
+// At device deconstructor context gets destroyed. To work around this the
+// collected data is synced target -> host after every kernel deinit.
+void clockSync(DeviceTy *Device) {
+  AsyncInfoTy async_info(*Device);
+  void *tgt_ptr = nullptr;
+  void *hst_ptr = &omptarget_device_time_profile;
+  Device->getTimeProfilePtr(&tgt_ptr);
+  Device->retrieveData(hst_ptr, tgt_ptr, sizeof(DeviceTimeProfileTy),
+                       async_info);
+  int32_t Idx = omptarget_device_time_profile.Idx;
+  assert(Idx > 0 && Idx < OMP_ARRAY_SIZE && "Out of beyond bounds");
+}
+
 /// passes data from the target, releases target memory and destroys
 /// the host-target mapping (top entry from the stack of data maps)
 /// created by the last __tgt_target_data_begin.
 EXTERN void __tgt_target_data_end_mapper(ident_t *Loc, int64_t DeviceId,
                                          int32_t ArgNum, void **ArgsBase,
                                          void **Args, int64_t *ArgSizes,
                                          int64_t *ArgTypes,
                                          map_var_info_t *ArgNames,
for (int I = 0; I < ArgNum; ++I) {
     DP("Entry %2d: Base=" DPxMOD ", Begin=" DPxMOD ", Size=%" PRId64
        ", Type=0x%" PRIx64 ", Name=%s\n",
        I, DPxPTR(ArgsBase[I]), DPxPTR(Args[I]), ArgSizes[I], ArgTypes[I],
        (ArgNames) ? getNameFromMapping(ArgNames[I]).c_str() : "unknown");
   }
 #endif
 
   AsyncInfoTy AsyncInfo(Device);
+  CLOCKSYNC();
   int Rc = targetDataEnd(Loc, Device, ArgNum, ArgsBase, Args, ArgSizes,
                          ArgTypes, ArgNames, ArgMappers, AsyncInfo);
   if (Rc == OFFLOAD_SUCCESS)
     Rc = AsyncInfo.synchronize();
   handleTargetOutcome(Rc == OFFLOAD_SUCCESS, Loc);
 }
 
 EXTERN void __tgt_target_data_end_nowait_mapper(

openmp/libomptarget/src/rtl.cpp

#endif
   *((void **)&RTL.destroy_event) =
       DynLibrary->getAddressOfSymbol("__tgt_rtl_destroy_event");
   *((void **)&RTL.release_async_info) =
       DynLibrary->getAddressOfSymbol("__tgt_rtl_release_async_info");
   *((void **)&RTL.init_async_info) =
       DynLibrary->getAddressOfSymbol("__tgt_rtl_init_async_info");
   *((void **)&RTL.init_device_info) =
       DynLibrary->getAddressOfSymbol("__tgt_rtl_init_device_info");
+  *((void **)&RTL.time_profile_ptr) =
+      DynLibrary->getAddressOfSymbol("__tgt_rtl_get_timeprofile_data");
 
   RTL.LibraryHandler = std::move(DynLibrary);
 
   // Successfully loaded
   return true;
 }
 
 ////////////////////////////////////////////////////////////////////////////////