Add implicit map for a list item appears in a reduction clause.

[jyu2]

Summary

A new rule is added in 5.0:
If a list item appears in a reduction, lastprivate or linear clause
on a combined target construct then it is treated as if it also appears
in a map clause with a map-type of tofrom.

Currently map clauses for all capture variables are added implicitly.
But missing for list item of expression for array elements or array
sections.

The change is to add implicit map clause for array of elements used in
reduction clause. Skip adding map clause if the expression is not
mappable.
Noted: For linear and lastprivate, since only variable name is
accepted, the map has been added though capture variables.

To do so:
During the mappable checking, if error, ignore diagnose and skip
adding implicit map clause.

The changes:
1> Add code to generate implicit map in ActOnOpenMPExecutableDirective,

for omp 5.0 and up.

2> Add extra default parameter NoDiagnose in ActOnOpenMPMapClause:
Use that to skip error as well as skip adding implicit map during the
mappable checking.

Note: there are only two places need to be check for NoDiagnose. Rest
of them either the check is for < omp 5.0 or the error already generated for
reduction clause.

[ABataev]

Why it can not be performed in codegen?

[ABataev]

Also, would be good to see a runtime test, which reveals the issue.

[jyu2]

Hi ABataev,
Thanks for reviedw.

In D108132#2946927, @ABataev wrote:

Why it can not be performed in codegen?

I am not sure I can do that. Do you mean when generate map adding coding code to look though reduction clause and generate map for it?

Here is the runtime test, I am trying to find way on how to add runtime test in clang. But in my added test reduction_implicit_map.cpp, I did checked IR for this.

The command line: without may change:
cmplrllvm-25845>clang -fopenmp -fopenmp-targets=x86_64-pc-linux-gnu o.cpp -g
cmplrllvm-25845>./a.out
Segmentation fault (core dumped)
with may change:
cmplrllvm-25845> ./a.out
Result=5050

test

extern "C" int printf(const char *,...);
void sum(int* input, int size, int* output)
{

#pragma omp target teams distribute parallel for reduction(+:output[0]) map(to:input[0:size]) //map(output[0])
for(int i=0; i<size; i++)
    output[0] += input[i];

}
int main()
{

const int size = 100;
int *array = new int[size];
int result = 0;
for (int i = 0; i < size; i++)
    array[i] = i + 1;
sum(array, size, &result);
printf("Result=%d\n", result);
delete[] array;
return 0;

}

[ABataev]

In D108132#2947053, @jyu2 wrote:

Hi ABataev,
Thanks for reviedw.

In D108132#2946927, @ABataev wrote:

Why it can not be performed in codegen?

I am not sure I can do that. Do you mean when generate map adding coding code to look though reduction clause and generate map for it?

Yes, exactly.

Here is the runtime test, I am trying to find way on how to add runtime test in clang. But in my added test reduction_implicit_map.cpp, I did checked IR for this.

Add it to libomptarget.

The command line: without may change:
cmplrllvm-25845>clang -fopenmp -fopenmp-targets=x86_64-pc-linux-gnu o.cpp -g
cmplrllvm-25845>./a.out
Segmentation fault (core dumped)
with may change:
cmplrllvm-25845> ./a.out
Result=5050

test

extern "C" int printf(const char *,...);
void sum(int* input, int size, int* output)
{

#pragma omp target teams distribute parallel for reduction(+:output[0]) map(to:input[0:size]) //map(output[0])
for(int i=0; i<size; i++)
    output[0] += input[i];

}
int main()
{

const int size = 100;
int *array = new int[size];
int result = 0;
for (int i = 0; i < size; i++)
    array[i] = i + 1;
sum(array, size, &result);
printf("Result=%d\n", result);
delete[] array;
return 0;

}

[jyu2]

I am not sure I can do that. Do you mean when generate map adding coding code to look though reduction clause and generate map for it?

Yes, exactly.

We are missing mappable checking for example:

#pragma omp target parallel for reduction(task, +: b[0:2][2:4][1])

In this, we should not add map clause, since the section is not contiguous storage.

[ABataev]

In D108132#2947080, @jyu2 wrote:

