diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index 121de42248e3b..44ba72c5c76c7 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -2000,22 +2000,29 @@ void CGOpenMPRuntime::emitCriticalRegion(CodeGenFunction &CGF, // Prepare arguments and build a call to __kmpc_critical if (!CGF.HaveInsertPoint()) return; + llvm::FunctionCallee RuntimeFcn = OMPBuilder.getOrCreateRuntimeFunction( + CGM.getModule(), + Hint ? OMPRTL___kmpc_critical_with_hint : OMPRTL___kmpc_critical); + llvm::Value *LockVar = getCriticalRegionLock(CriticalName); + unsigned LockVarArgIdx = 2; + if (cast(LockVar)->getAddressSpace() != + RuntimeFcn.getFunctionType() + ->getParamType(LockVarArgIdx) + ->getPointerAddressSpace()) + LockVar = CGF.Builder.CreateAddrSpaceCast( + LockVar, RuntimeFcn.getFunctionType()->getParamType(LockVarArgIdx)); llvm::Value *Args[] = {emitUpdateLocation(CGF, Loc), getThreadID(CGF, Loc), - getCriticalRegionLock(CriticalName)}; + LockVar}; llvm::SmallVector EnterArgs(std::begin(Args), std::end(Args)); if (Hint) { EnterArgs.push_back(CGF.Builder.CreateIntCast( CGF.EmitScalarExpr(Hint), CGM.Int32Ty, /*isSigned=*/false)); } - CommonActionTy Action( - OMPBuilder.getOrCreateRuntimeFunction( - CGM.getModule(), - Hint ? OMPRTL___kmpc_critical_with_hint : OMPRTL___kmpc_critical), - EnterArgs, - OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(), - OMPRTL___kmpc_end_critical), - Args); + CommonActionTy Action(RuntimeFcn, EnterArgs, + OMPBuilder.getOrCreateRuntimeFunction( + CGM.getModule(), OMPRTL___kmpc_end_critical), + Args); CriticalOpGen.setAction(Action); emitInlinedDirective(CGF, OMPD_critical, CriticalOpGen); } diff --git a/clang/test/OpenMP/force-usm.c b/clang/test/OpenMP/force-usm.c index 5c63a9a5e7004..45c0e28b525da 100644 --- a/clang/test/OpenMP/force-usm.c +++ b/clang/test/OpenMP/force-usm.c @@ -46,7 +46,7 @@ int main(void) { // CHECK-USM-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]] // CHECK-USM: user_code.entry: // CHECK-USM-NEXT: store i32 1, ptr [[TMP0]], align 4 -// CHECK-USM-NEXT: [[TMP2:%.*]] = load ptr, ptr @pGI_decl_tgt_ref_ptr, align 8 +// CHECK-USM-NEXT: [[TMP2:%.*]] = load ptr, ptr addrspace(1) @pGI_decl_tgt_ref_ptr, align 8 // CHECK-USM-NEXT: [[TMP3:%.*]] = load ptr, ptr [[TMP2]], align 8 // CHECK-USM-NEXT: store i32 2, ptr [[TMP3]], align 4 // CHECK-USM-NEXT: call void @__kmpc_target_deinit() diff --git a/clang/test/OpenMP/spirv_target_codegen_basic.cpp b/clang/test/OpenMP/spirv_target_codegen_basic.cpp index fb2810e88c063..6e029fb93644d 100644 --- a/clang/test/OpenMP/spirv_target_codegen_basic.cpp +++ b/clang/test/OpenMP/spirv_target_codegen_basic.cpp @@ -6,12 +6,18 @@ // CHECK: @__omp_offloading_{{.*}}_dynamic_environment = weak_odr protected addrspace(1) global %struct.DynamicEnvironmentTy zeroinitializer // CHECK: @__omp_offloading_{{.*}}_kernel_environment = weak_odr protected addrspace(1) constant %struct.KernelEnvironmentTy +// CHECK: @"_gomp_critical_user_$var" = common addrspace(1) global [8 x i32] zeroinitializer, align 8 + // CHECK: define weak_odr protected spir_kernel void @__omp_offloading_{{.*}} +// CHECK: call spir_func addrspace(9) void @__kmpc_critical(ptr addrspace(4) addrspacecast (ptr addrspace(1) @{{.*}} to ptr addrspace(4)), i32 %{{.*}}, ptr addrspace(4) addrspacecast (ptr addrspace(1) @"_gomp_critical_user_$var" to ptr addrspace(4))) +// CHECK: call spir_func addrspace(9) void @__kmpc_end_critical(ptr addrspace(4) addrspacecast (ptr addrspace(1) @{{.*}} to ptr addrspace(4)), i32 %{{.*}}, ptr addrspace(4) addrspacecast (ptr addrspace(1) @"_gomp_critical_user_$var" to ptr addrspace(4))) + int main() { int ret = 0; #pragma omp target for(int i = 0; i < 5; i++) + #pragma omp critical ret++; return ret; } diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h index b3d7ab4acf303..fd6b9729658c1 100644 --- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h +++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h @@ -3654,7 +3654,7 @@ class OpenMPIRBuilder { /// \param Name Name of the variable. LLVM_ABI GlobalVariable * getOrCreateInternalVariable(Type *Ty, const StringRef &Name, - unsigned AddressSpace = 0); + std::optional AddressSpace = {}); }; /// Class to represented the control flow structure of an OpenMP canonical loop. diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp index fff9a815e5368..7dc32fda0eed6 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -8460,9 +8460,8 @@ OpenMPIRBuilder::createPlatformSpecificName(ArrayRef Parts) const { Config.separator()); } -GlobalVariable * -OpenMPIRBuilder::getOrCreateInternalVariable(Type *Ty, const StringRef &Name, - unsigned AddressSpace) { +GlobalVariable *OpenMPIRBuilder::getOrCreateInternalVariable( + Type *Ty, const StringRef &Name, std::optional AddressSpace) { auto &Elem = *InternalVars.try_emplace(Name, nullptr).first; if (Elem.second) { assert(Elem.second->getValueType() == Ty && @@ -8472,16 +8471,18 @@ OpenMPIRBuilder::getOrCreateInternalVariable(Type *Ty, const StringRef &Name, // variable for possibly changing that to internal or private, or maybe // create different versions of the function for different OMP internal // variables. + const DataLayout &DL = M.getDataLayout(); + unsigned AddressSpaceVal = + AddressSpace ? *AddressSpace : DL.getDefaultGlobalsAddressSpace(); auto Linkage = this->M.getTargetTriple().getArch() == Triple::wasm32 ? GlobalValue::InternalLinkage : GlobalValue::CommonLinkage; auto *GV = new GlobalVariable(M, Ty, /*IsConstant=*/false, Linkage, Constant::getNullValue(Ty), Elem.first(), /*InsertBefore=*/nullptr, - GlobalValue::NotThreadLocal, AddressSpace); - const DataLayout &DL = M.getDataLayout(); + GlobalValue::NotThreadLocal, AddressSpaceVal); const llvm::Align TypeAlign = DL.getABITypeAlign(Ty); - const llvm::Align PtrAlign = DL.getPointerABIAlignment(AddressSpace); + const llvm::Align PtrAlign = DL.getPointerABIAlignment(AddressSpaceVal); GV->setAlignment(std::max(TypeAlign, PtrAlign)); Elem.second = GV; }