Skip to content

Commit c2aa543

Browse files
alexey-bataevalexey-bataev
committed
[OPENMP50]Codegen for array shaping expression in map clauses.
Added codegen support for array shaping operations in map/to/from clauses.
1 parent a3fab31 commit c2aa543

File tree

7 files changed

+336
-123
lines changed

7 files changed

+336
-123
lines changed

‎clang/lib/CodeGen/CGOpenMPRuntime.cpp

Lines changed: 32 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -7448,6 +7448,20 @@ class MappableExprsHandler {
74487448
llvm::Value *getExprTypeSize(const Expr *E) const {
74497449
QualType ExprTy = E->getType().getCanonicalType();
74507450

7451+
// Calculate the size for array shaping expression.
7452+
if (const auto *OAE = dyn_cast<OMPArrayShapingExpr>(E)) {
7453+
llvm::Value *Size =
7454+
CGF.getTypeSize(OAE->getBase()->getType()->getPointeeType());
7455+
for (const Expr *SE : OAE->getDimensions()) {
7456+
llvm::Value *Sz = CGF.EmitScalarExpr(SE);
7457+
Sz = CGF.EmitScalarConversion(Sz, SE->getType(),
7458+
CGF.getContext().getSizeType(),
7459+
SE->getExprLoc());
7460+
Size = CGF.Builder.CreateNUWMul(Size, Sz);
7461+
}
7462+
return Size;
7463+
}
7464+
74517465
// Reference types are ignored for mapping purposes.
74527466
if (const auto *RefTy = ExprTy->getAs<ReferenceType>())
74537467
ExprTy = RefTy->getPointeeType().getCanonicalType();
@@ -7779,6 +7793,7 @@ class MappableExprsHandler {
77797793
const Expr *AssocExpr = I->getAssociatedExpression();
77807794
const auto *AE = dyn_cast<ArraySubscriptExpr>(AssocExpr);
77817795
const auto *OASE = dyn_cast<OMPArraySectionExpr>(AssocExpr);
7796+
const auto *OAShE = dyn_cast<OMPArrayShapingExpr>(AssocExpr);
77827797

77837798
if (isa<MemberExpr>(AssocExpr)) {
77847799
// The base is the 'this' pointer. The content of the pointer is going
@@ -7788,6 +7803,11 @@ class MappableExprsHandler {
77887803
(OASE &&
77897804
isa<CXXThisExpr>(OASE->getBase()->IgnoreParenImpCasts()))) {
77907805
BP = CGF.EmitOMPSharedLValue(AssocExpr).getAddress(CGF);
7806+
} else if (OAShE &&
7807+
isa<CXXThisExpr>(OAShE->getBase()->IgnoreParenCasts())) {
7808+
BP = Address(
7809+
CGF.EmitScalarExpr(OAShE->getBase()),
7810+
CGF.getContext().getTypeAlignInChars(OAShE->getBase()->getType()));
77917811
} else {
77927812
// The base is the reference to the variable.
77937813
// BP = &Var.
@@ -7870,9 +7890,12 @@ class MappableExprsHandler {
78707890
// types.
78717891
const auto *OASE =
78727892
dyn_cast<OMPArraySectionExpr>(I->getAssociatedExpression());
7893+
const auto *OAShE =
7894+
dyn_cast<OMPArrayShapingExpr>(I->getAssociatedExpression());
78737895
const auto *UO = dyn_cast<UnaryOperator>(I->getAssociatedExpression());
78747896
const auto *BO = dyn_cast<BinaryOperator>(I->getAssociatedExpression());
78757897
bool IsPointer =
7898+
OAShE ||
78767899
(OASE && OMPArraySectionExpr::getBaseOriginalType(OASE)
78777900
.getCanonicalType()
78787901
->isAnyPointerType()) ||
@@ -7890,8 +7913,15 @@ class MappableExprsHandler {
78907913
isa<BinaryOperator>(Next->getAssociatedExpression())) &&
78917914
"Unexpected expression");
78927915

7893-
Address LB = CGF.EmitOMPSharedLValue(I->getAssociatedExpression())
7894-
.getAddress(CGF);
7916+
Address LB = Address::invalid();
7917+
if (OAShE) {
7918+
LB = Address(CGF.EmitScalarExpr(OAShE->getBase()),
7919+
CGF.getContext().getTypeAlignInChars(
7920+
OAShE->getBase()->getType()));
7921+
} else {
7922+
LB = CGF.EmitOMPSharedLValue(I->getAssociatedExpression())
7923+
.getAddress(CGF);
7924+
}
78957925

78967926
// If this component is a pointer inside the base struct then we don't
78977927
// need to create any entry for it - it will be combined with the object

‎clang/lib/Sema/SemaOpenMP.cpp

Lines changed: 28 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -1943,7 +1943,8 @@ bool Sema::isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level,
19431943

19441944
if (isa<ArraySubscriptExpr>(EI->getAssociatedExpression()) ||
19451945
isa<OMPArraySectionExpr>(EI->getAssociatedExpression()) ||
1946-
isa<MemberExpr>(EI->getAssociatedExpression())) {
1946+
isa<MemberExpr>(EI->getAssociatedExpression()) ||
1947+
isa<OMPArrayShapingExpr>(EI->getAssociatedExpression())) {
19471948
IsVariableAssociatedWithSection = true;
19481949
// There is nothing more we need to know about this variable.
19491950
return true;
@@ -3225,7 +3226,7 @@ class DSAAttrChecker final : public StmtVisitor<DSAAttrChecker, void> {
32253226
StackComponents,
32263227
OpenMPClauseKind) {
32273228
// Variable is used if it has been marked as an array, array
3228-
// section or the variable iself.
3229+
// section, array shaping or the variable iself.
32293230
return StackComponents.size() == 1 ||
32303231
std::all_of(
32313232
std::next(StackComponents.rbegin()),
@@ -3236,6 +3237,8 @@ class DSAAttrChecker final : public StmtVisitor<DSAAttrChecker, void> {
32363237
nullptr &&
32373238
(isa<OMPArraySectionExpr>(
32383239
MC.getAssociatedExpression()) ||
3240+
isa<OMPArrayShapingExpr>(
3241+
MC.getAssociatedExpression()) ||
32393242
isa<ArraySubscriptExpr>(
32403243
MC.getAssociatedExpression()));
32413244
});
@@ -3393,8 +3396,10 @@ class DSAAttrChecker final : public StmtVisitor<DSAAttrChecker, void> {
33933396
// Do both expressions have the same kind?
33943397
if (CCI->getAssociatedExpression()->getStmtClass() !=
33953398
SC.getAssociatedExpression()->getStmtClass())
3396-
if (!(isa<OMPArraySectionExpr>(
3397-
SC.getAssociatedExpression()) &&
3399+
if (!((isa<OMPArraySectionExpr>(
3400+
SC.getAssociatedExpression()) ||
3401+
isa<OMPArrayShapingExpr>(
3402+
SC.getAssociatedExpression())) &&
33983403
isa<ArraySubscriptExpr>(
33993404
CCI->getAssociatedExpression())))
34003405
return false;
@@ -16284,6 +16289,15 @@ class MapBaseChecker final : public StmtVisitor<MapBaseChecker, bool> {
1628416289
Components.emplace_back(OASE, nullptr);
1628516290
return RelevantExpr || Visit(E);
1628616291
}
16292+
bool VisitOMPArrayShapingExpr(OMPArrayShapingExpr *E) {
16293+
Expr *Base = E->getBase();
16294+
16295+
// Record the component - we don't have any declaration associated.
16296+
Components.emplace_back(E, nullptr);
16297+
16298+
return Visit(Base->IgnoreParenImpCasts());
16299+
}
16300+
1628716301
bool VisitUnaryOperator(UnaryOperator *UO) {
1628816302
if (SemaRef.getLangOpts().OpenMP < 50 || !UO->isLValue() ||
1628916303
UO->getOpcode() != UO_Deref) {
@@ -16409,9 +16423,11 @@ static bool checkMapConflicts(
1640916423
// variable in map clauses of the same construct.
1641016424
if (CurrentRegionOnly &&
1641116425
(isa<ArraySubscriptExpr>(CI->getAssociatedExpression()) ||
16412-
isa<OMPArraySectionExpr>(CI->getAssociatedExpression())) &&
16426+
isa<OMPArraySectionExpr>(CI->getAssociatedExpression()) ||
16427+
isa<OMPArrayShapingExpr>(CI->getAssociatedExpression())) &&
1641316428
(isa<ArraySubscriptExpr>(SI->getAssociatedExpression()) ||
16414-
isa<OMPArraySectionExpr>(SI->getAssociatedExpression()))) {
16429+
isa<OMPArraySectionExpr>(SI->getAssociatedExpression()) ||
16430+
isa<OMPArrayShapingExpr>(SI->getAssociatedExpression()))) {
1641516431
SemaRef.Diag(CI->getAssociatedExpression()->getExprLoc(),
1641616432
diag::err_omp_multiple_array_items_in_map_clause)
1641716433
<< CI->getAssociatedExpression()->getSourceRange();
@@ -16443,6 +16459,9 @@ static bool checkMapConflicts(
1644316459
const Expr *E = OASE->getBase()->IgnoreParenImpCasts();
1644416460
Type =
1644516461
OMPArraySectionExpr::getBaseOriginalType(E).getCanonicalType();
16462+
} else if (const auto *OASE = dyn_cast<OMPArrayShapingExpr>(
16463+
SI->getAssociatedExpression())) {
16464+
Type = OASE->getBase()->getType()->getPointeeType();
1644616465
}
1644716466
if (Type.isNull() || Type->isAnyPointerType() ||
1644816467
checkArrayExpressionDoesNotReferToWholeSize(
@@ -16905,6 +16924,7 @@ static void checkMappableExpressionList(
1690516924
QualType Type;
1690616925
auto *ASE = dyn_cast<ArraySubscriptExpr>(VE->IgnoreParens());
1690716926
auto *OASE = dyn_cast<OMPArraySectionExpr>(VE->IgnoreParens());
16927+
auto *OAShE = dyn_cast<OMPArrayShapingExpr>(VE->IgnoreParens());
1690816928
if (ASE) {
1690916929
Type = ASE->getType().getNonReferenceType();
1691016930
} else if (OASE) {
@@ -16915,6 +16935,8 @@ static void checkMappableExpressionList(
1691516935
else
1691616936
Type = BaseType->getPointeeType();
1691716937
Type = Type.getNonReferenceType();
16938+
} else if (OAShE) {
16939+
Type = OAShE->getBase()->getType()->getPointeeType();
1691816940
} else {
1691916941
Type = VE->getType();
1692016942
}

‎clang/test/OpenMP/target_data_ast_print.cpp

Lines changed: 8 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -1,10 +1,10 @@
1-
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -ast-print %s | FileCheck %s
2-
// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -emit-pch -o %t %s
3-
// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
1+
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -ast-print %s | FileCheck %s
2+
// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
3+
// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
44

5-
// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -ast-print %s | FileCheck %s
6-
// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -emit-pch -o %t %s
7-
// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
5+
// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=50 -ast-print %s | FileCheck %s
6+
// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
7+
// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
88
// expected-no-diagnostics
99

1010
#ifndef HEADER
@@ -139,6 +139,8 @@ int main (int argc, char **argv) {
139139
static int a;
140140
// CHECK: static int a;
141141

142+
#pragma omp target data map(to: ([argc][3][a])argv)
143+
// CHECK: #pragma omp target data map(to: ([argc][3][a])argv)
142144
#pragma omp target data map(to: c)
143145
// CHECK: #pragma omp target data map(to: c)
144146
a=2;

‎clang/test/OpenMP/target_map_codegen.cpp

Lines changed: 76 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -5353,5 +5353,81 @@ void explicit_maps_single (int ii){
53535353
// CK31: define {{.+}}[[CALL00]]
53545354
// CK31: define {{.+}}[[CALL01]]
53555355

5356+
#endif
5357+
///==========================================================================///
5358+
// RUN: %clang_cc1 -DCK32 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK32 --check-prefix CK32-64
5359+
// RUN: %clang_cc1 -DCK32 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
5360+
// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -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 CK32 --check-prefix CK32-64
5361+
// RUN: %clang_cc1 -DCK32 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK32 --check-prefix CK32-32
5362+
// RUN: %clang_cc1 -DCK32 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
5363+
// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -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 CK32 --check-prefix CK32-32
5364+
5365+
// RUN: %clang_cc1 -DCK32 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY32 %s
5366+
// RUN: %clang_cc1 -DCK32 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
5367+
// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY32 %s
5368+
// RUN: %clang_cc1 -DCK32 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY32 %s
5369+
// RUN: %clang_cc1 -DCK32 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
5370+
// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY32 %s
5371+
// SIMD-ONLY32-NOT: {{__kmpc|__tgt}}
5372+
#ifdef CK32
5373+
5374+
// CK32-DAG: [[MTYPE_TO:@.+]] = {{.+}}constant [1 x i64] [i64 33]
5375+
// CK32-DAG: [[MTYPE_FROM:@.+]] = {{.+}}constant [1 x i64] [i64 34]
5376+
5377+
void array_shaping(float *f, int sa) {
5378+
5379+
// CK32-DAG: call i32 @__tgt_target(i64 -1, i8* @{{.+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE_TO]]{{.+}})
5380+
// CK32-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
5381+
// CK32-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
5382+
// CK32-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
5383+
5384+
// CK32-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
5385+
// CK32-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
5386+
// CK32-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
5387+
5388+
// CK32-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float**
5389+
// CK32-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to float**
5390+
5391+
// CK32-DAG: store float* [[F1:%.+]], float** [[BPC0]],
5392+
// CK32-DAG: store float* [[F2:%.+]], float** [[PC0]],
5393+
// CK32-DAG: store i64 [[SIZE:%.+]], i64* [[S0]],
5394+
5395+
// CK32-DAG: [[F1]] = load float*, float** [[F_ADDR:%.+]],
5396+
// CK32-DAG: [[F2]] = load float*, float** [[F_ADDR]],
5397+
// CK32-64-DAG: [[SIZE]] = mul nuw i64 [[SZ1:%.+]], 4
5398+
// CK32-64-DAG: [[SZ1]] = mul nuw i64 12, %{{.+}}
5399+
// CK32-32-DAG: [[SIZE]] = sext i32 [[SZ1:%.+]] to i64
5400+
// CK32-32-DAG: [[SZ1]] = mul nuw i32 [[SZ2:%.+]], 4
5401+
// CK32-32-DAG: [[SZ2]] = mul nuw i32 12, %{{.+}}
5402+
#pragma omp target map(to:([3][sa][4])f)
5403+
f[0] = 1;
5404+
sa = 1;
5405+
// CK32-DAG: call i32 @__tgt_target(i64 -1, i8* @{{.+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE_FROM]]{{.+}})
5406+
// CK32-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
5407+
// CK32-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
5408+
// CK32-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
5409+
5410+
// CK32-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
5411+
// CK32-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
5412+
// CK32-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
5413+
5414+
// CK32-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float**
5415+
// CK32-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to float**
5416+
5417+
// CK32-DAG: store float* [[F1:%.+]], float** [[BPC0]],
5418+
// CK32-DAG: store float* [[F2:%.+]], float** [[PC0]],
5419+
// CK32-DAG: store i64 [[SIZE:%.+]], i64* [[S0]],
5420+
5421+
// CK32-DAG: [[F1]] = load float*, float** [[F_ADDR:%.+]],
5422+
// CK32-DAG: [[F2]] = load float*, float** [[F_ADDR]],
5423+
// CK32-64-DAG: [[SIZE]] = mul nuw i64 [[SZ1:%.+]], 5
5424+
// CK32-64-DAG: [[SZ1]] = mul nuw i64 4, %{{.+}}
5425+
// CK32-32-DAG: [[SIZE]] = sext i32 [[SZ1:%.+]] to i64
5426+
// CK32-32-DAG: [[SZ1]] = mul nuw i32 [[SZ2:%.+]], 5
5427+
// CK32-32-DAG: [[SZ2]] = mul nuw i32 4, %{{.+}}
5428+
#pragma omp target map(from: ([sa][5])f)
5429+
f[0] = 1;
5430+
}
5431+
53565432
#endif
53575433
#endif

0 commit comments

Comments
 (0)