I am not sure I can do that. Do you mean when generate map adding coding code to look though reduction clause and generate map for it?

Yes, exactly.

We are missing mappable checking for example:

#pragma omp target parallel for reduction(task, +: b[0:2][2:4][1])

In this, we should not add map clause, since the section is not contiguous storage.

.
OK, and what's the problem to check for this in codegen? Also, we can map non-contiguous storage, at least in some cases.

[jyu2]

Fix format and add new runtime test.

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

5825

isa. Also, what if this is a MemberExpr?

18475–18476

It still might be good to emit the warning if the reduction type is mappable. Also, need to extend this check in general and avoid emitting it if the user defined mapper for the type is provided (in a separate patch, not here).

19430–19432

You can remove /*NoDiagnose=*/ here.

19478–19479

Hmm, should not we still diagnose this case?

openmp/libomptarget/test/mapping/reduction_implicit_map.cpp

5

Can we make UNSUPPORTED instead of XFAIL?

[jyu2]

Address Alex's comments. Thanks Alex.

[jyu2]

clang/lib/Sema/SemaOpenMP.cpp

5825

Yes isa. Changed. Thanks.

For reduction MemeberExpr is not allowed, it is only allowed variable name, array element or array section.

18475–18476

Thanks. Make sense, in this case, implicit map will be add, so warning should be emit. Changed.

19430–19432

Changed. Thanks.

19478–19479

The rule is skip the error as well as skip adding implicit map clause. So that we don't get regression for old code.

openmp/libomptarget/test/mapping/reduction_implicit_map.cpp

5

changed. Thanks.

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

5825

I would also do this check if it is a not template parsing/analysis mode + also would do this for E->IgnoreParensImpCasts()

19478–19479

I think we already have the check for it for the reduction clause, so I think we can skip this check here.

[jyu2]

clang/lib/Sema/SemaOpenMP.cpp

5825

The code already under non DependentContext(). I add IgnoreParenImpCasts(). Thanks.

if (AStmt && !CurContext->isDependentContext() && Kind != OMPD_atomic &&

Kind != OMPD_critical && Kind != OMPD_section && Kind != OMPD_master &&
Kind != OMPD_masked && !isOpenMPLoopTransformationDirective(Kind))

19478–19479

In general yes. But not for some special case, please look clang/test/OpenMP/reduction_implicit_map.cpp, in the first test, if I remove this change the error will be emit. In CUDA device code, the local variable is implicitly threat as thridprivate.

#if defined(CUDA)
// expected-no-diagnostics

int foo(int n) {

double *e;
//no error and no implicit map generated for e[:1]
#pragma omp target parallel reduction(+: e[:1])
  *e=10;
;
return 0;

}

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

19478–19479

I just was not quite clear, I agree with your previous answer ;)

[jyu2]

Address Alex's comment.

[ABataev]

LG

[jyu2]

Thank you so much for Alex's review!!!

clang/include/clang/Sema/Sema.h

ActOnOpenMPDependClause(Expr *DepModifier, OpenMPDependClauseKind DepKind,
                           SourceLocation LParenLoc, SourceLocation EndLoc);
   /// Called on well-formed 'device' clause.
   OMPClause *ActOnOpenMPDeviceClause(OpenMPDeviceClauseModifier Modifier,
                                      Expr *Device, SourceLocation StartLoc,
                                      SourceLocation LParenLoc,
                                      SourceLocation ModifierLoc,
                                      SourceLocation EndLoc);
   /// Called on well-formed 'map' clause.
