[OPENMP]Fix PR49777: Clang should not try to specialize orphaned directives in device codegen.

[ABataev]

Summary

Compiler supports generic code emission, but in some cases may
erroneously consider the function context as SPMD context or Non-SPMD
parallel context. Need to clear/restore context upon entrance/exit
to/from function to avoid incorrect codegen.

[ABataev]

Rebase

[ABataev]

Rebase

[jdoerfert]

The bug has been fixed by D95976 already, I'll update the bug report now.
Also, the tracking of Generic/SPMD mode in clang is about to be finally removed, new code depending on that is short-lived.

clang/lib/CodeGen/CGOpenMPRuntimeGPU.h

private:
   /// region is L0 for sure.
   bool IsInTargetMasterThreadRegion = false;
   /// true if currently emitting code for target/teams/distribute region, false
   /// - otherwise.
   bool IsInTTDRegion = false;
   /// true if we're definitely in the parallel region.
   bool IsInParallelRegion = false;
 
+  struct StateMode {
+    ExecutionMode SavedExecutionMode = EM_Unknown;
+    bool SavedIsInTargetMasterThreadRegion = false;
+    bool SavedIsInTTDRegion = false;
+    bool SavedIsInParallelRegion = false;
+    StateMode(ExecutionMode SavedExecutionMode,
+              bool SavedIsInTargetMasterThreadRegion, bool SavedIsInTTDRegion,
+              bool SavedIsInParallelRegion)
+        : SavedExecutionMode(SavedExecutionMode),
+          SavedIsInTargetMasterThreadRegion(SavedIsInTargetMasterThreadRegion),
+          SavedIsInTTDRegion(SavedIsInTTDRegion),
+          SavedIsInParallelRegion(SavedIsInParallelRegion) {}
+  };
+  llvm::DenseMap<CanonicalDeclPtr<const Decl>, StateMode> SavedExecutionModes;
   /// Map between an outlined function and its wrapper.
   llvm::DenseMap<llvm::Function *, llvm::Function *> WrapperFunctionsMap;
 
   /// Emit function which wraps the outline parallel region
   /// and controls the parameters which are passed to this function.
   /// The wrapper ensures that the outlined function is called
   /// with the correct arguments when data is shared.
   llvm::Function *createParallelDataSharingWrapper(

clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp

llvm::Function *CGOpenMPRuntimeGPU::createParallelDataSharingWrapper(
 
   emitOutlinedFunctionCall(CGF, D.getBeginLoc(), OutlinedParallelFn, Args);
   CGF.FinishFunction();
   return Fn;
 }
 
 void CGOpenMPRuntimeGPU::emitFunctionProlog(CodeGenFunction &CGF,
                                               const Decl *D) {
-  if (getDataSharingMode(CGM) != CGOpenMPRuntimeGPU::Generic)
-    return;
-
   assert(D && "Expected function or captured|block decl.");
   assert(FunctionGlobalizedDecls.count(CGF.CurFn) == 0 &&
          "Function is registered already.");
   assert((!TeamAndReductions.first || TeamAndReductions.first == D) &&
          "Team is set but not processed.");
   const Stmt *Body = nullptr;
   bool NeedToDelayGlobalization = false;
   if (const auto *FD = dyn_cast<FunctionDecl>(D)) {
     Body = FD->getBody();
   } else if (const auto *BD = dyn_cast<BlockDecl>(D)) {
     Body = BD->getBody();
   } else if (const auto *CD = dyn_cast<CapturedDecl>(D)) {
     Body = CD->getBody();
     NeedToDelayGlobalization = CGF.CapturedStmtInfo->getKind() == CR_OpenMP;
     if (NeedToDelayGlobalization &&
         getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD)
       return;
   }
+  if (!NeedToDelayGlobalization) {
+    SavedExecutionModes.try_emplace(D, CurrentExecutionMode,
+                                    IsInTargetMasterThreadRegion, IsInTTDRegion,
+                                    IsInParallelRegion);
+    CurrentExecutionMode = EM_Unknown;
+    IsInTargetMasterThreadRegion = false;
+    IsInTTDRegion = false;
+    IsInParallelRegion = false;
+  }
+  if (getDataSharingMode(CGM) != CGOpenMPRuntimeGPU::Generic)
+    return;
+
   if (!Body)
     return;
   CheckVarsEscapingDeclContext VarChecker(CGF, TeamAndReductions.second);
   VarChecker.Visit(Body);
   const RecordDecl *GlobalizedVarsRecord =
       VarChecker.getGlobalizedRecord(IsInTTDRegion);
   TeamAndReductions.first = nullptr;
   TeamAndReductions.second.clear();
for (specific_attr_iterator<OMPReferencedVarAttr> IT(VD->attr_begin()),
         return VDI->second.PrivateAddr;
     }
   }
 
   return Address::invalid();
 }
 
 void CGOpenMPRuntimeGPU::functionFinished(CodeGenFunction &CGF) {
+  auto I = SavedExecutionModes.find(CGF.CurCodeDecl);
+  if (I != SavedExecutionModes.end()) {
+    CurrentExecutionMode = I->second.SavedExecutionMode;
+    IsInTargetMasterThreadRegion = I->second.SavedIsInTargetMasterThreadRegion;
+    IsInTTDRegion = I->second.SavedIsInTTDRegion;
+    IsInParallelRegion = I->second.SavedIsInParallelRegion;
+    SavedExecutionModes.erase(CGF.CurCodeDecl);
+  }
   FunctionGlobalizedDecls.erase(CGF.CurFn);
   CGOpenMPRuntime::functionFinished(CGF);
 }
 
 void CGOpenMPRuntimeGPU::getDefaultDistScheduleAndChunk(
     CodeGenFunction &CGF, const OMPLoopDirective &S,
     OpenMPDistScheduleClauseKind &ScheduleKind,
     llvm::Value *&Chunk) const {

clang/test/OpenMP/declare_target_codegen_globalization.cpp

 // CHECK: ret void
 
 // CHECK: define {{.*}}[[FOO]](i32* nonnull align {{[0-9]+}} dereferenceable{{.*}})
 // CHECK-NOT: @__kmpc_data_sharing_coalesced_push_stack
 
 // CHECK: define {{.*}}[[BAR]]()
 // CHECK: alloca i32,
 // CHECK: [[A_LOCAL_ADDR:%.+]] = alloca i32,
+// CHECK: [[PL:%.+]] = call i16 @__kmpc_parallel_level(
+// CHECK: [[IS_IN_PARALLEL:%.+]] = icmp eq i16 [[PL]], 0
 // CHECK: [[RES:%.+]] = call i8 @__kmpc_is_spmd_exec_mode()
 // CHECK: [[IS_SPMD:%.+]] = icmp ne i8 [[RES]], 0
 // CHECK: br i1 [[IS_SPMD]], label
 // CHECK: br label
-// CHECK: [[RES:%.+]] = call i8* @__kmpc_data_sharing_coalesced_push_stack(i64 128, i16 0)
+// CHECK: [[SZ:%.+]] = select i1 [[IS_IN_PARALLEL]], i64 4, i64 128
+// CHECK: [[RES:%.+]] = call i8* @__kmpc_data_sharing_coalesced_push_stack(i64 [[SZ]], i16 0)
 // CHECK: [[GLOBALS:%.+]] = bitcast i8* [[RES]] to [[GLOBAL_ST:%.+]]*
 // CHECK: br label
 // CHECK: [[ITEMS:%.+]] = phi [[GLOBAL_ST]]* [ null, {{.+}} ], [ [[GLOBALS]], {{.+}} ]
 // CHECK: [[A_ADDR:%.+]] = getelementptr inbounds [[GLOBAL_ST]], [[GLOBAL_ST]]* [[ITEMS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
 // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
 // CHECK: [[LID:%.+]] = and i32 [[TID]], 31
 // CHECK: [[A_GLOBAL_ADDR:%.+]] = getelementptr inbounds [32 x i32], [32 x i32]* [[A_ADDR]], i32 0, i32 [[LID]]
-// CHECK: [[A_ADDR:%.+]] = select i1 [[IS_SPMD]], i32* [[A_LOCAL_ADDR]], i32* [[A_GLOBAL_ADDR]]
+// CHECK: [[A_GLOBAL_PARALLEL_ADDR:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 0
+// CHECK: [[A_PARALLEL_ADDR:%.+]] = select i1 [[IS_IN_PARALLEL]], i32* [[A_GLOBAL_PARALLEL_ADDR]], i32* [[A_GLOBAL_ADDR]]
+// CHECK: [[A_ADDR:%.+]] = select i1 [[IS_SPMD]], i32* [[A_LOCAL_ADDR]], i32* [[A_PARALLEL_ADDR]]
 // CHECK: call {{.*}}[[FOO]](i32* nonnull align {{[0-9]+}} dereferenceable{{.*}} [[A_ADDR]])
 // CHECK: br i1 [[IS_SPMD]], label
 // CHECK: [[BC:%.+]] = bitcast [[GLOBAL_ST]]* [[ITEMS]] to i8*
 // CHECK: call void @__kmpc_data_sharing_pop_stack(i8* [[BC]])
 // CHECK: br label
 // CHECK: ret i32

clang/test/OpenMP/remarks_parallel_in_multiple_target_state_machines.c

#pragma omp parallel
   {
   }
 
 #pragma omp target teams distribute parallel for
   for (int i = 0; i < 100; ++i) {
   }
 }
 
-// all-remark@* 3 {{OpenMP runtime call __kmpc_global_thread_num moved to}}
-// all-remark@* 3 {{OpenMP runtime call __kmpc_global_thread_num deduplicated}}
+// all-remark@* 5 {{OpenMP runtime call __kmpc_global_thread_num moved to}}
+// all-remark@* 12 {{OpenMP runtime call __kmpc_global_thread_num deduplicated}}

clang/test/OpenMP/remarks_parallel_in_target_state_machine.c

#pragma omp parallel
   }
 
 #pragma omp target teams distribute parallel for
   for (int i = 0; i < 100; ++i) {
   }
 }
 
 // expected-remark@* {{OpenMP runtime call __kmpc_global_thread_num moved to}}
-// expected-remark@* {{OpenMP runtime call __kmpc_global_thread_num deduplicated}}
+// expected-remark@* 2 {{OpenMP runtime call __kmpc_global_thread_num deduplicated}}