diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp --- a/openmp/libomptarget/src/omptarget.cpp +++ b/openmp/libomptarget/src/omptarget.cpp @@ -1062,7 +1062,10 @@ TgtBaseOffset = 0; } else if (ArgTypes[I] & OMP_TGT_MAPTYPE_PRIVATE) { TgtBaseOffset = (intptr_t)HstPtrBase - (intptr_t)HstPtrBegin; - const bool IsFirstPrivate = ArgTypes[I] & OMP_TGT_MAPTYPE_TO; + // Can be marked for optimization if the next argument(s) do(es) not + // depend from this one. + const bool IsFirstPrivate = + (I >= ArgNum - 1 || !(ArgTypes[I + 1] & OMP_TGT_MAPTYPE_MEMBER_OF)); Ret = PrivateArgumentManager.addArg(HstPtrBegin, ArgSizes[I], TgtBaseOffset, IsFirstPrivate, TgtPtrBegin, TgtArgs.size()); diff --git a/openmp/libomptarget/test/mapping/lambda_mapping.cpp b/openmp/libomptarget/test/mapping/lambda_mapping.cpp new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/mapping/lambda_mapping.cpp @@ -0,0 +1,53 @@ +// RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu +// RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu +// RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu +// RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu +// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda + +#include + +template +inline void forall(int Begin, int End, LOOP_BODY LoopBody) { +#pragma omp target parallel for schedule(static) + for (int I = Begin; I < End; ++I) { + LoopBody(I); + } +} + +#define N (1000) + +// +// Demonstration of the RAJA abstraction using lambdas +// Requires data mapping onto the target section +// +int main() { + double A[N], B[N], C[N]; + + for (int I = 0; I < N; I++) { + A[I] = I + 1; + B[I] = -I; + C[I] = -9; + } + +#pragma omp target data map(tofrom : C [0:N]) map(to : A [0:N], B [0:N]) + { + forall(0, N, [&](int I) { C[I] += A[I] + B[I]; }); + } + + int Fail = 0; + for (int I = 0; I < N; I++) { + if (C[I] != -8) { + std::cout << "Failed at " << I << " with val " << C[I] << std::endl; + Fail = 1; + } + } + + // CHECK: Succeeded + if (Fail) { + std::cout << "Failed" << std::endl; + } else { + std::cout << "Succeeded" << std::endl; + } + + return 0; +}