[libomptarget] [amdgpu] Foundation for OMPT target callback support

[dhruvachak]

Summary

OMPT callback support for target task, target data op, and target submit

Patch from John Mellor-Crummey <[e-mail]>
(With contributions from Dhruva Chakrabarti <[e-mail]>)

o Connect up libomp, libomptarget and AMDGPU device rtl
o Provide a callback class used to interact with a tool
o Update omp-tool.h.var from LLVM main to ensure it is fresh
o Update libomp sources for two purposes

1. Adapt to new macros in omp-tools.h
2. Add API used by libomptarget to extract task_data and target_task_data when needed to support OMPT EMI calls

o Integrate support for EMI and NONEMI target callbacks
o Target, data operation, and target submit callbacks added

[protze.joachim]

I mainly checked the libomp code and put some inline comments.

openmp/libomptarget/CMakeLists.txt

72–95

This code does not reproduce the libomp control flow for in-llvm-tree and standalone build.

To use cmake variables from openmp/runtime, just make them visible in parent scope:

set(LIBOMP_HAVE_OMPT_SUPPORT ${LIBOMP_HAVE_OMPT_SUPPORT} PARENT_SCOPE)

see https://github.com/llvm/llvm-project/commit/3bc8ce5dd718beef0031bf4b070ac4026e6910d7#diff-01127988bc2b1599723b049e94ada3025063f4e0a7bac326aaf37b892a118fd1R332

openmp/runtime/src/ompt-general.cpp

504–509

This change seems unrelated and it is not obvious whether this change is correct.

[protze.joachim]

openmp/libomptarget/include/ompt_device_callbacks.h

64

Why are the functions virtual?

openmp/runtime/src/ompt-internal.h

16–37

These changes seem unrelated.

121–137

These are unrelated changes.

130

From __builtin_frame_address reference:

Calling this function with a nonzero argument can have unpredictable effects, including crashing the calling program. As a result, calls that are considered unsafe are diagnosed when the -Wframe-address option is in effect. Such calls should only be made in debugging situations.

How should we deal with this restriction?

openmp/runtime/src/ompt-specific.cpp

359

@AndreyChurbanov what do you think about the use of thread_local in libomp?

If a thread-local value is sufficient (which I doubt from OMPT perspective), this value could be stored in a new entry of ompt_thread_info_t:
__kmp_threads[__kmp_get_gtid()]->th.ompt_thread_info.target_data

openmp/runtime/src/ompt-specific.h

87–113

These are unrelated changes.

[dhruvachak]

Thanks, Joachim, for your comments. I removed the unrelated changes and made the cmake change in the next version.

openmp/libomptarget/CMakeLists.txt

72–95

Added the parent scope.

openmp/libomptarget/include/ompt_device_callbacks.h

64

I think the intent was to have plugin specific derived classes but that's not the case at the moment. So removed virtual.

openmp/runtime/src/ompt-internal.h

16–37

OMPT_FRAME_SET is called by ompt_set_frame_enter_internal which is part of target region support. ompt_set_frame_enter_internal is called by ompt_set_frame_enter which is used by ompt_callback.

130

Since it's unrelated, I removed it.

openmp/runtime/src/ompt-specific.cpp

359

Made the suggested change.

[dhruvachak]

Addressed feedback: made the suggested cmake change and removed unrelated changes

[protze.joachim]

openmp/runtime/src/ompt-general.cpp

512–520

I think, we need to reorder the conditions to make sure that target finalize is called even if the tool does not register a finalize callback. We should check the callback pointer in any case, even without OMPD support.

[dhruvachak]

openmp/runtime/src/ompt-general.cpp

512–520

Made the change in the next version. I think we should check for NULL-ness unconditionally, so removed the OMPD_SUPPORT ifdef.

[dhruvachak]

Addressed feedback: libomptarget finalize should not depend on OMPD

[JonChesterfield]

Skimmed through this patch. Lots of code that looks like it was copy + paste + modify which is usually a bug farm. Quite a lot of global variables. No associated tests.

Lack of tests seems the most significant problem here.

openmp/libomptarget/include/ompt-connector.h

54

I don't understand this note

78

error handling missing here

openmp/libomptarget/include/omptarget.h

356–357

I thought these functions were ABI stable, i.e. we can't change them

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

119

initialization order hazard here. When is this supposed to be constructed relative to the other globals?

openmp/libomptarget/src/ompt_callback.cpp

43

I can't remember if private inheritance of std containers is well defined or not. This class doesn't seem to use any of the members of std::list. What's going on here?

92

Why are these all global variables?

243

Can we factor this so copy&paste typos are less likely?

408

This sort of thing makes the library awkward to construct safely. I guess 102 is related to the other function tagged constructor, but we're better off exposing a single library_init function which is called by libomptarget at the appropriate point in time

openmp/runtime/src/ompt-specific.cpp

269

?

[JonChesterfield]

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

43

This is an unguarded if statement, will interact badly with other if statements next to it in code. Macros like this are traditionally wrapped in do { } while (0) to avoid that syntax hazard

[jmellorcrummey]

openmp/libomptarget/include/ompt-connector.h

54

This class is used to communicate between an OMPT implementation in libomptarget and libomp. It is also used to communicate between an OMPT implementation in a device-specific plugin. The decision whether OMPT is enabled or not needs to be made when the library is loaded before any functions in the library are invoked. For that reason, this class is intended to be declared in the constructor for libomptarget or a plugin so that the decision about whether OMPT is supposed to be enabled is known before any interface function in the library is invoked.

78

If the dlsym doesn't return a non-zero value, then the OMPT interface is not enabled. The check for a zero value is in the connect method. If library_ompt_connect is NULL because dlsym returned NULL, then the code in the client library (the one calling connect) won't try to call this function to connect with its server library.

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

119

The only requirement here is that this initialization must occur in a constructor so that libomp can use the pointer set up in lines 123 to initialize OMPT support when the call on line 130 is made so that OMPT support is either initialized or not before any entry point to libomptarget is invoked.

openmp/libomptarget/src/ompt_callback.cpp

92

These are global (but thread local) variables because they hold persistent state used by a thread operating in a target region.
When you enter a target region these variables are initialized. The operation id is adjusted every call to a target region operation. ompt_task_data and ompt_target_task_data are set when the region is entered and persist throughout a target region being performed by a thread.

Note that these variables are outside the ompt_interface class because they are manipulated by functions passed to methods such as ompt_device_callbacks.ompt_callback_target_data_op_emi.

243

The commonality is between begin/end pairs. What would you propose? Factoring out an EMI call into a function F that can be called with a few less arguments so that global variables like ompt_target_task_data, ompt_target_data don't need to be passed to be passed to F? I don't think this adds much simplification.

408

While this could be a global function called at the end of the other constructor in rtl.cpp, this would be the only reason any OMPT header has to be included in rtl.cpp.

The other option is to declare all of the priorities passed to constructors in one place and make the number for this constructor the successor to the other constructor, e.g.

#define RTL_CONSTRUCTOR_PRIORITY 101
#define OMPT_CONSTRUCTOR_PRIORITY (RTL_CONSTRUCTOR_PRIORITY + 1)

Using these defines means that no OMPT header file needs to be included in rtl.cpp. Or perhaps you'd prefer a constructor in its own file that calls the "constructor" functions in rtl.cpp and ompt_callback.cpp and the constructor can include headers for rtl.cpp and ompt_callback.cpp that define the functions that need to be called by the single constructor.

openmp/runtime/src/ompt-specific.cpp

269

I'm not sure what the question is.

The enter_frame_flags and exit_frame_flags used by the runtime entry points for inspecting frame objects provide information about how to locate the frame pointer of a procedure frame that called into or out of the runtime.

When setting up a lightweight task, these flags are cleared.

They will be set later to a logical or

either ompt_frame_runtime or ompt_frame_application

ored together with either ompt_frame_framepointer or ompt_frame_cfa

[dhruvachak]

openmp/libomptarget/include/ompt-connector.h

78

I deleted the default constructor in the next version so as to prevent its accidental use in new code. This will ensure proper initialization, no other error handling required here.

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

43

Fixed in the next version.

openmp/libomptarget/src/ompt_callback.cpp

43

Removed in the next version.

[dhruvachak]

Addressed feedback

Updated implementation status of target callbacks
Removed unused static var and unused function
Added do-while to macros
Removed unused inheritance from std::list
Deleted the default constructor to ensure proper initialization

[JonChesterfield]

I've added a bunch of reviewers from git shortlog -sn openmp to get more eyes on this as it's hard to guess who will be most interested.

I think a reasonable path forward here is to check the ompt hooks don't upset the existing machinery and then iterate in tree. @dhruvachak is committing this as-is to the amd internal rocm branch to get some data on that.

[dhruvachak]

Added new tests for OMPT target callbacks for amdgpu

[dhruvachak]

Guard ompt_init constructors with ifdef OMPT_SUPPORT, added comments to tests

[jdoerfert]

I commented on the most severe issues I could find. Since I don't know enough about the OMPT inner workings it is mostly style and general code comments.
That said, the lack of (doxygen) documentation is not a great sign. There is not much being explained in the code either. I fear we are piling up more and more code less and less people understand without reasonable means for new people to get into it. For now, I'm fine with this but @dhruvachak you should consider spending some time soon documenting the methods and non-trivial code pieces.

openmp/libomptarget/include/ompt-connector.h

38

This is not something we should expose in global namespace especially as it is not used here and has a rather generic name.

56

Nit: Try to make comments proper sentences incl. punctuation and such.

93

I would have used the LLVM coding convention in these new parts, it's a missed opportunity IMHO.

Early exists should be preferred over indention, e.g., if (is_initialized) return;

We should at least remove ; after methods (here and elsewhere).

Clang format all the code.

openmp/libomptarget/include/ompt_device_callbacks.h

3

Nit: Style

openmp/libomptarget/include/omptarget.h

356–357

What Jon said. Also, this method is effectively deprecated anyway, the one below is used even though we should not explicitly push the trip count in the first place.

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

128

These stray ; appear a few times. Will trigger warnings and should be removed.

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

1855

Clang format these things, consider OMPT_IF_ENABLED({ ... }).
Also all the others that span more than one line/statement.

openmp/libomptarget/src/device.cpp

502

These functions use LLVM coding style, please don't mix it with the C style other places use.

The pattern here is repetitive and I am unsure if it would not make it much easier to use a RAII object instead.

Comments apply to the rest of the file too.

openmp/libomptarget/src/ompt_callback.cpp

59

assertions should have messages, use nullptr for a pointer not 0.

121

what are those 0s? Use /* enter_frame */ nullptr, ...

381

Please no assert(0). In release builds this will silently do something while it arguably should not.
print the error, then explicitly abort, exit, trap, ...

423

technically this is not a reserved name. I doubt it will clash with the user names but it could for all we know. Any chance we can make it a reserved name instead? ompx_ omp_ (ompt_?) __libomptarget ...

openmp/libomptarget/src/ompt_callback.h

106

Yet another naming scheme?

openmp/libomptarget/src/omptarget.cpp

33

This seems something worth refactoring, third time I see it.

openmp/libomptarget/test/ompt/callbacks.h

72

Please no assert(0). In release builds this will silently do something while it arguably should not.
print the error, then explicitly abort, exit, trap, ...

126

Please no assert(0). In release builds this will silently do something while it arguably should not.
print the error, then explicitly abort, exit, trap, ...

openmp/runtime/src/ompt-internal.h

98

Is this something we need to keep stable?

openmp/runtime/src/ompt-specific.cpp

271

Here and above, stray ; again. Will cause warnings (=problems) eventually.

[dhruvachak]

Addressed feedback.

[dhruvachak]

I addressed most of the feedback and updated this review. I will work on splitting up this patch into multiple pieces so that they can be reviewed more easily.

openmp/libomptarget/include/ompt_device_callbacks.h

3

clang-format makes this change and I left it alone.

[jdoerfert]

openmp/libomptarget/include/ompt_device_callbacks.h

3

Your line is too long. "Target independent OpenMP target RTL" can be shortened, maybe even replaced with something that explains what this file does. I have no idea what "Target independent OpenMP target RTL" means wrt. OMPT callbacks. My guess is this was copied from somewhere else.

openmp/libomptarget/src/device.cpp

500

For these RAIIs, you could put OMPT_IF_BUILT into the RAII code and make them a no-op if OMPT is not used. That way you don't need any macro here, the RAII would always be there and removed when build w/o OMPT.

536

No need for the extra indention.

562

same as above.

openmp/libomptarget/src/interface.cpp

130

same as above, I'll stop pointing it out now.

[dhruvachak]

openmp/libomptarget/src/interface.cpp

130

Yes, I thought about the additional scope. But I was worried about future new code possibly messing up where exactly I want the RAII object dtor called.

[dhruvachak]

Abandoning this revision in favor of https://reviews.llvm.org/D127372 which is the top of a stack of reviews. Comments from this revision have been incorporated in that stack.

openmp/libomptarget/CMakeLists.txt

if(LIBOMPTARGET_CMAKE_BUILD_TYPE MATCHES debug)
   option(LIBOMPTARGET_ENABLE_DEBUG "Allow debug output with the environment variable LIBOMPTARGET_DEBUG=1" ON)
 else()
   option(LIBOMPTARGET_ENABLE_DEBUG "Allow debug output with the environment variable LIBOMPTARGET_DEBUG=1" OFF)
 endif()
 if(LIBOMPTARGET_ENABLE_DEBUG)
   add_definitions(-DOMPTARGET_DEBUG)
 endif()
 
+# OMPT support for libomptarget
+set(OMPT_TARGET_DEFAULT FALSE)
+if ((LIBOMP_HAVE_OMPT_SUPPORT) AND (NOT WIN32))
+  set (OMPT_TARGET_DEFAULT TRUE)
+endif()
+set(LIBOMPTARGET_OMPT_SUPPORT ${OMPT_TARGET_DEFAULT} CACHE BOOL "OMPT-target-support?")
+if (LIBOMPTARGET_OMPT_SUPPORT)
+  add_definitions(-DOMPT_SUPPORT=1)
+  message(STATUS "OMPT target enabled")
+else()
+  message(STATUS "OMPT target disabled")
+endif()
+
+pythonize_bool(LIBOMPTARGET_OMPT_SUPPORT)
+
 set(LIBOMPTARGET_INCLUDE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/include)
-include_directories(${LIBOMPTARGET_INCLUDE_DIR})
+set(LIBOMP_RUNTIME_SRC_BINARY_DIR ${CMAKE_BINARY_DIR}/openmp/runtime/src)
+include_directories(${LIBOMPTARGET_INCLUDE_DIR} ${LIBOMP_RUNTIME_SRC_BINARY_DIR})
 
 # Build target agnostic offloading library.
 set(LIBOMPTARGET_SRC_DIR ${CMAKE_CURRENT_SOURCE_DIR}/src)
 add_subdirectory(${LIBOMPTARGET_SRC_DIR})
 
 # Definitions for testing, for reuse when testing libomptarget-nvptx.
 set(LIBOMPTARGET_OPENMP_HEADER_FOLDER "${LIBOMP_INCLUDE_DIR}" CACHE STRING
   "Path to folder containing omp.h")
 set(LIBOMPTARGET_OPENMP_HOST_RTL_FOLDER "${LIBOMP_LIBRARY_DIR}" CACHE STRING
   "Path to folder containing libomp.so, and libLLVMSupport.so with profiling enabled")
 
 
 # Build offloading plugins and device RTLs if they are available.
 add_subdirectory(plugins)
 add_subdirectory(DeviceRTL)
 add_subdirectory(tools)
 
 # Add tests.
 add_subdirectory(test)

openmp/libomptarget/include/ompt-connector.h

+//=== ompt-connector.h - Target independent OpenMP target RTL -- 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
+//
+//===----------------------------------------------------------------------===//
+//
+// Support used by OMPT implementation to establish communication between
+// various OpenMP runtime libraries: host openmp library, target-independent
+// runtime library, and device-dependent runtime libraries.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _OMPT_CONNECTOR_H
+#define _OMPT_CONNECTOR_H
+
+//****************************************************************************
+// global includes
+//****************************************************************************
+
+#include <dlfcn.h>
+#include <string>
+
+//****************************************************************************
+// local includes
+//****************************************************************************
+
+#include <Debug.h>
+#include <omp-tools.h>
+#include <omptarget.h>
+
+//****************************************************************************
+// macros
+//****************************************************************************
+
+#define LIBOMPTARGET_STRINGIFY(s) #s
+
+#define LIBOMPTARGET_GET_TARGET_OPID libomptarget_get_target_opid
+
+//****************************************************************************
+// type declarations
+//****************************************************************************
+
+/// Type for the function to be invoked for connecting two libraries.
+typedef void (*OmptConnectRtnTy)(ompt_start_tool_result_t *result);
+
+/// Establish connection between openmp runtime libraries
+///
+/// This class is used to communicate between an OMPT implementation in
+/// libomptarget and libomp. It is also used to communicate between an
+/// OMPT implementation in a device-specific plugin and
+/// libomptarget. The decision whether OMPT is enabled or not needs to
+/// be made when the library is loaded before any functions in the
+/// library are invoked. For that reason, an instance of this class is
+/// intended to be defined in the constructor for libomptarget or a
+/// plugin so that the decision about whether OMPT is supposed to be
+/// enabled is known before any interface function in the library is
+/// invoked.
+class OmptLibraryConnectorTy {
+public:
+  /// Use \p LibName as the prefix of the global function used for connecting
+  /// two libraries, the source indicated by \p LibName and the destination
+  /// being the one that creates this object.
+  OmptLibraryConnectorTy(const char *LibName) {
+    LibConnRtn.append(LibName);
+    LibConnRtn.append("_connect");
+    IsInitialized = false;
+  }
+  OmptLibraryConnectorTy() = delete;
+  /// Use \p OmptResult init to connect the two libraries denoted by this
+  /// object. The init function of \p OmptResult will be used during connection
+  /// and the fini function of \p OmptResult will be used during teardown.
+  void connect(ompt_start_tool_result_t *OmptResult) {
+    initialize();
+    if (!LibConnHandle)
+      return;
+    // Call the function provided by the source library for connect
+    LibConnHandle(OmptResult);
+  }
+
+private:
+  void initialize() {
+    if (IsInitialized)
+      return;
+
+    DP("OMPT: Library connection routine = %s\n", LibConnRtn.c_str());
+
+    void *VPtr = dlsym(NULL, LibConnRtn.c_str());
+    // If dlsym fails, the handle will be null. connect() checks
+    // for this condition
+    LibConnHandle =
+        reinterpret_cast<OmptConnectRtnTy>(reinterpret_cast<long>(VPtr));
+    DP("OMPT: Library connection handle = %p\n", LibConnHandle);
+    IsInitialized = true;
+  }
+
+private:
+  /// Ensure initialization occurs only once
+  bool IsInitialized;
+  /// Handle of connect routine provided by source library
+  OmptConnectRtnTy LibConnHandle;
+  /// Name of connect routine provided by source library
+  std::string LibConnRtn;
+};
+
+#endif

openmp/libomptarget/include/ompt_device_callbacks.h

