maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[OPENMP]Fix behaviour of defaultmap for OpenMP 4.5. 

In OpenMP 4.5 pointers also must be considered as scalar types and
defaultmap(tofrom:scalar) clause must affect mapping of the pointers
too.
This commit is contained in:
Alexey Bataev 2019-11-22 11:09:33 -05:00
parent 1a58be2ac5
commit 6f7c8760a5
2 changed files with 19 additions and 13 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

18
clang/lib/Sema/SemaOpenMP.cpp
View File

@ -1763,7 +1763,12 @@ void Sema::checkOpenMPDeviceExpr(const Expr *E) {
}
static OpenMPDefaultmapClauseKind
getVariableCategoryFromDecl(const ValueDecl *VD) {
getVariableCategoryFromDecl(const LangOptions &LO, const ValueDecl *VD) {
if (LO.OpenMP <= 45) {
if (VD->getType().getNonReferenceType()->isScalarType())
return OMPC_DEFAULTMAP_scalar;
return OMPC_DEFAULTMAP_aggregate;
}
if (VD->getType().getNonReferenceType()->isAnyPointerType())
return OMPC_DEFAULTMAP_pointer;
if (VD->getType().getNonReferenceType()->isScalarType())
@ -1894,8 +1899,8 @@ bool Sema::isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level,
(DSAStack->isForceCaptureByReferenceInTargetExecutable() &&
!Ty->isAnyPointerType()) ||
!Ty->isScalarType() ||
DSAStack->isDefaultmapCapturedByRef(Level,
getVariableCategoryFromDecl(D)) ||
DSAStack->isDefaultmapCapturedByRef(
Level, getVariableCategoryFromDecl(LangOpts, D)) ||
DSAStack->hasExplicitDSA(
D, [](OpenMPClauseKind K) { return K == OMPC_reduction; }, Level);
}
@ -2112,8 +2117,8 @@ void Sema::setOpenMPCaptureKind(FieldDecl *FD, const ValueDecl *D,
if (DSAStack->hasExplicitDirective(isOpenMPTargetExecutionDirective,
NewLevel)) {
OMPC = OMPC_map;
if (DSAStack->mustBeFirstprivateAtLevel(NewLevel,
getVariableCategoryFromDecl(D)))
if (DSAStack->mustBeFirstprivateAtLevel(
NewLevel, getVariableCategoryFromDecl(LangOpts, D)))
OMPC = OMPC_firstprivate;
break;
}
@ -2944,7 +2949,8 @@ public:
// data-haring attribute clause (including a data-sharing attribute
// clause on a combined construct where target. is one of the
// constituent constructs), or an is_device_ptr clause.
OpenMPDefaultmapClauseKind ClauseKind = getVariableCategoryFromDecl(VD);
OpenMPDefaultmapClauseKind ClauseKind =
getVariableCategoryFromDecl(SemaRef.getLangOpts(), VD);
if (SemaRef.getLangOpts().OpenMP >= 50) {
bool IsModifierNone = Stack->getDefaultmapModifier(ClauseKind) ==
OMPC_DEFAULTMAP_MODIFIER_none;

14
clang/test/OpenMP/target_map_codegen.cpp
View File

@ -656,15 +656,15 @@ void implicit_maps_pointer (){
// CK11-LABEL: @.__omp_offloading_{{.*}}implicit_maps_double_complex{{.*}}_l678.region_id = weak constant i8 0
// CK11-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
// CK11-DAG: [[SIZES:@.+]] = {{.+}}constant [2 x i64] [i64 16, i64 {{8|4}}]
// Map types: OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 547
// CK11-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 547]
// CK11-DAG: [[TYPES:@.+]] = {{.+}}constant [2 x i64] [i64 547, i64 547]
// CK11-LABEL: implicit_maps_double_complex{{.*}}(
void implicit_maps_double_complex (int a){
void implicit_maps_double_complex (int a, int *b){
double _Complex dc = (double)a;
// CK11-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
// CK11-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 2, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
// CK11-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
// CK11-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
// CK11-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
@ -674,14 +674,14 @@ void implicit_maps_double_complex (int a){
// CK11-DAG: store { double, double }* [[PTR:%[^,]+]], { double, double }** [[CBP1]]
// CK11-DAG: store { double, double }* [[PTR]], { double, double }** [[CP1]]
// CK11: call void [[KERNEL:@.+]]({ double, double }* [[PTR]])
// CK11: call void [[KERNEL:@.+]]({ double, double }* [[PTR]], i32** %{{.+}})
#pragma omp target defaultmap(tofrom:scalar)
{
dc *= dc;
dc *= dc; *b = 1;
}
}
// CK11: define internal void [[KERNEL]]({ double, double }* {{.*}}[[ARG:%.+]])
// CK11: define internal void [[KERNEL]]({ double, double }* {{.*}}[[ARG:%.+]], i32** {{.*}})
// CK11: [[ADDR:%.+]] = alloca { double, double }*,
// CK11: store { double, double }* [[ARG]], { double, double }** [[ADDR]],
// CK11: [[REF:%.+]] = load { double, double }*, { double, double }** [[ADDR]],