[Openmp-commits] [llvm] [openmp] [Offload][OMPT] Properly implement `ompt_get_target_info` (PR #202392)
Jan André Reuter via Openmp-commits
openmp-commits at lists.llvm.org
Tue Jun 9 01:32:47 PDT 2026
https://github.com/Thyre updated https://github.com/llvm/llvm-project/pull/202392
>From db32ce551a41b6279fdb272bc6d4c8cd7921c458 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Jan=20Andr=C3=A9=20Reuter?= <jan at zyten.de>
Date: Mon, 8 Jun 2026 18:49:18 +0200
Subject: [PATCH 1/3] [Offload][OMPT] Properly implement `ompt_get_target_info`
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit
While `ompt_get_target_info` was present for quite a while, the implementation
always incorrectly returned that the encountering thread was not in a target
region. Hence, a tool was not able to retrieve the `target_id`, `host_op_id`
or `device_num`.
While the usability of the function is questionable, since the thread should
already have access to these variables via registered callbacks, still
provide a proper implementation for tools.
For this, `libomp` provides storage for a function pointer set by
`libomptarget` during the initialization. `libomp` can use this function
pointer to then query the information stored already.
If no function is provided e.g., because `libomptarget` was not linked, the
function behaves as before, but sets the values to `ompt_id_none` and
`omp_invalid_device`.
Signed-off-by: Jan André Reuter <jan at zyten.de>
---
offload/include/OpenMP/OMPT/Interface.h | 17 ++++++-
offload/libomptarget/OpenMP/OMPT/Callback.cpp | 51 ++++++++++++++++---
openmp/runtime/src/ompt-general.cpp | 23 ++++++++-
3 files changed, 82 insertions(+), 9 deletions(-)
diff --git a/offload/include/OpenMP/OMPT/Interface.h b/offload/include/OpenMP/OMPT/Interface.h
index 6961641769b76..080104a8036df 100644
--- a/offload/include/OpenMP/OMPT/Interface.h
+++ b/offload/include/OpenMP/OMPT/Interface.h
@@ -31,6 +31,9 @@
/// 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)();
+/// Callback from libomp -> libomptarget to retrieve the data
+/// required for ompt_get_target_info entry point.
+typedef void (*ompt_set_target_info_t)(ompt_get_target_info_t);
namespace llvm {
namespace omp {
@@ -41,6 +44,7 @@ namespace ompt {
/// 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_set_target_info_t ompt_set_target_info_fn;
/// Used to maintain execution state for this thread
class Interface {
@@ -221,16 +225,21 @@ class Interface {
void setTargetDataValue(uint64_t DataValue) { TargetData.value = DataValue; }
void setTargetDataPtr(void *DataPtr) { TargetData.ptr = DataPtr; }
void setHostOpId(ompt_id_t OpId) { HostOpId = OpId; }
+ void setDeviceNum(uint64_t Id) { DeviceNum = Id; }
/// Getters for target region and target operation correlation ids
uint64_t getTargetDataValue() { return TargetData.value; }
void *getTargetDataPtr() { return TargetData.ptr; }
ompt_id_t getHostOpId() { return HostOpId; }
+ uint64_t getDeviceNum() { return DeviceNum; }
private:
/// Target operations id
ompt_id_t HostOpId = 0;
+ /// Target device id associated with current target region
+ int32_t DeviceNum = /*omp_invalid_device*/ -2;
+
/// Target region data
ompt_data_t TargetData = ompt_data_none;
@@ -246,8 +255,14 @@ class Interface {
/// Used for marking end of a data operation
void endTargetDataOperation();
+ /// Used for marking begin of a submit operation
+ void beginTargetSubmitOperation();
+
+ /// Used for marking end of a submit operation
+ void endTargetSubmitOperation();
+
/// Used for marking begin of a target region
- void beginTargetRegion();
+ void beginTargetRegion(int64_t DeviceId);
/// Used for marking end of a target region
void endTargetRegion();
diff --git a/offload/libomptarget/OpenMP/OMPT/Callback.cpp b/offload/libomptarget/OpenMP/OMPT/Callback.cpp
index 1e03f1455d1b2..07ac09e0762aa 100644
--- a/offload/libomptarget/OpenMP/OMPT/Callback.cpp
+++ b/offload/libomptarget/OpenMP/OMPT/Callback.cpp
@@ -56,6 +56,7 @@ 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_set_target_info_t ompt_set_target_info_fn = nullptr;
/// Unique correlation id
static std::atomic<uint64_t> IdCounter(1);
@@ -77,6 +78,20 @@ static uint64_t createRegionId() {
return NewId;
}
+static int ompt_get_target_info_impl(uint64_t *DeviceNum, ompt_id_t *TargetId,
+ ompt_id_t *HostOpId) {
+ if (RegionInterface.getTargetDataValue() == ompt_id_none)
+ return 0;
+
+ if (DeviceNum)
+ *DeviceNum = RegionInterface.getDeviceNum();
+ if (TargetId)
+ *TargetId = RegionInterface.getTargetDataValue();
+ if (HostOpId)
+ *HostOpId = RegionInterface.getHostOpId();
+ return 1;
+}
+
void Interface::beginTargetDataAlloc(int64_t DeviceId, void *HstPtrBegin,
void **TgtPtrBegin, size_t Size,
void *Code) {
@@ -225,6 +240,7 @@ void Interface::endTargetDataRetrieve(int64_t SrcDeviceId, void *SrcPtrBegin,
}
void Interface::beginTargetSubmit(unsigned int NumTeams) {
+ beginTargetSubmitOperation();
if (ompt_callback_target_submit_emi_fn) {
// HostOpId is set by the tool. Invoke the tool supplied target submit EMI
// callback
@@ -245,10 +261,11 @@ void Interface::endTargetSubmit(unsigned int NumTeams) {
ompt_callback_target_submit_emi_fn(ompt_scope_end, &TargetData, &HostOpId,
NumTeams);
}
+ endTargetSubmitOperation();
}
void Interface::beginTargetDataEnter(int64_t DeviceId, void *Code) {
- beginTargetRegion();
+ beginTargetRegion(DeviceId);
if (ompt_callback_target_emi_fn) {
// Invoke the tool supplied target EMI callback
ompt_callback_target_emi_fn(ompt_target_enter_data, ompt_scope_begin,
@@ -276,7 +293,7 @@ void Interface::endTargetDataEnter(int64_t DeviceId, void *Code) {
}
void Interface::beginTargetDataExit(int64_t DeviceId, void *Code) {
- beginTargetRegion();
+ beginTargetRegion(DeviceId);
if (ompt_callback_target_emi_fn) {
// Invoke the tool supplied target EMI callback
ompt_callback_target_emi_fn(ompt_target_exit_data, ompt_scope_begin,
@@ -304,7 +321,7 @@ void Interface::endTargetDataExit(int64_t DeviceId, void *Code) {
}
void Interface::beginTargetUpdate(int64_t DeviceId, void *Code) {
- beginTargetRegion();
+ beginTargetRegion(DeviceId);
if (ompt_callback_target_emi_fn) {
// Invoke the tool supplied target EMI callback
ompt_callback_target_emi_fn(ompt_target_update, ompt_scope_begin, DeviceId,
@@ -415,7 +432,7 @@ void Interface::endTargetMemset(int64_t DeviceId, void *HostPtrBegin,
}
void Interface::beginTarget(int64_t DeviceId, void *Code) {
- beginTargetRegion();
+ beginTargetRegion(DeviceId);
if (ompt_callback_target_emi_fn) {
// Invoke the tool supplied target EMI callback
ompt_callback_target_emi_fn(ompt_target, ompt_scope_begin, DeviceId,
@@ -442,16 +459,28 @@ void Interface::endTarget(int64_t DeviceId, void *Code) {
}
void Interface::beginTargetDataOperation() {
- ODBG(ODT_Tool) << "in ompt_target_region_begin (TargetRegionId = "
+ ODBG(ODT_Tool) << "in ompt_target_data_region_begin (TargetRegionId = "
<< TargetData.value << ")";
}
void Interface::endTargetDataOperation() {
- ODBG(ODT_Tool) << "in ompt_target_region_end (TargetRegionId = "
+ ODBG(ODT_Tool) << "in ompt_target_data_region_end (TargetRegionId = "
+ << TargetData.value << ")";
+ HostOpId = ompt_id_none;
+}
+
+void Interface::beginTargetSubmitOperation() {
+ ODBG(ODT_Tool) << "in ompt_target_submit_region_begin (TargetRegionId = "
+ << TargetData.value << ")";
+}
+
+void Interface::endTargetSubmitOperation() {
+ ODBG(ODT_Tool) << "in ompt_target_submit_region_end (TargetRegionId = "
<< TargetData.value << ")";
+ HostOpId = ompt_id_none;
}
-void Interface::beginTargetRegion() {
+void Interface::beginTargetRegion(int64_t DeviceId) {
// Set up task state
assert(ompt_get_task_data_fn && "Calling a null task data function");
TaskData = ompt_get_task_data_fn();
@@ -459,6 +488,7 @@ void Interface::beginTargetRegion() {
assert(ompt_get_target_task_data_fn &&
"Calling a null target task data function");
TargetTaskData = ompt_get_target_task_data_fn();
+ DeviceNum = DeviceId;
// Target state will be set later
TargetData = ompt_data_none;
}
@@ -466,6 +496,7 @@ void Interface::beginTargetRegion() {
void Interface::endTargetRegion() {
TaskData = 0;
TargetTaskData = 0;
+ DeviceNum = /*omp_invalid_device*/ -2;
TargetData = ompt_data_none;
}
@@ -506,6 +537,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_set_target_info, ompt_set_target_info_fn);
#undef bindOmptFunctionName
// Store pointer of 'ompt_libomp_target_fn_lookup' for use by libomptarget
@@ -518,6 +550,9 @@ int llvm::omp::target::ompt::initializeLibrary(ompt_function_lookup_t lookup,
"ompt_get_target_task_data_fn should be non-null");
assert(LibraryFinalizer == nullptr &&
"LibraryFinalizer should not be initialized yet");
+ // Not required for tool functionality, hence no assertion.
+ if (ompt_set_target_info_fn)
+ ompt_set_target_info_fn(ompt_get_target_info_impl);
LibraryFinalizer = new LibomptargetRtlFinalizer();
@@ -528,6 +563,8 @@ int llvm::omp::target::ompt::initializeLibrary(ompt_function_lookup_t lookup,
void llvm::omp::target::ompt::finalizeLibrary(ompt_data_t *data) {
ODBG(ODT_Tool) << "Executing finalizeLibrary";
+ if (ompt_set_target_info_fn)
+ ompt_set_target_info_fn(nullptr);
// Before disabling OMPT, call the (plugin) finalizations that were registered
// with this library
LibraryFinalizer->finalize();
diff --git a/openmp/runtime/src/ompt-general.cpp b/openmp/runtime/src/ompt-general.cpp
index 959457d380d03..d86c08bf3d87b 100644
--- a/openmp/runtime/src/ompt-general.cpp
+++ b/openmp/runtime/src/ompt-general.cpp
@@ -81,6 +81,8 @@ enum tool_setting_e {
omp_tool_enabled
};
+typedef void (*ompt_set_target_info_t)(ompt_get_target_info_t);
+
/*****************************************************************************
* global variables
****************************************************************************/
@@ -102,6 +104,7 @@ kmp_mutex_impl_info_t kmp_mutex_impl_info[] = {
ompt_callbacks_internal_t ompt_callbacks;
static ompt_start_tool_result_t *ompt_start_tool_result = NULL;
+static ompt_get_target_info_t ompt_get_target_info_impl = NULL;
#if KMP_OS_WINDOWS
static HMODULE ompt_tool_module = NULL;
@@ -861,7 +864,20 @@ OMPT_API_ROUTINE void ompt_finalize_tool(void) { __kmp_internal_end_atexit(); }
OMPT_API_ROUTINE int ompt_get_target_info(uint64_t *device_num,
ompt_id_t *target_id,
ompt_id_t *host_op_id) {
- return 0; // thread is not in a target region
+ if (!ompt_enabled.enabled)
+ return 0;
+
+ if (device_num)
+ *device_num = -2; /* omp_invalid_device */
+ if (target_id)
+ *target_id = ompt_id_none;
+ if (host_op_id)
+ *host_op_id = ompt_id_none;
+
+ if (ompt_get_target_info_impl)
+ return ompt_get_target_info_impl(device_num, target_id, host_op_id);
+
+ return 0;
}
OMPT_API_ROUTINE int ompt_get_num_devices(void) {
@@ -892,6 +908,10 @@ static ompt_data_t *ompt_get_target_task_data() {
return __ompt_get_target_task_data();
}
+static void ompt_set_target_info(ompt_get_target_info_t fn) {
+ ompt_get_target_info_impl = fn;
+}
+
/// Lookup function to query libomp callbacks registered by the tool
static ompt_interface_fn_t ompt_libomp_target_fn_lookup(const char *s) {
#define provide_fn(fn) \
@@ -901,6 +921,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_set_target_info);
#undef provide_fn
#define ompt_interface_fn(fn, type, code) \
>From 407bfb505b75c01d369d453912dc887a738addc1 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Jan=20Andr=C3=A9=20Reuter?= <jan at zyten.de>
Date: Mon, 8 Jun 2026 18:57:56 +0200
Subject: [PATCH 2/3] [Offload][OMPT] Add test for `ompt_get_target_info`
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit
Signed-off-by: Jan André Reuter <jan at zyten.de>
---
offload/test/ompt/get_target_info.c | 184 ++++++++++++++++++++++++++++
1 file changed, 184 insertions(+)
create mode 100644 offload/test/ompt/get_target_info.c
diff --git a/offload/test/ompt/get_target_info.c b/offload/test/ompt/get_target_info.c
new file mode 100644
index 0000000000000..0e85f023a3aa2
--- /dev/null
+++ b/offload/test/ompt/get_target_info.c
@@ -0,0 +1,184 @@
+// clang-format off
+// RUN: %libomptarget-compile-run-and-check-generic
+// REQUIRES: ompt
+// REQUIRES: gpu
+// clang-format on
+
+/*
+ * Example OpenMP program that checks if get_target_info entry point is
+ * correctly implemented.
+ */
+
+#include <assert.h>
+#include <stdio.h>
+
+#include <omp-tools.h>
+#include <omp.h>
+
+// From openmp/runtime/test/ompt/callback.h
+#define register_ompt_callback_t(name, type) \
+ do { \
+ type f_##name = &on_##name; \
+ if (ompt_set_callback(name, (ompt_callback_t)f_##name) == ompt_set_never) \
+ printf("0: Could not register callback '" #name "'\n"); \
+ } while (0)
+
+#define register_ompt_callback(name) register_ompt_callback_t(name, name##_t)
+
+// OMPT entry point handles
+static ompt_set_callback_t ompt_set_callback = 0;
+static ompt_get_target_info_t ompt_get_target_info = 0;
+static ompt_get_unique_id_t ompt_get_unique_id = 0;
+
+// Callback implementation
+static void on_ompt_callback_device_initialize(int device_num, const char *type,
+ ompt_device_t *device,
+ ompt_function_lookup_t lookup,
+ const char *documentation) {
+ printf("Callback Device Init\n");
+}
+
+static void on_ompt_callback_target_emi(ompt_target_t kind,
+ ompt_scope_endpoint_t endpoint,
+ int device_num, ompt_data_t *task_data,
+ ompt_data_t *target_task_data,
+ ompt_data_t *target_data,
+ const void *codeptr_ra) {
+ printf("Callback Target EMI\n");
+
+ if (endpoint == ompt_scope_begin)
+ target_data->value = ompt_get_unique_id();
+ else if (endpoint == ompt_scope_end) {
+ uint64_t retrieved_device_num;
+ ompt_id_t retrieved_target_id;
+ ompt_id_t retrieved_host_op_id;
+ int in_target_region = ompt_get_target_info(
+ &retrieved_device_num, &retrieved_target_id, &retrieved_host_op_id);
+
+ assert(in_target_region && "ompt_get_target_info did not return 1");
+ assert(retrieved_device_num == device_num && "device_num does not match!");
+ assert(retrieved_target_id == target_data->value &&
+ "target_id does not match!");
+ assert(retrieved_host_op_id == ompt_id_none &&
+ "host_op_id should not be set!");
+ }
+}
+
+static void on_ompt_callback_target_data_op_emi(
+ ompt_scope_endpoint_t endpoint, ompt_data_t *target_task_data,
+ ompt_data_t *target_data, ompt_id_t *host_op_id,
+ ompt_target_data_op_t optype, void *src_addr, int src_device_num,
+ void *dest_addr, int dest_device_num, size_t bytes,
+ const void *codeptr_ra) {
+ printf("Callback DataOp EMI\n");
+
+ if (endpoint == ompt_scope_begin)
+ *host_op_id = ompt_get_unique_id();
+ else if (endpoint == ompt_scope_end) {
+ ompt_id_t retrieved_target_id;
+ ompt_id_t retrieved_host_op_id;
+ int in_target_region =
+ ompt_get_target_info(NULL, &retrieved_target_id, &retrieved_host_op_id);
+
+ // Values are only valid if thread is inside of a target region.
+ // This is not always the case for data_op_emi callbacks though, e.g.
+ // during omp_target_memcpy or while transferring the device image on
+ // runtime startup.
+ if (in_target_region) {
+ assert(retrieved_target_id == target_data->value &&
+ "target_id does not match!");
+ assert(retrieved_host_op_id == *host_op_id &&
+ "host_op_id does not match!");
+ }
+ }
+}
+
+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\n");
+
+ if (endpoint == ompt_scope_begin)
+ *host_op_id = ompt_get_unique_id();
+ else if (endpoint == ompt_scope_end) {
+ ompt_id_t retrieved_target_id;
+ ompt_id_t retrieved_host_op_id;
+ int in_target_region =
+ ompt_get_target_info(NULL, &retrieved_target_id, &retrieved_host_op_id);
+
+ assert(in_target_region && "ompt_get_target_info did not return 1");
+ assert(retrieved_target_id == target_data->value &&
+ "target_id does not match!");
+ assert(retrieved_host_op_id == *host_op_id && "host_op_id does not match!");
+ }
+}
+
+// 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
+
+ ompt_get_unique_id = (ompt_get_unique_id_t)lookup("ompt_get_unique_id");
+ if (!ompt_get_unique_id)
+ return 0; // failed
+
+ ompt_get_target_info = (ompt_get_target_info_t)lookup("ompt_get_target_info");
+ if (!ompt_get_target_info)
+ return 0; // failed
+
+ register_ompt_callback(ompt_callback_device_initialize);
+ 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) {}
+
+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;
+}
+
+int main(void) {
+ int numDevices = omp_get_num_devices();
+ int N = 100000;
+
+ int a[N];
+ int b[N];
+
+ int i;
+
+ for (i = 0; i < N; i++)
+ a[i] = 0;
+
+ for (i = 0; i < N; i++)
+ b[i] = i;
+
+ for (int dev = 0; dev < numDevices; ++dev) {
+#pragma omp target parallel for device(dev)
+ {
+ for (int j = 0; j < N; j++)
+ a[j] = b[j];
+ }
+ }
+}
+
+// clang-format off
+
+/// CHECK: Callback Device Init
+/// CHECK: Callback Target EMI
+/// CHECK: Callback DataOp EMI
+/// CHECK: Callback DataOp EMI
+/// CHECK: Callback Submit EMI
+/// CHECK: Callback Submit EMI
+/// CHECK: Callback DataOp EMI
+/// CHECK: Callback DataOp EMI
+/// CHECK: Callback Target EMI
+
+// clang-format on
>From 911fdf655cb584c7ec45f1ed1b123dece8d3eded Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Jan=20Andr=C3=A9=20Reuter?= <j.reuter at fz-juelich.de>
Date: Tue, 9 Jun 2026 10:25:56 +0200
Subject: [PATCH 3/3] [OpenMP][OMPT] Add host only test for get_target_info
entry point
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>
---
.../runtime/test/ompt/misc/get_target_info.c | 76 +++++++++++++++++++
1 file changed, 76 insertions(+)
create mode 100644 openmp/runtime/test/ompt/misc/get_target_info.c
diff --git a/openmp/runtime/test/ompt/misc/get_target_info.c b/openmp/runtime/test/ompt/misc/get_target_info.c
new file mode 100644
index 0000000000000..b950f33a135e1
--- /dev/null
+++ b/openmp/runtime/test/ompt/misc/get_target_info.c
@@ -0,0 +1,76 @@
+// clang-format off
+// RUN: %libomp-compile && %libomp-run | FileCheck %s
+// REQUIRES: ompt
+// clang-format on
+
+/*
+ * Example OpenMP program that checks if get_target_info entry point is
+ * correctly implemented.
+ */
+
+#include <stdio.h>
+
+#include <omp-tools.h>
+#include <omp.h>
+
+// From openmp/runtime/test/ompt/callback.h
+#define register_ompt_callback_t(name, type) \
+ do { \
+ type f_##name = &on_##name; \
+ if (ompt_set_callback(name, (ompt_callback_t)f_##name) == ompt_set_never) \
+ printf("0: Could not register callback '" #name "'\n"); \
+ } while (0)
+
+#define register_ompt_callback(name) register_ompt_callback_t(name, name##_t)
+
+// OMPT entry point handles
+static ompt_get_target_info_t ompt_get_target_info = 0;
+
+// Init functions
+int ompt_initialize(ompt_function_lookup_t lookup, int initial_device_num,
+ ompt_data_t *tool_data) {
+ ompt_get_target_info = (ompt_get_target_info_t)lookup("ompt_get_target_info");
+ if (!ompt_get_target_info)
+ return 0; // failed
+
+ return 1; // success
+}
+
+void ompt_finalize(ompt_data_t *tool_data) {}
+
+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;
+}
+
+int main(void) {
+ // Ensure OMPT is initialized
+#pragma omp parallel
+ {
+ }
+
+ uint64_t retrieved_device_num;
+ ompt_id_t retrieved_target_id;
+ ompt_id_t retrieved_host_op_id;
+ int in_target_region = ompt_get_target_info(
+ &retrieved_device_num, &retrieved_target_id, &retrieved_host_op_id);
+
+ printf("omp_invalid_device: %lu\n", (uint64_t)omp_invalid_device);
+ printf("ompt_id_none: %lu\n", ompt_id_none);
+ printf("in target: %d | device_num: %lu | target_id: %lu | host_op_id: %lu\n",
+ in_target_region, retrieved_device_num, retrieved_target_id,
+ retrieved_host_op_id);
+}
+
+// clang-format off
+
+/// CHECK: omp_invalid_device: [[OMP_INVALID_DEVICE:[0-9]+]]
+/// CHECK: ompt_id_none: [[OMPT_ID_NONE:[0-9]+]]
+/// CHECK: in target: 0
+/// CHECK-SAME: device_num: [[OMP_INVALID_DEVICE]]
+/// CHECK-SAME: target_id: [[OMPT_ID_NONE]]
+/// CHECK-SAME: host_op_id: [[OMPT_ID_NONE]]
+
+// clang-format on
More information about the Openmp-commits
mailing list