[Openmp-commits] [llvm] [openmp] [OMPT][Offload][OpenMP] Fixes for OMPT data used by libomptarget (PR #156020)

Fri Aug 29 06:05:23 PDT 2025

https://github.com/kaloyan-ignatov created https://github.com/llvm/llvm-project/pull/156020

These commits fix issues regarding storage of tool data within libomptarget. Both libomp and libomptarget have been modified to accommodate this. We differentiate between two cases depending on the type of the target region:

- merged target regions (default, without `nowait` clause): behavior remains unchanged, tool data is stored in the thread local RegionInterface class within libomptarget.
- deferred target regions (using `nowait` clause): tool data is moved to `ompt_task_info_t` struct within libomp, as `RegionInterface` is thread local and its data is lost whenever another task is scheduled on the thread, which happens with deferred target regions.

In the new implementation, `RegionInterface` receives pointers to `ompt_task_info_t` within libomp which are handled transparently within libomptarget. Thus, the problem of tool data getting lost when a thread receives a new task is resolved: `target_data` and `target_task_data` remain set. 

Another issue was the value of `task_data` which is supposed to belong to the generating task of the region according to the OpenMP standard, but instead had been set to the `task_data` of the target task itself until now.

Test cases have been added which check both of these fixes.

>From 8bcdbcf2a72d4acdc25433efb2596af4c687655b Mon Sep 17 00:00:00 2001
From: Kaloyan Ignatov <kaloyan.ignatov at rwth-aachen.de>
Date: Sat, 16 Aug 2025 01:59:45 +0200
Subject: [PATCH 1/2] [OpenMP][OMPT] Change handling of target-related tool
 data in libomp

- store target_data in ompt_task_info_t to prevent data loss across
  scheduling of target region (both for deferred and undeferred target
  tasks)
- target_task_data is already in ompt_task_info_t, replace previous
  implementation which returned the wrong value
- combine queries for target_data and target_task_data and directly
  return ompt_task_info_t
- return correct task_data, OpenMP standard defines task_data in target
  callbacks as belonging to the generating (encountering) task
---
 openmp/runtime/src/ompt-general.cpp  | 10 ++++++----
 openmp/runtime/src/ompt-internal.h   |  5 ++++-
 openmp/runtime/src/ompt-specific.cpp | 26 +++++++++++++++++++++-----
 openmp/runtime/src/ompt-specific.h   |  4 ++--
 openmp/runtime/test/ompt/callback.h  |  3 +++
 5 files changed, 36 insertions(+), 12 deletions(-)

diff --git a/openmp/runtime/src/ompt-general.cpp b/openmp/runtime/src/ompt-general.cpp
index 1a778e4ecac3a..a529f9976f9fc 100644
--- a/openmp/runtime/src/ompt-general.cpp
+++ b/openmp/runtime/src/ompt-general.cpp
@@ -886,10 +886,12 @@ static ompt_interface_fn_t ompt_fn_lookup(const char *s) {
   return NULL;
 }
 
-static ompt_data_t *ompt_get_task_data() { return __ompt_get_task_data(); }
+static ompt_data_t *ompt_get_task_data() {
+  return __ompt_get_generating_task();
+}
 
-static ompt_data_t *ompt_get_target_task_data() {
-  return __ompt_get_target_task_data();
+static ompt_task_info_t *ompt_get_task_info_target() {
+  return __ompt_get_task_info_target();
 }
 
 /// Lookup function to query libomp callbacks registered by the tool
@@ -900,7 +902,7 @@ static ompt_interface_fn_t ompt_libomp_target_fn_lookup(const char *s) {
 
   provide_fn(ompt_get_callback);
   provide_fn(ompt_get_task_data);
-  provide_fn(ompt_get_target_task_data);
+  provide_fn(ompt_get_task_info_target);
 #undef provide_fn
 
 #define ompt_interface_fn(fn, type, code)                                      \
diff --git a/openmp/runtime/src/ompt-internal.h b/openmp/runtime/src/ompt-internal.h
index 36b45f7a91ea2..6a9d6776e27e9 100644
--- a/openmp/runtime/src/ompt-internal.h
+++ b/openmp/runtime/src/ompt-internal.h
@@ -57,8 +57,11 @@ typedef struct ompt_callbacks_active_s {
       (info->td_flags.merged_if0 ? ompt_task_mergeable : 0x0)
 
 typedef struct {
-  ompt_frame_t frame;
+  // liboffload only uses task_data and target_data. They must be the first
+  // elements!
   ompt_data_t task_data;
+  ompt_data_t target_data;
+  ompt_frame_t frame;
   struct kmp_taskdata *scheduling_parent;
   int thread_num;
   ompt_dispatch_chunk_t dispatch_chunk;
diff --git a/openmp/runtime/src/ompt-specific.cpp b/openmp/runtime/src/ompt-specific.cpp
index 94ae2e5293875..fea428cb38344 100644
--- a/openmp/runtime/src/ompt-specific.cpp
+++ b/openmp/runtime/src/ompt-specific.cpp
@@ -346,14 +346,30 @@ void __ompt_lw_taskteam_unlink(kmp_info_t *thr) {
 // task support
 //----------------------------------------------------------
 
-ompt_data_t *__ompt_get_task_data() {
+ompt_data_t *__ompt_get_generating_task() {
   kmp_info_t *thr = ompt_get_thread();
-  ompt_data_t *task_data = thr ? OMPT_CUR_TASK_DATA(thr) : NULL;
-  return task_data;
+  if (thr) {
+    kmp_taskdata_t *taskdata = thr->th.th_current_task;
+    if (taskdata == NULL)
+      return NULL;
+    if (taskdata->td_flags.target)
+      return &(taskdata->td_parent->ompt_task_info.task_data);
+    else
+      return &(taskdata->ompt_task_info.task_data);
+  }
+  return NULL;
 }
 
-ompt_data_t *__ompt_get_target_task_data() {
-  return &__kmp_threads[__kmp_get_gtid()]->th.ompt_thread_info.target_task_data;
+ompt_task_info_t *__ompt_get_task_info_target() {
+  kmp_info_t *thr = ompt_get_thread();
+  if (thr) {
+    kmp_taskdata_t *taskdata = thr->th.th_current_task;
+    if (taskdata == NULL)
+      return NULL;
+    if (taskdata->td_flags.target)
+      return &taskdata->ompt_task_info;
+  }
+  return NULL;
 }
 
 int __ompt_get_task_info_internal(int ancestor_level, int *type,
diff --git a/openmp/runtime/src/ompt-specific.h b/openmp/runtime/src/ompt-specific.h
index b7eb140458b40..6dad1d76274db 100644
--- a/openmp/runtime/src/ompt-specific.h
+++ b/openmp/runtime/src/ompt-specific.h
@@ -37,9 +37,9 @@ 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_generating_task();
 
-ompt_data_t *__ompt_get_target_task_data();
+ompt_task_info_t *__ompt_get_task_info_target();
 
 ompt_task_info_t *__ompt_get_task_info_object(int depth);
 
diff --git a/openmp/runtime/test/ompt/callback.h b/openmp/runtime/test/ompt/callback.h
index cd8acb57ee2f7..1e5283856c5b5 100644
--- a/openmp/runtime/test/ompt/callback.h
+++ b/openmp/runtime/test/ompt/callback.h
@@ -1018,6 +1018,8 @@ static void on_ompt_callback_error(ompt_severity_t severity,
          codeptr_ra);
 }
 
+#ifndef SKIP_CALLBACK_REGISTRATION
+
 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");
@@ -1094,6 +1096,7 @@ ompt_start_tool_result_t *ompt_start_tool(unsigned int omp_version,
 #ifdef __cplusplus
 }
 #endif
+#endif // ifndef SKIP_CALLBACK_REGISTRATION
 #endif // ifndef USE_PRIVATE_TOOL
 #ifdef _OMPT_TESTS
 #undef _OMPT_TESTS

>From 1235b74f30592eed17b9a4ee4b34286f0d4f6208 Mon Sep 17 00:00:00 2001
From: Kaloyan Ignatov <kaloyan.ignatov at outlook.com>
Date: Fri, 22 Aug 2025 11:16:48 +0200
Subject: [PATCH 2/2] [Offload][OMPT] Change handling of target_data and
 target_task_data

- store target_data and target_task_data in OMP runtime
- redefine ompt_task_info_t from libomp with only first two fields -
  target_task_data and target_data
- replace get_target_task_data with ompt_get_task_info_target
- define macro for transparent handling of task_data and
  target_task_data
- provide a local ompt_task_info_t struct for merged target regions
- provide tests to ensure correct values of target_data and
  target_task_data
---
 offload/include/OpenMP/OMPT/Interface.h       |  28 ++-
 offload/libomptarget/OpenMP/OMPT/Callback.cpp |  21 +--
 offload/test/ompt/callbacks.h                 |  67 +++++--
 offload/test/ompt/omp_api.c                   |  10 +-
 offload/test/ompt/register_with_host.h        |  68 +++++++
 offload/test/ompt/target_memcpy.c             |  12 +-
 offload/test/ompt/target_memcpy_emi.c         |  22 ++-
 offload/test/ompt/target_tool_data.c          | 152 +++++++++++++++
 offload/test/ompt/target_tool_data_nowait.c   | 176 ++++++++++++++++++
 .../ompt/target_tool_data_nowait_nodepend.c   |  63 +++++++
 offload/test/ompt/veccopy.c                   |  43 ++---
 offload/test/ompt/veccopy_data.c              |  67 +++----
 offload/test/ompt/veccopy_disallow_both.c     |  75 ++++----
 offload/test/ompt/veccopy_emi.c               |  83 +++++----
 offload/test/ompt/veccopy_emi_map.c           |  83 +++++----
 offload/test/ompt/veccopy_map.c               |  46 ++---
 offload/test/ompt/veccopy_no_device_init.c    |  42 +++--
 offload/test/ompt/veccopy_wrong_return.c      |  43 +++--
 18 files changed, 814 insertions(+), 287 deletions(-)
 create mode 100644 offload/test/ompt/register_with_host.h
 create mode 100644 offload/test/ompt/target_tool_data.c
 create mode 100644 offload/test/ompt/target_tool_data_nowait.c
 create mode 100644 offload/test/ompt/target_tool_data_nowait_nodepend.c

diff --git a/offload/include/OpenMP/OMPT/Interface.h b/offload/include/OpenMP/OMPT/Interface.h
index 43fb193bc75a6..b45c953dd2edb 100644
--- a/offload/include/OpenMP/OMPT/Interface.h
+++ b/offload/include/OpenMP/OMPT/Interface.h
@@ -25,12 +25,22 @@
 
 #define OMPT_IF_BUILT(stmt) stmt
 
+#define TargetTaskData                                                         \
+  ((OmptTaskInfoPtr == &OmptTaskInfo) ? nullptr                                \
+                                      : (&(OmptTaskInfoPtr->task_data)))
+#define TargetData (OmptTaskInfoPtr->target_data)
+
+typedef struct ompt_task_info_t {
+  ompt_data_t task_data;
+  ompt_data_t target_data;
+} ompt_task_info_t;
+
 /// Callbacks for target regions require task_data representing the
 /// encountering task.
 /// Callbacks for target regions and target data ops require
 /// target_task_data representing the target task region.
 typedef ompt_data_t *(*ompt_get_task_data_t)();
-typedef ompt_data_t *(*ompt_get_target_task_data_t)();
+typedef ompt_task_info_t *(*ompt_get_task_info_target_t)();
 
 namespace llvm {
 namespace omp {
@@ -40,7 +50,7 @@ namespace ompt {
 /// Function pointers that will be used to track task_data and
 /// target_task_data.
 static ompt_get_task_data_t ompt_get_task_data_fn;
-static ompt_get_target_task_data_t ompt_get_target_task_data_fn;
+static ompt_get_task_info_target_t ompt_get_task_info_target_fn;
 
 /// Used to maintain execution state for this thread
 class Interface {
@@ -216,16 +226,16 @@ class Interface {
 
 private:
   /// Target operations id
-  ompt_id_t HostOpId = 0;
-
-  /// Target region data
-  ompt_data_t TargetData = ompt_data_none;
+  ompt_id_t HostOpId{0};
 
   /// Task data representing the encountering task
-  ompt_data_t *TaskData = nullptr;
+  ompt_data_t *TaskData{nullptr};
+
+  /// TaskInfo contains target_data and task_data
+  ompt_task_info_t OmptTaskInfo{ompt_data_none, ompt_data_none};
 
-  /// Target task data representing the target task region
-  ompt_data_t *TargetTaskData = nullptr;
+  /// Ptr to TaskInfo in OpenMP runtime in case of deferred target tasks
+  ompt_task_info_t *OmptTaskInfoPtr{&OmptTaskInfo};
 
   /// Used for marking begin of a data operation
   void beginTargetDataOperation();
diff --git a/offload/libomptarget/OpenMP/OMPT/Callback.cpp b/offload/libomptarget/OpenMP/OMPT/Callback.cpp
index ab0942ed4fd3f..b59fe72ae514c 100644
--- a/offload/libomptarget/OpenMP/OMPT/Callback.cpp
+++ b/offload/libomptarget/OpenMP/OMPT/Callback.cpp
@@ -50,8 +50,8 @@ bool llvm::omp::target::ompt::Initialized = false;
 
 ompt_get_callback_t llvm::omp::target::ompt::lookupCallbackByCode = nullptr;
 ompt_function_lookup_t llvm::omp::target::ompt::lookupCallbackByName = nullptr;
-ompt_get_target_task_data_t ompt_get_target_task_data_fn = nullptr;
 ompt_get_task_data_t ompt_get_task_data_fn = nullptr;
+ompt_get_task_info_target_t ompt_get_task_info_target_fn = nullptr;
 
 /// Unique correlation id
 static std::atomic<uint64_t> IdCounter(1);
@@ -421,18 +421,17 @@ void Interface::beginTargetRegion() {
   // Set up task state
   assert(ompt_get_task_data_fn && "Calling a null task data function");
   TaskData = ompt_get_task_data_fn();
-  // Set up target task state
-  assert(ompt_get_target_task_data_fn &&
-         "Calling a null target task data function");
-  TargetTaskData = ompt_get_target_task_data_fn();
-  // Target state will be set later
-  TargetData = ompt_data_none;
+  // Set up target task and target state
+  assert(ompt_get_task_info_target_fn &&
+         "Calling a null target task info function");
+  if (ompt_task_info_t *TempTaskInfo = ompt_get_task_info_target_fn())
+    OmptTaskInfoPtr = TempTaskInfo;
 }
 
 void Interface::endTargetRegion() {
   TaskData = 0;
-  TargetTaskData = 0;
-  TargetData = ompt_data_none;
+  OmptTaskInfo = {ompt_data_none, ompt_data_none};
+  OmptTaskInfoPtr = &OmptTaskInfo;
 }
 
 /// Used to maintain the finalization functions that are received
@@ -471,7 +470,7 @@ int llvm::omp::target::ompt::initializeLibrary(ompt_function_lookup_t lookup,
 
   bindOmptFunctionName(ompt_get_callback, lookupCallbackByCode);
   bindOmptFunctionName(ompt_get_task_data, ompt_get_task_data_fn);
-  bindOmptFunctionName(ompt_get_target_task_data, ompt_get_target_task_data_fn);
+  bindOmptFunctionName(ompt_get_task_info_target, ompt_get_task_info_target_fn);
 #undef bindOmptFunctionName
 
   // Store pointer of 'ompt_libomp_target_fn_lookup' for use by libomptarget
@@ -480,8 +479,6 @@ int llvm::omp::target::ompt::initializeLibrary(ompt_function_lookup_t lookup,
   assert(lookupCallbackByCode && "lookupCallbackByCode should be non-null");
   assert(lookupCallbackByName && "lookupCallbackByName should be non-null");
   assert(ompt_get_task_data_fn && "ompt_get_task_data_fn should be non-null");
-  assert(ompt_get_target_task_data_fn &&
-         "ompt_get_target_task_data_fn should be non-null");
   assert(LibraryFinalizer == nullptr &&
          "LibraryFinalizer should not be initialized yet");
 
diff --git a/offload/test/ompt/callbacks.h b/offload/test/ompt/callbacks.h
index 95437d9cdcfb1..2e7763f0abbac 100644
--- a/offload/test/ompt/callbacks.h
+++ b/offload/test/ompt/callbacks.h
@@ -5,6 +5,37 @@
 // Tool related code below
 #include <omp-tools.h>
 
+static const char *ompt_target_data_op_t_values[] = {
+    "",
+    "ompt_target_data_alloc",
+    "ompt_target_data_transfer_to_device",
+    "ompt_target_data_transfer_from_device",
+    "ompt_target_data_delete",
+    "ompt_target_data_associate",
+    "ompt_target_data_disassociate",
+    "ompt_target_data_alloc_async",
+    "ompt_target_data_transfer_to_device_async",
+    "ompt_target_data_transfer_from_device_async",
+    "ompt_target_data_delete_async"};
+
+static const char *ompt_scope_endpoint_t_values[] = {
+    "", "ompt_scope_begin", "ompt_scope_end", "ompt_scope_beginend"};
+
+static const char *ompt_target_t_values[] = {"",
+                                             "ompt_target",
+                                             "ompt_target_enter_data",
+                                             "ompt_target_exit_data",
+                                             "ompt_target_update",
+                                             "",
+                                             "",
+                                             "",
+                                             "",
+                                             "",
+                                             "ompt_target_nowait",
+                                             "ompt_target_enter_data_nowait",
+                                             "ompt_target_exit_data_nowait",
+                                             "ompt_target_update_nowait"};
+
 // For EMI callbacks
 ompt_id_t next_op_id = 0x8000000000000001;
 
@@ -38,11 +69,11 @@ static void on_ompt_callback_target_data_op(
     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");
-  printf("  Callback DataOp: target_id=%lu host_op_id=%lu optype=%d src=%p "
+  printf("  Callback DataOp: target_id=%lu host_op_id=%lu optype=%s 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);
+         target_id, host_op_id, ompt_target_data_op_t_values[optype], src_addr,
+         src_device_num, dest_addr, dest_device_num, bytes, codeptr_ra);
 }
 
 static void on_ompt_callback_target(ompt_target_t kind,
@@ -51,9 +82,10 @@ static void on_ompt_callback_target(ompt_target_t kind,
                                     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 "
+  printf("Callback Target: target_id=%lu kind=%s endpoint=%s device_num=%d "
          "code=%p\n",
-         target_id, kind, endpoint, device_num, codeptr_ra);
+         target_id, ompt_target_t_values[kind],
+         ompt_scope_endpoint_t_values[endpoint], device_num, codeptr_ra);
 }
 
 static void on_ompt_callback_target_submit(ompt_id_t target_id,
@@ -84,13 +116,15 @@ static void on_ompt_callback_target_data_op_emi(
   // target_task_data may be null, avoid dereferencing it
   uint64_t target_task_data_value =
       (target_task_data) ? target_task_data->value : 0;
-  printf("  Callback DataOp EMI: endpoint=%d optype=%d target_task_data=%p "
+  printf("  Callback DataOp EMI: endpoint=%s optype=%s 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);
+         ompt_scope_endpoint_t_values[endpoint],
+         ompt_target_data_op_t_values[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,
@@ -102,20 +136,21 @@ static void on_ompt_callback_target_emi(ompt_target_t kind,
   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 "
+  printf("Callback Target EMI: kind=%s endpoint=%s 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);
+         ompt_target_t_values[kind], ompt_scope_endpoint_t_values[endpoint],
+         device_num, task_data, task_data ? task_data->value : 0,
+         target_task_data, target_task_data ? target_task_data->value : 0,
+         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 "
+  printf("  Callback Submit EMI: endpoint=%s  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);
+         ompt_scope_endpoint_t_values[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,
diff --git a/offload/test/ompt/omp_api.c b/offload/test/ompt/omp_api.c
index a16ef7a64aa7d..5fb2098f0ce79 100644
--- a/offload/test/ompt/omp_api.c
+++ b/offload/test/ompt/omp_api.c
@@ -1,6 +1,8 @@
+// clang-format off
 // RUN: %libomptarget-compile-run-and-check-generic
 // REQUIRES: ompt
 // REQUIRES: gpu
+// clang-format on
 
 #include "omp.h"
 #include <stdlib.h>
@@ -32,8 +34,8 @@ int main(int argc, char **argv) {
 
 // clang-format off
 /// CHECK: Callback Init:
-/// 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=5
-/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=6
-/// 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=ompt_target_data_alloc
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_associate
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_disassociate
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_delete
 /// CHECK: Callback Fini:
diff --git a/offload/test/ompt/register_with_host.h b/offload/test/ompt/register_with_host.h
new file mode 100644
index 0000000000000..5e97f2c0b751a
--- /dev/null
+++ b/offload/test/ompt/register_with_host.h
@@ -0,0 +1,68 @@
+#define SKIP_CALLBACK_REGISTRATION 1
+
+#include "../../../openmp/runtime/test/ompt/callback.h"
+#include "callbacks.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)
+
+// 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
+
+  // host runtime functions
+  ompt_get_unique_id = (ompt_get_unique_id_t)lookup("ompt_get_unique_id");
+  ompt_get_thread_data = (ompt_get_thread_data_t)lookup("ompt_get_thread_data");
+  ompt_get_task_info = (ompt_get_task_info_t)lookup("ompt_get_task_info");
+
+  ompt_get_unique_id();
+
+  // host callbacks
+  register_ompt_callback(ompt_callback_sync_region);
+  register_ompt_callback_t(ompt_callback_sync_region_wait,
+                           ompt_callback_sync_region_t);
+  register_ompt_callback_t(ompt_callback_reduction,
+                           ompt_callback_sync_region_t);
+  register_ompt_callback(ompt_callback_implicit_task);
+  register_ompt_callback(ompt_callback_parallel_begin);
+  register_ompt_callback(ompt_callback_parallel_end);
+  register_ompt_callback(ompt_callback_task_create);
+  register_ompt_callback(ompt_callback_task_schedule);
+
+  // device callbacks
+  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
diff --git a/offload/test/ompt/target_memcpy.c b/offload/test/ompt/target_memcpy.c
index f244e0f418ed6..f769995579f50 100644
--- a/offload/test/ompt/target_memcpy.c
+++ b/offload/test/ompt/target_memcpy.c
@@ -1,6 +1,8 @@
+// clang-format off
 // RUN: %libomptarget-compile-run-and-check-generic
 // REQUIRES: ompt
 // REQUIRES: gpu
+// clang-format on
 
 /*
  * Verify that for the target OpenMP APIs, the return address is non-null and
@@ -46,26 +48,26 @@ int main() {
 }
 
 // clang-format off
-/// 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=ompt_target_data_alloc
 /// CHECK-SAME: src_device_num=[[HOST:[0-9]+]]
 /// CHECK-SAME: dest_device_num=[[DEVICE:[0-9]+]]
 /// CHECK-NOT: code=(nil)
 /// CHECK: code=[[CODE1:0x[0-f]+]]
-/// 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=ompt_target_data_transfer_to_device
 /// CHECK-SAME: src_device_num=[[HOST]] {{.+}} dest_device_num=[[DEVICE]]
 /// CHECK-NOT: code=(nil)
 /// CHECK-NOT: code=[[CODE1]]
 /// CHECK: code=[[CODE2:0x[0-f]+]]
-/// 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=ompt_target_data_transfer_from_device
 /// CHECK-SAME: src_device_num=[[DEVICE]] {{.+}} dest_device_num=[[DEVICE]]
 /// CHECK-NOT: code=(nil)
 /// CHECK-NOT: code=[[CODE2]]
 /// CHECK: code=[[CODE3:0x[0-f]+]]
-/// 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=ompt_target_data_transfer_from_device
 /// CHECK-SAME: src_device_num=[[DEVICE]] {{.+}} dest_device_num=[[HOST]]
 /// CHECK-NOT: code=(nil)
 /// CHECK-NOT: code=[[CODE3]]
 /// CHECK: code=[[CODE4:0x[0-f]+]]
-/// 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=ompt_target_data_delete
 /// CHECK-NOT: code=(nil)
 /// CHECK-NOT: code=[[CODE4]]
diff --git a/offload/test/ompt/target_memcpy_emi.c b/offload/test/ompt/target_memcpy_emi.c
index 934caba6efab3..39f262a366f94 100644
--- a/offload/test/ompt/target_memcpy_emi.c
+++ b/offload/test/ompt/target_memcpy_emi.c
@@ -1,6 +1,8 @@
+// clang-format off
 // RUN: %libomptarget-compile-run-and-check-generic
 // REQUIRES: ompt
 // REQUIRES: gpu
+// clang-format on
 
 /*
  * Verify all three data transfer directions: H2D, D2D and D2H
@@ -54,28 +56,28 @@ int main(void) {
 /// CHECK: Callback Init:
 
 /// CHECK: Allocating Memory on Device
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_alloc
 /// CHECK-SAME: src_device_num=[[HOST:[0-9]+]]
 /// CHECK-SAME: dest_device_num=[[DEVICE:[0-9]+]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=1 {{.+}} src_device_num=[[HOST]] {{.+}} dest_device_num=[[DEVICE]]
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_alloc {{.+}} src_device_num=[[HOST]] {{.+}} dest_device_num=[[DEVICE]]
 
 /// CHECK: Testing: Host to Device
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=2 {{.+}} src_device_num=[[HOST]] {{.+}} dest_device_num=[[DEVICE]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=2 {{.+}} src_device_num=[[HOST]] {{.+}} dest_device_num=[[DEVICE]]
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_to_device {{.+}} src_device_num=[[HOST]] {{.+}} dest_device_num=[[DEVICE]]
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_to_device {{.+}} src_device_num=[[HOST]] {{.+}} dest_device_num=[[DEVICE]]
 
 /// CHECK: Testing: Device to Device
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=3 {{.+}} src_device_num=[[DEVICE]] {{.+}} dest_device_num=[[DEVICE]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=3 {{.+}} src_device_num=[[DEVICE]] {{.+}} dest_device_num=[[DEVICE]]
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_from_device {{.+}} src_device_num=[[DEVICE]] {{.+}} dest_device_num=[[DEVICE]]
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_from_device {{.+}} src_device_num=[[DEVICE]] {{.+}} dest_device_num=[[DEVICE]]
 
 /// CHECK: Testing: Device to Host
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=3 {{.+}} src_device_num=[[DEVICE]] {{.+}} dest_device_num=[[HOST]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=3 {{.+}} src_device_num=[[DEVICE]] {{.+}} dest_device_num=[[HOST]]
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_from_device {{.+}} src_device_num=[[DEVICE]] {{.+}} dest_device_num=[[HOST]]
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_from_device {{.+}} src_device_num=[[DEVICE]] {{.+}} dest_device_num=[[HOST]]
 
 /// CHECK: Checking Correctness
 
 /// CHECK: Freeing Memory on Device
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=4 {{.+}} src_device_num=[[DEVICE]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=4 {{.+}} src_device_num=[[DEVICE]]
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_delete {{.+}} src_device_num=[[DEVICE]]
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_delete {{.+}} src_device_num=[[DEVICE]]
 
 /// CHECK: Callback Fini:
 
diff --git a/offload/test/ompt/target_tool_data.c b/offload/test/ompt/target_tool_data.c
new file mode 100644
index 0000000000000..ceec476cef375
--- /dev/null
+++ b/offload/test/ompt/target_tool_data.c
@@ -0,0 +1,152 @@
+// clang-format off
+// RUN: env LIBOMP_NUM_HIDDEN_HELPER_THREADS=1 %libomptarget-compile-run-and-check-generic
+// REQUIRES: ompt
+// clang-format on
+
+#include <inttypes.h>
+#include <omp-tools.h>
+#include <omp.h>
+#include <stdio.h>
+#include <string.h>
+
+#include "register_with_host.h"
+
+#define N 1000000
+#define M 1000
+
+int main() {
+  float *x = malloc(N * sizeof(float));
+  float *y = malloc(N * sizeof(float));
+
+  for (int i = 0; i < N; i++) {
+    x[i] = 1;
+    y[i] = 1;
+  }
+
+#pragma omp target enter data map(to : x [0:N]) map(alloc : y [0:N])
+#pragma omp target teams distribute parallel for
+  for (int i = 0; i < N; i++) {
+    for (int j = 0; j < M; j++) {
+      y[i] += 3 * x[i];
+    }
+  }
+
+#pragma omp target teams distribute parallel for
+  for (int i = 0; i < N; i++) {
+    for (int j = 0; j < M; j++) {
+      y[i] += 3 * x[i];
+    }
+  }
+
+#pragma omp target exit data map(release : x [0:N]) map(from : y [0:N])
+
+  printf("%f, %f\n", x[0], y[0]);
+
+  free(x);
+  free(y);
+  return 0;
+}
+
+// clang-format off
+/// CHECK: ompt_event_initial_task_begin
+/// CHECK-SAME: task_id=[[ENCOUNTERING_TASK:[0-f]+]]
+
+/// CHECK: Callback Target EMI: kind=ompt_target_enter_data endpoint=ompt_scope_begin
+/// CHECK-SAME: task_data=0x{{[0-f]+}} (0x[[ENCOUNTERING_TASK]])
+/// CHECK-SAME: target_task_data=(nil) (0x0)
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA:[0-f]+]])
+
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin
+/// CHECK-SAME: target_task_data=(nil) (0x0)
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end
+/// CHECK-SAME: target_task_data=(nil) (0x0)
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin
+/// CHECK-SAME: target_task_data=(nil) (0x0)
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end
+/// CHECK-SAME: target_task_data=(nil) (0x0)
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin
+/// CHECK-SAME: target_task_data=(nil) (0x0)
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end
+/// CHECK-SAME: target_task_data=(nil) (0x0)
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback Target EMI: kind=ompt_target_enter_data endpoint=ompt_scope_end
+/// CHECK-SAME: task_data=0x{{[0-f]+}} (0x[[ENCOUNTERING_TASK]])
+/// CHECK-SAME: target_task_data=(nil) (0x0)
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_begin
+/// CHECK-SAME: task_data=0x{{[0-f]+}} (0x[[ENCOUNTERING_TASK]])
+/// CHECK-SAME: target_task_data=(nil) (0x0)
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA:[0-f]+]])
+
+/// CHECK: Callback Submit EMI: endpoint=ompt_scope_begin
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback Submit EMI: endpoint=ompt_scope_end
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_end
+/// CHECK-SAME: task_data=0x{{[0-f]+}} (0x[[ENCOUNTERING_TASK]])
+/// CHECK-SAME: target_task_data=(nil) (0x0)
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_begin
+/// CHECK-SAME: task_data=0x{{[0-f]+}} (0x[[ENCOUNTERING_TASK]])
+/// CHECK-SAME: target_task_data=(nil) (0x0)
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA:[0-f]+]])
+
+/// CHECK: Callback Submit EMI: endpoint=ompt_scope_begin
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback Submit EMI: endpoint=ompt_scope_end
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_end
+/// CHECK-SAME: task_data=0x{{[0-f]+}} (0x[[ENCOUNTERING_TASK]])
+/// CHECK-SAME: target_task_data=(nil) (0x0)
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback Target EMI: kind=ompt_target_exit_data endpoint=ompt_scope_begin
+/// CHECK-SAME: task_data=0x{{[0-f]+}} (0x[[ENCOUNTERING_TASK]])
+/// CHECK-SAME: target_task_data=(nil) (0x0)
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA:[0-f]+]])
+
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin
+/// CHECK-SAME: target_task_data=(nil) (0x0)
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end
+/// CHECK-SAME: target_task_data=(nil) (0x0)
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin
+/// CHECK-SAME: target_task_data=(nil) (0x0)
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end
+/// CHECK-SAME: target_task_data=(nil) (0x0)
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin
+/// CHECK-SAME: target_task_data=(nil) (0x0)
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end
+/// CHECK-SAME: target_task_data=(nil) (0x0)
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback Target EMI: kind=ompt_target_exit_data endpoint=ompt_scope_end
+/// CHECK-SAME: task_data=0x{{[0-f]+}} (0x[[ENCOUNTERING_TASK]])
+/// CHECK-SAME: target_task_data=(nil) (0x0)
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
diff --git a/offload/test/ompt/target_tool_data_nowait.c b/offload/test/ompt/target_tool_data_nowait.c
new file mode 100644
index 0000000000000..3ef6a495aaf80
--- /dev/null
+++ b/offload/test/ompt/target_tool_data_nowait.c
@@ -0,0 +1,176 @@
+// clang-format off
+// RUN: env LIBOMP_NUM_HIDDEN_HELPER_THREADS=1 %libomptarget-compile-run-and-check-generic
+// REQUIRES: ompt
+// clang-format on
+
+#include <inttypes.h>
+#include <omp-tools.h>
+#include <omp.h>
+#include <stdio.h>
+#include <string.h>
+
+#include "register_with_host.h"
+
+#define N 1000000
+#define M 1000
+
+int main() {
+  float *x = malloc(N * sizeof(float));
+  float *y = malloc(N * sizeof(float));
+
+  for (int i = 0; i < N; i++) {
+    x[i] = 1;
+    y[i] = 1;
+  }
+
+#pragma omp target enter data map(to                                           \
+                                  : x [0:N]) map(alloc                         \
+                                                 : y [0:N])                    \
+    nowait depend(inout                                                        \
+                  : x)
+#pragma omp target teams distribute parallel for nowait depend(inout : x)
+  for (int i = 0; i < N; i++) {
+    for (int j = 0; j < M; j++) {
+      y[i] += 3 * x[i];
+    }
+  }
+
+#pragma omp target teams distribute parallel for nowait depend(inout : x)
+  for (int i = 0; i < N; i++) {
+    for (int j = 0; j < M; j++) {
+      y[i] += 3 * x[i];
+    }
+  }
+
+#pragma omp target exit data map(release                                       \
+                                 : x [0:N]) map(from                           \
+                                                : y [0:N]) nowait depend(inout \
+                                                                         : x)
+#pragma omp taskwait
+
+  printf("%f, %f\n", x[0], y[0]);
+
+  free(x);
+  free(y);
+  return 0;
+}
+
+// clang-format off
+/// CHECK: ompt_event_initial_task_begin
+/// CHECK-SAME: task_id=[[ENCOUNTERING_TASK:[0-f]+]]
+
+/// CHECK: ompt_event_task_create
+/// CHECK-SAME: new_task_id=[[TARGET_TASK_1:[0-f]+]]
+/// CHECK-SAME: task_type=ompt_task_target
+
+/// CHECK: ompt_event_task_create
+/// CHECK-SAME: new_task_id=[[TARGET_TASK_2:[0-f]+]]
+/// CHECK-SAME: task_type=ompt_task_target
+
+/// CHECK: ompt_event_task_create
+/// CHECK-SAME: new_task_id=[[TARGET_TASK_3:[0-f]+]]
+/// CHECK-SAME: task_type=ompt_task_target
+
+/// CHECK: ompt_event_task_create
+/// CHECK-SAME: new_task_id=[[TARGET_TASK_4:[0-f]+]]
+/// CHECK-SAME: task_type=ompt_task_target
+
+/// CHECK: Callback Target EMI: kind=ompt_target_enter_data endpoint=ompt_scope_begin
+/// CHECK-SAME: task_data=0x{{[0-f]+}} (0x[[ENCOUNTERING_TASK]])
+/// CHECK-SAME: target_task_data=0x{{[0-f]+}} (0x[[TARGET_TASK_1]])
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA:[0-f]+]])
+
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin
+/// CHECK-SAME: target_task_data=0x{{[0-f]+}} (0x[[TARGET_TASK_1]])
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end
+/// CHECK-SAME: target_task_data=0x{{[0-f]+}} (0x[[TARGET_TASK_1]])
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin
+/// CHECK-SAME: target_task_data=0x{{[0-f]+}} (0x[[TARGET_TASK_1]])
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end
+/// CHECK-SAME: target_task_data=0x{{[0-f]+}} (0x[[TARGET_TASK_1]])
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin
+/// CHECK-SAME: target_task_data=0x{{[0-f]+}} (0x[[TARGET_TASK_1]])
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end
+/// CHECK-SAME: target_task_data=0x{{[0-f]+}} (0x[[TARGET_TASK_1]])
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback Target EMI: kind=ompt_target_enter_data endpoint=ompt_scope_end
+/// CHECK-SAME: task_data=0x{{[0-f]+}} (0x[[ENCOUNTERING_TASK]])
+/// CHECK-SAME: target_task_data=0x{{[0-f]+}} (0x[[TARGET_TASK_1]])
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_begin
+/// CHECK-SAME: task_data=0x{{[0-f]+}} (0x[[ENCOUNTERING_TASK]])
+/// CHECK-SAME: target_task_data=0x{{[0-f]+}} (0x[[TARGET_TASK_2]])
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA:[0-f]+]])
+
+/// CHECK: Callback Submit EMI: endpoint=ompt_scope_begin
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback Submit EMI: endpoint=ompt_scope_end
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_end
+/// CHECK-SAME: task_data=0x{{[0-f]+}} (0x[[ENCOUNTERING_TASK]])
+/// CHECK-SAME: target_task_data=0x{{[0-f]+}} (0x[[TARGET_TASK_2]])
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_begin
+/// CHECK-SAME: task_data=0x{{[0-f]+}} (0x[[ENCOUNTERING_TASK]])
+/// CHECK-SAME: target_task_data=0x{{[0-f]+}} (0x[[TARGET_TASK_3]])
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA:[0-f]+]])
+
+/// CHECK: Callback Submit EMI: endpoint=ompt_scope_begin
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback Submit EMI: endpoint=ompt_scope_end
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_end
+/// CHECK-SAME: task_data=0x{{[0-f]+}} (0x[[ENCOUNTERING_TASK]])
+/// CHECK-SAME: target_task_data=0x{{[0-f]+}} (0x[[TARGET_TASK_3]])
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback Target EMI: kind=ompt_target_exit_data endpoint=ompt_scope_begin
+/// CHECK-SAME: task_data=0x{{[0-f]+}} (0x[[ENCOUNTERING_TASK]])
+/// CHECK-SAME: target_task_data=0x{{[0-f]+}} (0x[[TARGET_TASK_4]])
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA:[0-f]+]])
+
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin
+/// CHECK-SAME: target_task_data=0x{{[0-f]+}} (0x[[TARGET_TASK_4]])
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end
+/// CHECK-SAME: target_task_data=0x{{[0-f]+}} (0x[[TARGET_TASK_4]])
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin
+/// CHECK-SAME: target_task_data=0x{{[0-f]+}} (0x[[TARGET_TASK_4]])
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end
+/// CHECK-SAME: target_task_data=0x{{[0-f]+}} (0x[[TARGET_TASK_4]])
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin
+/// CHECK-SAME: target_task_data=0x{{[0-f]+}} (0x[[TARGET_TASK_4]])
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end
+/// CHECK-SAME: target_task_data=0x{{[0-f]+}} (0x[[TARGET_TASK_4]])
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
+
+/// CHECK: Callback Target EMI: kind=ompt_target_exit_data endpoint=ompt_scope_end
+/// CHECK-SAME: task_data=0x{{[0-f]+}} (0x[[ENCOUNTERING_TASK]])
+/// CHECK-SAME: target_task_data=0x{{[0-f]+}} (0x[[TARGET_TASK_4]])
+/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
diff --git a/offload/test/ompt/target_tool_data_nowait_nodepend.c b/offload/test/ompt/target_tool_data_nowait_nodepend.c
new file mode 100644
index 0000000000000..15b3971383535
--- /dev/null
+++ b/offload/test/ompt/target_tool_data_nowait_nodepend.c
@@ -0,0 +1,63 @@
+// clang-format off
+// RUN: env LIBOMP_NUM_HIDDEN_HELPER_THREADS=1 %libomptarget-compile-run-and-check-generic
+// REQUIRES: ompt
+// clang-format on
+
+#include <inttypes.h>
+#include <omp-tools.h>
+#include <omp.h>
+#include <stdio.h>
+#include <string.h>
+
+#include "register_with_host.h"
+
+#define N 1000000
+#define M 1000
+
+int main() {
+  float *x = malloc(N * sizeof(float));
+  float *y = malloc(N * sizeof(float));
+  float *a = malloc(N * sizeof(float));
+  float *b = malloc(N * sizeof(float));
+
+  for (int i = 0; i < N; i++) {
+    x[i] = 1;
+    y[i] = 1;
+    a[i] = 1;
+    b[i] = 1;
+  }
+
+#pragma omp target teams distribute parallel for nowait map(to                 \
+                                                            : x [0:N])         \
+    map(from                                                                   \
+        : y [0:N])
+  for (int i = 0; i < N; i++) {
+    for (int j = 0; j < M; j++) {
+      y[i] += 3 * x[i];
+    }
+  }
+
+#pragma omp target teams distribute parallel for nowait map(to                 \
+                                                            : a [0:N])         \
+    map(from                                                                   \
+        : b [0:N])
+  for (int i = 0; i < N; i++) {
+    for (int j = 0; j < M; j++) {
+      b[i] += 3 * a[i];
+    }
+  }
+
+#pragma omp taskwait
+
+  printf("%f, %f, %f, %f\n", x[0], y[0], a[0], b[0]);
+
+  free(x);
+  free(y);
+  free(a);
+  free(b);
+  return 0;
+}
+
+// clang-format off
+/// CHECK-NOT: target_task_data=(nil) (0x0)
+/// CHECK-NOT: target_data=(nil) (0x0)
diff --git a/offload/test/ompt/veccopy.c b/offload/test/ompt/veccopy.c
index f28d94f524bb8..24d7363e65599 100644
--- a/offload/test/ompt/veccopy.c
+++ b/offload/test/ompt/veccopy.c
@@ -1,6 +1,8 @@
+// clang-format off
 // RUN: %libomptarget-compile-run-and-check-generic
 // REQUIRES: ompt
 // REQUIRES: gpu
+// clang-format on
 
 /*
  * Example OpenMP program that registers non-EMI callbacks
@@ -54,48 +56,47 @@ int main() {
 // clang-format off
 /// CHECK: Callback Init:
 /// CHECK: Callback Load:
-/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=1 device_num=[[DEVICE_NUM:[0-9]+]]
+/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=ompt_target endpoint=ompt_scope_begin device_num=[[DEVICE_NUM:[0-9]+]]
 /// CHECK-NOT: code=(nil)
 /// CHECK: code=[[CODE1:.*]]
-/// 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=ompt_target_data_alloc
 /// CHECK: code=[[CODE1]]
-/// 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=ompt_target_data_transfer_to_device
 /// CHECK: code=[[CODE1]]
-/// 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=ompt_target_data_alloc
 /// CHECK: code=[[CODE1]]
-/// 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=ompt_target_data_transfer_to_device
 /// CHECK: code=[[CODE1]]
 /// 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=ompt_target_data_transfer_from_device
 /// CHECK: code=[[CODE1]]
-/// 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=ompt_target_data_transfer_from_device
 /// CHECK: code=[[CODE1]]
-/// 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=ompt_target_data_delete
 /// CHECK: code=[[CODE1]]
-/// 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=ompt_target_data_delete
 /// CHECK: code=[[CODE1]]
-/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=2 device_num=[[DEVICE_NUM]] code=[[CODE1]]
+/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=ompt_target endpoint=ompt_scope_end device_num=[[DEVICE_NUM]] code=[[CODE1]]
 
-/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=1
-/// device_num=[[DEVICE_NUM]]
+/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=ompt_target endpoint=ompt_scope_begin device_num=[[DEVICE_NUM]]
 /// CHECK-NOT: code=(nil)
 /// CHECK: code=[[CODE2:.*]]
-/// 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=ompt_target_data_alloc
 /// CHECK: code=[[CODE2]]
-/// 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=ompt_target_data_transfer_to_device
 /// CHECK: code=[[CODE2]]
-/// 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=ompt_target_data_alloc
 /// CHECK: code=[[CODE2]]
-/// 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=ompt_target_data_transfer_to_device
 /// CHECK: code=[[CODE2]]
 /// 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=ompt_target_data_transfer_from_device
 /// CHECK: code=[[CODE2]]
-/// 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=ompt_target_data_transfer_from_device
 /// CHECK: code=[[CODE2]]
-/// 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=ompt_target_data_delete
 /// CHECK: code=[[CODE2]]
-/// 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=ompt_target_data_delete
 /// CHECK: code=[[CODE2]]
-/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=2 device_num=[[DEVICE_NUM]] code=[[CODE2]]
+/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=ompt_target endpoint=ompt_scope_end device_num=[[DEVICE_NUM]] code=[[CODE2]]
 /// CHECK: Callback Fini:
diff --git a/offload/test/ompt/veccopy_data.c b/offload/test/ompt/veccopy_data.c
index 059ca97c3cde3..9df5374193e94 100644
--- a/offload/test/ompt/veccopy_data.c
+++ b/offload/test/ompt/veccopy_data.c
@@ -1,6 +1,8 @@
+// clang-format off
 // RUN: %libomptarget-compile-run-and-check-generic
 // REQUIRES: ompt
 // REQUIRES: gpu
+// clang-format on
 
 /*
  * Example OpenMP program that registers EMI callbacks.
@@ -73,85 +75,86 @@ int main() {
   return rc;
 }
 
+// clang-format off
 /// CHECK-NOT: Callback Target EMI:
 /// CHECK-NOT: device_num=-1
 /// CHECK: Callback Init:
 /// CHECK: Callback Load:
-/// CHECK: Callback Target EMI: kind=2 endpoint=1
+/// CHECK: Callback Target EMI: kind=ompt_target_enter_data endpoint=ompt_scope_begin
 /// CHECK-NOT: device_num=-1
 /// CHECK-NOT: code=(nil)
 /// CHECK: code=[[CODE1:.*]]
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_alloc
 /// CHECK: code=[[CODE1]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_alloc
 /// CHECK-NOT: dest=(nil)
 /// CHECK: code=[[CODE1]]
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_to_device
 /// CHECK: code=[[CODE1]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_to_device
 /// CHECK: code=[[CODE1]]
-/// CHECK: Callback Target EMI: kind=2 endpoint=2
+/// CHECK: Callback Target EMI: kind=ompt_target_enter_data endpoint=ompt_scope_end
 /// CHECK-NOT: device_num=-1
 /// CHECK: code=[[CODE1]]
-/// CHECK: Callback Target EMI: kind=1 endpoint=1
+/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_begin
 /// CHECK-NOT: device_num=-1
 /// CHECK-NOT: code=(nil)
 /// CHECK: code=[[CODE2:.*]]
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_alloc
 /// CHECK: code=[[CODE2]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_alloc
 /// CHECK-NOT: dest=(nil)
 /// CHECK: code=[[CODE2]]
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_to_device
 /// CHECK: code=[[CODE2]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_to_device
 /// CHECK: code=[[CODE2]]
-/// 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 Submit EMI: endpoint=ompt_scope_begin  req_num_teams=1
+/// CHECK: Callback Submit EMI: endpoint=ompt_scope_end  req_num_teams=1
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_from_device
 /// CHECK: code=[[CODE2]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_from_device
 /// CHECK: code=[[CODE2]]
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_delete
 /// CHECK: code=[[CODE2]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_delete
 /// CHECK: code=[[CODE2]]
-/// CHECK: Callback Target EMI: kind=1 endpoint=2
+/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_end
 /// CHECK-NOT: device_num=-1
 /// CHECK: code=[[CODE2]]
-/// CHECK: Callback Target EMI: kind=3 endpoint=1
+/// CHECK: Callback Target EMI: kind=ompt_target_exit_data endpoint=ompt_scope_begin
 /// CHECK-NOT: device_num=-1
 /// CHECK-NOT: code=(nil)
 /// CHECK: code=[[CODE3:.*]]
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_from_device
 /// CHECK: code=[[CODE3]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_from_device
 /// CHECK: code=[[CODE3]]
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_delete
 /// CHECK: code=[[CODE3]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_delete
 /// CHECK: code=[[CODE3]]
-/// CHECK: Callback Target EMI: kind=3 endpoint=2
+/// CHECK: Callback Target EMI: kind=ompt_target_exit_data endpoint=ompt_scope_end
 /// CHECK-NOT: device_num=-1
 /// CHECK: code=[[CODE3]]
-/// CHECK: Callback Target EMI: kind=1 endpoint=1
+/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_begin
 /// CHECK-NOT: device_num=-1
 /// CHECK-NOT: code=(nil)
 /// CHECK: code=[[CODE4:.*]]
-/// CHECK: Callback Submit EMI: endpoint=1  req_num_teams=1
-/// CHECK: Callback Submit EMI: endpoint=2  req_num_teams=1
-/// CHECK: Callback Target EMI: kind=1 endpoint=2
+/// CHECK: Callback Submit EMI: endpoint=ompt_scope_begin  req_num_teams=1
+/// CHECK: Callback Submit EMI: endpoint=ompt_scope_end  req_num_teams=1
+/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_end
 /// CHECK-NOT: device_num=-1
 /// CHECK: code=[[CODE4]]
-/// CHECK: Callback Target EMI: kind=4 endpoint=1
+/// CHECK: Callback Target EMI: kind=ompt_target_update endpoint=ompt_scope_begin
 /// CHECK-NOT: device_num=-1
 /// CHECK-NOT: code=(nil)
 /// CHECK: code=[[CODE5:.*]]
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_from_device
 /// CHECK: code=[[CODE5]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_from_device
 /// CHECK: code=[[CODE5]]
-/// CHECK: Callback Target EMI: kind=4 endpoint=2
+/// CHECK: Callback Target EMI: kind=ompt_target_update endpoint=ompt_scope_end
 /// CHECK-NOT: device_num=-1
 /// CHECK: code=[[CODE5]]
 /// CHECK: Callback Fini:
diff --git a/offload/test/ompt/veccopy_disallow_both.c b/offload/test/ompt/veccopy_disallow_both.c
index b531a628803e4..bfc67c5f4d274 100644
--- a/offload/test/ompt/veccopy_disallow_both.c
+++ b/offload/test/ompt/veccopy_disallow_both.c
@@ -1,6 +1,8 @@
+// clang-format off
 // RUN: %libomptarget-compile-run-and-check-generic
 // REQUIRES: ompt
 // REQUIRES: gpu
+// clang-format on
 
 /*
  * Example OpenMP program that shows that both EMI and non-EMI
@@ -54,48 +56,49 @@ int main() {
   return rc;
 }
 
+// clang-format off
 /// 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 Target EMI: kind=ompt_target endpoint=ompt_scope_begin
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_alloc
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_alloc
 /// CHECK-NOT: dest=(nil)
-/// 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=ompt_scope_begin optype=ompt_target_data_transfer_to_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_to_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_alloc
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_alloc
 /// CHECK-NOT: dest=(nil)
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=2
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_to_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_to_device
 /// 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=ompt_scope_begin optype=ompt_target_data_transfer_from_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_from_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_from_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_from_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_delete
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_delete
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_delete
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_delete
+/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_end
+/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_begin
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_alloc
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_alloc
 /// CHECK-NOT: dest=(nil)
-/// 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=ompt_scope_begin optype=ompt_target_data_transfer_to_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_to_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_alloc
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_alloc
 /// CHECK-NOT: dest=(nil)
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=2
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_to_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_to_device
 /// 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 DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_from_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_from_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_from_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_from_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_delete
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_delete
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_delete
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_delete
+/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_end
 /// CHECK: Callback Fini:
diff --git a/offload/test/ompt/veccopy_emi.c b/offload/test/ompt/veccopy_emi.c
index 2c57a85c14756..a1427b86a58fa 100644
--- a/offload/test/ompt/veccopy_emi.c
+++ b/offload/test/ompt/veccopy_emi.c
@@ -1,6 +1,8 @@
+// clang-format off
 // RUN: %libomptarget-compile-run-and-check-generic
 // REQUIRES: ompt
 // REQUIRES: gpu
+// clang-format on
 
 /*
  * Example OpenMP program that registers EMI callbacks
@@ -52,89 +54,90 @@ int main() {
   return rc;
 }
 
+// clang-format off
 /// CHECK: Callback Init:
 /// CHECK: Callback Load:
-/// CHECK: Callback Target EMI: kind=1 endpoint=1
+/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_begin
 /// CHECK-NOT: code=(nil)
 /// CHECK: code=[[CODE1:.*]]
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_alloc
 /// CHECK: code=[[CODE1]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_alloc
 /// CHECK-NOT: dest=(nil)
 /// CHECK: code=[[CODE1]]
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_to_device
 /// CHECK: code=[[CODE1]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_to_device
 /// CHECK: code=[[CODE1]]
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_alloc
 /// CHECK: code=[[CODE1]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_alloc
 /// CHECK-NOT: dest=(nil)
 /// CHECK: code=[[CODE1]]
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_to_device
 /// CHECK: code=[[CODE1]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_to_device
 /// CHECK: code=[[CODE1]]
-/// 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 Submit EMI: endpoint=ompt_scope_begin req_num_teams=1
+/// CHECK: Callback Submit EMI: endpoint=ompt_scope_end req_num_teams=1
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_from_device
 /// CHECK: code=[[CODE1]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_from_device
 /// CHECK: code=[[CODE1]]
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_from_device
 /// CHECK: code=[[CODE1]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_from_device
 /// CHECK: code=[[CODE1]]
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_delete
 /// CHECK: code=[[CODE1]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_delete
 /// CHECK: code=[[CODE1]]
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_delete
 /// CHECK: code=[[CODE1]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_delete
 /// CHECK: code=[[CODE1]]
-/// CHECK: Callback Target EMI: kind=1 endpoint=2
+/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_end
 /// CHECK: code=[[CODE1]]
 
-/// CHECK: Callback Target EMI: kind=1 endpoint=1
+/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_begin
 /// CHECK-NOT: code=(nil)
 /// CHECK: code=[[CODE2:.*]]
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_alloc
 /// CHECK: code=[[CODE2]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_alloc
 /// CHECK-NOT: dest=(nil)
 /// CHECK: code=[[CODE2]]
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_to_device
 /// CHECK: code=[[CODE2]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_to_device
 /// CHECK: code=[[CODE2]]
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_alloc
 /// CHECK: code=[[CODE2]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=1
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_alloc
 /// CHECK-NOT: dest=(nil)
 /// CHECK: code=[[CODE2]]
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_to_device
 /// CHECK: code=[[CODE2]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=2
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_to_device
 /// CHECK: code=[[CODE2]]
-/// 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 Submit EMI: endpoint=ompt_scope_begin req_num_teams=0
+/// CHECK: Callback Submit EMI: endpoint=ompt_scope_end req_num_teams=0
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_from_device
 /// CHECK: code=[[CODE2]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_from_device
 /// CHECK: code=[[CODE2]]
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_from_device
 /// CHECK: code=[[CODE2]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=3
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_from_device
 /// CHECK: code=[[CODE2]]
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_delete
 /// CHECK: code=[[CODE2]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_delete
 /// CHECK: code=[[CODE2]]
-/// CHECK: Callback DataOp EMI: endpoint=1 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_delete
 /// CHECK: code=[[CODE2]]
-/// CHECK: Callback DataOp EMI: endpoint=2 optype=4
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_delete
 /// CHECK: code=[[CODE2]]
-/// CHECK: Callback Target EMI: kind=1 endpoint=2
+/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_end
 /// CHECK: code=[[CODE2]]
 /// CHECK: Callback Fini:
diff --git a/offload/test/ompt/veccopy_emi_map.c b/offload/test/ompt/veccopy_emi_map.c
index fa18a43cd8a50..450faa1f28b0e 100644
--- a/offload/test/ompt/veccopy_emi_map.c
+++ b/offload/test/ompt/veccopy_emi_map.c
@@ -1,6 +1,8 @@
+// clang-format off
 // RUN: %libomptarget-compile-run-and-check-generic
 // REQUIRES: ompt
 // REQUIRES: gpu
+// clang-format on
 
 /*
  * Example OpenMP program that shows that map-EMI callbacks are not supported.
@@ -52,51 +54,52 @@ int main() {
   return rc;
 }
 
+// clang-format off
 /// 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 Target EMI: kind=ompt_target endpoint=ompt_scope_begin
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_alloc
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_alloc
 /// CHECK-NOT: dest=(nil)
-/// 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=ompt_scope_begin optype=ompt_target_data_transfer_to_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_to_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_alloc
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_alloc
 /// CHECK-NOT: dest=(nil)
-/// 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=ompt_scope_begin optype=ompt_target_data_transfer_to_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_to_device
+/// CHECK: Callback Submit EMI: endpoint=ompt_scope_begin req_num_teams=1
+/// CHECK: Callback Submit EMI: endpoint=ompt_scope_end req_num_teams=1
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_from_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_from_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_from_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_from_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_delete
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_delete
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_delete
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_delete
+/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_end
+/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_begin
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_alloc
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_alloc
 /// CHECK-NOT: dest=(nil)
-/// 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=ompt_scope_begin optype=ompt_target_data_transfer_to_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_to_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_alloc
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_alloc
 /// CHECK-NOT: dest=(nil)
-/// 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 DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_to_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_to_device
+/// CHECK: Callback Submit EMI: endpoint=ompt_scope_begin req_num_teams=0
+/// CHECK: Callback Submit EMI: endpoint=ompt_scope_end req_num_teams=0
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_from_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_from_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_from_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_from_device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_delete
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_delete
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_delete
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_delete
+/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_end
 /// CHECK: Callback Fini:
diff --git a/offload/test/ompt/veccopy_map.c b/offload/test/ompt/veccopy_map.c
index 2e817d328e59f..12e141ea74d07 100644
--- a/offload/test/ompt/veccopy_map.c
+++ b/offload/test/ompt/veccopy_map.c
@@ -1,6 +1,8 @@
+// clang-format off
 // RUN: %libomptarget-compile-run-and-check-generic
 // REQUIRES: ompt
 // REQUIRES: gpu
+// clang-format on
 
 /*
  * Example OpenMP program that shows that map callbacks are not supported.
@@ -51,31 +53,31 @@ int main() {
   return rc;
 }
 
-
+// clang-format off
 /// 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 Target: target_id=[[TARGET_ID:[0-9]+]] kind=ompt_target endpoint=ompt_scope_begin
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_alloc
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_transfer_to_device
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_alloc
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_transfer_to_device
 /// 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 DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_transfer_from_device
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_transfer_from_device
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_delete
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_delete
+/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=ompt_target endpoint=ompt_scope_end
+
+/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=ompt_target endpoint=ompt_scope_begin
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_alloc
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_transfer_to_device
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_alloc
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_transfer_to_device
 /// 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 DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_transfer_from_device
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_transfer_from_device
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_delete
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_delete
+/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=ompt_target endpoint=ompt_scope_end
 /// CHECK: Callback Fini:
diff --git a/offload/test/ompt/veccopy_no_device_init.c b/offload/test/ompt/veccopy_no_device_init.c
index 8ee8243281187..ade06fcc92290 100644
--- a/offload/test/ompt/veccopy_no_device_init.c
+++ b/offload/test/ompt/veccopy_no_device_init.c
@@ -1,6 +1,7 @@
 // clang-format off
 // RUN: %libomptarget-compile-run-and-check-generic
 // REQUIRES: ompt
+// clang-format on
 
 /*
  * Example OpenMP program that shows that if no device init callback
@@ -51,30 +52,31 @@ int main() {
 
   return rc;
 }
+
 // clang-format off
 /// CHECK-NOT: 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 Target: target_id=[[TARGET_ID:[0-9]+]] kind=ompt_target endpoint=ompt_scope_begin
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_alloc
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_transfer_to_device
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_alloc
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_transfer_to_device
 /// 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 DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_transfer_from_device
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_transfer_from_device
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_delete
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_delete
+/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=ompt_target endpoint=ompt_scope_end
 
-/// 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 Target: target_id=[[TARGET_ID:[0-9]+]] kind=ompt_target endpoint=ompt_scope_begin
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_alloc
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_transfer_to_device
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_alloc
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_transfer_to_device
 /// 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 DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_transfer_from_device
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_transfer_from_device
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_delete
+/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_delete
+/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=ompt_target endpoint=ompt_scope_end
 /// CHECK-NOT: Callback Fini:
diff --git a/offload/test/ompt/veccopy_wrong_return.c b/offload/test/ompt/veccopy_wrong_return.c
index 2d07b4e1bf04a..17327f3553817 100644
--- a/offload/test/ompt/veccopy_wrong_return.c
+++ b/offload/test/ompt/veccopy_wrong_return.c
@@ -1,5 +1,7 @@
+// clang-format off
 // RUN: %libomptarget-compile-run-and-check-generic
 // REQUIRES: ompt
+// clang-format on
 
 /*
  * Example OpenMP program that shows that if the initialize function
@@ -51,29 +53,30 @@ int main() {
   return rc;
 }
 
+// clang-format off
 /// 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 Target: target_id=[[TARGET_ID:[0-9]+]] kind=ompt_target endpoint=ompt_scope_begin
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_alloc
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_transfer_to_device
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_alloc
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_transfer_to_device
 /// 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 DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_transfer_from_device
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_transfer_from_device
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_delete
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_delete
+/// CHECK-NOT: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=ompt_target endpoint=ompt_scope_end
 
-/// 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 Target: target_id=[[TARGET_ID:[0-9]+]] kind=ompt_target endpoint=ompt_scope_begin
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_alloc
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_transfer_to_device
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_alloc
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_transfer_to_device
 /// 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 DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_transfer_from_device
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_transfer_from_device
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_delete
+/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=ompt_target_data_delete
+/// CHECK-NOT: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=ompt_target endpoint=ompt_scope_end
 /// CHECK-NOT: Callback Fini

