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

[Jan André Reuter via llvm-commits]

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

>From 4764abc0aa46662685446ef415217eab417032b3 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 1/2] [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 | 99 ++++++++++++++++++++++++++++++
 1 file changed, 99 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..6616e7f04ad46
--- /dev/null
+++ b/offload/test/ompt/host_op_id_emi.c
@@ -0,0 +1,99 @@
+// 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 {{.+}} host_op_id=[[HOSTOP1:0x[[:xdigit:]]+ \(0x[[:xdigit:]]+\)]] {{.+}}
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_alloc {{.+}} host_op_id=[[HOSTOP1]] {{.+}}
+
+/// CHECK: Testing: Host to Device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_to_device {{.+}} 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]] {{.+}}
+
+/// CHECK: Testing: Device to Device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_from_device {{.+}} 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]] {{.+}}
+
+/// CHECK: Testing: Device to Host
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_transfer_from_device {{.+}} 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]] {{.+}}
+
+/// CHECK: Checking Correctness
+
+/// CHECK: Setting Device Memory to 0
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_memset {{.+}} host_op_id=[[HOSTOP5:0x[[:xdigit:]]+ \(0x[[:xdigit:]]+\)]] {{.+}}
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end optype=ompt_target_data_memset {{.+}} host_op_id=[[HOSTOP5]] {{.+}}
+
+/// CHECK: Submitting a Kernel
+/// CHECK: Callback Submit EMI: endpoint=ompt_scope_begin {{.+}} host_op_id=[[HOSTOP6:0x[[:xdigit:]]+ \(0x[[:xdigit:]]+\)]]
+/// CHECK: Callback Submit EMI: endpoint=ompt_scope_end {{.+}} host_op_id=[[HOSTOP6]]
+
+/// CHECK: Freeing Memory on Device
+/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin optype=ompt_target_data_delete {{.+}} 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]] {{.+}}
+
+/// CHECK: Callback Fini:
+
+// clang-format on

>From 42c6e29b4fb6d09df5e48e4e63ac80018a508de1 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Jan=20Andr=C3=A9=20Reuter?= <j.reuter at fz-juelich.de>
Date: Fri, 5 Jun 2026 23:57:01 +0200
Subject: [PATCH 2/2] [OMPT][Offload] Explicitly pass `is_device_ptr` for
 host_op_id_emi test
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

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

diff --git a/offload/test/ompt/host_op_id_emi.c b/offload/test/ompt/host_op_id_emi.c
index 6616e7f04ad46..c96e602cda1f5 100644
--- a/offload/test/ompt/host_op_id_emi.c
+++ b/offload/test/ompt/host_op_id_emi.c
@@ -50,7 +50,7 @@ int main(void) {
   assert(DevMemsetPtr == DevPtr && "Memset returned incorrect pointer.\n");
 
   printf("Submitting a Kernel\n");
-#pragma omp target
+#pragma omp target is_device_ptr(DevPtr)
   *DevPtr = 42;
 
   printf("Freeing Memory on Device\n");