[clang] 6f7c876 - [OPENMP]Fix behaviour of defaultmap for OpenMP 4.5.

[Alexey Bataev via cfe-commits]

Author: Alexey Bataev
Date: 2019-11-22T11:26:09-05:00
New Revision: 6f7c8760a5c01edaa8e9c8e48940e1b8a7aa4128

URL: https://github.com/llvm/llvm-project/commit/6f7c8760a5c01edaa8e9c8e48940e1b8a7aa4128
DIFF: https://github.com/llvm/llvm-project/commit/6f7c8760a5c01edaa8e9c8e48940e1b8a7aa4128.diff

LOG: [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.

Added: 
    

Modified: 
    clang/lib/Sema/SemaOpenMP.cpp
    clang/test/OpenMP/target_map_codegen.cpp

Removed: 
    


################################################################################
diff  --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 45cac595004f..53840d400164 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -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 @@ class DSAAttrChecker final : public StmtVisitor<DSAAttrChecker, void> {
       // 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;

diff  --git a/clang/test/OpenMP/target_map_codegen.cpp b/clang/test/OpenMP/target_map_codegen.cpp
index 0a8198a90417..cd1b5e49174e 100644
--- a/clang/test/OpenMP/target_map_codegen.cpp
+++ b/clang/test/OpenMP/target_map_codegen.cpp
@@ -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]],