-  OMPClause *
-  ActOnOpenMPMapClause(ArrayRef<OpenMPMapModifierKind> MapTypeModifiers,
-                       ArrayRef<SourceLocation> MapTypeModifiersLoc,
-                       CXXScopeSpec &MapperIdScopeSpec,
-                       DeclarationNameInfo &MapperId,
-                       OpenMPMapClauseKind MapType, bool IsMapTypeImplicit,
-                       SourceLocation MapLoc, SourceLocation ColonLoc,
-                       ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs,
-                       ArrayRef<Expr *> UnresolvedMappers = llvm::None);
+  OMPClause *ActOnOpenMPMapClause(
+      ArrayRef<OpenMPMapModifierKind> MapTypeModifiers,
+      CXXScopeSpec &MapperIdScopeSpec, DeclarationNameInfo &MapperId,
+      SourceLocation MapLoc, SourceLocation ColonLoc, ArrayRef<Expr *> VarList,
+      const OMPVarListLocTy &Locs, bool NoDiagnose = false,
   /// Called on well-formed 'num_teams' clause.
   OMPClause *ActOnOpenMPNumTeamsClause(Expr *NumTeams, SourceLocation StartLoc,
                                        SourceLocation LParenLoc,
                                        SourceLocation EndLoc);
   /// Called on well-formed 'thread_limit' clause.
   OMPClause *ActOnOpenMPThreadLimitClause(Expr *ThreadLimit,
                                           SourceLocation StartLoc,
                                           SourceLocation LParenLoc,

clang/lib/Sema/SemaOpenMP.cpp

if (!ImplicitFirstprivates.empty()) {
               SourceLocation())) {
         ClausesWithImplicit.push_back(Implicit);
         ErrorFound = cast<OMPFirstprivateClause>(Implicit)->varlist_size() !=
                      ImplicitFirstprivates.size();
       } else {
         ErrorFound = true;
       }
     }
