Index: lib/CodeGen/CGOpenMPRuntime.cpp =================================================================== --- lib/CodeGen/CGOpenMPRuntime.cpp +++ lib/CodeGen/CGOpenMPRuntime.cpp @@ -5271,15 +5271,13 @@ // 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()) + Ty->castAs()) .getPointer(); // We do not need to generate individual map information for the Index: lib/CodeGen/CGStmtOpenMP.cpp =================================================================== --- lib/CodeGen/CGStmtOpenMP.cpp +++ lib/CodeGen/CGStmtOpenMP.cpp @@ -264,7 +264,17 @@ // 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; Index: test/OpenMP/target_map_codegen.cpp =================================================================== --- test/OpenMP/target_map_codegen.cpp +++ test/OpenMP/target_map_codegen.cpp @@ -60,12 +60,15 @@ // 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 @@ -85,6 +88,25 @@ { ++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:%.+]]) @@ -99,6 +121,14 @@ // 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 @@ -4471,4 +4501,67 @@ // 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