[OpenMP] Avoid emitting maps for target link variables when unified memory is used

[gtbercea]

Summary

This patch avoids the emission of maps for target link variables when unified memory is present.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

7956 ↗

(On Diff #195798)

Hmm, what about regular declare target declarations? Also, this is not enough. You also should not capture such variables in the captured regions.

[gtbercea]

• Fix function call.

[gtbercea]

lib/CodeGen/CGOpenMPRuntime.cpp

7956 ↗

(On Diff #195798)

For now, only variables under the link clause are considered for unified memory usage.

[gtbercea]

• Include sema check.

[ABataev]

Tests?

lib/CodeGen/CGOpenMPRuntime.h

1628 ↗

(On Diff #200818)

Do we need to make it virtual?

[gtbercea]

• Add test.

[ABataev]

Just one question after looking at the test: how we're going to link the device variable to its host copy? I think the address should be passed to the runtime, no?

lib/CodeGen/CGOpenMPRuntime.h

1628 ↗

(On Diff #200818)

Still not answered

[gtbercea]

In D60883#1516693, @ABataev wrote:

Just one question after looking at the test: how we're going to link the device variable to its host copy? I think the address should be passed to the runtime, no?

The runtime will take care of it by performing a copy of the host address into the device pointer.

[gtbercea]

• Remove virtual.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.h

1628 ↗

(On Diff #204780)

You can make this function const.

[gtbercea]

• Make function const.

cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h

public:
   /// Perform check on requires decl to ensure that target architecture
   /// supports unified addressing
   virtual void checkArchForUnifiedAddressing(const OMPRequiresDecl *D);
 
   /// Checks if the variable has associated OMPAllocateDeclAttr attribute with
   /// the predefined allocator and translates it into the corresponding address
   /// space.
   virtual bool hasAllocateAttributeForGlobalVar(const VarDecl *VD, LangAS &AS);
+
+  /// Return whether the unified_shared_memory has been specified.
+  bool hasRequiresUnifiedSharedMemory() const;
 };
 
 /// Class supports emissionof SIMD-only code.
 class CGOpenMPSIMDRuntime final : public CGOpenMPRuntime {
 public:
   explicit CGOpenMPSIMDRuntime(CodeGenModule &CGM) : CGOpenMPRuntime(CGM) {}
   ~CGOpenMPSIMDRuntime() override {}
 

cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp

for (const auto *C : this->CurDir.getClausesOfKind<OMPMapClause>()) {
       for (const auto &L : C->component_lists()) {
         if (!L.first)
           continue;
         const auto *VD = dyn_cast<VarDecl>(L.first);
         if (!VD)
           continue;
         llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
             OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
-        if (!Res || *Res != OMPDeclareTargetDeclAttr::MT_Link)
+        if (CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory() ||
+            !Res || *Res != OMPDeclareTargetDeclAttr::MT_Link)
           continue;
         StructRangeInfoTy PartialStruct;
         generateInfoForComponentList(
             C->getMapType(), C->getMapTypeModifiers(), L.second, BasePointers,
             Pointers, Sizes, Types, PartialStruct,
             /*IsFirstComponentList=*/true, C->isImplicit());
         assert(!PartialStruct.Base.isValid() &&
                "No partial structs for declare target link expected.");
case OMPAllocateDeclAttr::OMPPTeamMemAlloc:
     return true;
   case OMPAllocateDeclAttr::OMPUserDefinedMemAlloc:
     llvm_unreachable("Expected predefined allocator for the variables with the "
                      "static storage.");
   }
   return false;
 }
 
+bool CGOpenMPRuntime::hasRequiresUnifiedSharedMemory() const {
+  return HasRequiresUnifiedSharedMemory;
+}
+
 CGOpenMPRuntime::DisableAutoDeclareTargetRAII::DisableAutoDeclareTargetRAII(
     CodeGenModule &CGM)
     : CGM(CGM) {
   if (CGM.getLangOpts().OpenMPIsDevice) {
     SavedShouldMarkAsGlobal = CGM.getOpenMPRuntime().ShouldMarkAsGlobal;
     CGM.getOpenMPRuntime().ShouldMarkAsGlobal = false;
   }
 }

cfe/trunk/lib/Sema/SemaOpenMP.cpp

if (auto *VD = dyn_cast<VarDecl>(E->getDecl())) {
       // Check if the variable has explicit DSA set and stop analysis if it so.
       if (DVar.RefExpr || !ImplicitDeclarations.insert(VD).second)
         return;
 
       // Skip internally declared static variables.
       llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
           OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
       if (VD->hasGlobalStorage() && CS && !CS->capturesVariable(VD) &&
-          (!Res || *Res != OMPDeclareTargetDeclAttr::MT_Link))
+          (Stack->hasRequiresDeclWithClause<OMPUnifiedSharedMemoryClause>() ||
+           !Res || *Res != OMPDeclareTargetDeclAttr::MT_Link))
         return;
 
       SourceLocation ELoc = E->getExprLoc();
       OpenMPDirectiveKind DKind = Stack->getCurrentDirective();
       // The default(none) clause requires that each variable that is referenced
       // in the construct, and does not have a predetermined data-sharing
       // attribute, must have its data-sharing attribute explicitly determined
       // by being listed in a data-sharing attribute clause.

cfe/trunk/test/OpenMP/nvptx_target_requires_unified_shared_memory.cpp

#pragma omp target
   }
 
   return sum;
 }
 
 // CHECK: [[VAR:@.+]] = global double 1.000000e+01
 // CHECK: [[VAR_DECL_TGT_LINK_PTR:@.+]] = global double* [[VAR]]
 
-// CHECK: [[OFFLOAD_SIZES:@.+]] = private unnamed_addr constant [3 x i64] [i64 4, i64 8, i64 8]
-// CHECK: [[OFFLOAD_MAPTYPES:@.+]] = private unnamed_addr constant [3 x i64] [i64 800, i64 800, i64 531]
+// CHECK: [[OFFLOAD_SIZES:@.+]] = private unnamed_addr constant [2 x i64] [i64 4, i64 8]
+// CHECK: [[OFFLOAD_MAPTYPES:@.+]] = private unnamed_addr constant [2 x i64] [i64 800, i64 800]
 
 // CHECK: [[N_CASTED:%.+]] = alloca i64
 // CHECK: [[SUM_CASTED:%.+]] = alloca i64
 
-// CHECK: [[OFFLOAD_BASEPTRS:%.+]] = alloca [3 x i8*]
-// CHECK: [[OFFLOAD_PTRS:%.+]] = alloca [3 x i8*]
+// CHECK: [[OFFLOAD_BASEPTRS:%.+]] = alloca [2 x i8*]
+// CHECK: [[OFFLOAD_PTRS:%.+]] = alloca [2 x i8*]
 
 // CHECK: [[LOAD1:%.+]] = load i64, i64* [[N_CASTED]]
 // CHECK: [[LOAD2:%.+]] = load i64, i64* [[SUM_CASTED]]
 
-// CHECK: [[BPTR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[OFFLOAD_BASEPTRS]], i32 0, i32 0
+// CHECK: [[BPTR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[OFFLOAD_BASEPTRS]], i32 0, i32 0
 // CHECK: [[BCAST1:%.+]] = bitcast i8** [[BPTR1]] to i64*
 // CHECK: store i64 [[LOAD1]], i64* [[BCAST1]]
-// CHECK: [[BPTR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[OFFLOAD_PTRS]], i32 0, i32 0
+// CHECK: [[BPTR2:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[OFFLOAD_PTRS]], i32 0, i32 0
 // CHECK: [[BCAST2:%.+]] = bitcast i8** [[BPTR2]] to i64*
 // CHECK: store i64 [[LOAD1]], i64* [[BCAST2]]
 
-// CHECK: [[BPTR3:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[OFFLOAD_BASEPTRS]], i32 0, i32 1
+// CHECK: [[BPTR3:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[OFFLOAD_BASEPTRS]], i32 0, i32 1
 // CHECK: [[BCAST3:%.+]] = bitcast i8** [[BPTR3]] to i64*
 // CHECK: store i64 [[LOAD2]], i64* [[BCAST3]]
-// CHECK: [[BPTR4:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[OFFLOAD_PTRS]], i32 0, i32 1
+// CHECK: [[BPTR4:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[OFFLOAD_PTRS]], i32 0, i32 1
 // CHECK: [[BCAST4:%.+]] = bitcast i8** [[BPTR4]] to i64*
 // CHECK: store i64 [[LOAD2]], i64* [[BCAST4]]
 
-// CHECK: [[BPTR5:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[OFFLOAD_BASEPTRS]], i32 0, i32 2
-// CHECK: [[BCAST5:%.+]] = bitcast i8** [[BPTR5]] to double***
-// CHECK: store double** [[VAR_DECL_TGT_LINK_PTR]], double*** [[BCAST5]]
-// CHECK: [[BPTR6:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[OFFLOAD_PTRS]], i32 0, i32 2
-// CHECK: [[BCAST6:%.+]] = bitcast i8** [[BPTR6]] to double**
-// CHECK: store double* [[VAR]], double** [[BCAST6]]
+// CHECK: [[BPTR7:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[OFFLOAD_BASEPTRS]], i32 0, i32 0
+// CHECK: [[BPTR8:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[OFFLOAD_PTRS]], i32 0, i32 0
 
-// CHECK: [[BPTR7:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[OFFLOAD_BASEPTRS]], i32 0, i32 0
-// CHECK: [[BPTR8:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[OFFLOAD_PTRS]], i32 0, i32 0
-
-// CHECK: call i32 @__tgt_target(i64 -1, i8* @{{.*}}.region_id, i32 3, i8** [[BPTR7]], i8** [[BPTR8]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[OFFLOAD_SIZES]], i32 0, i32 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[OFFLOAD_MAPTYPES]], i32 0, i32 0))
+// CHECK: call i32 @__tgt_target(i64 -1, i8* @{{.*}}.region_id, i32 2, i8** [[BPTR7]], i8** [[BPTR8]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[OFFLOAD_SIZES]], i32 0, i32 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[OFFLOAD_MAPTYPES]], i32 0, i32 0))
 
 #endif