+    // OpenMP 5.0 [2.19.7]
+    // If a list item appears in a reduction, lastprivate or linear
+    // clause on a combined target construct then it is treated as
+    // if it also appears in a map clause with a map-type of tofrom
+    if (getLangOpts().OpenMP >= 50 && Kind != OMPD_target &&
+        isOpenMPTargetExecutionDirective(Kind)) {
+      SmallVector<Expr *, 4> ImplicitExprs;
+      for (OMPClause *C : Clauses) {
+        if (auto *RC = dyn_cast<OMPReductionClause>(C))
+          for (Expr *E : RC->varlists())
+            if (!isa<DeclRefExpr>(E->IgnoreParenImpCasts()))
+              ImplicitExprs.emplace_back(E);
+      }
+      if (!ImplicitExprs.empty()) {
+        ArrayRef<Expr *> Exprs = ImplicitExprs;
+        CXXScopeSpec MapperIdScopeSpec;
+        DeclarationNameInfo MapperId;
+        if (OMPClause *Implicit = ActOnOpenMPMapClause(
+                OMPC_MAP_MODIFIER_unknown, SourceLocation(), MapperIdScopeSpec,
+                MapperId, OMPC_MAP_tofrom,
+                /*IsMapTypeImplicit=*/true, SourceLocation(), SourceLocation(),
+                Exprs, OMPVarListLocTy(), /*NoDiagnose=*/true))
+          ClausesWithImplicit.emplace_back(Implicit);
+      }
+    }
     for (unsigned I = 0, E = DefaultmapKindNum; I < E; ++I) {
       int ClauseKindCnt = -1;
       for (ArrayRef<Expr *> ImplicitMap : ImplicitMaps[I]) {
         ++ClauseKindCnt;
         if (ImplicitMap.empty())
           continue;
         CXXScopeSpec MapperIdScopeSpec;
         DeclarationNameInfo MapperId;

 
 static bool checkTypeMappable(SourceLocation SL, SourceRange SR, Sema &SemaRef,
                               DSAStackTy *Stack, QualType QTy,
                               bool FullCheck = true) {
   if (SemaRef.RequireCompleteType(SL, QTy, diag::err_incomplete_type))
     return false;
   if (FullCheck && !SemaRef.CurContext->isDependentContext() &&
       !QTy.isTriviallyCopyableType(SemaRef.Context))
     SemaRef.Diag(SL, diag::warn_omp_non_trivial_type_mapped) << QTy << SR;
   return true;
 }
 
 /// Return true if it can be proven that the provided array expression
 /// (array section or array subscript) does NOT specify the whole size of the
 /// array whose base type is \a BaseQTy.
 static bool checkArrayExpressionDoesNotReferToWholeSize(Sema &SemaRef,
                                                         const Expr *E,
                                                         QualType BaseQTy) {
bool VisitArraySubscriptExpr(ArraySubscriptExpr *AE) {
 
     // Record the component - we don't have any declaration associated.
     Components.emplace_back(AE, nullptr, IsNonContiguous);
 
     return RelevantExpr || Visit(E);
   }
 
   bool VisitOMPArraySectionExpr(OMPArraySectionExpr *OASE) {
-    assert(!NoDiagnose && "Array sections cannot be implicitly mapped.");
+    // After OMP 5.0  Array section in reduction clause will be implicitly
+    // mapped
+    assert(!(SemaRef.getLangOpts().OpenMP < 50 && NoDiagnose) &&
+           "Array sections cannot be implicitly mapped.");
     Expr *E = OASE->getBase()->IgnoreParenImpCasts();
     QualType CurType =
       OMPArraySectionExpr::getBaseOriginalType(E).getCanonicalType();
 
     // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C++, p.1]
     //  If the type of a list item is a reference to a type T then the type
     //  will be considered to be T for all purposes of this clause.
     if (CurType->isReferenceType())
if (AllowWholeSizeArraySection) {
       if (IsPointer && !AllowAnotherPtr)
         SemaRef.Diag(ELoc, diag::err_omp_section_length_undefined)
             << /*array of unknown bound */ 1;
       else
         IsNonContiguous = true;
     } else if (AllowUnitySizeArraySection && NotUnity) {
       // A unity or whole array section is not allowed and that is not
       // compatible with the properties of the current array section.
+      if (NoDiagnose)
+        return false;
       SemaRef.Diag(
         ELoc, diag::err_array_section_does_not_specify_contiguous_storage)
         << OASE->getSourceRange();
       return false;
     }
 
     if (IsPointer)
       AllowAnotherPtr = false;

 // UDMapperList in MVLI. \a MapType, \a IsMapTypeImplicit, \a MapperIdScopeSpec,
 // and \a MapperId are expected to be valid if the clause kind is 'map'.
 static void checkMappableExpressionList(
     Sema &SemaRef, DSAStackTy *DSAS, OpenMPClauseKind CKind,
     MappableVarListInfo &MVLI, SourceLocation StartLoc,
     CXXScopeSpec &MapperIdScopeSpec, DeclarationNameInfo MapperId,
     ArrayRef<Expr *> UnresolvedMappers,
     OpenMPMapClauseKind MapType = OMPC_MAP_unknown,
-    bool IsMapTypeImplicit = false) {
+    bool IsMapTypeImplicit = false, bool NoDiagnose = false) {
   // We only expect mappable expressions in 'to', 'from', and 'map' clauses.
   assert((CKind == OMPC_map || CKind == OMPC_to || CKind == OMPC_from) &&
          "Unexpected clause kind with mappable expressions!");
 
   // If the identifier of user-defined mapper is not specified, it is "default".
   // We do not change the actual name in this clause to distinguish whether a
   // mapper is specified explicitly, i.e., it is not explicitly specified when
   // MapperId.getName() is empty.
if (!RE->isLValue()) {
       continue;
     }
 
     OMPClauseMappableExprCommon::MappableExprComponentList CurComponents;
     ValueDecl *CurDeclaration = nullptr;
 
     // Obtain the array or member expression bases if required. Also, fill the
     // components array with all the components identified in the process.
-    const Expr *BE = checkMapClauseExpressionBase(
-        SemaRef, SimpleExpr, CurComponents, CKind, DSAS->getCurrentDirective(),
-        /*NoDiagnose=*/false);
+    const Expr *BE =
+        checkMapClauseExpressionBase(SemaRef, SimpleExpr, CurComponents, CKind,
+                                     DSAS->getCurrentDirective(), NoDiagnose);
     if (!BE)
       continue;
 
     assert(!CurComponents.empty() &&
            "Invalid mappable expression information.");
 
     if (const auto *TE = dyn_cast<CXXThisExpr>(BE)) {
       // Add store "this" pointer to class in DSAStackTy for future checking
for (Expr *RE : MVLI.VarList) {
     assert((VD || FD) && "Only variables or fields are expected here!");
     (void)FD;
 
     // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.10]
     // threadprivate variables cannot appear in a map clause.
     // OpenMP 4.5 [2.10.5, target update Construct]
     // threadprivate variables cannot appear in a from clause.
     if (VD && DSAS->isThreadPrivate(VD)) {
+      if (NoDiagnose)
+        continue;
       DSAStackTy::DSAVarData DVar = DSAS->getTopDSA(VD, /*FromParent=*/false);
       SemaRef.Diag(ELoc, diag::err_omp_threadprivate_in_clause)
           << getOpenMPClauseName(CKind);
       reportOriginalDsa(SemaRef, DSAS, VD, DVar);
       continue;
     }
 
     // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.9]
if (ASE) {
       Type = VE->getType();
     }
 
     // OpenMP 4.5 [2.10.5, target update Construct, Restrictions, p.4]
     // A list item in a to or from clause must have a mappable type.
     // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.9]
     //  A list item must have a mappable type.
     if (!checkTypeMappable(VE->getExprLoc(), VE->getSourceRange(), SemaRef,
-                           DSAS, Type))
+                           DSAS, Type, /*FullCheck=*/true))
       continue;
 
     if (CKind == OMPC_map) {
       // target enter data
       // OpenMP [2.10.2, Restrictions, p. 99]
       // A map-type must be specified in all map clauses and must be either
       // to or alloc.
       OpenMPDirectiveKind DKind = DSAS->getCurrentDirective();

 }
 
 OMPClause *Sema::ActOnOpenMPMapClause(
     ArrayRef<OpenMPMapModifierKind> MapTypeModifiers,
     ArrayRef<SourceLocation> MapTypeModifiersLoc,
     CXXScopeSpec &MapperIdScopeSpec, DeclarationNameInfo &MapperId,
     OpenMPMapClauseKind MapType, bool IsMapTypeImplicit, SourceLocation MapLoc,
     SourceLocation ColonLoc, ArrayRef<Expr *> VarList,
-    const OMPVarListLocTy &Locs, ArrayRef<Expr *> UnresolvedMappers) {
+    const OMPVarListLocTy &Locs, bool NoDiagnose,
+    ArrayRef<Expr *> UnresolvedMappers) {
   OpenMPMapModifierKind Modifiers[] = {
       OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown,
       OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown};
   SourceLocation ModifiersLoc[NumberOfOMPMapClauseModifiers];
 
   // Process map-type-modifiers, flag errors for duplicate modifiers.
   unsigned Count = 0;
   for (unsigned I = 0, E = MapTypeModifiers.size(); I < E; ++I) {
     if (MapTypeModifiers[I] != OMPC_MAP_MODIFIER_unknown &&
         llvm::find(Modifiers, MapTypeModifiers[I]) != std::end(Modifiers)) {
       Diag(MapTypeModifiersLoc[I], diag::err_omp_duplicate_map_type_modifier);
       continue;
     }
     assert(Count < NumberOfOMPMapClauseModifiers &&
            "Modifiers exceed the allowed number of map type modifiers");
     Modifiers[Count] = MapTypeModifiers[I];
     ModifiersLoc[Count] = MapTypeModifiersLoc[I];
     ++Count;
   }
 
   MappableVarListInfo MVLI(VarList);
   checkMappableExpressionList(*this, DSAStack, OMPC_map, MVLI, Locs.StartLoc,
                               MapperIdScopeSpec, MapperId, UnresolvedMappers,
-                              MapType, IsMapTypeImplicit);
+                              MapType, IsMapTypeImplicit, NoDiagnose);
 
   // We need to produce a map clause even if we don't have variables so that
   // other diagnostics related with non-existing map clauses are accurate.
   return OMPMapClause::Create(Context, Locs, MVLI.ProcessedVarList,
                               MVLI.VarBaseDeclarations, MVLI.VarComponents,
                               MVLI.UDMapperList, Modifiers, ModifiersLoc,
                               MapperIdScopeSpec.getWithLocInContext(Context),
                               MapperId, MapType, IsMapTypeImplicit, MapLoc);

clang/lib/Sema/TreeTransform.h

   /// Subclasses may override this routine to provide different behavior.
   OMPClause *RebuildOMPMapClause(
       ArrayRef<OpenMPMapModifierKind> MapTypeModifiers,
       ArrayRef<SourceLocation> MapTypeModifiersLoc,
       CXXScopeSpec MapperIdScopeSpec, DeclarationNameInfo MapperId,
       OpenMPMapClauseKind MapType, bool IsMapTypeImplicit,
       SourceLocation MapLoc, SourceLocation ColonLoc, ArrayRef<Expr *> VarList,
       const OMPVarListLocTy &Locs, ArrayRef<Expr *> UnresolvedMappers) {
-    return getSema().ActOnOpenMPMapClause(MapTypeModifiers, MapTypeModifiersLoc,
-                                          MapperIdScopeSpec, MapperId, MapType,
-                                          IsMapTypeImplicit, MapLoc, ColonLoc,
-                                          VarList, Locs, UnresolvedMappers);
+    return getSema().ActOnOpenMPMapClause(
+        MapTypeModifiers, MapTypeModifiersLoc, MapperIdScopeSpec, MapperId,
+        MapType, IsMapTypeImplicit, MapLoc, ColonLoc, VarList, Locs,
+        /*NoDiagnose=*/false, UnresolvedMappers);
   }
 
   /// Build a new OpenMP 'allocate' clause.
   ///
   /// By default, performs semantic analysis to build the new OpenMP clause.
   /// Subclasses may override this routine to provide different behavior.
   OMPClause *RebuildOMPAllocateClause(Expr *Allocate, ArrayRef<Expr *> VarList,
                                       SourceLocation StartLoc,

clang/test/OpenMP/reduction_implicit_map.cpp

+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ \
+// RUN:  -triple powerpc64le-unknown-unknown -DCUDA \
+// RUN:  -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o \
+// RUN:  %t-ppc-host.bc
+
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ \
+// RUN:  -triple nvptx64-unknown-unknown -DCUA \
+// RUN:  -fopenmp-targets=nvptx64-nvidia-cuda -DCUDA -emit-llvm %s \
+// RUN:  -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc \
+// RUN:  -o - | FileCheck %s --check-prefix CHECK
+
+// RUN: %clang_cc1 -verify -fopenmp -x c++ \
+// RUN:   -triple powerpc64le-unknown-unknown -DDIAG\
+// RUN:   -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm \
+// RUN:   %s -o - | FileCheck  %s \
+// RUN:   --check-prefix=CHECK1
+
+// RUN: %clang_cc1 -verify -fopenmp -x c++ \
+// RUN:   -triple i386-unknown-unknown \
+// RUN:   -fopenmp-targets=i386-pc-linux-gnu -emit-llvm \
+// RUN:   %s -o - | FileCheck  %s \
+// RUN:   --check-prefix=CHECK2
+
+
+#if defined(CUDA)
+// expected-no-diagnostics
+
+int foo(int n) {
+  double *e;
+  //no error and no implicit map generated for e[:1]
+  #pragma omp target parallel reduction(+: e[:1])
+    *e=10;
+  ;
+  return 0;
+}
+// CHECK-NOT @.offload_maptypes
+// CHECK: call void @__kmpc_nvptx_end_reduce_nowait(
+#elif defined(DIAG)
+class S2 {
+  mutable int a;
+public:
+  S2():a(0) { }
+  S2(S2 &s2):a(s2.a) { }
+  S2 &operator +(S2 &s);
+};
+int bar() {
+ S2 o[5];
+  //warnig "copyable and not guaranteed to be mapped correctly" and
+  //implicit map generated.
+#pragma omp target parallel reduction(+:o[0]) //expected-warning {{Type 'S2' is not trivially copyable and not guaranteed to be mapped correctly}}
+  for (int i = 0; i < 10; i++);
+  double b[10][10][10];
+  //no error no implicit map generated, the map for b is generated but not
+  //for b[0:2][2:4][1].
+#pragma omp target parallel for reduction(task, +: b[0:2][2:4][1])
+  for (long long i = 0; i < 10; ++i);
+  return 0;
+}
+// map for variable o
+// CHECK1: offload_sizes = private unnamed_addr constant [1 x i64] [i64 4]
+// CHECK1: offload_maptypes = private unnamed_addr constant [1 x i64] [i64 547]
+// map for b:
+// CHECK1: @.offload_sizes{{.*}} = private unnamed_addr constant [1 x i64] [i64 8000]
+// CHECK1: @.offload_maptypes{{.*}} = private unnamed_addr constant [1 x i64] [i64 547]
+#else
+// expected-no-diagnostics
+
+// generate implicit map for array elements or array sections in reduction
+// clause. In following case: the implicit map is generate for output[0]
+// with map size 4 and output[:3] with map size 12.
+void sum(int* input, int size, int* output)
+{
+#pragma omp target teams distribute parallel for reduction(+: output[0]) \
+                                                 map(to: input [0:size])
+  for (int i = 0; i < size; i++)
+    output[0] += input[i];
+#pragma omp target teams distribute parallel for reduction(+: output[:3])  \
+                                                 map(to: input [0:size])
+  for (int i = 0; i < size; i++)
+    output[0] += input[i];
+  int a[10];
+#pragma omp target parallel reduction(+: a[:2])
+  for (int i = 0; i < size; i++)
+    ;
+#pragma omp target parallel reduction(+: a[3])
+  for (int i = 0; i < size; i++)
+    ;
+}
+//CHECK2: @.offload_sizes = private unnamed_addr constant [2 x i64] [i64 4, i64 8]
+//CHECK2: @.offload_maptypes.10 = private unnamed_addr constant [2 x i64] [i64 800, i64 547]
+//CHECK2: @.offload_sizes.13 = private unnamed_addr constant [2 x i64] [i64 4, i64 4]
+//CHECK2: @.offload_maptypes.14 = private unnamed_addr constant [2 x i64] [i64 800, i64 547]
+//CHECK2: define dso_local void @_Z3sumPiiS_
+//CHECK2-NEXT: entry
+//CHECK2-NEXT: [[INP:%.*]] = alloca i32*
+//CHECK2-NEXT: [[SIZE:%.*]] = alloca i32
+//CHECK2-NEXT: [[OUTP:%.*]] = alloca i32*
+//CHECK2:      [[OFFSIZE:%.*]] = alloca [3 x i64]
+//CHECK2:      [[OFFSIZE10:%.*]] = alloca [3 x i64]
+//CHECK2:      [[T15:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[OFFSIZE]], i32 0, i32 0
+//CHECK2-NEXT: store i64 4, i64* [[T15]]
+//CHECK2:      [[T21:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[OFFSIZE]], i32 0, i32 1
+//CHECK2-NEXT: store i64 4, i64* [[T21]]
+//CHECK2:     [[T53:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[OFFSIZE10]], i32 0, i32 0
+//CHECK2-NEXT: store i64 4, i64* [[T53]]
+//CHECK2:     [[T59:%.*]] = getelementptr inbounds [3 x i64], [3 x i64]* [[OFFSIZE10]], i32 0, i32 1
+//CHECK2-NEXT: store i64 12, i64* [[T59]]
+#endif
+int main()
+{
+#if defined(CUDA)
+  int a = foo(10);
+#elif defined(DIAG)
+  int a = bar();
+#else
+  const int size = 100;
+  int *array = new int[size];
+  int result = 0;
+  sum(array, size, &result);
+#endif
+  return 0;
+}

openmp/libomptarget/test/mapping/reduction_implicit_map.cpp

+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// amdgcn does not have printf definition
+// UNSUPPORTED: amdgcn-amd-amdhsa
+
+#include <stdio.h>
+
+void sum(int* input, int size, int* output)
+{
+#pragma omp target teams distribute parallel for reduction(+:output[0]) \
+                                                 map(to:input[0:size])
+  for (int i = 0; i < size; i++)
+    output[0] += input[i];
+}
+int main()
+{
+  const int size = 100;
+  int *array = new int[size];
+  int result = 0;
+  for (int i = 0; i < size; i++)
+    array[i] = i + 1;
+  sum(array, size, &result);
+  // CHECK: Result=5050
+  printf("Result=%d\n", result);
+  delete[] array;
+  return 0;
+}
+