shylie/llvm-project
1
0
Fork 0
Code Issues Pull Requests Actions 6 Packages Projects Releases Wiki Activity
llvm-project/offload/test/ompt/target_tool_data.c
Kaloyan Ignatov 726452720c
[OMPT][Offload][OpenMP] Fixes for OMPT data used by libomptarget (#156020) 
These commits fix issues regarding storage of tool data within
libomptarget. Both libomp and libomptarget have been modified to
accommodate this. We differentiate between two cases depending on the
type of the target region:

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

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

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

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

---------

Co-authored-by: Joachim <jenke@itc.rwth-aachen.de>
2025-12-18 10:20:14 +01:00

153 lines
5.5 KiB
C
Raw Blame History

// clang-format off
// RUN: env LIBOMP_NUM_HIDDEN_HELPER_THREADS=1 %libomptarget-compile-run-and-check-generic
// REQUIRES: ompt
// clang-format on
#include <inttypes.h>
#include <omp-tools.h>
#include <omp.h>
#include <stdio.h>
#include <string.h>
#include "register_with_host.h"
#define N 1000000
#define M 1000
int main() {
float *x = malloc(N * sizeof(float));
float *y = malloc(N * sizeof(float));
for (int i = 0; i < N; i++) {
x[i] = 1;
y[i] = 1;
}
#pragma omp target enter data map(to : x[0 : N]) map(alloc : y[0 : N])
#pragma omp target teams distribute parallel for
for (int i = 0; i < N; i++) {
for (int j = 0; j < M; j++) {
y[i] += 3 * x[i];
}
}
#pragma omp target teams distribute parallel for
for (int i = 0; i < N; i++) {
for (int j = 0; j < M; j++) {
y[i] += 3 * x[i];
}
}
#pragma omp target exit data map(release : x[0 : N]) map(from : y[0 : N])
printf("%f, %f\n", x[0], y[0]);
free(x);
free(y);
return 0;
}
// clang-format off
/// CHECK: ompt_event_initial_task_begin
/// CHECK-SAME: task_id=[[ENCOUNTERING_TASK:[0-f]+]]
/// CHECK: Callback Target EMI: kind=ompt_target_enter_data endpoint=ompt_scope_begin
/// CHECK-SAME: task_data=0x{{[0-f]+}} (0x[[ENCOUNTERING_TASK]])
/// CHECK-SAME: target_task_data=(nil) (0x0)
/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA:[0-f]+]])
/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin
/// CHECK-SAME: target_task_data=(nil) (0x0)
/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end
/// CHECK-SAME: target_task_data=(nil) (0x0)
/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin
/// CHECK-SAME: target_task_data=(nil) (0x0)
/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end
/// CHECK-SAME: target_task_data=(nil) (0x0)
/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin
/// CHECK-SAME: target_task_data=(nil) (0x0)
/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end
/// CHECK-SAME: target_task_data=(nil) (0x0)
/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
/// CHECK: Callback Target EMI: kind=ompt_target_enter_data endpoint=ompt_scope_end
/// CHECK-SAME: task_data=0x{{[0-f]+}} (0x[[ENCOUNTERING_TASK]])
/// CHECK-SAME: target_task_data=(nil) (0x0)
/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_begin
/// CHECK-SAME: task_data=0x{{[0-f]+}} (0x[[ENCOUNTERING_TASK]])
/// CHECK-SAME: target_task_data=(nil) (0x0)
/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA:[0-f]+]])
/// CHECK: Callback Submit EMI: endpoint=ompt_scope_begin
/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
/// CHECK: Callback Submit EMI: endpoint=ompt_scope_end
/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_end
/// CHECK-SAME: task_data=0x{{[0-f]+}} (0x[[ENCOUNTERING_TASK]])
/// CHECK-SAME: target_task_data=(nil) (0x0)
/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_begin
/// CHECK-SAME: task_data=0x{{[0-f]+}} (0x[[ENCOUNTERING_TASK]])
/// CHECK-SAME: target_task_data=(nil) (0x0)
/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA:[0-f]+]])
/// CHECK: Callback Submit EMI: endpoint=ompt_scope_begin
/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
/// CHECK: Callback Submit EMI: endpoint=ompt_scope_end
/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
/// CHECK: Callback Target EMI: kind=ompt_target endpoint=ompt_scope_end
/// CHECK-SAME: task_data=0x{{[0-f]+}} (0x[[ENCOUNTERING_TASK]])
/// CHECK-SAME: target_task_data=(nil) (0x0)
/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
/// CHECK: Callback Target EMI: kind=ompt_target_exit_data endpoint=ompt_scope_begin
/// CHECK-SAME: task_data=0x{{[0-f]+}} (0x[[ENCOUNTERING_TASK]])
/// CHECK-SAME: target_task_data=(nil) (0x0)
/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA:[0-f]+]])
/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin
/// CHECK-SAME: target_task_data=(nil) (0x0)
/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end
/// CHECK-SAME: target_task_data=(nil) (0x0)
/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin
/// CHECK-SAME: target_task_data=(nil) (0x0)
/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end
/// CHECK-SAME: target_task_data=(nil) (0x0)
/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_begin
/// CHECK-SAME: target_task_data=(nil) (0x0)
/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
/// CHECK: Callback DataOp EMI: endpoint=ompt_scope_end
/// CHECK-SAME: target_task_data=(nil) (0x0)
/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
/// CHECK: Callback Target EMI: kind=ompt_target_exit_data endpoint=ompt_scope_end
/// CHECK-SAME: task_data=0x{{[0-f]+}} (0x[[ENCOUNTERING_TASK]])
/// CHECK-SAME: target_task_data=(nil) (0x0)
/// CHECK-SAME: target_data=0x{{[0-f]+}} (0x[[TARGET_DATA]])
Reference in New Issue View Git Blame Copy Permalink