[llvm] [Offload][OMPT][NFC] Add test for `host_op_id` consistency (PR #201144)

[Jan André Reuter via llvm-commits]

https://github.com/Thyre created https://github.com/llvm/llvm-project/pull/201144

Add test to ensure that `host_op_id` is correctly passed for external monitoring interface callbacks. A set `host_op_id` during `ompt_scope_begin` should show up during `ompt_scope_end`.

This catches cases like 2c8ca9679e85e466e8d7105027df322d03c45e3d, where a nested `data_op` callback causes an incorrect `host_op_id` on the outer `ompt_scope_end` event.

>From 34c2d55f96c234f3ea1747f381d4a11f914a87a6 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Jan=20Andr=C3=A9=20Reuter?= <j.reuter at fz-juelich.de>
Date: Tue, 2 Jun 2026 16:35:07 +0200
Subject: [PATCH] [Offload][OMPT][NFC] Add test for `host_op_id` consistency
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

Add test to ensure that `host_op_id` is correctly passed for external
monitoring interface callbacks. A set `host_op_id` during
`ompt_scope_begin` should show up during `ompt_scope_end`.

This catches cases like 2c8ca9679e85e466e8d7105027df322d03c45e3d, where
a nested `data_op` callback causes an incorrect `host_op_id` on the
outer `ompt_scope_end` event.

Signed-off-by: Jan André Reuter <j.reuter at fz-juelich.de>
---
 offload/test/ompt/host_op_id_emi.c | 108 +++++++++++++++++++++++++++++
 1 file changed, 108 insertions(+)
 create mode 100644 offload/test/ompt/host_op_id_emi.c

diff --git a/offload/test/ompt/host_op_id_emi.c b/offload/test/ompt/host_op_id_emi.c
new file mode 100644
index 0000000000000..e45119086628d
--- /dev/null
+++ b/offload/test/ompt/host_op_id_emi.c
@@ -0,0 +1,108 @@
+// clang-format off
+// RUN: %libomptarget-compile-run-and-check-generic
+// REQUIRES: ompt
+// REQUIRES: gpu
+// clang-format on
+
+/*
+ * Example OpenMP program that checks if EMI callbacks
+ * correctly pass host_op_id between ompt_scope_begin and ompt_scope_end
+ */
+
+#include <assert.h>
+#include <omp.h>
+#include <stdio.h>
+
+#include "callbacks.h"
+#include "register_emi.h"
+
+int main(void) {
+  int NumDevices = omp_get_num_devices();
+  assert(NumDevices > 0 && "No device(s) present.");
+  int Device = omp_get_default_device();
+  int Host = omp_get_initial_device();
+  // Note: Zero value depicts an OFFLOAD_SUCCESS
+  int Status;
+
+  printf("Allocating Memory on Device\n");
+  int *DevPtr = (int *)omp_target_alloc(sizeof(int), Device);
+  assert(DevPtr && "Could not allocate memory on device.");
+  int *HstPtr = (int *)malloc(sizeof(int));
+  *HstPtr = 42;
+
+  printf("Testing: Host to Device\n");
+  Status = omp_target_memcpy(DevPtr, HstPtr, sizeof(int), 0, 0, Device, Host);
+  assert(Status == 0 && "H2D memory copy operation failed.\n");
+
+  printf("Testing: Device to Device\n");
+  Status = omp_target_memcpy(DevPtr, DevPtr, sizeof(int), 0, 0, Device, Device);
+  assert(Status == 0 && "D2D memory copy operation failed.\n");
+
+  printf("Testing: Device to Host\n");
+  Status = omp_target_memcpy(HstPtr, DevPtr, sizeof(int), 0, 0, Host, Device);
+  assert(Status == 0 && "D2H memory copy operation failed.\n");
+
+  printf("Checking Correctness\n");
+  assert(*HstPtr == 42);
+
+  printf("Setting Device Memory to 0\n");
+  int *DevMemsetPtr = omp_target_memset(DevPtr, 0, sizeof(int), Device);
+  assert(DevMemsetPtr == DevPtr && "Memset returned incorrect pointer.\n");
+
+  printf("Submitting a Kernel\n");
+#pragma omp target
+  *DevPtr = 42;
+
+  printf("Freeing Memory on Device\n");
+  free(HstPtr);
+  omp_target_free(DevPtr, Device);
+
+  return 0;
+}
+
+// clang-format off
+
+/// CHECK: Callback Init:
+
+/// CHECK: Allocating Memory on Device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_alloc
+/// CHECK-SAME: host_op_id=[[HOSTOP1:0x[[:xdigit:]]+ \(0x[[:xdigit:]]+\)]]
+/// CHECK-SAME: src_device_num=[[HOST:-1]]
+/// CHECK-SAME: dest_device_num=[[DEVICE:[0-9]+]]
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_alloc {{.+}} host_op_id=[[HOSTOP1]] {{.+}} src_device_num=[[HOST]] {{.+}} dest_device_num=[[DEVICE]]
+
+/// CHECK: Testing: Host to Device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_to_device
+/// CHECK-SAME: host_op_id=[[HOSTOP2:0x[[:xdigit:]]+ \(0x[[:xdigit:]]+\)]]
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_to_device {{.+}} host_op_id=[[HOSTOP2]] {{.+}} src_device_num=[[HOST]] {{.+}} dest_device_num=[[DEVICE]]
+
+/// CHECK: Testing: Device to 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-SAME: host_op_id=[[HOSTOP3:0x[[:xdigit:]]+ \(0x[[:xdigit:]]+\)]]
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_from_device {{.+}} host_op_id=[[HOSTOP3]] {{.+}} src_device_num=[[DEVICE]] {{.+}} dest_device_num=[[DEVICE]]
+
+/// CHECK: Testing: Device to 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-SAME: host_op_id=[[HOSTOP4:0x[[:xdigit:]]+ \(0x[[:xdigit:]]+\)]]
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_from_device {{.+}} host_op_id=[[HOSTOP4]] {{.+}} src_device_num=[[DEVICE]] {{.+}} dest_device_num=[[HOST]]
+
+/// CHECK: Checking Correctness
+
+/// CHECK: Setting Device Memory to 0
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_to_device {{.+}} src_device_num=[[HOST]] {{.+}} dest_device_num=[[DEVICE]]
+/// CHECK-SAME: host_op_id=[[HOSTOP5:0x[[:xdigit:]]+ \(0x[[:xdigit:]]+\)]]
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_to_device {{.+}} host_op_id=[[HOSTOP5]] {{.+}} src_device_num=[[HOST]] {{.+}} dest_device_num=[[DEVICE]]
+
+/// CHECK: Submitting a Kernel
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_to_device {{.+}} src_device_num=[[HOST]] {{.+}} dest_device_num=[[DEVICE]]
+/// CHECK-SAME: host_op_id=[[HOSTOP6:0x[[:xdigit:]]+ \(0x[[:xdigit:]]+\)]]
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_transfer_to_device {{.+}} host_op_id=[[HOSTOP6]] {{.+}} src_device_num=[[HOST]] {{.+}} dest_device_num=[[DEVICE]]
+
+/// CHECK: Freeing Memory on Device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_delete {{.+}} src_device_num=[[DEVICE]]
+/// CHECK-SAME: host_op_id=[[HOSTOP7:0x[[:xdigit:]]+ \(0x[[:xdigit:]]+\)]]
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_delete {{.+}} host_op_id=[[HOSTOP7]] {{.+}} src_device_num=[[DEVICE]]
+
+/// CHECK: Callback Fini:
+
+// clang-format on
\ No newline at end of file