From 8bcdbcf2a72d4acdc25433efb2596af4c687655b Mon Sep 17 00:00:00 2001 From: Kaloyan Ignatov 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 2b58c602f05687ebf71ce42eeeb4b6b39604e5dd Mon Sep 17 00:00:00 2001 From: Kaloyan Ignatov 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 | 171 ++++++++++++++++++ .../ompt/target_tool_data_nowait_nodepend.c | 59 ++++++ 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, 805 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 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 +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 @@ -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 + +// 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..7407e106700ba --- /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 +#include +#include +#include +#include + +#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..c5e20bec1bd55 --- /dev/null +++ b/offload/test/ompt/target_tool_data_nowait.c @@ -0,0 +1,171 @@ +// clang-format off +// RUN: env LIBOMP_NUM_HIDDEN_HELPER_THREADS=1 %libomptarget-compile-run-and-check-generic +// REQUIRES: ompt +// clang-format on + +#include +#include +#include +#include +#include + +#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..b28184bc304db --- /dev/null +++ b/offload/test/ompt/target_tool_data_nowait_nodepend.c @@ -0,0 +1,59 @@ +// clang-format off +// RUN: env LIBOMP_NUM_HIDDEN_HELPER_THREADS=1 %libomptarget-compile-run-and-check-generic +// REQUIRES: ompt +// clang-format on + +#include +#include +#include +#include +#include + +#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