+//=== ompt_device_callbacks.h - Target independent OpenMP target RTL -- 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
+//
+//===----------------------------------------------------------------------===//
+//
+// Interface used by both target-independent and device-dependent runtimes
+// to coordinate registration and invocation of OMPT callbacks
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _OMPT_DEVICE_CALLBACKS_H
+#define _OMPT_DEVICE_CALLBACKS_H
+
+//****************************************************************************
+// local includes
+//****************************************************************************
+
+#include <atomic>
+#include <string.h>
+
+#include <omp-tools.h>
+
+#define FOREACH_OMPT_TARGET_CALLBACK(macro)                                    \
+  FOREACH_OMPT_DEVICE_EVENT(macro)                                             \
+  FOREACH_OMPT_NOEMI_EVENT(macro)                                              \
+  FOREACH_OMPT_EMI_EVENT(macro)
+
+typedef uint64_t (*IdInterfaceTy)();
+
+/// Internal representation for OMPT device
+class OmptDeviceTy {
+public:
+  OmptDeviceTy() { Enabled.store(false); }
+  bool initialize() {
+    bool Old = false;
+    return Enabled.compare_exchange_strong(Old, true);
+  }
+  bool finalize() {
+    bool Old = true;
+    return Enabled.compare_exchange_strong(Old, false);
+  }
+
+private:
+  std::atomic<bool> Enabled;
+};
+
+/// Internal representation for OMPT device callback functions. An EMI callback
+/// passes a pointer to a correlation-id (HostOpId) to be set by the tool. For a
+/// non-EMI callback, the correlation-id (HostOpId) is set by the runtime and
+/// provided to the tool.
+class OmptDeviceCallbacksTy {
+public:
+  /// Invoked when a device is initialized
+  void OmptCallbackDeviceInitialize(int DeviceNum, const char *Type) {
+    if (ompt_callback_device_initialize_fn) {
+      OmptDeviceTy *Device = lookupDevice(DeviceNum);
+      if (Device->initialize()) {
+        ompt_callback_device_initialize_fn(
+            DeviceNum, Type, (ompt_device_t *)Device, doLookup, Documentation);
+      }
+    }
+  }
+
+  /// Invoked when a device is finalized
+  void OmptCallbackDeviceFinalize(int DeviceNum) {
+    if (ompt_callback_device_finalize_fn) {
+      OmptDeviceTy *Device = lookupDevice(DeviceNum);
+      if (Device->finalize()) {
+        ompt_callback_device_finalize_fn(DeviceNum);
+      }
+    }
+  }
+
+  /// Invoked when a device image is loaded
+  void OmptCallbackDeviceLoad(int DeviceNum, const char *Filename,
+                              int64_t OffsetInFile, void *VmaInFile,
+                              size_t Bytes, void *HostAddr, void *DeviceAddr,
+                              uint64_t ModuleId) {
+    if (ompt_callback_device_load_fn) {
+      ompt_callback_device_load_fn(DeviceNum, Filename, OffsetInFile, VmaInFile,
+                                   Bytes, HostAddr, DeviceAddr, ModuleId);
+    }
+  }
+
+  /// Invoked when a device image is unloaded
+  void OmptCallbackDeviceUnload(int DeviceNum, uint64_t ModuleId) {
+    if (ompt_callback_device_unload_fn) {
+      ompt_callback_device_unload_fn(DeviceNum, ModuleId);
+    }
+  }
+
+  /// Invoked when a data transfer is initiated
+  void OmptCallbackTargetDataOpEmi(
+      ompt_scope_endpoint_t EndPoint, ompt_data_t *TargetTaskData,
+      ompt_data_t *TargetData, ompt_target_data_op_t OpType, void *SrcAddr,
+      int SrcDeviceNum, void *DestAddr, int DestDeviceNum, size_t Bytes,
+      const void *CodePtrRA, IdInterfaceTy IdInterface, ompt_id_t *HostOpId) {
+    if (ompt_callback_target_data_op_emi_fn) {
+      ompt_callback_target_data_op_emi_fn(
+          EndPoint, TargetTaskData, TargetData, HostOpId, OpType, SrcAddr,
+          SrcDeviceNum, DestAddr, DestDeviceNum, Bytes, CodePtrRA);
+    } else if (EndPoint == ompt_scope_begin) {
+      OmptCallbackTargetDataOp(TargetData->value, OpType, SrcAddr, SrcDeviceNum,
+                               DestAddr, DestDeviceNum, Bytes, CodePtrRA,
+                               IdInterface, HostOpId);
+    }
+  }
+
+  /// Invoked when a data transfer is initiated
+  void OmptCallbackTargetDataOp(
+      ompt_id_t TargetId, ompt_target_data_op_t OpType, void *SrcAddr,
+      int SrcDeviceNum, void *DestAddr, int DestDeviceNum, size_t Bytes,
+      const void *CodePtrRA, IdInterfaceTy IdInterface, ompt_id_t *HostOpId) {
+    if (ompt_callback_target_data_op_fn) {
+      *HostOpId = IdInterface();
+      ompt_callback_target_data_op_fn(TargetId, *HostOpId, OpType, SrcAddr,
+                                      SrcDeviceNum, DestAddr, DestDeviceNum,
+                                      Bytes, CodePtrRA);
+    }
+  }
+
+  /// Invoked when a target region is executed
+  void OmptCallbackTargetEmi(ompt_target_t Kind, ompt_scope_endpoint_t EndPoint,
+                             int DeviceNum, ompt_data_t *TaskData,
+                             ompt_data_t *TargetTaskData,
+                             ompt_data_t *TargetData, const void *CodePtrRA,
+                             IdInterfaceTy IdInterface) {
+    if (ompt_callback_target_emi_fn) {
+      ompt_callback_target_emi_fn(Kind, EndPoint, DeviceNum, TaskData,
+                                  TargetTaskData, TargetData, CodePtrRA);
+    } else {
+      OmptCallbackTarget(Kind, EndPoint, DeviceNum, TaskData, CodePtrRA,
+                         TargetData, IdInterface);
+    }
+  }
+
+  /// Invoked when a target region is executed
+  void OmptCallbackTarget(ompt_target_t Kind, ompt_scope_endpoint_t EndPoint,
+                          int DeviceNum, ompt_data_t *TaskData,
+                          const void *CodePtrRA, ompt_data_t *TargetData,
+                          IdInterfaceTy IdInterface) {
+    // if we reach this point, ompt_callback_target_emi was not
+    // invoked so a tool didn't provide a target id. thus, we must
+    // unconditionally get an id here. even if there is no
+    // OMPT callback target, we need to have an id for use by other
+    // callbacks.
+    // note:
+    // on a scope_begin callback, IdInterface() will generate an id.
+    // on a scope_end callback, IdInterface() will return the existing
+    // id. it is safe to do the assignment again.
+    TargetData->value = IdInterface();
+    if (ompt_callback_target_fn) {
+      ompt_callback_target_fn(Kind, EndPoint, DeviceNum, TaskData,
+                              TargetData->value, CodePtrRA);
+    }
+  }
+
+  /// Invoked when a target map clause is executed
+  void OmptCallbackTargetMapEmi(ompt_data_t *TargetData, unsigned int NItems,
+                                void **HostAddr, void **DeviceAddr,
+                                size_t *Bytes, unsigned int *MappingFlags,
+                                const void *CodePtrRA) {
+    if (ompt_callback_target_map_emi_fn) {
+      ompt_callback_target_map_emi_fn(TargetData, NItems, HostAddr, DeviceAddr,
+                                      Bytes, MappingFlags, CodePtrRA);
+    } else {
+      OmptCallbackTargetMap(TargetData->value, NItems, HostAddr, DeviceAddr,
+                            Bytes, MappingFlags, CodePtrRA);
+    }
+  }
+
+  /// Invoked when a target map clause is executed
+  void OmptCallbackTargetMap(ompt_id_t TargetId, unsigned int NItems,
+                             void **HostAddr, void **DeviceAddr, size_t *Bytes,
+                             unsigned int *MappingFlags,
+                             const void *CodePtrRA) {
+    if (ompt_callback_target_map_fn) {
+      ompt_callback_target_map_fn(TargetId, NItems, HostAddr, DeviceAddr, Bytes,
+                                  MappingFlags, CodePtrRA);
+    }
+  }
+
+  /// Invoked when a target submit is executed
+  void OmptCallbackTargetSubmitEmi(ompt_scope_endpoint_t EndPoint,
+                                   ompt_data_t *TargetData,
+                                   unsigned int RequestedNumTeams,
+                                   IdInterfaceTy IdInterface,
+                                   ompt_id_t *HostOpId) {
+    if (ompt_callback_target_submit_emi_fn) {
+      ompt_callback_target_submit_emi_fn(EndPoint, TargetData, HostOpId,
+                                         RequestedNumTeams);
+    } else if (EndPoint == ompt_scope_begin) {
+      return OmptCallbackTargetSubmit(TargetData->value, RequestedNumTeams,
+                                      IdInterface, HostOpId);
+    }
+  }
+
+  /// Invoked when a target submit is executed
+  void OmptCallbackTargetSubmit(ompt_id_t TargetId,
+                                unsigned int RequestedNumTeams,
+                                IdInterfaceTy IdInterface,
+                                ompt_id_t *HostOpId) {
+    if (ompt_callback_target_submit_fn) {
+      *HostOpId = IdInterface();
+      ompt_callback_target_submit_fn(TargetId, *HostOpId, RequestedNumTeams);
+    }
+  }
+
+  /// Initialize the enabled flag and all the callbacks
+  void init() {
+    Enabled = false;
+#define initName(Name, Type, Code) Name##_fn = 0;
+    FOREACH_OMPT_TARGET_CALLBACK(initName)
+#undef initName
+  }
+
+  bool isEnabled() { return Enabled; }
+  void prepareDevices(int NumDevices) { resize(NumDevices); }
+
+  /// Used to register callbacks. \p Lookup is used to query a give callback
+  /// and the result is assigned to the callback of this object.
+  void registerCallbacks(ompt_function_lookup_t Lookup) {
+    Enabled = true;
+#define OmptBindCallback(Name, Type, Code)                                     \
+  Name##_fn = (Name##_t)Lookup(#Name);                                         \
+  DP("OMPT: class bound %s=%p\n", #Name, ((void *)(uint64_t)Name##_fn));
+
+    FOREACH_OMPT_TARGET_CALLBACK(OmptBindCallback);
+#undef OmptBindCallback
+  }
+
+  /// Used to find a callback given its name
+  ompt_interface_fn_t lookupCallback(const char *InterfaceFunctionName) {
+#define OmptLookup(Name, Type, Code)                                           \
+  if (strcmp(InterfaceFunctionName, #Name) == 0)                               \
+    return (ompt_interface_fn_t)Name##_fn;
+
+    FOREACH_OMPT_TARGET_CALLBACK(OmptLookup);
+#undef OmptLookup
+    return (ompt_interface_fn_t)0;
+  }
+
+  /// Wrapper function to find a callback given its name
+  static ompt_interface_fn_t doLookup(const char *InterfaceFunctionName);
+
+private:
+  bool Enabled;
+
+#define DeclareName(Name, Type, Code) Name##_t Name##_fn;
+  FOREACH_OMPT_TARGET_CALLBACK(DeclareName)
+#undef DeclareName
+
+  /// Allocate devices
+  static void resize(int NumDevices);
+  /// Find a device given its id
+  static OmptDeviceTy *lookupDevice(int DeviceNum);
+  /// Documentation based on omp-tools
+  static const char *Documentation;
+};
+
+extern OmptDeviceCallbacksTy ompt_device_callbacks;
+
+#endif

openmp/libomptarget/include/omptarget.h

int __tgt_target_teams_mapper(ident_t *loc, int64_t device_id, void *host_ptr,
                               int32_t num_teams, int32_t thread_limit);
 int __tgt_target_teams_nowait_mapper(
     ident_t *loc, int64_t device_id, void *host_ptr, int32_t arg_num,
     void **args_base, void **args, int64_t *arg_sizes, int64_t *arg_types,
     map_var_info_t *arg_names, void **arg_mappers, int32_t num_teams,
     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(ident_t *loc, 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);
 
 int __tgt_print_device_info(int64_t device_id);
 #ifdef __cplusplus

openmp/libomptarget/plugins/amdgpu/CMakeLists.txt

 add_library(omptarget.rtl.amdgpu SHARED
       impl/impl.cpp
       impl/interop_hsa.cpp
       impl/data.cpp
       impl/get_elf_mach_gfx_name.cpp
       impl/system.cpp
       impl/msgpack.cpp
       src/rtl.cpp
+      src/ompt_callback.cpp
       ${LIBOMPTARGET_EXTRA_SOURCE}
       )
 add_dependencies(omptarget.rtl.amdgpu omptarget.devicertl.amdgpu)
 
 # Install plugin under the lib destination folder.
 # When we build for debug, OPENMP_LIBDIR_SUFFIX get set to -debug
 install(TARGETS omptarget.rtl.amdgpu LIBRARY DESTINATION "lib${OPENMP_LIBDIR_SUFFIX}")
 set_property(TARGET omptarget.rtl.amdgpu PROPERTY INSTALL_RPATH_USE_LINK_PATH ON)

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

+//===------ ompt_callback.cpp - Target RTLs Implementation -------- 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
+//
+//===----------------------------------------------------------------------===//
+//
+// OMPT support for AMDGPU
+//
+//===----------------------------------------------------------------------===//
+
+//****************************************************************************
+// global includes
+//****************************************************************************
+
+#include <atomic>
+#include <vector>
+
+#include <string.h>
+
+//****************************************************************************
+// debug macro needed by include files
+//****************************************************************************
+
+#ifndef DEBUG_PREFIX
+#define DEBUG_PREFIX "Target AMDGPU RTL"
+#endif
+
+//****************************************************************************
+// local includes
+//****************************************************************************
+
+#include <Debug.h>
+#include <ompt-connector.h>
+#include <ompt_device_callbacks.h>
+
+//****************************************************************************
+// macros
+//****************************************************************************
+
+#define FOREACH_TARGET_FN(macro)
+
+#define fnptr_to_ptr(x) ((void *)(uint64_t)x)
+
+//****************************************************************************
+// global data
+//****************************************************************************
+
+OmptDeviceCallbacksTy ompt_device_callbacks;
+
+//****************************************************************************
+// private data
+//****************************************************************************
+
+const char *OmptDeviceCallbacksTy::Documentation = 0;
+
+static OmptDeviceTy *devices = 0;
+
+//****************************************************************************
+// private operations
+//****************************************************************************
+
+void OmptDeviceCallbacksTy::resize(int NumDevices) {
+  devices = new OmptDeviceTy[NumDevices];
+}
+
+OmptDeviceTy *OmptDeviceCallbacksTy::lookupDevice(int DeviceNum) {
+  return &devices[DeviceNum];
+}
+
+ompt_interface_fn_t
+OmptDeviceCallbacksTy::doLookup(const char *InterfaceFunctionName) {
+#define macro(fn)                                                              \
+  if (strcmp(InterfaceFunctionName, #fn) == 0)                                 \
+    return (ompt_interface_fn_t)fn;
+
+  FOREACH_TARGET_FN(macro);
+#undef macro
+  return (ompt_interface_fn_t)0;
+}
+
+#ifdef OMPT_SUPPORT
+
+static bool ompt_enabled = false;
+static ompt_get_target_info_t LIBOMPTARGET_GET_TARGET_OPID;
+
+static int ompt_device_init(ompt_function_lookup_t lookup,
+                            int initial_device_num, ompt_data_t *tool_data) {
+  DP("OMPT: Enter ompt_device_init\n");
+
+  ompt_enabled = true;
+
+  LIBOMPTARGET_GET_TARGET_OPID = (ompt_get_target_info_t)lookup(
+      LIBOMPTARGET_STRINGIFY(LIBOMPTARGET_GET_TARGET_OPID));
+
+  DP("OMPT: libomptarget_get_target_info = %p\n",
+     fnptr_to_ptr(LIBOMPTARGET_GET_TARGET_OPID));
+
+  ompt_device_callbacks.registerCallbacks(lookup);
+
+  DP("OMPT: Exit ompt_device_init\n");
+
+  return 0;
+}
+
+static void ompt_device_fini(ompt_data_t *tool_data) {
+  DP("OMPT: executing amdgpu_ompt_device_fini\n");
+}
+
+//****************************************************************************
+// constructor
+//****************************************************************************
+/// Used to initialize callbacks implemented by the tool. This interface will
+/// lookup the callbacks table in libomptarget and assign them to the callbacks
+/// table maintained in this plugin library.
+__attribute__((constructor)) static void ompt_init(void) {
+  DP("OMPT: Entering ompt_init\n");
+  /// Connect with libomptarget
+  static OmptLibraryConnectorTy LibomptargetConnector("ompt_libomptarget");
+  static ompt_start_tool_result_t OmptResult;
+
+  // Initialize OmptResult with the init and fini functions that will be
+  // called by the connector
+  OmptResult.initialize = ompt_device_init;
+  OmptResult.finalize = ompt_device_fini;
+  OmptResult.tool_data.value = 0;
+
+  // Initialize the device callbacks first
+  ompt_device_callbacks.init();
+
+  // Now call connect that causes the above init/fini functions to be called
+  LibomptargetConnector.connect(&OmptResult);
+  DP("OMPT: Exiting ompt_init\n");
+}
+#endif

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

 #include "DeviceEnvironment.h"
 #include "get_elf_mach_gfx_name.h"
 #include "omptargetplugin.h"
 #include "print_tracing.h"
 
 #include "llvm/Frontend/OpenMP/OMPConstants.h"
 #include "llvm/Frontend/OpenMP/OMPGridValues.h"
 
+#ifdef OMPT_SUPPORT
+#include <ompt_device_callbacks.h>
+#define OMPT_IF_BUILT_AND_ENABLED(stmts)                                       \
+  do {                                                                         \
+    if (ompt_device_callbacks.isEnabled()) {                                   \
+      stmts;                                                                   \
+    }                                                                          \
+  } while (0)
+#else
+#define OMPT_IF_BUILT_AND_ENABLED(stmts)
+#endif
+
 // hostrpc interface, FIXME: consider moving to its own include these are
 // statically linked into amdgpu/plugin if present from hostrpc_services.a,
 // linked as --whole-archive to override the weak symbols that are used to
 // implement a fallback for toolchains that do not yet have a hostrpc library.
 extern "C" {
 uint64_t hostrpc_assign_buffer(hsa_agent_t agent, hsa_queue_t *this_Q,
                                uint32_t device_id);
 hsa_status_t hostrpc_init();
static int readEnvElseMinusOne(const char *Env) {
     int res = -1;
     if (envStr) {
       res = std::stoi(envStr);
       DP("Parsed %s=%d\n", Env, res);
     }
     return res;
   }
 
+#ifdef OMPT_SUPPORT
+  void doOmptDeviceLoad(int32_t DeviceId, const char *FileName,
+                        int64_t OffsetInFile, void *VmaInFile, size_t Bytes,
+                        void *HostAddr, void *DeviceAddr, uint64_t ModuleId) {
+    ompt_device_callbacks.OmptCallbackDeviceLoad(
+        DeviceId, FileName, OffsetInFile, VmaInFile, Bytes, HostAddr,
+        DeviceAddr, ModuleId);
+  }
+
+  void doOmptDeviceInitialize(int32_t DeviceId, const char *InfoName) {
+    std::string OmptGPUType("AMD ");
+    OmptGPUType += InfoName;
+    ompt_device_callbacks.OmptCallbackDeviceInitialize(DeviceId,
+                                                       OmptGPUType.c_str());
+  }
+
+  void doOmptDeviceFinalize() {
+    for (int i = 0; i < NumberOfDevices; i++)
+      ompt_device_callbacks.OmptCallbackDeviceFinalize(i);
+  }
+#endif
+
   RTLDeviceInfoTy() {
     DP("Start initializing " GETNAME(TARGET_NAME) "\n");
 
     // LIBOMPTARGET_KERNEL_TRACE provides a kernel launch trace to stderr
     // anytime. You do not need a debug library build.
     //  0 => no tracing
     //  1 => tracing dispatch only
     // >1 => verbosity increase
RTLDeviceInfoTy() {
     }
 
     err = setupHostMemoryPools(CPUAgents);
     if (err != HSA_STATUS_SUCCESS) {
       DP("Setup for Host Memory Pools failed\n");
       return;
     }
 
+#ifdef OMPT_SUPPORT
+    // TODO ompt_device_callbacks.enabled is not yet set since
+    // register_callbacks on the plugin instance is not yet
+    // called. Hence, unconditionally prepare devices.
+    ompt_device_callbacks.prepareDevices(NumberOfDevices);
+#endif
+
     for (int i = 0; i < NumberOfDevices; i++) {
       uint32_t queue_size = 0;
       {
         hsa_status_t err = hsa_agent_get_info(
             HSAAgents[i], HSA_AGENT_INFO_QUEUE_MAX_SIZE, &queue_size);
         if (err != HSA_STATUS_SUCCESS) {
           DP("HSA query QUEUE_MAX_SIZE failed for agent %d\n", i);
           return;
#endif
   }
 
   ~RTLDeviceInfoTy() {
     DP("Finalizing the " GETNAME(TARGET_NAME) " DeviceInfo.\n");
     if (!HSAInitSuccess()) {
       // Then none of these can have been set up and they can't be torn down
       return;
     }
+
+    OMPT_IF_BUILT_AND_ENABLED(doOmptDeviceFinalize());
+
     // Run destructors on types that use HSA before
     // impl_finalize removes access to it
     deviceStateStore.clear();
     KernelArgPoolMap.clear();
     // Terminate hostrpc before finalizing hsa
     hostrpc_terminate();
 
     hsa_status_t Err;
DP("Device %d: default limit for groupsPerDevice %d & threadsPerGroup %d\n",
      DeviceInfo.ThreadsPerGroup[device_id]);
 
   DP("Device %d: wavefront size %d, total threads %d x %d = %d\n", device_id,
      DeviceInfo.WarpSize[device_id], DeviceInfo.ThreadsPerGroup[device_id],
      DeviceInfo.GroupsPerDevice[device_id],
      DeviceInfo.GroupsPerDevice[device_id] *
          DeviceInfo.ThreadsPerGroup[device_id]);
 
+  OMPT_IF_BUILT_AND_ENABLED(
+      DeviceInfo.doOmptDeviceInitialize(device_id, GetInfoName));
+
   return OFFLOAD_SUCCESS;
 }
 
 static __tgt_target_table *
 __tgt_rtl_load_binary_locked(int32_t device_id, __tgt_device_image *image);
 
 __tgt_target_table *__tgt_rtl_load_binary(int32_t device_id,
                                           __tgt_device_image *image) {
if (err != HSA_STATUS_SUCCESS) {
            " compiler flag: -march=<gpu>\n",
            DeviceName, ElfName);
       } else {
         DP("Error loading image onto GPU: %s\n", get_error_string(err));
       }
 
       return NULL;
     }
-
     err = env.after_loading();
     if (err != HSA_STATUS_SUCCESS) {
       return NULL;
     }
   }
 
   DP("AMDGPU module successfully loaded!\n");
 
+  OMPT_IF_BUILT_AND_ENABLED(DeviceInfo.doOmptDeviceLoad(
+      device_id, nullptr /* FileName */, 0 /* OffsetInFile */,
+      nullptr /* VmaInFile */, img_size, image->ImageStart,
+      nullptr /* DeviceAddr */, 0 /* FIXME ModuleId */));
   {
     // the device_State array is either large value in bss or a void* that
     // needs to be assigned to a pointer to an array of size device_state_bytes
     // If absent, it has been deadstripped and needs no setup.
 
     void *state_ptr;
     uint32_t state_ptr_size;
     auto &SymbolInfoMap = DeviceInfo.SymbolInfoTable[device_id];

openmp/libomptarget/src/CMakeLists.txt

 libomptarget_say("Building offloading runtime library libomptarget.")
 
 set(LIBOMPTARGET_SRC_FILES
   ${CMAKE_CURRENT_SOURCE_DIR}/api.cpp
   ${CMAKE_CURRENT_SOURCE_DIR}/device.cpp
   ${CMAKE_CURRENT_SOURCE_DIR}/interface.cpp
   ${CMAKE_CURRENT_SOURCE_DIR}/interop.cpp
   ${CMAKE_CURRENT_SOURCE_DIR}/omptarget.cpp
+  ${CMAKE_CURRENT_SOURCE_DIR}/ompt_callback.cpp
   ${CMAKE_CURRENT_SOURCE_DIR}/rtl.cpp
 )
 
 set(LIBOMPTARGET_SRC_FILES ${LIBOMPTARGET_SRC_FILES} PARENT_SCOPE)
 
 # Build libomptarget library with libdl dependency.
 add_library(omptarget SHARED ${LIBOMPTARGET_SRC_FILES})
 set_target_properties(omptarget PROPERTIES INSTALL_RPATH "$ORIGIN" BUILD_RPATH "$ORIGIN")

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 "ompt_callback.h"
 #include "omptarget.h"
 #include "private.h"
 #include "rtl.h"
 
 #include <cassert>
 #include <climits>
 #include <cstdint>
 #include <cstdio>

 // Load binary to device.
 __tgt_target_table *DeviceTy::load_binary(void *Img) {
   std::lock_guard<decltype(RTL->Mtx)> LG(RTL->Mtx);
   __tgt_target_table *rc = RTL->load_binary(RTLDeviceID, Img);
   return rc;
 }
 
 void *DeviceTy::allocData(int64_t Size, void *HstPtr, int32_t Kind) {
+  OMPT_IF_BUILT(OmptInterfaceTargetDataOpRAII TgtDataAlloc(
+      RTLDeviceID, Size, HstPtr, nullptr /* TgtPtr */, ompt_target_data_alloc));
   return RTL->data_alloc(RTLDeviceID, Size, HstPtr, Kind);
 }
 
 int32_t DeviceTy::deleteData(void *TgtPtrBegin) {
+  OMPT_IF_BUILT(OmptInterfaceTargetDataOpRAII TgtDataDelete(
+      RTLDeviceID, 0 /* Size */, nullptr /* HostPtr */, TgtPtrBegin,
+      ompt_target_data_delete));
   return RTL->data_delete(RTLDeviceID, TgtPtrBegin);
 }
 
 // Submit data to device
 int32_t DeviceTy::submitData(void *TgtPtrBegin, void *HstPtrBegin, 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,
          "Copying data from host to device, HstPtr=" DPxMOD ", TgtPtr=" DPxMOD
          ", Size=%" PRId64 ", Name=%s\n",
          DPxPTR(HstPtrBegin), DPxPTR(TgtPtrBegin), Size,
          (HT && HT->HstPtrName) ? getNameFromMapping(HT->HstPtrName).c_str()
                                 : "unknown");
   }
-
+  {
+    OMPT_IF_BUILT(OmptInterfaceTargetDataOpRAII TargetDataSubmitRAII(
+        RTLDeviceID, Size, HstPtrBegin, TgtPtrBegin,
+        ompt_target_data_transfer_to_device));
     if (!AsyncInfo || !RTL->data_submit_async || !RTL->synchronize)
       return RTL->data_submit(RTLDeviceID, TgtPtrBegin, HstPtrBegin, Size);
     else
       return RTL->data_submit_async(RTLDeviceID, TgtPtrBegin, HstPtrBegin, Size,
                                     AsyncInfo);
   }
+}
 
 // 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,
          "Copying data from device to host, TgtPtr=" DPxMOD ", HstPtr=" DPxMOD
          ", Size=%" PRId64 ", Name=%s\n",
          DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin), Size,
          (HT && HT->HstPtrName) ? getNameFromMapping(HT->HstPtrName).c_str()
                                 : "unknown");
   }
-
+  {
+    OMPT_IF_BUILT(OmptInterfaceTargetDataOpRAII TargetDataRetrieve(
+        RTLDeviceID, Size, HstPtrBegin, TgtPtrBegin,
+        ompt_target_data_transfer_from_device));
     if (!RTL->data_retrieve_async || !RTL->synchronize)
       return RTL->data_retrieve(RTLDeviceID, HstPtrBegin, TgtPtrBegin, Size);
     else
-    return RTL->data_retrieve_async(RTLDeviceID, HstPtrBegin, TgtPtrBegin, Size,
-                                    AsyncInfo);
+      return RTL->data_retrieve_async(RTLDeviceID, HstPtrBegin, TgtPtrBegin,
+                                      Size, AsyncInfo);
+  }
 }
 
 // Copy data from current device to destination device directly
 int32_t DeviceTy::dataExchange(void *SrcPtr, DeviceTy &DstDev, void *DstPtr,
                                int64_t Size, AsyncInfoTy &AsyncInfo) {
   if (!AsyncInfo || !RTL->data_exchange_async || !RTL->synchronize) {
     assert(RTL->data_exchange && "RTL->data_exchange is nullptr");
     return RTL->data_exchange(RTLDeviceID, SrcPtr, DstDev.RTLDeviceID, DstPtr,
                               Size);

openmp/libomptarget/src/exports

global:
     omp_get_interop_str;
     omp_get_interop_int;
     omp_get_interop_name;
     omp_get_interop_type_desc;
     omp_get_interop_rc_desc;
     __tgt_interop_init;
     __tgt_interop_use;
     __tgt_interop_destroy;
+    ompt_libomptarget_connect;
   local:
     *;
 };
 

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 <omptarget.h>
+
 #include "device.h"
 #include "omptarget.h"
 #include "private.h"
 #include "rtl.h"
 
 #include <cassert>
 #include <cstdio>
 #include <cstdlib>
 #include <mutex>
 
+#include "ompt_callback.h"
+
 ////////////////////////////////////////////////////////////////////////////////
 /// adds requires flags
 EXTERN void __tgt_register_requires(int64_t flags) {
   TIMESCOPE();
   PM->RTLs.RegisterRequires(flags);
 }
 
 ////////////////////////////////////////////////////////////////////////////////
EXTERN void __tgt_target_data_begin_mapper(ident_t *loc, int64_t device_id,
                                            int32_t arg_num, void **args_base,
                                            void **args, int64_t *arg_sizes,
                                            int64_t *arg_types,
                                            map_var_info_t *arg_names,
                                            void **arg_mappers) {
   TIMESCOPE_WITH_IDENT(loc);
   DP("Entering data begin region for device %" PRId64 " with %d mappings\n",
      device_id, arg_num);
+
   if (checkDeviceAndCtors(device_id, loc)) {
     DP("Not offloading to device %" PRId64 "\n", device_id);
     return;
   }
 
   DeviceTy &Device = *PM->Devices[device_id];
 
   if (getInfoLevel() & OMP_INFOTYPE_KERNEL_ARGS)
     printKernelArguments(loc, device_id, arg_num, arg_sizes, arg_types,
                          arg_names, "Entering OpenMP data region");
 #ifdef OMPTARGET_DEBUG
   for (int i = 0; i < arg_num; ++i) {
     DP("Entry %2d: Base=" DPxMOD ", Begin=" DPxMOD ", Size=%" PRId64
        ", Type=0x%" PRIx64 ", Name=%s\n",
        i, DPxPTR(args_base[i]), DPxPTR(args[i]), arg_sizes[i], arg_types[i],
        (arg_names) ? getNameFromMapping(arg_names[i]).c_str() : "unknown");
   }
 #endif
 
+  {
+    OMPT_IF_BUILT(OmptInterfaceTargetRAII TargetDataBeginRAII(
+        device_id, ompt_target_enter_data));
     AsyncInfoTy AsyncInfo(Device);
     int rc = targetDataBegin(loc, Device, arg_num, args_base, args, arg_sizes,
                              arg_types, arg_names, arg_mappers, AsyncInfo);
     if (rc == OFFLOAD_SUCCESS)
       rc = AsyncInfo.synchronize();
     handleTargetOutcome(rc == OFFLOAD_SUCCESS, loc);
   }
+}
 
 EXTERN void __tgt_target_data_begin_nowait_mapper(
     ident_t *loc, int64_t device_id, int32_t arg_num, void **args_base,
     void **args, int64_t *arg_sizes, int64_t *arg_types,
     map_var_info_t *arg_names, void **arg_mappers, int32_t depNum,
     void *depList, int32_t noAliasDepNum, void *noAliasDepList) {
   TIMESCOPE_WITH_IDENT(loc);
 
#ifdef OMPTARGET_DEBUG
   for (int i = 0; i < arg_num; ++i) {
     DP("Entry %2d: Base=" DPxMOD ", Begin=" DPxMOD ", Size=%" PRId64
        ", Type=0x%" PRIx64 ", Name=%s\n",
        i, DPxPTR(args_base[i]), DPxPTR(args[i]), arg_sizes[i], arg_types[i],
        (arg_names) ? getNameFromMapping(arg_names[i]).c_str() : "unknown");
   }
 #endif
 
+  {
+    OMPT_IF_BUILT(OmptInterfaceTargetRAII TargetDataEndRAII(
+        device_id, ompt_target_exit_data));
     AsyncInfoTy AsyncInfo(Device);
     int rc = targetDataEnd(loc, Device, arg_num, args_base, args, arg_sizes,
                            arg_types, arg_names, arg_mappers, AsyncInfo);
     if (rc == OFFLOAD_SUCCESS)
       rc = AsyncInfo.synchronize();
     handleTargetOutcome(rc == OFFLOAD_SUCCESS, loc);
   }
+}
 
 EXTERN void __tgt_target_data_end_nowait_mapper(
     ident_t *loc, int64_t device_id, int32_t arg_num, void **args_base,
     void **args, int64_t *arg_sizes, int64_t *arg_types,
     map_var_info_t *arg_names, void **arg_mappers, int32_t depNum,
     void *depList, int32_t noAliasDepNum, void *noAliasDepList) {
   TIMESCOPE_WITH_IDENT(loc);
 

 EXTERN void __tgt_target_data_update_mapper(ident_t *loc, int64_t device_id,
                                             int32_t arg_num, void **args_base,
                                             void **args, int64_t *arg_sizes,
                                             int64_t *arg_types,
                                             map_var_info_t *arg_names,
                                             void **arg_mappers) {
   TIMESCOPE_WITH_IDENT(loc);
   DP("Entering data update with %d mappings\n", arg_num);
+
   if (checkDeviceAndCtors(device_id, loc)) {
     DP("Not offloading to device %" PRId64 "\n", device_id);
     return;
   }
 
   if (getInfoLevel() & OMP_INFOTYPE_KERNEL_ARGS)
     printKernelArguments(loc, device_id, arg_num, arg_sizes, arg_types,
                          arg_names, "Updating OpenMP data");
 
   DeviceTy &Device = *PM->Devices[device_id];
+  {
+    OMPT_IF_BUILT(OmptInterfaceTargetRAII TargetDataUpdateRAII(
+        device_id, ompt_target_update));
     AsyncInfoTy AsyncInfo(Device);
     int rc = targetDataUpdate(loc, Device, arg_num, args_base, args, arg_sizes,
                               arg_types, arg_names, arg_mappers, AsyncInfo);
     if (rc == OFFLOAD_SUCCESS)
       rc = AsyncInfo.synchronize();
     handleTargetOutcome(rc == OFFLOAD_SUCCESS, loc);
   }
+}
 
 EXTERN void __tgt_target_data_update_nowait_mapper(
     ident_t *loc, int64_t device_id, int32_t arg_num, void **args_base,
     void **args, int64_t *arg_sizes, int64_t *arg_types,
     map_var_info_t *arg_names, void **arg_mappers, int32_t depNum,
     void *depList, int32_t noAliasDepNum, void *noAliasDepList) {
   TIMESCOPE_WITH_IDENT(loc);
 
for (int i = 0; i < arg_num; ++i) {
     DP("Entry %2d: Base=" DPxMOD ", Begin=" DPxMOD ", Size=%" PRId64
        ", Type=0x%" PRIx64 ", Name=%s\n",
        i, DPxPTR(args_base[i]), DPxPTR(args[i]), arg_sizes[i], arg_types[i],
        (arg_names) ? getNameFromMapping(arg_names[i]).c_str() : "unknown");
   }
 #endif
 
   DeviceTy &Device = *PM->Devices[device_id];
+
+  {
+    OMPT_IF_BUILT(OmptInterfaceTargetRAII TargetRAII(device_id, ompt_target));
     AsyncInfoTy AsyncInfo(Device);
     int rc = target(loc, Device, host_ptr, arg_num, args_base, args, arg_sizes,
                     arg_types, arg_names, arg_mappers, 0, 0, false /*team*/,
                     AsyncInfo);
     if (rc == OFFLOAD_SUCCESS)
       rc = AsyncInfo.synchronize();
     handleTargetOutcome(rc == OFFLOAD_SUCCESS, loc);
     assert(rc == OFFLOAD_SUCCESS && "__tgt_target_mapper unexpected failure!");
+  }
   return OMP_TGT_SUCCESS;
 }
 
 EXTERN int __tgt_target_nowait_mapper(
     ident_t *loc, int64_t device_id, void *host_ptr, int32_t arg_num,
     void **args_base, void **args, int64_t *arg_sizes, int64_t *arg_types,
     map_var_info_t *arg_names, void **arg_mappers, int32_t depNum,
     void *depList, int32_t noAliasDepNum, void *noAliasDepList) {
for (int i = 0; i < arg_num; ++i) {
     DP("Entry %2d: Base=" DPxMOD ", Begin=" DPxMOD ", Size=%" PRId64
        ", Type=0x%" PRIx64 ", Name=%s\n",
        i, DPxPTR(args_base[i]), DPxPTR(args[i]), arg_sizes[i], arg_types[i],
        (arg_names) ? getNameFromMapping(arg_names[i]).c_str() : "unknown");
   }
 #endif
 
   DeviceTy &Device = *PM->Devices[device_id];
+
+  {
+    OMPT_IF_BUILT(OmptInterfaceTargetRAII TargetRAII(device_id, ompt_target));
     AsyncInfoTy AsyncInfo(Device);
     int rc = target(loc, Device, host_ptr, arg_num, args_base, args, arg_sizes,
                     arg_types, arg_names, arg_mappers, team_num, thread_limit,
                     true /*team*/, AsyncInfo);
     if (rc == OFFLOAD_SUCCESS)
       rc = AsyncInfo.synchronize();
     handleTargetOutcome(rc == OFFLOAD_SUCCESS, loc);
     assert(rc == OFFLOAD_SUCCESS &&
            "__tgt_target_teams_mapper unexpected failure!");
+  }
   return OMP_TGT_SUCCESS;
 }
 
 EXTERN int __tgt_target_teams_nowait_mapper(
     ident_t *loc, int64_t device_id, void *host_ptr, int32_t arg_num,
     void **args_base, void **args, int64_t *arg_sizes, int64_t *arg_types,
     map_var_info_t *arg_names, void **arg_mappers, int32_t team_num,
     int32_t thread_limit, int32_t depNum, void *depList, int32_t noAliasDepNum,
DP("__tgt_push_mapper_component(Handle=" DPxMOD
      ", Type=0x%" PRIx64 ", Name=%s).\n",
      DPxPTR(rt_mapper_handle), DPxPTR(base), DPxPTR(begin), size, type,
      (name) ? getNameFromMapping(name).c_str() : "unknown");
   auto *MapperComponentsPtr = (struct MapperComponentsTy *)rt_mapper_handle;
   MapperComponentsPtr->Components.push_back(
       MapComponentInfoTy(base, begin, size, type, name));
 }
 
-EXTERN void __kmpc_push_target_tripcount(int64_t device_id,
+EXTERN void __kmpc_push_target_tripcount(ident_t *loc, int64_t device_id,
                                          uint64_t loop_tripcount) {
-  __kmpc_push_target_tripcount_mapper(nullptr, device_id, loop_tripcount);
+  __kmpc_push_target_tripcount_mapper(loc, device_id, loop_tripcount);
 }
 
 EXTERN void __kmpc_push_target_tripcount_mapper(ident_t *loc, int64_t device_id,
                                                 uint64_t loop_tripcount) {
   TIMESCOPE_WITH_IDENT(loc);
   if (checkDeviceAndCtors(device_id, loc)) {
     DP("Not offloading to device %" PRId64 "\n", device_id);
     return;

openmp/libomptarget/src/ompt_callback.h

+//===----------- device.h - 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
+//
+//===----------------------------------------------------------------------===//
+//
+// Declarations for OpenMP Tool callback dispatchers
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _OMPTARGET_OMPT_CALLBACK_H
+#define _OMPTARGET_OMPT_CALLBACK_H
+
+// If target OMPT support is compiled in
+#ifdef OMPT_SUPPORT
+#define OMPT_IF_BUILT(stmt) stmt
+#else
+#define OMPT_IF_BUILT(stmt)
+#endif
+
+#if (__PPC64__ | __arm__)
+#define OMPT_GET_FRAME_ADDRESS(level) __builtin_frame_address(level)
+#define OMPT_FRAME_POSITION_DEFAULT ompt_frame_cfa
+#else
+#define OMPT_GET_FRAME_ADDRESS(level) __builtin_frame_address(level)
+#define OMPT_FRAME_POSITION_DEFAULT ompt_frame_framepointer
+#endif
+
+#define OMPT_FRAME_FLAGS (ompt_frame_runtime | OMPT_FRAME_POSITION_DEFAULT)
+#define OMPT_GET_RETURN_ADDRESS(level) __builtin_return_address(level)
+
+#include <omp-tools.h>
+
+/// Used to maintain execution state for this thread
+class OmptInterface {
+public:
+  OmptInterface()
+      : EnterFrame(NULL), Codeptr(NULL), InterfaceState(ompt_state_idle) {}
+
+  void setState(void *Frame, void *Codeptr);
+
+  void clearState();
+
+  // target op callbacks
+  void beginTargetDataAlloc(int64_t DeviceId, void *TgtPtrBegin, size_t Size,
+                            void *Code);
+
+  void endTargetDataAlloc(int64_t DeviceId, void *TgtPtrBegin, size_t Size,
+                          void *Code);
+
+  void beginTargetDataSubmit(int64_t DeviceId, void *HstPtrBegin,
+                             void *TgtPtrBegin, size_t Size, void *Code);
+
+  void endTargetDataSubmit(int64_t DeviceId, void *HstPtrBegin,
+                           void *TgtPtrBegin, size_t Size, void *Code);
+
+  void beginTargetDataDelete(int64_t DeviceId, void *TgtPtrBegin, void *Code);
+
+  void endTargetDataDelete(int64_t DeviceId, void *TgtPtrBegin, void *Code);
+
+  void beginTargetDataRetrieve(int64_t DeviceId, void *HstPtrBegin,
+                               void *TgtPtrBegin, size_t Size, void *Code);
+
+  void endTargetDataRetrieve(int64_t DeviceId, void *HstPtrBegin,
+                             void *TgtPtrBegin, size_t Size, void *Code);
+
+  void beginTargetSubmit(unsigned int NumTeams = 1);
+
+  void endTargetSubmit(unsigned int NumTeams = 1);
+
+  // target region callbacks
+  void beginTargetDataEnter(int64_t DeviceId, void *Code);
+
+  void endTargetDataEnter(int64_t DeviceId, void *Code);
+
+  void beginTargetDataExit(int64_t DeviceId, void *Code);
+
+  void endTargetDataExit(int64_t DeviceId, void *Code);
+
+  void beginTargetUpdate(int64_t DeviceId, void *Code);
+
+  void endTargetUpdate(int64_t DeviceId, void *Code);
+
+  void beginTarget(int64_t DeviceId, void *Code);
+
+  void endTarget(int64_t DeviceId, void *Code);
+
+private:
+  void setStateHelper(void *Frame, void *Code, int Flags, int State);
+
+  // begin/end target op marks
+  void beginTargetOperation();
+
+  void endTargetOperation();
+
+  // begin/end target region marks
+  void beginTargetRegion();
+
+  void endTargetRegion();
+
+  void announceTargetRegion(const char *Name);
+
+private:
+  void *EnterFrame;
+  void *Codeptr;
+  int InterfaceState;
+};
+
+extern thread_local OmptInterface ompt_interface;
+extern bool ompt_enabled;
+
+// RAII objects used by the runtime entry points
+struct OmptInterfaceTargetRAII {
+  OmptInterfaceTargetRAII(int64_t Id, ompt_target_t Op)
+      : CodePtr{nullptr}, DeviceId{Id}, TgtOp{Op} {
+    if (!ompt_enabled)
+      return;
+    CodePtr = OMPT_GET_RETURN_ADDRESS(0);
+    ompt_interface.setState(OMPT_GET_FRAME_ADDRESS(0), CodePtr);
+    switch (TgtOp) {
+    case ompt_target_enter_data:
+    case ompt_target_enter_data_nowait:
+      ompt_interface.beginTargetDataEnter(DeviceId, CodePtr);
+      break;
+    case ompt_target_exit_data:
+    case ompt_target_exit_data_nowait:
+      ompt_interface.beginTargetDataExit(DeviceId, CodePtr);
+      break;
+    case ompt_target_update:
+    case ompt_target_update_nowait:
+      ompt_interface.beginTargetUpdate(DeviceId, CodePtr);
+      break;
+    case ompt_target:
+    case ompt_target_nowait:
+      ompt_interface.beginTarget(DeviceId, CodePtr);
+      break;
+    }
+  }
+  ~OmptInterfaceTargetRAII() {
+    if (!ompt_enabled)
+      return;
+    switch (TgtOp) {
+    case ompt_target_enter_data:
+    case ompt_target_enter_data_nowait:
+      ompt_interface.endTargetDataEnter(DeviceId, CodePtr);
+      break;
+    case ompt_target_exit_data:
+    case ompt_target_exit_data_nowait:
+      ompt_interface.endTargetDataExit(DeviceId, CodePtr);
+      break;
+    case ompt_target_update:
+    case ompt_target_update_nowait:
+      ompt_interface.endTargetUpdate(DeviceId, CodePtr);
+      break;
+    case ompt_target:
+    case ompt_target_nowait:
+      ompt_interface.endTarget(DeviceId, CodePtr);
+      break;
+    }
+    ompt_interface.clearState();
+  }
+
+private:
+  void *CodePtr;
+  int64_t DeviceId;
+  ompt_target_t TgtOp;
+};
+
+struct OmptInterfaceTargetSubmitRAII {
+  OmptInterfaceTargetSubmitRAII(int32_t Teams) : NumTeams{Teams} {
+    if (!ompt_enabled)
+      return;
+    ompt_interface.setState(OMPT_GET_FRAME_ADDRESS(0),
+                            OMPT_GET_RETURN_ADDRESS(0));
+    ompt_interface.beginTargetSubmit(NumTeams);
+  }
+  ~OmptInterfaceTargetSubmitRAII() {
+    if (!ompt_enabled)
+      return;
+    ompt_interface.endTargetSubmit(NumTeams);
+    ompt_interface.clearState();
+  }
+
+private:
+  int32_t NumTeams;
+};
+
+struct OmptInterfaceTargetDataOpRAII {
+  OmptInterfaceTargetDataOpRAII(int32_t Id, int64_t Sz, void *HPtr, void *TPtr,
+                                ompt_target_data_op_t Op)
+      : DeviceId{Id}, Size{Sz}, CodePtr{nullptr}, HostPtr{HPtr}, TgtPtr{TPtr},
+        TgtDataOp{Op} {
+    if (!ompt_enabled)
+      return;
+    CodePtr = OMPT_GET_RETURN_ADDRESS(0);
+    ompt_interface.setState(OMPT_GET_FRAME_ADDRESS(0), CodePtr);
+    switch (TgtDataOp) {
+    case ompt_target_data_alloc:
+    case ompt_target_data_alloc_async:
+      ompt_interface.beginTargetDataAlloc(DeviceId, HostPtr, Size, CodePtr);
+      break;
+    case ompt_target_data_delete:
+    case ompt_target_data_delete_async:
+      ompt_interface.beginTargetDataDelete(DeviceId, TgtPtr, CodePtr);
+      break;
+    case ompt_target_data_transfer_to_device:
+    case ompt_target_data_transfer_to_device_async:
+      ompt_interface.beginTargetDataSubmit(DeviceId, TgtPtr, HostPtr, Size,
+                                           CodePtr);
+      break;
+    case ompt_target_data_transfer_from_device:
+    case ompt_target_data_transfer_from_device_async:
+      ompt_interface.beginTargetDataRetrieve(DeviceId, HostPtr, TgtPtr, Size,
+                                             CodePtr);
+      break;
+    default:
+      break;
+    }
+  }
+  ~OmptInterfaceTargetDataOpRAII() {
+    if (!ompt_enabled)
+      return;
+    switch (TgtDataOp) {
+    case ompt_target_data_alloc:
+    case ompt_target_data_alloc_async:
+      ompt_interface.endTargetDataAlloc(DeviceId, HostPtr, Size, CodePtr);
+      break;
+    case ompt_target_data_delete:
+    case ompt_target_data_delete_async:
+      ompt_interface.endTargetDataDelete(DeviceId, TgtPtr, CodePtr);
+      break;
+    case ompt_target_data_transfer_to_device:
+    case ompt_target_data_transfer_to_device_async:
+      ompt_interface.endTargetDataSubmit(DeviceId, TgtPtr, HostPtr, Size,
+                                         CodePtr);
+      break;
+    case ompt_target_data_transfer_from_device:
+    case ompt_target_data_transfer_from_device_async:
+      ompt_interface.endTargetDataRetrieve(DeviceId, HostPtr, TgtPtr, Size,
+                                           CodePtr);
+      break;
+    default:
+      break;
+    }
+    ompt_interface.clearState();
+  }
+
+private:
+  int32_t DeviceId;
+  int64_t Size;
+  void *CodePtr;
+  void *HostPtr;
+  void *TgtPtr;
+  ompt_target_data_op_t TgtDataOp;
+};
+
+#endif // _OMPTARGET_OMPT_CALLBACK_H

openmp/libomptarget/src/ompt_callback.cpp

+//===-- ompt_callback.cpp - Target independent OpenMP target RTL -- 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
+//
+//===----------------------------------------------------------------------===//
+//
+// Implementation of OMPT callback interfaces for target independent layer
+//
+//===----------------------------------------------------------------------===//
+
+#include <assert.h>
+#include <atomic>
+#include <cstdlib>
+#include <cstring>
+#include <dlfcn.h>
+
+//****************************************************************************
+// local include files
+//****************************************************************************
+
+#include <omp-tools.h>
+
+#include "ompt_callback.h"
+#include "private.h"
+
+#include <ompt-connector.h>
+#include <ompt_device_callbacks.h>
+
+/*************************************************************************
+ * macros
+ ************************************************************************/
+
+#define fnptr_to_ptr(x) ((void *)(uint64_t)x)
+
+/*****************************************************************************
+ * global data
+ *****************************************************************************/
+
+/// Used to indicate whether OMPT was enabled for this library
+bool ompt_enabled = false;
+/// Object maintaining all the callbacks for this library
+OmptDeviceCallbacksTy ompt_device_callbacks;
+/// Thread local state during execution for target region and associated
+/// metadata
+thread_local OmptInterface ompt_interface;
+
+namespace {
+
+/*******************************************************************************
+ * type declarations
+ *******************************************************************************/
+
+/// Used to maintain the init and fini functions
+class LibomptargetRtlFinalizer {
+public:
+  LibomptargetRtlFinalizer() : Fn(nullptr) {}
+  void registerRtl(ompt_finalize_t _fn) {
+    assert(Fn == nullptr && "Finalizer already registered");
+    Fn = _fn;
+  }
+  void finalize() {
+    if (Fn)
+      Fn(NULL);
+    Fn = nullptr;
+  }
+
+private:
+  ompt_finalize_t Fn;
+};
+
+typedef int (*ompt_set_frame_enter_t)(void *addr, int flags, int state);
+typedef ompt_data_t *(*ompt_get_task_data_t)();
+typedef ompt_data_t *(*ompt_get_target_task_data_t)();
+
+} // namespace
+
+/*****************************************************************************
+ * private data
+ *****************************************************************************/
+
+static ompt_set_frame_enter_t ompt_set_frame_enter_fn = nullptr;
+static ompt_get_task_data_t ompt_get_task_data_fn = nullptr;
+static ompt_get_target_task_data_t ompt_get_target_task_data_fn = nullptr;
+
+static LibomptargetRtlFinalizer libomptarget_rtl_finalizer;
+
+static std::atomic<uint64_t> unique_id_ticket(1);
+
+/*****************************************************************************
+ * Thread local data
+ *****************************************************************************/
+
+/// Correlation id that is incremented with target operations
+static thread_local uint64_t ompt_target_region_opid = 1;
+// The following are thread local data structures that are passed to callbacks
+static thread_local ompt_data_t ompt_target_data = ompt_data_none;
+static thread_local ompt_data_t *ompt_task_data = nullptr;
+static thread_local ompt_data_t *ompt_target_task_data = nullptr;
+static thread_local ompt_id_t host_op_id = 0;
+
+/*****************************************************************************
+ * OMPT private operations
+ *****************************************************************************/
+
+/// Used to create a new correlation id
+static uint64_t id_create() { return unique_id_ticket.fetch_add(1); }
+
+static uint64_t opid_create() {
+  host_op_id = id_create();
+  return host_op_id;
+}
+
+static uint64_t opid_get() { return host_op_id; }
+
+static uint64_t regionid_create() {
+  ompt_target_data.value = id_create();
+  return ompt_target_data.value;
+}
+
+static uint64_t regionid_get() { return ompt_target_data.value; }
+
+static void LIBOMPTARGET_GET_TARGET_OPID(uint64_t *device_num,
+                                         ompt_id_t *target_id,
+                                         ompt_id_t *host_op_id) {
+  *host_op_id = ompt_target_region_opid;
+}
+
+/*****************************************************************************
+ * OMPT callbacks
+ *****************************************************************************/
+
+void OmptInterface::setStateHelper(void *Frame, void *Code, int Flags,
+                                   int State) {
+  EnterFrame = Frame;
+  Codeptr = Code;
+  if (ompt_set_frame_enter_fn)
+    InterfaceState = ompt_set_frame_enter_fn(Frame, Flags, State);
+}
+
+void OmptInterface::setState(void *Frame, void *Codeptr) {
+  setStateHelper(Frame, Codeptr, OMPT_FRAME_FLAGS, ompt_state_work_parallel);
+}
+
+void OmptInterface::clearState() {
+  setStateHelper(nullptr /* Frame */, nullptr /* Code */, 0 /* Flags */,
+                 InterfaceState);
+}
+
+void OmptInterface::beginTargetRegion() {
+  // set up task region state
+  ompt_task_data = ompt_get_task_data_fn();
+  ompt_target_task_data = ompt_get_target_task_data_fn();
+
+  *ompt_task_data = ompt_data_none;
+  *ompt_target_task_data = ompt_data_none;
+  ompt_target_data = ompt_data_none;
+}
+
+void OmptInterface::endTargetRegion() {
+  ompt_task_data = 0;
+  ompt_target_task_data = 0;
+  ompt_target_data = ompt_data_none;
+}
+
+void OmptInterface::announceTargetRegion(const char *Name) {
+  DP("in OmptInterface::target_region_%s target_id=%lu\n", Name,
+     ompt_target_data.value);
+}
+
+void OmptInterface::beginTargetOperation() {
+  DP("in ompt_target_region_begin (ompt_target_region_opid = %lu)\n",
+     ompt_target_data.value);
+}
+
+void OmptInterface::endTargetOperation() {
+  DP("in ompt_target_region_end (ompt_target_region_opid = %lu)\n",
+     ompt_target_data.value);
+}
+
+/*****************************************************************************
+ * OMPT public operations
+ *****************************************************************************/
+
+// FIXME: optional implementation of target map?
+
+void OmptInterface::beginTargetDataAlloc(int64_t DeviceId, void *HstPtrBegin,
+                                         size_t Size, void *Code) {
+  ompt_device_callbacks.OmptCallbackTargetDataOpEmi(
+      ompt_scope_begin, ompt_target_task_data, &ompt_target_data,
+      ompt_target_data_alloc, HstPtrBegin, DeviceId, NULL, 0, Size, Code,
+      opid_create, &ompt_target_region_opid);
+  beginTargetOperation();
+}
+
+void OmptInterface::endTargetDataAlloc(int64_t DeviceId, void *HstPtrBegin,
+                                       size_t Size, void *Code) {
+  ompt_device_callbacks.OmptCallbackTargetDataOpEmi(
+      ompt_scope_end, ompt_target_task_data, &ompt_target_data,
+      ompt_target_data_alloc, HstPtrBegin, DeviceId, NULL, 0, Size, Code,
+      opid_get, &ompt_target_region_opid);
+  endTargetOperation();
+}
+
+void OmptInterface::beginTargetDataSubmit(int64_t DeviceId, void *TgtPtrBegin,
+                                          void *HstPtrBegin, size_t Size,
+                                          void *Code) {
+  ompt_device_callbacks.OmptCallbackTargetDataOpEmi(
+      ompt_scope_begin, ompt_target_task_data, &ompt_target_data,
+      ompt_target_data_transfer_to_device, HstPtrBegin, 0, TgtPtrBegin,
+      DeviceId, Size, Code, opid_create, &ompt_target_region_opid);
+  beginTargetOperation();
+}
+
+void OmptInterface::endTargetDataSubmit(int64_t DeviceId, void *TgtPtrBegin,
+                                        void *HstPtrBegin, size_t Size,
+                                        void *Code) {
+  ompt_device_callbacks.OmptCallbackTargetDataOpEmi(
+      ompt_scope_end, ompt_target_task_data, &ompt_target_data,
+      ompt_target_data_transfer_to_device, HstPtrBegin, 0 /* DeviceId */,
+      TgtPtrBegin, DeviceId, Size, Code, opid_get, &ompt_target_region_opid);
+  endTargetOperation();
+}
+
+void OmptInterface::beginTargetDataDelete(int64_t DeviceId, void *TgtPtrBegin,
+                                          void *Code) {
+  ompt_device_callbacks.OmptCallbackTargetDataOpEmi(
+      ompt_scope_begin, ompt_target_task_data, &ompt_target_data,
+      ompt_target_data_delete, TgtPtrBegin, DeviceId, NULL /* DstAddr */,
+      0 /* DestDeviceNum */, 0 /* Bytes */, Code, opid_create,
+      &ompt_target_region_opid);
+  beginTargetOperation();
+}
+
+void OmptInterface::endTargetDataDelete(int64_t DeviceId, void *TgtPtrBegin,
+                                        void *Code) {
+  ompt_device_callbacks.OmptCallbackTargetDataOpEmi(
+      ompt_scope_end, ompt_target_task_data, &ompt_target_data,
+      ompt_target_data_delete, TgtPtrBegin, DeviceId, NULL /* DstAddr */,
+      0 /* DestDeviceNum */, 0 /* Bytes */, Code, opid_get,
+      &ompt_target_region_opid);
+  endTargetOperation();
+}
+
+void OmptInterface::beginTargetDataRetrieve(int64_t DeviceId, void *HstPtrBegin,
+                                            void *TgtPtrBegin, size_t Size,
+                                            void *Code) {
+  ompt_device_callbacks.OmptCallbackTargetDataOpEmi(
+      ompt_scope_begin, ompt_target_task_data, &ompt_target_data,
+      ompt_target_data_transfer_from_device, TgtPtrBegin, DeviceId, HstPtrBegin,
+      0 /* DeviceId */, Size, Code, opid_create, &ompt_target_region_opid);
+  beginTargetOperation();
+}
+
+void OmptInterface::endTargetDataRetrieve(int64_t DeviceId, void *HstPtrBegin,
+                                          void *TgtPtrBegin, size_t Size,
+                                          void *Code) {
+  ompt_device_callbacks.OmptCallbackTargetDataOpEmi(
+      ompt_scope_end, ompt_target_task_data, &ompt_target_data,
+      ompt_target_data_transfer_from_device, TgtPtrBegin, DeviceId, HstPtrBegin,
+      0 /* Device Id */, Size, Code, opid_get, &ompt_target_region_opid);
+  endTargetOperation();
+}
+
+void OmptInterface::beginTargetSubmit(unsigned int numTeams) {
+  ompt_device_callbacks.OmptCallbackTargetSubmitEmi(
+      ompt_scope_begin, &ompt_target_data, numTeams, opid_create,
+      &ompt_target_region_opid);
+}
+
+void OmptInterface::endTargetSubmit(unsigned int numTeams) {
+  ompt_device_callbacks.OmptCallbackTargetSubmitEmi(
+      ompt_scope_end, &ompt_target_data, numTeams, opid_get,
+      &ompt_target_region_opid);
+}
+
+void OmptInterface::beginTargetDataEnter(int64_t DeviceId, void *Code) {
+  beginTargetRegion();
+  ompt_device_callbacks.OmptCallbackTargetEmi(
+      ompt_target_enter_data, ompt_scope_begin, DeviceId, ompt_task_data,
+      ompt_target_task_data, &ompt_target_data, Code, regionid_create);
+}
+
+void OmptInterface::endTargetDataEnter(int64_t DeviceId, void *Code) {
+  ompt_device_callbacks.OmptCallbackTargetEmi(
+      ompt_target_enter_data, ompt_scope_end, DeviceId, ompt_task_data,
+      ompt_target_task_data, &ompt_target_data, Code, regionid_get);
+  endTargetRegion();
+}
+
+void OmptInterface::beginTargetDataExit(int64_t DeviceId, void *Code) {
+  beginTargetRegion();
+  ompt_device_callbacks.OmptCallbackTargetEmi(
+      ompt_target_exit_data, ompt_scope_begin, DeviceId, ompt_task_data,
+      ompt_target_task_data, &ompt_target_data, Code, regionid_create);
+  announceTargetRegion("begin");
+}
+
+void OmptInterface::endTargetDataExit(int64_t DeviceId, void *Code) {
+  ompt_device_callbacks.OmptCallbackTargetEmi(
+      ompt_target_exit_data, ompt_scope_end, DeviceId, ompt_task_data,
+      ompt_target_task_data, &ompt_target_data, Code, regionid_get);
+  endTargetRegion();
+}
+
+void OmptInterface::beginTargetUpdate(int64_t DeviceId, void *Code) {
+  beginTargetRegion();
+  ompt_device_callbacks.OmptCallbackTargetEmi(
+      ompt_target_update, ompt_scope_begin, DeviceId, ompt_task_data,
+      ompt_target_task_data, &ompt_target_data, Code, regionid_create);
+  announceTargetRegion("begin");
+}
+
+void OmptInterface::endTargetUpdate(int64_t DeviceId, void *Code) {
+  ompt_device_callbacks.OmptCallbackTargetEmi(
+      ompt_target_update, ompt_scope_end, DeviceId, ompt_task_data,
+      ompt_target_task_data, &ompt_target_data, Code, regionid_get);
+  endTargetRegion();
+}
+
+void OmptInterface::beginTarget(int64_t DeviceId, void *Code) {
+  beginTargetRegion();
+  ompt_device_callbacks.OmptCallbackTargetEmi(
+      ompt_target, ompt_scope_begin, DeviceId, ompt_task_data,
+      ompt_target_task_data, &ompt_target_data, Code, regionid_create);
+  announceTargetRegion("begin");
+}
+
+void OmptInterface::endTarget(int64_t DeviceId, void *Code) {
+  ompt_device_callbacks.OmptCallbackTargetEmi(
+      ompt_target, ompt_scope_end, DeviceId, ompt_task_data,
+      ompt_target_task_data, &ompt_target_data, Code, regionid_get);
+  endTargetRegion();
+}
+
+/*****************************************************************************
+ * OMPT device callback operations specific to libomptarget
+ *****************************************************************************/
+
+const char *OmptDeviceCallbacksTy::Documentation = 0;
+
+// Today, this is not called from libomptarget
+OmptDeviceTy *OmptDeviceCallbacksTy::lookupDevice(int DeviceNnum) {
+  DP("Lookup device should be invoked in the plugin\n");
+  abort();
+  return nullptr;
+}
+
+ompt_interface_fn_t
+OmptDeviceCallbacksTy::doLookup(const char *InterfaceFunctionName) {
+  if (strcmp(InterfaceFunctionName,
+             LIBOMPTARGET_STRINGIFY(LIBOMPTARGET_GET_TARGET_OPID)) == 0)
+    return (ompt_interface_fn_t)LIBOMPTARGET_GET_TARGET_OPID;
+
+  return ompt_device_callbacks.lookupCallback(InterfaceFunctionName);
+}
+
+#ifdef OMPT_SUPPORT
+static int ompt_libomptarget_initialize(ompt_function_lookup_t lookup,
+                                        int initial_device_num,
+                                        ompt_data_t *tool_data) {
+  DP("enter ompt_libomptarget_initialize!\n");
+  ompt_enabled = true;
+
+#define ompt_bind_name(fn)                                                     \
+  fn##_fn = (fn##_t)lookup(#fn);                                               \
+  DP("%s=%p\n", #fn, fnptr_to_ptr(fn##_fn));
+
+  ompt_bind_name(ompt_set_frame_enter);
+  ompt_bind_name(ompt_get_task_data);
+  ompt_bind_name(ompt_get_target_task_data);
+#undef ompt_bind_name
+
+  ompt_device_callbacks.registerCallbacks(lookup);
+  DP("exit ompt_libomptarget_initialize!\n");
+  return 0;
+}
+
+static void ompt_libomptarget_finalize(ompt_data_t *data) {
+  DP("enter ompt_libomptarget_finalize!\n");
+  libomptarget_rtl_finalizer.finalize();
+  ompt_enabled = false;
+  DP("exit ompt_libomptarget_finalize!\n");
+}
+
+/*****************************************************************************
+ * constructor
+ *****************************************************************************/
+/// Used to initialize callbacks implemented by the tool. This interface
+/// will lookup the callbacks table in libomp and assign them to the callbacks
+/// maintained in libomptarget.
+__attribute__((constructor(102))) static void ompt_init(void) {
+  /// Connect with libomp
+  static OmptLibraryConnectorTy LibompConnector("ompt_libomp");
+  static ompt_start_tool_result_t OmptResult;
+
+  // Initialize OmptResult with the init and fini functions that will be
+  // called by the connector
+  OmptResult.initialize = ompt_libomptarget_initialize;
+  OmptResult.finalize = ompt_libomptarget_finalize;
+  OmptResult.tool_data.value = 0;
+
+  // Initialize the device callbacks first
+  ompt_device_callbacks.init();
+
+  // Now call connect that causes the above init/fini functions to be called
+  LibompConnector.connect(&OmptResult);
+  DP("OMPT: Exit ompt_init\n");
+}
+#endif
+
+extern "C" {
+/// Used for connecting libomptarget with a plugin
+void ompt_libomptarget_connect(ompt_start_tool_result_t *result) {
+  DP("OMPT: Enter libomptarget_ompt_connect\n");
+  if (ompt_enabled && result) {
+    // Cache the fini function so that it can be invoked on exit
+    libomptarget_rtl_finalizer.registerRtl(result->finalize);
+    // Invoke the provided init function with the lookup function maintained
+    // in this library so that callbacks maintained by this library are obtained
+    result->initialize(OmptDeviceCallbacksTy::doLookup,
+                       0 /* initial_device_num */, nullptr /* tool_data */);
+  }
+  DP("OMPT: Leave libomptarget_ompt_connect\n");
+}
+}

openmp/libomptarget/src/omptarget.cpp

 #include "device.h"
 #include "private.h"
 #include "rtl.h"
 
 #include <cassert>
 #include <cstdint>
 #include <vector>
 
+#include "ompt_callback.h"
+
 int AsyncInfoTy::synchronize() {
   int Result = OFFLOAD_SUCCESS;
   if (AsyncInfo.Queue) {
     // If we have a queue we need to synchronize it now.
     Result = Device.synchronize(*this);
     assert(AsyncInfo.Queue == nullptr &&
            "The device plugin should have nulled the queue to indicate there "
            "are no outstanding actions!");
   }
   return Result;
 }
 
 void *&AsyncInfoTy::getVoidPtrLocation() {
   BufferLocations.push_back(nullptr);
   return BufferLocations.back();
 }
 
int target(ident_t *loc, DeviceTy &Device, void *HostPtr, int32_t ArgNum,
   // Launch device execution.
   void *TgtEntryPtr = TargetTable->EntriesBegin[TM->Index].addr;
   DP("Launching target execution %s with pointer " DPxMOD " (index=%d).\n",
      TargetTable->EntriesBegin[TM->Index].name, DPxPTR(TgtEntryPtr), TM->Index);
 
   {
     TIMESCOPE_WITH_NAME_AND_IDENT(
         IsTeamConstruct ? "runTargetTeamRegion" : "runTargetRegion", loc);
+
+    OMPT_IF_BUILT(OmptInterfaceTargetSubmitRAII TargetSubmitRAII(TeamNum));
+
     if (IsTeamConstruct)
       Ret = Device.runTeamRegion(TgtEntryPtr, &TgtArgs[0], &TgtOffsets[0],
                                  TgtArgs.size(), TeamNum, ThreadLimit,
                                  getLoopTripCount(DeviceId), AsyncInfo);
     else
       Ret = Device.runRegion(TgtEntryPtr, &TgtArgs[0], &TgtOffsets[0],
                              TgtArgs.size(), AsyncInfo);
   }

openmp/libomptarget/test/lit.cfg

 # Allow REQUIRES / UNSUPPORTED / XFAIL to work
 config.target_triple = [ ]
 for feature in config.test_compiler_features:
     config.available_features.add(feature)
 
 if config.libomptarget_debug:
   config.available_features.add('libomptarget-debug')
 
+if config.has_libomptarget_ompt:
+  config.available_features.add('ompt')
+  
 config.available_features.add(config.libomptarget_current_target)
 
 # Determine whether the test system supports unified memory.
 # For CUDA, this is the case with compute capability 70 (Volta) or higher.
 # For all other targets, we currently assume it is.
 supports_unified_shared_memory = True
 if config.libomptarget_current_target.startswith('nvptx'):
   try:

openmp/libomptarget/test/lit.site.cfg.in

 config.omp_header_directory = "@LIBOMPTARGET_OPENMP_HEADER_FOLDER@"
 config.omp_host_rtl_directory = "@LIBOMPTARGET_OPENMP_HOST_RTL_FOLDER@"
 config.operating_system = "@CMAKE_SYSTEM_NAME@"
 config.libomptarget_all_targets = "@LIBOMPTARGET_ALL_TARGETS@".split()
 config.libomptarget_current_target = "@CURRENT_TARGET@"
 config.libomptarget_filecheck = "@OPENMP_FILECHECK_EXECUTABLE@"
 config.libomptarget_not = "@OPENMP_NOT_EXECUTABLE@"
 config.libomptarget_debug = @LIBOMPTARGET_DEBUG@
+config.has_libomptarget_ompt = @LIBOMPTARGET_OMPT_SUPPORT@
 
 # Let the main config do the real work.
 lit_config.load_config(config, "@CMAKE_CURRENT_SOURCE_DIR@/lit.cfg")

openmp/libomptarget/test/ompt/callbacks.h

+#include <assert.h>
+#include <stdlib.h>
+
+// Tool related code below
+#include <omp-tools.h>
+
+// For EMI callbacks
+ompt_id_t next_op_id = 0x8000000000000001;
+
+// OMPT callbacks
+
+// Synchronous callbacks
+static void on_ompt_callback_device_initialize(int device_num, const char *type,
+                                               ompt_device_t *device,
+                                               ompt_function_lookup_t lookup,
+                                               const char *documentation) {
+  printf("Callback Init: device_num=%d type=%s device=%p lookup=%p doc=%p\n",
+         device_num, type, device, lookup, documentation);
+}
+
+static void on_ompt_callback_device_finalize(int device_num) {
+  printf("Callback Fini: device_num=%d\n", device_num);
+}
+
+static void on_ompt_callback_device_load(int device_num, const char *filename,
+                                         int64_t offset_in_file,
+                                         void *vma_in_file, size_t bytes,
+                                         void *host_addr, void *device_addr,
+                                         uint64_t module_id) {
+  printf("Callback Load: device_num:%d module_id:%lu filename:%s host_adddr:%p "
+         "device_addr:%p bytes:%lu\n",
+         device_num, module_id, filename, host_addr, device_addr, bytes);
+}
+
+static void on_ompt_callback_target_data_op(
+    ompt_id_t target_id, ompt_id_t host_op_id, ompt_target_data_op_t optype,
+    void *src_addr, int src_device_num, void *dest_addr, int dest_device_num,
+    size_t bytes, const void *codeptr_ra) {
+  assert(codeptr_ra != 0 && "Unexpected null codeptr");
+  // Both src and dest must not be null
+  assert((src_addr != 0 || dest_addr != 0) && "Both src and dest addr null");
+  printf("  Callback DataOp: target_id=%lu host_op_id=%lu optype=%d src=%p "
+         "src_device_num=%d "
+         "dest=%p dest_device_num=%d bytes=%lu code=%p\n",
+         target_id, host_op_id, optype, src_addr, src_device_num, dest_addr,
+         dest_device_num, bytes, codeptr_ra);
+}
+
+static void on_ompt_callback_target(ompt_target_t kind,
+                                    ompt_scope_endpoint_t endpoint,
+                                    int device_num, ompt_data_t *task_data,
+                                    ompt_id_t target_id,
+                                    const void *codeptr_ra) {
+  assert(codeptr_ra != 0 && "Unexpected null codeptr");
+  printf("Callback Target: target_id=%lu kind=%d endpoint=%d device_num=%d "
+         "code=%p\n",
+         target_id, kind, endpoint, device_num, codeptr_ra);
+}
+
+static void on_ompt_callback_target_submit(ompt_id_t target_id,
+                                           ompt_id_t host_op_id,
+                                           unsigned int requested_num_teams) {
+  printf("  Callback Submit: target_id=%lu host_op_id=%lu req_num_teams=%d\n",
+         target_id, host_op_id, requested_num_teams);
+}
+
+static void on_ompt_callback_target_map(ompt_id_t target_id,
+                                        unsigned int nitems, void **host_addr,
+                                        void **device_addr, size_t *bytes,
+                                        unsigned int *mapping_flags,
+                                        const void *codeptr_ra) {
+  printf("Target map callback is unimplemented\n");
+  abort();
+}
+
+static void on_ompt_callback_target_data_op_emi(
+    ompt_scope_endpoint_t endpoint, ompt_data_t *target_task_data,
+    ompt_data_t *target_data, ompt_id_t *host_op_id,
+    ompt_target_data_op_t optype, void *src_addr, int src_device_num,
+    void *dest_addr, int dest_device_num, size_t bytes,
+    const void *codeptr_ra) {
+  assert(codeptr_ra != 0 && "Unexpected null codeptr");
+  // Both src and dest must not be null
+  assert((src_addr != 0 || dest_addr != 0) && "Both src and dest addr null");
+  if (endpoint == ompt_scope_begin)
+    *host_op_id = next_op_id++;
+  printf("  Callback DataOp EMI: endpoint=%d optype=%d target_task_data=%p "
+         "(0x%lx) target_data=%p (0x%lx) host_op_id=%p (0x%lx) src=%p "
+         "src_device_num=%d "
+         "dest=%p dest_device_num=%d bytes=%lu code=%p\n",
+         endpoint, optype, target_task_data, target_task_data->value,
+         target_data, target_data->value, host_op_id, *host_op_id, src_addr,
+         src_device_num, dest_addr, dest_device_num, bytes, codeptr_ra);
+}
+
+static void on_ompt_callback_target_emi(ompt_target_t kind,
+                                        ompt_scope_endpoint_t endpoint,
+                                        int device_num, ompt_data_t *task_data,
+                                        ompt_data_t *target_task_data,
+                                        ompt_data_t *target_data,
+                                        const void *codeptr_ra) {
+  assert(codeptr_ra != 0 && "Unexpected null codeptr");
+  if (endpoint == ompt_scope_begin)
+    target_data->value = next_op_id++;
+  printf("Callback Target EMI: kind=%d endpoint=%d device_num=%d task_data=%p "
+         "(0x%lx) target_task_data=%p (0x%lx) target_data=%p (0x%lx) code=%p\n",
+         kind, endpoint, device_num, task_data, task_data->value,
+         target_task_data, target_task_data->value, target_data,
+         target_data->value, codeptr_ra);
+}
+
+static void on_ompt_callback_target_submit_emi(
+    ompt_scope_endpoint_t endpoint, ompt_data_t *target_data,
+    ompt_id_t *host_op_id, unsigned int requested_num_teams) {
+  printf("  Callback Submit EMI: endpoint=%d  req_num_teams=%d target_data=%p "
+         "(0x%lx) host_op_id=%p (0x%lx)\n",
+         endpoint, requested_num_teams, target_data, target_data->value,
+         host_op_id, *host_op_id);
+}
+
+static void on_ompt_callback_target_map_emi(ompt_data_t *target_data,
+                                            unsigned int nitems,
+                                            void **host_addr,
+                                            void **device_addr, size_t *bytes,
+                                            unsigned int *mapping_flags,
+                                            const void *codeptr_ra) {
+  printf("Target map emi callback is unimplemented\n");
+  abort();
+}

openmp/libomptarget/test/ompt/register_both.h

+#include <omp-tools.h>
+
+// From openmp/runtime/test/ompt/callback.h
+#define register_ompt_callback_t(name, type)                                   \
+  do {                                                                         \
+    type f_##name = &on_##name;                                                \
+    if (ompt_set_callback(name, (ompt_callback_t)f_##name) == ompt_set_never)  \
+      printf("0: Could not register callback '" #name "'\n");                  \
+  } while (0)
+
+#define register_ompt_callback(name) register_ompt_callback_t(name, name##_t)
+
+// OMPT entry point handles
+static ompt_set_callback_t ompt_set_callback = 0;
+
+// Init functions
+int ompt_initialize(ompt_function_lookup_t lookup, int initial_device_num,
+                    ompt_data_t *tool_data) {
+  ompt_set_callback = (ompt_set_callback_t)lookup("ompt_set_callback");
+
+  if (!ompt_set_callback)
+    return 0; // failed
+
+  register_ompt_callback(ompt_callback_device_initialize);
+  register_ompt_callback(ompt_callback_device_finalize);
+  register_ompt_callback(ompt_callback_device_load);
+  register_ompt_callback(ompt_callback_target_data_op_emi);
+  register_ompt_callback(ompt_callback_target_data_op);
+  register_ompt_callback(ompt_callback_target);
+  register_ompt_callback(ompt_callback_target_emi);
+  register_ompt_callback(ompt_callback_target_submit);
+
+  return 1; // success
+}
+
+void ompt_finalize(ompt_data_t *tool_data) {}
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+ompt_start_tool_result_t *ompt_start_tool(unsigned int omp_version,
+                                          const char *runtime_version) {
+  static ompt_start_tool_result_t ompt_start_tool_result = {&ompt_initialize,
+                                                            &ompt_finalize, 0};
+  return &ompt_start_tool_result;
+}
+#ifdef __cplusplus
+}
+#endif

openmp/libomptarget/test/ompt/register_emi.h

+#include <omp-tools.h>
+
+// From openmp/runtime/test/ompt/callback.h
+#define register_ompt_callback_t(name, type)                                   \
+  do {                                                                         \
+    type f_##name = &on_##name;                                                \
+    if (ompt_set_callback(name, (ompt_callback_t)f_##name) == ompt_set_never)  \
+      printf("0: Could not register callback '" #name "'\n");                  \
+  } while (0)
+
+#define register_ompt_callback(name) register_ompt_callback_t(name, name##_t)
+
+// OMPT entry point handles
+static ompt_set_callback_t ompt_set_callback = 0;
+
+// Init functions
+int ompt_initialize(ompt_function_lookup_t lookup, int initial_device_num,
+                    ompt_data_t *tool_data) {
+  ompt_set_callback = (ompt_set_callback_t)lookup("ompt_set_callback");
+
+  if (!ompt_set_callback)
+    return 0; // failed
+
+  register_ompt_callback(ompt_callback_device_initialize);
+  register_ompt_callback(ompt_callback_device_finalize);
+  register_ompt_callback(ompt_callback_device_load);
+  register_ompt_callback(ompt_callback_target_data_op_emi);
+  register_ompt_callback(ompt_callback_target_emi);
+  register_ompt_callback(ompt_callback_target_submit_emi);
+
+  return 1; // success
+}
+
+void ompt_finalize(ompt_data_t *tool_data) {}
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+ompt_start_tool_result_t *ompt_start_tool(unsigned int omp_version,
+                                          const char *runtime_version) {
+  static ompt_start_tool_result_t ompt_start_tool_result = {&ompt_initialize,
+                                                            &ompt_finalize, 0};
+  return &ompt_start_tool_result;
+}
+#ifdef __cplusplus
+}
+#endif

openmp/libomptarget/test/ompt/register_emi_map.h

+#include <omp-tools.h>
+
+// From openmp/runtime/test/ompt/callback.h
+#define register_ompt_callback_t(name, type)                                   \
+  do {                                                                         \
+    type f_##name = &on_##name;                                                \
+    if (ompt_set_callback(name, (ompt_callback_t)f_##name) == ompt_set_never)  \
+      printf("0: Could not register callback '" #name "'\n");                  \
+  } while (0)
+
+#define register_ompt_callback(name) register_ompt_callback_t(name, name##_t)
+
+// OMPT entry point handles
+static ompt_set_callback_t ompt_set_callback = 0;
+
+// Init functions
+int ompt_initialize(ompt_function_lookup_t lookup, int initial_device_num,
+                    ompt_data_t *tool_data) {
+  ompt_set_callback = (ompt_set_callback_t)lookup("ompt_set_callback");
+
+  if (!ompt_set_callback)
+    return 0; // failed
+
+  register_ompt_callback(ompt_callback_device_initialize);
+  register_ompt_callback(ompt_callback_device_finalize);
+  register_ompt_callback(ompt_callback_device_load);
+  register_ompt_callback(ompt_callback_target_data_op_emi);
+  register_ompt_callback(ompt_callback_target_emi);
+  register_ompt_callback(ompt_callback_target_submit_emi);
+  register_ompt_callback(ompt_callback_target_map_emi);
+
+  return 1; // success
+}
+
+void ompt_finalize(ompt_data_t *tool_data) {}
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+ompt_start_tool_result_t *ompt_start_tool(unsigned int omp_version,
+                                          const char *runtime_version) {
+  static ompt_start_tool_result_t ompt_start_tool_result = {&ompt_initialize,
+                                                            &ompt_finalize, 0};
+  return &ompt_start_tool_result;
+}
+#ifdef __cplusplus
+}
+#endif

openmp/libomptarget/test/ompt/register_no_device_init.h

+#include <omp-tools.h>
+
+// From openmp/runtime/test/ompt/callback.h
+#define register_ompt_callback_t(name, type)                                   \
+  do {                                                                         \
+    type f_##name = &on_##name;                                                \
+    if (ompt_set_callback(name, (ompt_callback_t)f_##name) == ompt_set_never)  \
+      printf("0: Could not register callback '" #name "'\n");                  \
+  } while (0)
+
+#define register_ompt_callback(name) register_ompt_callback_t(name, name##_t)
+
+// OMPT entry point handles
+static ompt_set_callback_t ompt_set_callback = 0;
+
+// Init functions
+int ompt_initialize(ompt_function_lookup_t lookup, int initial_device_num,
+                    ompt_data_t *tool_data) {
+  ompt_set_callback = (ompt_set_callback_t)lookup("ompt_set_callback");
+
+  if (!ompt_set_callback)
+    return 0; // failed
+
+  // If no device init callback is registered, the other callbacks won't be
+  // activated.
+  register_ompt_callback(ompt_callback_device_load);
+  register_ompt_callback(ompt_callback_target_data_op);
+  register_ompt_callback(ompt_callback_target);
+  register_ompt_callback(ompt_callback_target_submit);
+
+  return 1; // success
+}
+
+void ompt_finalize(ompt_data_t *tool_data) {}
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+ompt_start_tool_result_t *ompt_start_tool(unsigned int omp_version,
+                                          const char *runtime_version) {
+  static ompt_start_tool_result_t ompt_start_tool_result = {&ompt_initialize,
+                                                            &ompt_finalize, 0};
+  return &ompt_start_tool_result;
+}
+#ifdef __cplusplus
+}
+#endif

openmp/libomptarget/test/ompt/register_non_emi.h

+#include <omp-tools.h>
+
+// From openmp/runtime/test/ompt/callback.h
+#define register_ompt_callback_t(name, type)                                   \
+  do {                                                                         \
+    type f_##name = &on_##name;                                                \
+    if (ompt_set_callback(name, (ompt_callback_t)f_##name) == ompt_set_never)  \
+      printf("0: Could not register callback '" #name "'\n");                  \
+  } while (0)
+
+#define register_ompt_callback(name) register_ompt_callback_t(name, name##_t)
+
+// OMPT entry point handles
+static ompt_set_callback_t ompt_set_callback = 0;
+
+// Init functions
+int ompt_initialize(ompt_function_lookup_t lookup, int initial_device_num,
+                    ompt_data_t *tool_data) {
+  ompt_set_callback = (ompt_set_callback_t)lookup("ompt_set_callback");
+
+  if (!ompt_set_callback)
+    return 0; // failed
+
+  register_ompt_callback(ompt_callback_device_initialize);
+  register_ompt_callback(ompt_callback_device_finalize);
+  register_ompt_callback(ompt_callback_device_load);
+  register_ompt_callback(ompt_callback_target_data_op);
+  register_ompt_callback(ompt_callback_target);
+  register_ompt_callback(ompt_callback_target_submit);
+
+  return 1; // success
+}
+
+void ompt_finalize(ompt_data_t *tool_data) {}
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+ompt_start_tool_result_t *ompt_start_tool(unsigned int omp_version,
+                                          const char *runtime_version) {
+  static ompt_start_tool_result_t ompt_start_tool_result = {&ompt_initialize,
+                                                            &ompt_finalize, 0};
+  return &ompt_start_tool_result;
+}
+#ifdef __cplusplus
+}
+#endif

openmp/libomptarget/test/ompt/register_non_emi_map.h

+#include <omp-tools.h>
+
+// From openmp/runtime/test/ompt/callback.h
+#define register_ompt_callback_t(name, type)                                   \
+  do {                                                                         \
+    type f_##name = &on_##name;                                                \
+    if (ompt_set_callback(name, (ompt_callback_t)f_##name) == ompt_set_never)  \
+      printf("0: Could not register callback '" #name "'\n");                  \
+  } while (0)
+
+#define register_ompt_callback(name) register_ompt_callback_t(name, name##_t)
+
+// OMPT entry point handles
+static ompt_set_callback_t ompt_set_callback = 0;
+
+// Init functions
+int ompt_initialize(ompt_function_lookup_t lookup, int initial_device_num,
+                    ompt_data_t *tool_data) {
+  ompt_set_callback = (ompt_set_callback_t)lookup("ompt_set_callback");
+
+  if (!ompt_set_callback)
+    return 0; // failed
+
+  register_ompt_callback(ompt_callback_device_initialize);
+  register_ompt_callback(ompt_callback_device_finalize);
+  register_ompt_callback(ompt_callback_device_load);
+  register_ompt_callback(ompt_callback_target_data_op);
+  register_ompt_callback(ompt_callback_target);
+  register_ompt_callback(ompt_callback_target_submit);
+  register_ompt_callback(ompt_callback_target_map);
+
+  return 1; // success
+}
+
+void ompt_finalize(ompt_data_t *tool_data) {}
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+ompt_start_tool_result_t *ompt_start_tool(unsigned int omp_version,
+                                          const char *runtime_version) {
+  static ompt_start_tool_result_t ompt_start_tool_result = {&ompt_initialize,
+                                                            &ompt_finalize, 0};
+  return &ompt_start_tool_result;
+}
+#ifdef __cplusplus
+}
+#endif

openmp/libomptarget/test/ompt/register_wrong_return.h

+#include <omp-tools.h>
+
+// From openmp/runtime/test/ompt/callback.h
+#define register_ompt_callback_t(name, type)                                   \
+  do {                                                                         \
+    type f_##name = &on_##name;                                                \
+    if (ompt_set_callback(name, (ompt_callback_t)f_##name) == ompt_set_never)  \
+      printf("0: Could not register callback '" #name "'\n");                  \
+  } while (0)
+
+#define register_ompt_callback(name) register_ompt_callback_t(name, name##_t)
+
+// OMPT entry point handles
+static ompt_set_callback_t ompt_set_callback = 0;
+
+// Init functions
+int ompt_initialize(ompt_function_lookup_t lookup, int initial_device_num,
+                    ompt_data_t *tool_data) {
+  ompt_set_callback = (ompt_set_callback_t)lookup("ompt_set_callback");
+
+  if (!ompt_set_callback)
+    return 1; // failed but wrongly returning 1
+
+  register_ompt_callback(ompt_callback_device_initialize);
+  register_ompt_callback(ompt_callback_device_finalize);
+  register_ompt_callback(ompt_callback_device_load);
+  register_ompt_callback(ompt_callback_target_data_op);
+  register_ompt_callback(ompt_callback_target);
+  register_ompt_callback(ompt_callback_target_submit);
+
+  return 0; // success but should return 1 according to the spec
+}
+
+void ompt_finalize(ompt_data_t *tool_data) {}
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+ompt_start_tool_result_t *ompt_start_tool(unsigned int omp_version,
+                                          const char *runtime_version) {
+  static ompt_start_tool_result_t ompt_start_tool_result = {&ompt_initialize,
+                                                            &ompt_finalize, 0};
+  return &ompt_start_tool_result;
+}
+#ifdef __cplusplus
+}
+#endif

openmp/libomptarget/test/ompt/veccopy.c

+// RUN: %libomptarget-compile-run-and-check-generic
+// REQUIRES: ompt
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-newDriver
+// UNSUPPORTED: x86_64-pc-linux-gnu
+// UNSUPPORTED: x86_64-pc-linux-gnu-newDriver
+
+/*
+ * Example OpenMP program that registers non-EMI callbacks
+ */
+
+#include <omp.h>
+#include <stdio.h>
+
+#include "callbacks.h"
+#include "register_non_emi.h"
+
+int main() {
+  int N = 100000;
+
+  int a[N];
+  int b[N];
+
+  int i;
+
+  for (i = 0; i < N; i++)
+    a[i] = 0;
+
+  for (i = 0; i < N; i++)
+    b[i] = i;
+
+#pragma omp target parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+#pragma omp target teams distribute parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+  int rc = 0;
+  for (i = 0; i < N; i++)
+    if (a[i] != b[i]) {
+      rc++;
+      printf("Wrong value: a[%d]=%d\n", i, a[i]);
+    }
+
+  if (!rc)
+    printf("Success\n");
+
+  return rc;
+}
+
+/// CHECK: Callback Init:
+/// CHECK: Callback Load:
+/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=1
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2
+/// CHECK: Callback Submit: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] req_num_teams=1
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4
+/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=2
+
+/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=1
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2
+/// CHECK: Callback Submit: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] req_num_teams=0
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4
+/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=2
+/// CHECK: Callback Fini:

openmp/libomptarget/test/ompt/veccopy_disallow_both.c

+// RUN: %libomptarget-compile-run-and-check-generic
+// REQUIRES: ompt
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-newDriver
+// UNSUPPORTED: x86_64-pc-linux-gnu
+// UNSUPPORTED: x86_64-pc-linux-gnu-newDriver
+
+/*
+ * Example OpenMP program that shows that both EMI and non-EMI
+ * callbacks cannot be registered for the same type. In the
+ * current implementation, the EMI callback overrides the non-EMI
+ * callback.
+ */
+
+#include <omp.h>
+#include <stdio.h>
+
+#include "callbacks.h"
+#include "register_both.h"
+
+int main() {
+  int N = 100000;
+
+  int a[N];
+  int b[N];
+
+  int i;
+
+  for (i = 0; i < N; i++)
+    a[i] = 0;
+
+  for (i = 0; i < N; i++)
+    b[i] = i;
+
+#pragma omp target parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+#pragma omp target teams distribute parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+  int rc = 0;
+  for (i = 0; i < N; i++)
+    if (a[i] != b[i]) {
+      rc++;
+      printf("Wrong value: a[%d]=%d\n", i, a[i]);
+    }
+
+  if (!rc)
+    printf("Success\n");
+
+  return rc;
+}
+
+/// CHECK: Callback Init:
+/// CHECK: Callback Load:
+/// CHECK: Callback Target EMI: kind=1 endpoint=1
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=2
+/// CHECK: Callback Submit: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] req_num_teams=1
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=4
+/// CHECK: Callback Target EMI: kind=1 endpoint=2
+/// CHECK: Callback Target EMI: kind=1 endpoint=1
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=2
+/// CHECK: Callback Submit: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] req_num_teams=0
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=4
+/// CHECK: Callback Target EMI: kind=1 endpoint=2
+/// CHECK: Callback Fini:

openmp/libomptarget/test/ompt/veccopy_emi.c

+// RUN: %libomptarget-compile-run-and-check-generic
+// REQUIRES: ompt
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-newDriver
+// UNSUPPORTED: x86_64-pc-linux-gnu
+// UNSUPPORTED: x86_64-pc-linux-gnu-newDriver
+
+/*
+ * Example OpenMP program that registers EMI callbacks
+ */
+
+#include <assert.h>
+#include <omp.h>
+#include <stdio.h>
+
+#include "callbacks.h"
+#include "register_emi.h"
+
+int main() {
+  int N = 100000;
+
+  int a[N];
+  int b[N];
+
+  int i;
+
+  for (i = 0; i < N; i++)
+    a[i] = 0;
+
+  for (i = 0; i < N; i++)
+    b[i] = i;
+
+#pragma omp target parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+#pragma omp target teams distribute parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+  int rc = 0;
+  for (i = 0; i < N; i++)
+    if (a[i] != b[i]) {
+      rc++;
+      printf("Wrong value: a[%d]=%d\n", i, a[i]);
+    }
+
+  if (!rc)
+    printf("Success\n");
+
+  return rc;
+}
+
+/// CHECK: Callback Init:
+/// CHECK: Callback Load:
+/// CHECK: Callback Target EMI: kind=1 endpoint=1
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=2
+/// CHECK: Callback Submit EMI: endpoint=1 req_num_teams=1
+/// CHECK: Callback Submit EMI: endpoint=2 req_num_teams=1
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=4
+/// CHECK: Callback Target EMI: kind=1 endpoint=2
+/// CHECK: Callback Target EMI: kind=1 endpoint=1
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=2
+/// CHECK: Callback Submit EMI: endpoint=1 req_num_teams=0
+/// CHECK: Callback Submit EMI: endpoint=2 req_num_teams=0
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=4
+/// CHECK: Callback Target EMI: kind=1 endpoint=2
+/// CHECK: Callback Fini:

openmp/libomptarget/test/ompt/veccopy_emi_map.c

+// RUN: %libomptarget-compile-run-and-check-generic
+// REQUIRES: ompt
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-newDriver
+// UNSUPPORTED: x86_64-pc-linux-gnu
+// UNSUPPORTED: x86_64-pc-linux-gnu-newDriver
+
+/*
+ * Example OpenMP program that shows that map-EMI callbacks are not supported.
+ */
+
+#include <assert.h>
+#include <omp.h>
+#include <stdio.h>
+
+#include "callbacks.h"
+#include "register_emi_map.h"
+
+int main() {
+  int N = 100000;
+
+  int a[N];
+  int b[N];
+
+  int i;
+
+  for (i = 0; i < N; i++)
+    a[i] = 0;
+
+  for (i = 0; i < N; i++)
+    b[i] = i;
+
+#pragma omp target parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+#pragma omp target teams distribute parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+  int rc = 0;
+  for (i = 0; i < N; i++)
+    if (a[i] != b[i]) {
+      rc++;
+      printf("Wrong value: a[%d]=%d\n", i, a[i]);
+    }
+
+  if (!rc)
+    printf("Success\n");
+
+  return rc;
+}
+
+/// CHECK: 0: Could not register callback 'ompt_callback_target_map_emi'
+/// CHECK: Callback Init:
+/// CHECK: Callback Load:
+/// CHECK: Callback Target EMI: kind=1 endpoint=1
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=2
+/// CHECK: Callback Submit EMI: endpoint=1 req_num_teams=1
+/// CHECK: Callback Submit EMI: endpoint=2 req_num_teams=1
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=4
+/// CHECK: Callback Target EMI: kind=1 endpoint=2
+/// CHECK: Callback Target EMI: kind=1 endpoint=1
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=2
+/// CHECK: Callback Submit EMI: endpoint=1 req_num_teams=0
+/// CHECK: Callback Submit EMI: endpoint=2 req_num_teams=0
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=1 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=2 optype=4
+/// CHECK: Callback Target EMI: kind=1 endpoint=2
+/// CHECK: Callback Fini:

openmp/libomptarget/test/ompt/veccopy_map.c

+// RUN: %libomptarget-compile-run-and-check-generic
+// REQUIRES: ompt
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-newDriver
+// UNSUPPORTED: x86_64-pc-linux-gnu
+// UNSUPPORTED: x86_64-pc-linux-gnu-newDriver
+
+/*
+ * Example OpenMP program that shows that map callbacks are not supported.
+ */
+
+#include <omp.h>
+#include <stdio.h>
+
+#include "callbacks.h"
+#include "register_non_emi_map.h"
+
+int main() {
+  int N = 100000;
+
+  int a[N];
+  int b[N];
+
+  int i;
+
+  for (i = 0; i < N; i++)
+    a[i] = 0;
+
+  for (i = 0; i < N; i++)
+    b[i] = i;
+
+#pragma omp target parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+#pragma omp target teams distribute parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+  int rc = 0;
+  for (i = 0; i < N; i++)
+    if (a[i] != b[i]) {
+      rc++;
+      printf("Wrong value: a[%d]=%d\n", i, a[i]);
+    }
+
+  if (!rc)
+    printf("Success\n");
+
+  return rc;
+}
+
+/// CHECK: 0: Could not register callback 'ompt_callback_target_map'
+/// CHECK: Callback Init:
+/// CHECK: Callback Load:
+/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=1
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2
+/// CHECK: Callback Submit: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] req_num_teams=1
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4
+/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=2
+
+/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=1
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2
+/// CHECK: Callback Submit: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] req_num_teams=0
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4
+/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=2
+/// CHECK: Callback Fini:

openmp/libomptarget/test/ompt/veccopy_no_device_init.c

+// RUN: %libomptarget-compile-run-and-check-generic
+// REQUIRES: ompt
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-newDriver
+// UNSUPPORTED: x86_64-pc-linux-gnu
+// UNSUPPORTED: x86_64-pc-linux-gnu-newDriver
+
+/*
+ * Example OpenMP program that shows that if no device init callback
+ * is registered, the other callbacks won't be activated.
+ */
+
+#include <omp.h>
+#include <stdio.h>
+
+#include "callbacks.h"
+#include "register_no_device_init.h"
+
+int main() {
+  int N = 100000;
+
+  int a[N];
+  int b[N];
+
+  int i;
+
+  for (i = 0; i < N; i++)
+    a[i] = 0;
+
+  for (i = 0; i < N; i++)
+    b[i] = i;
+
+#pragma omp target parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+#pragma omp target teams distribute parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+  int rc = 0;
+  for (i = 0; i < N; i++)
+    if (a[i] != b[i]) {
+      rc++;
+      printf("Wrong value: a[%d]=%d\n", i, a[i]);
+    }
+
+  if (!rc)
+    printf("Success\n");
+
+  return rc;
+}
+
+/// CHECK-NOT: Callback Init:
+/// CHECK-NOT: Callback Load:
+/// CHECK-NOT: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=1
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2
+/// CHECK-NOT: Callback Submit: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] req_num_teams=1
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4
+/// CHECK-NOT: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=2
+
+/// CHECK-NOT: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=1
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2
+/// CHECK-NOT: Callback Submit: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] req_num_teams=0
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4
+/// CHECK-NOT: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=2
+/// CHECK-NOT: Callback Fini:

openmp/libomptarget/test/ompt/veccopy_wrong_return.c

+// RUN: %libomptarget-compile-run-and-check-generic
+// REQUIRES: ompt
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-newDriver
+// UNSUPPORTED: x86_64-pc-linux-gnu
+// UNSUPPORTED: x86_64-pc-linux-gnu-newDriver
+
+/*
+ * Example OpenMP program that shows that if the initialize function
+ * returns the wrong status code, the callbacks won't be activated.
+ */
+
+#include <omp.h>
+#include <stdio.h>
+
+#include "callbacks.h"
+#include "register_wrong_return.h"
+
+int main() {
+  int N = 100000;
+
+  int a[N];
+  int b[N];
+
+  int i;
+
+  for (i = 0; i < N; i++)
+    a[i] = 0;
+
+  for (i = 0; i < N; i++)
+    b[i] = i;
+
+#pragma omp target parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+#pragma omp target teams distribute parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+  int rc = 0;
+  for (i = 0; i < N; i++)
+    if (a[i] != b[i]) {
+      rc++;
+      printf("Wrong value: a[%d]=%d\n", i, a[i]);
+    }
+
+  if (!rc)
+    printf("Success\n");
+
+  return rc;
+}
+
+/// CHECK-NOT: Callback Init:
+/// CHECK-NOT: Callback Load:
+/// CHECK-NOT: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=1
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2
+/// CHECK-NOT: Callback Submit: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] req_num_teams=1
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4
+/// CHECK-NOT: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=2
+
+/// CHECK-NOT: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=1
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2
+/// CHECK-NOT: Callback Submit: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] req_num_teams=0
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4
+/// CHECK-NOT: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=2
+/// CHECK-NOT: Callback Fini:

openmp/runtime/cmake/config-ix.cmake

#      (LIBOMP_ARCH STREQUAL arm) OR
      AND # OS supported?
      ((WIN32 AND LIBOMP_HAVE_PSAPI) OR APPLE OR (NOT WIN32 AND LIBOMP_HAVE_WEAK_ATTRIBUTE)))
     set(LIBOMP_HAVE_OMPT_SUPPORT TRUE)
   else()
     set(LIBOMP_HAVE_OMPT_SUPPORT FALSE)
   endif()
 endif()
 
+set(LIBOMP_HAVE_OMPT_SUPPORT ${LIBOMP_HAVE_OMPT_SUPPORT} PARENT_SCOPE)
+
 # Check if HWLOC support is available
 if(${LIBOMP_USE_HWLOC})
   set(CMAKE_REQUIRED_INCLUDES ${LIBOMP_HWLOC_INSTALL_DIR}/include)
   check_include_file(hwloc.h LIBOMP_HAVE_HWLOC_H)
   set(CMAKE_REQUIRED_INCLUDES)
   find_library(LIBOMP_HWLOC_LIBRARY
     NAMES hwloc libhwloc
     HINTS ${LIBOMP_HWLOC_INSTALL_DIR}/lib)

openmp/runtime/src/exports_so.txt

global: # Exported symbols.
         #
         omp_*;     # Standard OpenMP functions.
         OMP_*;     # Standard OpenMP symbols.
 
         #
         # OMPT API
         #
         ompt_start_tool;     # OMPT start interface
+        ompt_libomp_connect; # OMPT libomptarget interface
 
         ompc_*;    # omp.h renames some standard functions to ompc_*.
         kmp_*;     # Intel extensions.
         kmpc_*;    # Intel extensions.
         __kmpc_*;  # Functions called by compiler-generated code.
         GOMP_*;    # GNU C compatibility functions.
 
         _You_must_link_with_*;     # Mutual detection/MS compatibility symbols.

openmp/runtime/src/include/omp-tools.h.var

 
 
 #define FOREACH_KMP_MUTEX_IMPL(macro)                                                \
     macro (kmp_mutex_impl_none, 0)         /* unknown implementation */              \
     macro (kmp_mutex_impl_spin, 1)         /* based on spin */                       \
     macro (kmp_mutex_impl_queuing, 2)      /* based on some fair policy */           \
     macro (kmp_mutex_impl_speculative, 3)  /* based on HW-supported speculation */
 
-#define FOREACH_OMPT_EVENT(macro)                                                                                        \
+#define FOREACH_OMPT_HOST_EVENT(macro)                                                                                   \
                                                                                                                          \
     /*--- Mandatory Events ---*/                                                                                         \
     macro (ompt_callback_thread_begin,      ompt_callback_thread_begin_t,       1) /* thread begin                    */ \
     macro (ompt_callback_thread_end,        ompt_callback_thread_end_t,         2) /* thread end                      */ \
                                                                                                                          \
     macro (ompt_callback_parallel_begin,    ompt_callback_parallel_begin_t,     3) /* parallel begin                  */ \
     macro (ompt_callback_parallel_end,      ompt_callback_parallel_end_t,       4) /* parallel end                    */ \
                                                                                                                          \
     macro (ompt_callback_task_create,       ompt_callback_task_create_t,        5) /* task begin                      */ \
     macro (ompt_callback_task_schedule,     ompt_callback_task_schedule_t,      6) /* task schedule                   */ \
     macro (ompt_callback_implicit_task,     ompt_callback_implicit_task_t,      7) /* implicit task                   */ \
                                                                                                                          \
-    macro (ompt_callback_target,            ompt_callback_target_t,             8) /* target                          */ \
-    macro (ompt_callback_target_data_op,    ompt_callback_target_data_op_t,     9) /* target data op                  */ \
-    macro (ompt_callback_target_submit,     ompt_callback_target_submit_t,     10) /* target  submit                  */ \
                                                                                                                          \
     macro (ompt_callback_control_tool,      ompt_callback_control_tool_t,      11) /* control tool                    */ \
                                                                                                                          \
-    macro (ompt_callback_device_initialize, ompt_callback_device_initialize_t, 12) /* device initialize               */ \
-    macro (ompt_callback_device_finalize,   ompt_callback_device_finalize_t,   13) /* device finalize                 */ \
-                                                                                                                         \
-    macro (ompt_callback_device_load,       ompt_callback_device_load_t,       14) /* device load                     */ \
-    macro (ompt_callback_device_unload,     ompt_callback_device_unload_t,     15) /* device unload                   */ \
                                                                                                                          \
     /* Optional Events */                                                                                                \
     macro (ompt_callback_sync_region_wait,  ompt_callback_sync_region_t,       16) /* sync region wait begin or end   */ \
                                                                                                                          \
     macro (ompt_callback_mutex_released,    ompt_callback_mutex_t,             17) /* mutex released                  */ \
                                                                                                                          \
     macro (ompt_callback_dependences,       ompt_callback_dependences_t,       18) /* report task dependences         */ \
     macro (ompt_callback_task_dependence,   ompt_callback_task_dependence_t,   19) /* report task dependence          */ \
                                                                                                                          \
     macro (ompt_callback_work,              ompt_callback_work_t,              20) /* task at work begin or end       */ \
                                                                                                                          \
     macro (ompt_callback_masked,            ompt_callback_masked_t,            21) /* task at masked begin or end     */ \
                                                                                                                          \
-    macro (ompt_callback_target_map,        ompt_callback_target_map_t,        22) /* target map                      */ \
                                                                                                                          \
     macro (ompt_callback_sync_region,       ompt_callback_sync_region_t,       23) /* sync region begin or end        */ \
                                                                                                                          \
     macro (ompt_callback_lock_init,         ompt_callback_mutex_acquire_t,     24) /* lock init                       */ \
     macro (ompt_callback_lock_destroy,      ompt_callback_mutex_t,             25) /* lock destroy                    */ \
                                                                                                                          \
     macro (ompt_callback_mutex_acquire,     ompt_callback_mutex_acquire_t,     26) /* mutex acquire                   */ \
     macro (ompt_callback_mutex_acquired,    ompt_callback_mutex_t,             27) /* mutex acquired                  */ \
                                                                                                                          \
     macro (ompt_callback_nest_lock,         ompt_callback_nest_lock_t,         28) /* nest lock                       */ \
                                                                                                                          \
     macro (ompt_callback_flush,             ompt_callback_flush_t,             29) /* after executing flush           */ \
                                                                                                                          \
     macro (ompt_callback_cancel,            ompt_callback_cancel_t,            30) /* cancel innermost binding region */ \
                                                                                                                          \
     macro (ompt_callback_reduction,         ompt_callback_sync_region_t,       31) /* reduction                       */ \
                                                                                                                          \
     macro (ompt_callback_dispatch,          ompt_callback_dispatch_t,          32) /* dispatch of work                */ \
+    macro (ompt_callback_error,             ompt_callback_error_t,             37) /* error                           */
+
+
+#define FOREACH_OMPT_DEVICE_EVENT(macro)                                                                                 \
+    macro (ompt_callback_device_initialize, ompt_callback_device_initialize_t, 12) /* device initialize               */ \
+    macro (ompt_callback_device_finalize,   ompt_callback_device_finalize_t,   13) /* device finalize                 */ \
+                                                                                                                         \
+    macro (ompt_callback_device_load,       ompt_callback_device_load_t,       14) /* device load                     */ \
+    macro (ompt_callback_device_unload,     ompt_callback_device_unload_t,     15) /* device unload                   */ \
+
+
+#define FOREACH_OMPT_NOEMI_EVENT(macro)                                                                                  \
+    macro (ompt_callback_target,            ompt_callback_target_t,             8) /* target                          */ \
+    macro (ompt_callback_target_data_op,    ompt_callback_target_data_op_t,     9) /* target data op                  */ \
+    macro (ompt_callback_target_submit,     ompt_callback_target_submit_t,     10) /* target  submit                  */ \
+    macro (ompt_callback_target_map,        ompt_callback_target_map_t,        22) /* target map                      */ \
+
+
+#define FOREACH_OMPT_EMI_EVENT(macro)                                                                                    \
     macro (ompt_callback_target_emi,        ompt_callback_target_emi_t,        33) /* target                          */ \
     macro (ompt_callback_target_data_op_emi,ompt_callback_target_data_op_emi_t,34) /* target data op                  */ \
     macro (ompt_callback_target_submit_emi, ompt_callback_target_submit_emi_t, 35) /* target submit                   */ \
     macro (ompt_callback_target_map_emi,    ompt_callback_target_map_emi_t,    36) /* target map                      */ \
-    macro (ompt_callback_error,             ompt_callback_error_t,             37) /* error                           */
+
+#define FOREACH_OMPT_50_TARGET_EVENT(macro)                                                                              \
+    FOREACH_OMPT_DEVICE_EVENT(macro)                                                                                     \
+    FOREACH_OMPT_NOEMI_EVENT(macro) 
+
+#define FOREACH_OMPT_51_TARGET_EVENT(macro)                                                                              \
+    FOREACH_OMPT_DEVICE_EVENT(macro)                                                                                     \
+    FOREACH_OMPT_EMI_EVENT(macro) 
+
+#define FOREACH_OMPT_EVENT(macro)                                                                                        \
+    FOREACH_OMPT_HOST_EVENT(macro)                                                                                       \
+    FOREACH_OMPT_DEVICE_EVENT(macro)                                                                                     \
+    FOREACH_OMPT_NOEMI_EVENT(macro)                                                                                      \
+    FOREACH_OMPT_EMI_EVENT(macro)
+
+#define FOREACH_OMPT_51_EVENT(macro)                                                                                     \
+    FOREACH_OMPT_HOST_EVENT(macro)                                                                                       \
+    FOREACH_OMPT_DEVICE_EVENT(macro)                                                                                     \
+    FOREACH_OMPT_EMI_EVENT(macro)
 
 /*****************************************************************************
  * implementation specific types
  *****************************************************************************/
 
 typedef enum kmp_mutex_impl_t {
 #define kmp_mutex_impl_macro(impl, code) impl = code,
     FOREACH_KMP_MUTEX_IMPL(kmp_mutex_impl_macro)

openmp/runtime/src/ompt-event-specific.h

 #define ompt_callback_parallel_begin_implemented ompt_event_MAY_ALWAYS
 #define ompt_callback_parallel_end_implemented ompt_event_MAY_ALWAYS
 
 #define ompt_callback_task_create_implemented ompt_event_MAY_ALWAYS
 #define ompt_callback_task_schedule_implemented ompt_event_MAY_ALWAYS
 
 #define ompt_callback_implicit_task_implemented ompt_event_MAY_ALWAYS
 
-#define ompt_callback_target_implemented ompt_event_UNIMPLEMENTED
-#define ompt_callback_target_emi_implemented ompt_event_UNIMPLEMENTED
-#define ompt_callback_target_data_op_implemented ompt_event_UNIMPLEMENTED
-#define ompt_callback_target_data_op_emi_implemented ompt_event_UNIMPLEMENTED
-#define ompt_callback_target_submit_implemented ompt_event_UNIMPLEMENTED
-#define ompt_callback_target_submit_emi_implemented ompt_event_UNIMPLEMENTED
+#define ompt_callback_target_implemented ompt_event_MAY_ALWAYS
+#define ompt_callback_target_emi_implemented ompt_event_MAY_ALWAYS
+#define ompt_callback_target_data_op_implemented ompt_event_MAY_ALWAYS
+#define ompt_callback_target_data_op_emi_implemented ompt_event_MAY_ALWAYS
+#define ompt_callback_target_submit_implemented ompt_event_MAY_ALWAYS
+#define ompt_callback_target_submit_emi_implemented ompt_event_MAY_ALWAYS
 
 #define ompt_callback_control_tool_implemented ompt_event_MAY_ALWAYS
 
-#define ompt_callback_device_initialize_implemented ompt_event_UNIMPLEMENTED
-#define ompt_callback_device_finalize_implemented ompt_event_UNIMPLEMENTED
+#define ompt_callback_device_initialize_implemented ompt_event_MAY_ALWAYS
+#define ompt_callback_device_finalize_implemented ompt_event_MAY_ALWAYS
 
-#define ompt_callback_device_load_implemented ompt_event_UNIMPLEMENTED
+#define ompt_callback_device_load_implemented ompt_event_MAY_ALWAYS
 #define ompt_callback_device_unload_implemented ompt_event_UNIMPLEMENTED
 
 /*----------------------------------------------------------------------------
  | Optional Events
  +--------------------------------------------------------------------------*/
 
 #define ompt_callback_sync_region_wait_implemented                             \
   ompt_event_MAY_ALWAYS_OPTIONAL

openmp/runtime/src/ompt-general.cpp

 
 #if KMP_OS_WINDOWS
 static HMODULE ompt_tool_module = NULL;
 #define OMPT_DLCLOSE(Lib) FreeLibrary(Lib)
 #else
 static void *ompt_tool_module = NULL;
 #define OMPT_DLCLOSE(Lib) dlclose(Lib)
 #endif
+static ompt_start_tool_result_t *libomptarget_ompt_result = NULL;
 
 /*****************************************************************************
  * forward declarations
  ****************************************************************************/
 
 static ompt_interface_fn_t ompt_fn_lookup(const char *s);
 
 OMPT_API_ROUTINE ompt_data_t *ompt_get_thread_data(void);
fprintf(stderr,
             "  legal values are (NULL,\"\",\"disabled\","
             "\"enabled\").\n",
             ompt_env_var);
     break;
   }
   if (verbose_init && verbose_file != stderr && verbose_file != stdout)
     fclose(verbose_file);
 #if OMPT_DEBUG
-  printf("ompt_pre_init(): ompt_enabled = %d\n", ompt_enabled);
+  printf("ompt_pre_init(): ompt_enabled = %d\n", ompt_enabled.enabled);
 #endif
 }
 
 extern "C" int omp_get_initial_device(void);
 
 void ompt_post_init() {
   //--------------------------------------------------
   // Execute the post-initialization logic only once.
if (ompt_start_tool_result) {
     kmp_info_t *root_thread = ompt_get_thread();
 
     ompt_set_thread_state(root_thread, ompt_state_overhead);
 
     if (ompt_enabled.ompt_callback_thread_begin) {
       ompt_callbacks.ompt_callback(ompt_callback_thread_begin)(
           ompt_thread_initial, __ompt_get_thread_data_internal());
     }
-    ompt_data_t *task_data;
-    ompt_data_t *parallel_data;
+    ompt_data_t *task_data = nullptr;
+    ompt_data_t *parallel_data = nullptr;
     __ompt_get_task_info_internal(0, NULL, &task_data, NULL, &parallel_data,
                                   NULL);
     if (ompt_enabled.ompt_callback_implicit_task) {
       ompt_callbacks.ompt_callback(ompt_callback_implicit_task)(
           ompt_scope_begin, parallel_data, task_data, 1, 1, ompt_task_initial);
     }
 
     ompt_set_thread_state(root_thread, ompt_state_work_serial);
   }
 }
 
 void ompt_fini() {
-  if (ompt_enabled.enabled
-#if OMPD_SUPPORT
-      && ompt_start_tool_result && ompt_start_tool_result->finalize
-#endif
-  ) {
+  if (ompt_enabled.enabled) {
+    if (ompt_start_tool_result && ompt_start_tool_result->finalize) {
       ompt_start_tool_result->finalize(&(ompt_start_tool_result->tool_data));
     }
+    if (libomptarget_ompt_result && libomptarget_ompt_result->finalize) {
+      libomptarget_ompt_result->finalize(NULL);
+    }
+  }
 
   if (ompt_tool_module)
     OMPT_DLCLOSE(ompt_tool_module);
   memset(&ompt_enabled, 0, sizeof(ompt_enabled));
 }
 
 /*****************************************************************************
  * interface operations

 
 #define ompt_interface_fn(fn)                                                  \
   fn##_t fn##_f = fn;                                                          \
   if (strcmp(s, #fn) == 0)                                                     \
     return (ompt_interface_fn_t)fn##_f;
 
   FOREACH_OMPT_INQUIRY_FN(ompt_interface_fn)
 
-  return NULL;
+#undef ompt_interface_fn
+
+  return (ompt_interface_fn_t)0;
+}
+
+static int ompt_set_frame_enter(void *addr, int flags, int state) {
+  return __ompt_set_frame_enter_internal(addr, flags, state);
+}
+
+static ompt_data_t *ompt_get_task_data() { return __ompt_get_task_data(); }
+
+static ompt_data_t *ompt_get_target_task_data() {
+  return __ompt_get_target_task_data();
+}
+
+static ompt_interface_fn_t ompt_libomp_target_fn_lookup(const char *s) {
+#define provide_fn(fn)                                                         \
+  if (strcmp(s, #fn) == 0)                                                     \
+    return (ompt_interface_fn_t)fn;
+
+  provide_fn(ompt_set_frame_enter);
+  provide_fn(ompt_get_task_data);
+  provide_fn(ompt_get_target_task_data);
+
+#define ompt_interface_fn(fn, type, code)                                      \
+  if (strcmp(s, #fn) == 0)                                                     \
+    return (ompt_interface_fn_t)ompt_callbacks.ompt_callback(fn);
+
+  FOREACH_OMPT_DEVICE_EVENT(ompt_interface_fn)
+  FOREACH_OMPT_EMI_EVENT(ompt_interface_fn)
+  FOREACH_OMPT_NOEMI_EVENT(ompt_interface_fn)
+
+#undef ompt_interface_fn
+
+  return (ompt_interface_fn_t)0;
+}
+
+_OMP_EXTERN void ompt_libomp_connect(ompt_start_tool_result_t *result) {
+  OMPT_VERBOSE_INIT_PRINT("libomp --> OMPT: Enter libomp_ompt_connect\n");
+
+  __ompt_force_initialization();
+
+  if (ompt_enabled.enabled &&
+      ompt_callbacks.ompt_callback(ompt_callback_device_initialize)) {
+    if (result) {
+      OMPT_VERBOSE_INIT_PRINT(
+          "libomp --> OMPT: Connecting with libomptarget\n");
+      result->initialize(ompt_libomp_target_fn_lookup,
+                         0 /* initial_device_num */, nullptr /* tool_data */);
+      libomptarget_ompt_result = result;
+    }
+  }
+  OMPT_VERBOSE_INIT_PRINT("libomp --> OMPT: Exit libomp_ompt_connect\n");
 }

openmp/runtime/src/ompt-internal.h

 /*
  * ompt-internal.h - header of OMPT internal data structures
  */
 
 //===----------------------------------------------------------------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef __OMPT_INTERNAL_H__
 #define __OMPT_INTERNAL_H__
 
+#include "kmp_platform.h"
+
 #include "ompt-event-specific.h"
 #include "omp-tools.h"
 
 #define OMPT_VERSION 1
 
 #define _OMP_EXTERN extern "C"
 
 #define OMPT_INVOKER(x)                                                        \
   ((x == fork_context_gnu) ? ompt_parallel_invoker_program                     \
                            : ompt_parallel_invoker_runtime)
 
+#define OMPT_FRAME_SET(frame, which, ptr_value, flags)                         \
+  {                                                                            \
+    frame->which##_frame.ptr = ptr_value;                                      \
+    frame->which##_frame_flags = flags;                                        \
+  }
+
+#define OMPT_FRAME_CLEAR(frame, which) OMPT_FRAME_SET(frame, which, 0, 0)
+
+#define OMPT_FRAME_SET_P(frame, which) (frame->which##_frame.ptr != NULL)
+
 #define ompt_callback(e) e##_callback
 
 typedef struct ompt_callbacks_internal_s {
 #define ompt_event_macro(event, callback, eventid)                             \
   callback ompt_callback(event);
 
   FOREACH_OMPT_EVENT(ompt_event_macro)
 
typedef struct ompt_lw_taskteam_s {
   int heap;
   struct ompt_lw_taskteam_s *parent;
 } ompt_lw_taskteam_t;
 
 typedef struct {
   ompt_data_t thread_data;
   ompt_data_t task_data; /* stored here from implicit barrier-begin until
                             implicit-task-end */
+  ompt_data_t target_task_data;
   void *return_address; /* stored here on entry of runtime */
   ompt_state_t state;
   ompt_wait_id_t wait_id;
   int ompt_task_yielded;
   int parallel_flags; // information for the last parallel region invoked
   void *idle_frame;
 } ompt_thread_info_t;
 
 extern ompt_callbacks_internal_t ompt_callbacks;
 
 #if OMPT_SUPPORT && OMPT_OPTIONAL
 #if USE_FAST_MEMORY
 #define KMP_OMPT_DEPS_ALLOC __kmp_fast_allocate
 #define KMP_OMPT_DEPS_FREE __kmp_fast_free
 #else
 #define KMP_OMPT_DEPS_ALLOC __kmp_thread_malloc
 #define KMP_OMPT_DEPS_FREE __kmp_thread_free
 #endif
 #endif /* OMPT_SUPPORT && OMPT_OPTIONAL */
 
 #ifdef __cplusplus
 extern "C" {
 #endif
 
 void ompt_pre_init(void);
 void ompt_post_init(void);
 void ompt_fini(void);
 
 #define OMPT_GET_RETURN_ADDRESS(level) __builtin_return_address(level)
 #define OMPT_GET_FRAME_ADDRESS(level) __builtin_frame_address(level)
 
 int __kmp_control_tool(uint64_t command, uint64_t modifier, void *arg);
 
 extern ompt_callbacks_active_t ompt_enabled;
 
 #if KMP_OS_WINDOWS
 #define UNLIKELY(x) (x)
 #define OMPT_NOINLINE __declspec(noinline)
 #else
 #define UNLIKELY(x) __builtin_expect(!!(x), 0)
 #define OMPT_NOINLINE __attribute__((noinline))
 #endif
 
 #ifdef __cplusplus
 }
 #endif
 
 #endif

openmp/runtime/src/ompt-specific.h

 
 #include "kmp.h"
 
 #if OMPT_SUPPORT
 /*****************************************************************************
  * forward declarations
  ****************************************************************************/
 
+void __ompt_force_initialization();
+
+int __ompt_set_frame_enter_internal(void *addr, int flags, int state);
+
 void __ompt_team_assign_id(kmp_team_t *team, ompt_data_t ompt_pid);
+
 void __ompt_thread_assign_wait_id(void *variable);
 
 void __ompt_lw_taskteam_init(ompt_lw_taskteam_t *lwt, kmp_info_t *thr, int gtid,
                              ompt_data_t *ompt_pid, void *codeptr);
 
 void __ompt_lw_taskteam_link(ompt_lw_taskteam_t *lwt, kmp_info_t *thr,
                              int on_heap, bool always = false);
 
 void __ompt_lw_taskteam_unlink(kmp_info_t *thr);
 
 ompt_team_info_t *__ompt_get_teaminfo(int depth, int *size);
 
+ompt_data_t *__ompt_get_task_data();
+
+ompt_data_t *__ompt_get_target_task_data();
+
 ompt_task_info_t *__ompt_get_task_info_object(int depth);
 
 int __ompt_get_parallel_info_internal(int ancestor_level,
                                       ompt_data_t **parallel_data,
                                       int *team_size);
 
 int __ompt_get_task_info_internal(int ancestor_level, int *type,
                                   ompt_data_t **task_data,
                                   ompt_frame_t **task_frame,
                                   ompt_data_t **parallel_data, int *thread_num);
 
 ompt_data_t *__ompt_get_thread_data_internal();
 
 /*
  * Unused currently
 static uint64_t __ompt_get_get_unique_id_internal();
 */
 
 ompt_sync_region_t __ompt_get_barrier_kind(enum barrier_type, kmp_info_t *);
 
 /*****************************************************************************
  * macros
  ****************************************************************************/
 
-#define OMPT_CUR_TASK_INFO(thr) (&(thr->th.th_current_task->ompt_task_info))
+#define OMPT_CUR_TASK_INFO(thr) (&((thr)->th.th_current_task->ompt_task_info))
 #define OMPT_CUR_TASK_DATA(thr)                                                \
-  (&(thr->th.th_current_task->ompt_task_info.task_data))
-#define OMPT_CUR_TEAM_INFO(thr) (&(thr->th.th_team->t.ompt_team_info))
+  (&((thr)->th.th_current_task->ompt_task_info.task_data))
+#define OMPT_CUR_TEAM_INFO(thr) (&((thr)->th.th_team->t.ompt_team_info))
 #define OMPT_CUR_TEAM_DATA(thr)                                                \
-  (&(thr->th.th_team->t.ompt_team_info.parallel_data))
+  (&((thr)->th.th_team->t.ompt_team_info.parallel_data))
 
 #define OMPT_HAVE_WEAK_ATTRIBUTE KMP_HAVE_WEAK_ATTRIBUTE
 #define OMPT_HAVE_PSAPI KMP_HAVE_PSAPI
 #define OMPT_STR_MATCH(haystack, needle) __kmp_str_match(haystack, 0, needle)
 
 inline void *__ompt_load_return_address(int gtid) {
   kmp_info_t *thr = __kmp_threads[gtid];
   void *return_address = thr->th.ompt_thread_info.return_address;
   thr->th.ompt_thread_info.return_address = NULL;
   return return_address;
 }
 
 /*#define OMPT_STORE_RETURN_ADDRESS(gtid) \
   if (ompt_enabled.enabled && gtid >= 0 && __kmp_threads[gtid] &&              \
       !__kmp_threads[gtid]->th.ompt_thread_info.return_address)                \
   __kmp_threads[gtid]->th.ompt_thread_info.return_address =                    \
-      __builtin_return_address(0)*/
+
 #define OMPT_STORE_RETURN_ADDRESS(gtid)                                        \
   OmptReturnAddressGuard ReturnAddressGuard{gtid, __builtin_return_address(0)};
 #define OMPT_LOAD_RETURN_ADDRESS(gtid) __ompt_load_return_address(gtid)
 #define OMPT_LOAD_OR_GET_RETURN_ADDRESS(gtid)                                  \
   ((ompt_enabled.enabled && gtid >= 0 && __kmp_threads[gtid] &&                \
     __kmp_threads[gtid]->th.ompt_thread_info.return_address)                   \
        ? __ompt_load_return_address(gtid)                                      \
        : __builtin_return_address(0))
 
 #define OMPT_GET_DISPATCH_CHUNK(chunk, lb, ub, incr)                           \
   do {                                                                         \
     if (incr > 0) {                                                            \
       chunk.start = static_cast<uint64_t>(lb);                                 \
       chunk.iterations = static_cast<uint64_t>(((ub) - (lb)) / (incr) + 1);    \
     } else {                                                                   \
       chunk.start = static_cast<uint64_t>(ub);                                 \
       chunk.iterations = static_cast<uint64_t>(((lb) - (ub)) / -(incr) + 1);   \
     }                                                                          \
   } while (0)
 
 //******************************************************************************
 // inline functions
 //******************************************************************************
 
 inline kmp_info_t *ompt_get_thread_gtid(int gtid) {
   return (gtid >= 0) ? __kmp_thread_from_gtid(gtid) : NULL;
 }
 
 inline kmp_info_t *ompt_get_thread() {

openmp/runtime/src/ompt-specific.cpp

ompt_task_info_t *__ompt_get_scheduling_taskinfo(int depth) {
   }
 
   return info;
 }
 
 //******************************************************************************
 // interface operations
 //******************************************************************************
+//----------------------------------------------------------
+// initialization support
+//----------------------------------------------------------
+
+void __ompt_force_initialization() { __kmp_serial_initialize(); }
 
 //----------------------------------------------------------
 // thread support
 //----------------------------------------------------------
 
 ompt_data_t *__ompt_get_thread_data_internal() {
   if (__kmp_get_gtid() >= 0) {
     kmp_info_t *thread = ompt_get_thread();

 void __ompt_lw_taskteam_init(ompt_lw_taskteam_t *lwt, kmp_info_t *thr, int gtid,
                              ompt_data_t *ompt_pid, void *codeptr) {
   // initialize parallel_data with input, return address to parallel_data on
   // exit
   lwt->ompt_team_info.parallel_data = *ompt_pid;
   lwt->ompt_team_info.master_return_address = codeptr;
   lwt->ompt_task_info.task_data.value = 0;
   lwt->ompt_task_info.frame.enter_frame = ompt_data_none;
+  lwt->ompt_task_info.frame.enter_frame_flags = 0;
   lwt->ompt_task_info.frame.exit_frame = ompt_data_none;
+  lwt->ompt_task_info.frame.exit_frame_flags = 0;
   lwt->ompt_task_info.scheduling_parent = NULL;
   lwt->heap = 0;
   lwt->parent = 0;
 }
 
 void __ompt_lw_taskteam_link(ompt_lw_taskteam_t *lwt, kmp_info_t *thr,
                              int on_heap, bool always) {
   ompt_lw_taskteam_t *link_lwt = lwt;
   if (always ||
#endif
   }
   //    return lwtask;
 }
 
 //----------------------------------------------------------
 // task support
 //----------------------------------------------------------
 
+ompt_data_t *__ompt_get_task_data() {
+  kmp_info_t *thr = ompt_get_thread();
+  ompt_data_t *task_data = thr ? OMPT_CUR_TASK_DATA(thr) : NULL;
+  return task_data;
+}
+
+ompt_data_t *__ompt_get_target_task_data() {
+  return &__kmp_threads[__kmp_get_gtid()]->th.ompt_thread_info.target_task_data;
+}
+
 int __ompt_get_task_info_internal(int ancestor_level, int *type,
                                   ompt_data_t **task_data,
                                   ompt_frame_t **task_frame,
                                   ompt_data_t **parallel_data,
                                   int *thread_num) {
   if (__kmp_get_gtid() < 0)
     return 0;
 
   if (ancestor_level < 0)
if (ret_size < 0)
     return 0;
 
   *addr = ret_addr;
   *size = (size_t)ret_size;
   return 1;
 }
 
 //----------------------------------------------------------
+// target region support
+//----------------------------------------------------------
+
+int __ompt_set_frame_enter_internal(void *addr, int flags, int state) {
+  int gtid = __kmp_entry_gtid();
+  kmp_info_t *thr = __kmp_threads[gtid];
+
+  ompt_frame_t *ompt_frame = &OMPT_CUR_TASK_INFO(thr)->frame;
+  OMPT_FRAME_SET(ompt_frame, enter, addr, flags);
+  int old_state = thr->th.ompt_thread_info.state;
+  thr->th.ompt_thread_info.state = ompt_state_work_parallel;
+  return old_state;
+}
+
+//----------------------------------------------------------
 // team support
 //----------------------------------------------------------
 
 void __ompt_team_assign_id(kmp_team_t *team, ompt_data_t ompt_pid) {
   team->t.ompt_team_info.parallel_data = ompt_pid;
 }
 
 //----------------------------------------------------------