diff --git a/openmp/libomptarget/src/interface.cpp b/openmp/libomptarget/src/interface.cpp --- a/openmp/libomptarget/src/interface.cpp +++ b/openmp/libomptarget/src/interface.cpp @@ -28,6 +28,35 @@ #ifdef OMPT_SUPPORT using namespace llvm::omp::target::ompt; + +typedef InterfaceRAII< + std::pair, + std::_Mem_fn>, + long, void *> *TargetDataRAIIPtr; + +/// Helper function for OMPT callbacks related to target data operations. +/// Ordering of callbacks as well as displayed data could be affected by the +/// fact that device CTOR is called lazily. This function helps to create RAII +/// objects at a later stage of the execution, when initialization is done. +/// Note: This will allocate an object, which has to be manually deleted. +inline TargetDataRAIIPtr +createTargetDataRAII(int64_t DeviceId, TargetDataFuncPtrTy TargetDataFunction) { + if (TargetDataFunction == targetDataBegin) + return new InterfaceRAII( + RegionInterface.getCallbacks(), DeviceId, + OMPT_GET_RETURN_ADDRESS(0)); + + if (TargetDataFunction == targetDataEnd) + return new InterfaceRAII( + RegionInterface.getCallbacks(), DeviceId, + OMPT_GET_RETURN_ADDRESS(0)); + + if (TargetDataFunction == targetDataUpdate) + return new InterfaceRAII(RegionInterface.getCallbacks(), + DeviceId, OMPT_GET_RETURN_ADDRESS(0)); + + return (TargetDataRAIIPtr) nullptr; +} #endif //////////////////////////////////////////////////////////////////////////////// @@ -108,6 +137,10 @@ TargetAsyncInfoTy TargetAsyncInfo(Device); AsyncInfoTy &AsyncInfo = TargetAsyncInfo; + /// RAII to establish tool anchors before and after data begin / end / update + OMPT_IF_BUILT(auto *TargetDataRAII = + createTargetDataRAII(DeviceId, TargetDataFunction);) + int Rc = OFFLOAD_SUCCESS; Rc = TargetDataFunction(Loc, Device, ArgNum, ArgsBase, Args, ArgSizes, ArgTypes, ArgNames, ArgMappers, AsyncInfo, @@ -116,6 +149,9 @@ if (Rc == OFFLOAD_SUCCESS) Rc = AsyncInfo.synchronize(); + // Destroy allocated RAII object + OMPT_IF_BUILT(delete TargetDataRAII;) + handleTargetOutcome(Rc == OFFLOAD_SUCCESS, Loc); } @@ -129,12 +165,6 @@ map_var_info_t *ArgNames, void **ArgMappers) { TIMESCOPE_WITH_IDENT(Loc); - /// RAII to establish tool anchors before and after data begin - OMPT_IF_BUILT(InterfaceRAII TargetDataEnterRAII( - RegionInterface.getCallbacks(), - DeviceId, - /* CodePtr */ OMPT_GET_RETURN_ADDRESS(0));) - targetDataMapper(Loc, DeviceId, ArgNum, ArgsBase, Args, ArgSizes, ArgTypes, ArgNames, ArgMappers, targetDataBegin, "Entering OpenMP data region", "begin"); @@ -161,12 +191,6 @@ map_var_info_t *ArgNames, void **ArgMappers) { TIMESCOPE_WITH_IDENT(Loc); - /// RAII to establish tool anchors before and after data end - OMPT_IF_BUILT(InterfaceRAII TargetDataExitRAII( - RegionInterface.getCallbacks(), - DeviceId, - /* CodePtr */ OMPT_GET_RETURN_ADDRESS(0));) - targetDataMapper(Loc, DeviceId, ArgNum, ArgsBase, Args, ArgSizes, ArgTypes, ArgNames, ArgMappers, targetDataEnd, "Exiting OpenMP data region", "end"); @@ -190,12 +214,6 @@ map_var_info_t *ArgNames, void **ArgMappers) { TIMESCOPE_WITH_IDENT(Loc); - /// RAII to establish tool anchors before and after data update - OMPT_IF_BUILT(InterfaceRAII TargetDataUpdateRAII( - RegionInterface.getCallbacks(), - DeviceId, - /* CodePtr */ OMPT_GET_RETURN_ADDRESS(0));) - targetDataMapper( Loc, DeviceId, ArgNum, ArgsBase, Args, ArgSizes, ArgTypes, ArgNames, ArgMappers, targetDataUpdate, "Updating OpenMP data", "update"); @@ -295,7 +313,8 @@ DeviceTy &Device = *PM->Devices[DeviceId]; TargetAsyncInfoTy TargetAsyncInfo(Device); AsyncInfoTy &AsyncInfo = TargetAsyncInfo; - OMPT_IF_BUILT(InterfaceRAII TargetDataAllocRAII( + /// RAII to establish tool anchors before and after target region + OMPT_IF_BUILT(InterfaceRAII TargetRAII( RegionInterface.getCallbacks(), DeviceId, /* CodePtr */ OMPT_GET_RETURN_ADDRESS(0));) @@ -386,7 +405,8 @@ return OMP_TGT_FAIL; } DeviceTy &Device = *PM->Devices[DeviceId]; - OMPT_IF_BUILT(InterfaceRAII TargetDataAllocRAII( + /// RAII to establish tool anchors before and after target region + OMPT_IF_BUILT(InterfaceRAII TargetRAII( RegionInterface.getCallbacks(), DeviceId, /* CodePtr */ OMPT_GET_RETURN_ADDRESS(0));) diff --git a/openmp/libomptarget/test/ompt/veccopy_data.c b/openmp/libomptarget/test/ompt/veccopy_data.c new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/ompt/veccopy_data.c @@ -0,0 +1,105 @@ +// RUN: %libomptarget-compile-run-and-check-generic +// REQUIRES: ompt +// UNSUPPORTED: aarch64-unknown-linux-gnu +// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO +// UNSUPPORTED: x86_64-pc-linux-gnu +// UNSUPPORTED: x86_64-pc-linux-gnu-oldDriver +// UNSUPPORTED: x86_64-pc-linux-gnu-LTO + +/* + * Example OpenMP program that registers non-EMI callbacks. + * Explicitly testing for an initialized device num and + * #pragma omp target [data enter / data exit / update] + * The latter with the addition of a nowait clause. + */ + +#include +#include + +#include "callbacks.h" +#include "register_non_emi.h" + +#define N 100000 + +#pragma omp declare target +int c[N]; +#pragma omp end declare target + +int main() { + int a[N]; + int b[N]; + + int i; + + for (i = 0; i < N; i++) + a[i] = 0; + + for (i = 0; i < N; i++) + b[i] = i; + + for (i = 0; i < N; i++) + c[i] = 0; + +#pragma omp target enter data map(to : a) +#pragma omp target parallel for + { + for (int j = 0; j < N; j++) + a[j] = b[j]; + } +#pragma omp target exit data map(from : a) + +#pragma omp target parallel for map(alloc : c) + { + for (int j = 0; j < N; j++) + c[j] = 2 * j + 1; + } +#pragma omp target update from(c) nowait +#pragma omp barrier + + int rc = 0; + for (i = 0; i < N; i++) { + if (a[i] != i) { + rc++; + printf("Wrong value: a[%d]=%d\n", i, a[i]); + } + } + + for (i = 0; i < N; i++) { + if (c[i] != 2 * i + 1) { + rc++; + printf("Wrong value: c[%d]=%d\n", i, c[i]); + } + } + + if (!rc) + printf("Success\n"); + + return rc; +} + +/// CHECK: Callback Init: +/// CHECK: Callback Load: +/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=2 endpoint=1 +/// CHECK-NOT: device_num=-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 Target: target_id=[[TARGET_ID:[0-9]+]] kind=2 endpoint=2 +/// CHECK-NOT: device_num=-1 +/// 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 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=4 +/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=2 +/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=3 endpoint=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=4 +/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=3 endpoint=2 +/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=1 +/// CHECK: Callback Submit: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] req_num_teams=1 +/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=2 +/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=4 endpoint=1 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3 +/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=4 endpoint=2 +/// CHECK: Callback Fini: