[OpenMP] Add support for mapping array sections through pointer references.

[sfantao]

Summary

This patch fixes a bug in the map of array sections whose base is a reference to a pointer. The existing mapping support was not prepared to deal with it, causing the compiler to crash.

Mapping a reference to a pointer enjoys the same characteristics of a regular pointer, i.e., it is passed by value. Therefore, the reference has to be materialized in the target region.

[ABataev]

LG with a nit

lib/CodeGen/CGOpenMPRuntime.cpp

5280

getAs -> castAs

[sfantao]

• Use castAs instead of getAs in pointer type.

lib/CodeGen/CGOpenMPRuntime.cpp

void generateInfoForComponentList(
     } else {
       // The base is the reference to the variable.
       // BP = &Var.
       BP = CGF.EmitLValue(cast<DeclRefExpr>(I->getAssociatedExpression()))
                .getPointer();
 
       // If the variable is a pointer and is being dereferenced (i.e. is not
       // the last component), the base has to be the pointer itself, not its
-      // reference.
-      if (I->getAssociatedDeclaration()->getType()->isAnyPointerType() &&
-          std::next(I) != CE) {
-        auto PtrAddr = CGF.MakeNaturalAlignAddrLValue(
-            BP, I->getAssociatedDeclaration()->getType());
+      // reference. References are ignored for mapping purposes.
+      QualType Ty =
+          I->getAssociatedDeclaration()->getType().getNonReferenceType();
+      if (Ty->isAnyPointerType() && std::next(I) != CE) {
+        auto PtrAddr = CGF.MakeNaturalAlignAddrLValue(BP, Ty);
         BP = CGF.EmitLoadOfPointerLValue(PtrAddr.getAddress(),
-                                         I->getAssociatedDeclaration()
-                                             ->getType()
-                                             ->getAs<PointerType>())
+                                         Ty->castAs<PointerType>())
                  .getPointer();
 
         // We do not need to generate individual map information for the
         // pointer, it can be associated with the combined storage.
         ++I;
       }
     }
 

lib/CodeGen/CGStmtOpenMP.cpp

CodeGenFunction::GenerateOpenMPCapturedStmtFunction(const CapturedStmt &S) {
   StartFunction(CD, Ctx.VoidTy, F, FuncInfo, Args, CD->getLocation(),
                 CD->getBody()->getLocStart());
   unsigned Cnt = CD->getContextParamPosition();
   I = S.captures().begin();
   for (auto *FD : RD->fields()) {
     // If we are capturing a pointer by copy we don't need to do anything, just
     // use the value that we get from the arguments.
     if (I->capturesVariableByCopy() && FD->getType()->isAnyPointerType()) {
-      setAddrOfLocalVar(I->getCapturedVar(), GetAddrOfLocalVar(Args[Cnt]));
+      const VarDecl *CurVD = I->getCapturedVar();
+      Address LocalAddr = GetAddrOfLocalVar(Args[Cnt]);
+      // If the variable is a reference we need to materialize it here.
+      if (CurVD->getType()->isReferenceType()) {
+        Address RefAddr = CreateMemTemp(CurVD->getType(), getPointerAlign(),
+                                        ".materialized_ref");
+        EmitStoreOfScalar(LocalAddr.getPointer(), RefAddr, /*Volatile=*/false,
+                          CurVD->getType());
+        LocalAddr = RefAddr;
+      }
+      setAddrOfLocalVar(CurVD, LocalAddr);
       ++Cnt;
       ++I;
       continue;
     }
 
     LValue ArgLVal =
         MakeAddrLValue(GetAddrOfLocalVar(Args[Cnt]), Args[Cnt]->getType(),
                        AlignmentSource::Decl);

test/OpenMP/target_map_codegen.cpp

 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-64
 // RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK2 --check-prefix CK2-64
 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK2 --check-prefix CK2-32
 // RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK2 --check-prefix CK2-32
 #ifdef CK2
 
-// CK2-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
+// CK2: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
 // Map types: OMP_MAP_PRIVATE_VAL | OMP_MAP_IS_FIRST = 288
-// CK2-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
+// CK2: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
+// CK2: [[SIZES2:@.+]] = {{.+}}constant [1 x i[[sz]]] zeroinitializer
+// Map types: OMP_MAP_IS_PTR = 32
+// CK2: [[TYPES2:@.+]] = {{.+}}constant [1 x i32] [i32 32]
 
-// CK2-LABEL: implicit_maps_integer_reference
-void implicit_maps_integer_reference (int a){
+// CK2-LABEL: implicit_maps_reference
+void implicit_maps_reference (int a, int *b){
   int &i = a;
   // CK2-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
   // CK2-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
   // CK2-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
   // CK2-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
   // CK2-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
   // CK2-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
   // CK2-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
   // CK2-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
   // CK2-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8*
   // CK2-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]],
   // CK2-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32*
   // CK2-64-DAG: store i32 {{.+}}, i32* [[CADDR]],
 
   // CK2: call void [[KERNEL:@.+]](i[[sz]] [[VAL]])
   #pragma omp target
   {
    ++i;
   }
+
+  int *&p = b;
+  // CK2-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES2]]{{.+}}, {{.+}}[[TYPES2]]{{.+}})
+  // CK2-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
+  // CK2-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
+  // CK2-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
+  // CK2-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
+  // CK2-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]],
+  // CK2-DAG: store i8* [[VALP:%.+]], i8** [[P1]],
+  // CK2-DAG: [[VALBP]] = bitcast i32* [[VAL:%.+]] to i8*
+  // CK2-DAG: [[VALP]] = bitcast i32* [[VAL]] to i8*
+  // CK2-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]],
+  // CK2-DAG: [[ADDR]] = load i32**, i32*** [[ADDR2:%.+]],
+
+  // CK2: call void [[KERNEL2:@.+]](i32* [[VAL]])
+  #pragma omp target
+  {
+   ++p;
+  }
 }
 
 // CK2: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]])
 // CK2: [[ADDR:%.+]] = alloca i[[sz]],
 // CK2: [[REF:%.+]] = alloca i32*,
 // CK2: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]],
 // CK2-64: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32*
 // CK2-64: store i32* [[CADDR]], i32** [[REF]],
 // CK2-64: [[RVAL:%.+]] = load i32*, i32** [[REF]],
 // CK2-64: {{.+}} = load i32, i32* [[RVAL]],
 // CK2-32: store i32* [[ADDR]], i32** [[REF]],
 // CK2-32: [[RVAL:%.+]] = load i32*, i32** [[REF]],
 // CK2-32: {{.+}} = load i32, i32* [[RVAL]],
 
+// CK2: define internal void [[KERNEL2]](i32* [[ARG:%.+]])
+// CK2: [[ADDR:%.+]] = alloca i32*,
+// CK2: [[REF:%.+]] = alloca i32**,
+// CK2: store i32* [[ARG]], i32** [[ADDR]],
+// CK2: store i32** [[ADDR]], i32*** [[REF]],
+// CK2: [[T:%.+]] = load i32**, i32*** [[REF]],
+// CK2: [[TT:%.+]] = load i32*, i32** [[T]],
+// CK2: getelementptr inbounds i32, i32* [[TT]], i32 1
 #endif
 ///==========================================================================///
 // RUN: %clang_cc1 -DCK3 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK3 --check-prefix CK3-64
 // RUN: %clang_cc1 -DCK3 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK3 --check-prefix CK3-64
 // RUN: %clang_cc1 -DCK3 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK3 --check-prefix CK3-32
 // RUN: %clang_cc1 -DCK3 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK3 --check-prefix CK3-32

 // CK27: define {{.+}}[[CALL01]]
 // CK27: define {{.+}}[[CALL02]]
 // CK27: define {{.+}}[[CALL03]]
 // CK27: define {{.+}}[[CALL04]]
 // CK27: define {{.+}}[[CALL05]]
 // CK27: define {{.+}}[[CALL06]]
 // CK27: define {{.+}}[[CALL07]]
 #endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK28 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK28 --check-prefix CK28-64
+// RUN: %clang_cc1 -DCK28 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK28 --check-prefix CK28-64
+// RUN: %clang_cc1 -DCK28 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK28 --check-prefix CK28-32
+// RUN: %clang_cc1 -DCK28 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK28 --check-prefix CK28-32
+#ifdef CK28
+
+// CK28: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] {{8|4}}]
+// CK28: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
+
+// CK28: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 400]
+// CK28: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 35]
+
+// CK28-LABEL: explicit_maps_pointer_references
+void explicit_maps_pointer_references (int *p){
+  int *&a = p;
+
+  // Region 00
+  // CK28-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // CK28-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK28-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK28-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK28-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK28-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK28-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK28-DAG: [[CBPVAL0]] = bitcast i32** [[VAR0:%.+]] to i8*
+  // CK28-DAG: [[CPVAL0]] = bitcast i32** [[VAR1:%.+]] to i8*
+  // CK28-DAG: [[VAR0]] = load i32**, i32*** [[VAR00:%.+]],
+  // CK28-DAG: [[VAR1]] = load i32**, i32*** [[VAR11:%.+]],
+
+  // CK28: call void [[CALL00:@.+]](i32** {{[^,]+}})
+  #pragma omp target map(a)
+  {
+    ++a;
+  }
+
+  // Region 01
+  // CK28-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
+  // CK28-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK28-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK28-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK28-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK28-DAG: store i8* [[CBPVAL0:%[^,]+]], i8** [[BP0]]
+  // CK28-DAG: store i8* [[CPVAL0:%[^,]+]], i8** [[P0]]
+  // CK28-DAG: [[CBPVAL0]] = bitcast i32* [[VAR0:%.+]] to i8*
+  // CK28-DAG: [[CPVAL0]] = bitcast i32* [[VAR1:%.+]] to i8*
+  // CK28-DAG: [[VAR0]] = load i32*, i32** [[VAR00:%.+]],
+  // CK28-DAG: [[VAR00]] = load i32**, i32*** [[VAR000:%.+]],
+  // CK28-DAG: [[VAR1]] = getelementptr inbounds i32, i32* [[VAR11:%.+]], i{{64|32}} 2
+  // CK28-DAG: [[VAR11]] = load i32*, i32** [[VAR111:%.+]],
+  // CK28-DAG: [[VAR111]] = load i32**, i32*** [[VAR1111:%.+]],
+
+  // CK28: call void [[CALL01:@.+]](i32* {{[^,]+}})
+  #pragma omp target map(a[2:100])
+  {
+    ++a;
+  }
+}
+#endif
 #endif