diff --git a/clang/lib/CodeGen/CGDecl.cpp b/clang/lib/CodeGen/CGDecl.cpp --- a/clang/lib/CodeGen/CGDecl.cpp +++ b/clang/lib/CodeGen/CGDecl.cpp @@ -2468,33 +2468,7 @@ (IPD->getParameterKind() == ImplicitParamDecl::ThreadPrivateVar); } - Address DeclPtr = Address::invalid(); - Address AllocaPtr = Address::invalid(); - bool DoStore = false; - bool IsScalar = hasScalarEvaluationKind(Ty); - // If we already have a pointer to the argument, reuse the input pointer. - if (Arg.isIndirect()) { - // If we have a prettier pointer type at this point, bitcast to that. - DeclPtr = Arg.getIndirectAddress(); - DeclPtr = Builder.CreateElementBitCast(DeclPtr, ConvertTypeForMem(Ty), - D.getName()); - // Indirect argument is in alloca address space, which may be different - // from the default address space. - auto AllocaAS = CGM.getASTAllocaAddressSpace(); - auto *V = DeclPtr.getPointer(); - AllocaPtr = DeclPtr; - auto SrcLangAS = getLangOpts().OpenCL ? LangAS::opencl_private : AllocaAS; - auto DestLangAS = - getLangOpts().OpenCL ? LangAS::opencl_private : LangAS::Default; - if (SrcLangAS != DestLangAS) { - assert(getContext().getTargetAddressSpace(SrcLangAS) == - CGM.getDataLayout().getAllocaAddrSpace()); - auto DestAS = getContext().getTargetAddressSpace(DestLangAS); - auto *T = DeclPtr.getElementType()->getPointerTo(DestAS); - DeclPtr = DeclPtr.withPointer(getTargetHooks().performAddrSpaceCast( - *this, V, SrcLangAS, DestLangAS, T, true)); - } - + auto PushCleanupIfNeeded = [this, Ty, &D](Address DeclPtr) { // Push a destructor cleanup for this parameter if the ABI requires it. // Don't push a cleanup in a thunk for a method that will also emit a // cleanup. @@ -2510,87 +2484,123 @@ EHStack.stable_begin(); } } - } else { - // Check if the parameter address is controlled by OpenMP runtime. - Address OpenMPLocalAddr = - getLangOpts().OpenMP - ? CGM.getOpenMPRuntime().getAddressOfLocalVariable(*this, &D) - : Address::invalid(); - if (getLangOpts().OpenMP && OpenMPLocalAddr.isValid()) { - DeclPtr = OpenMPLocalAddr; - AllocaPtr = DeclPtr; + }; + + Address DeclPtr = Address::invalid(); + Address AllocaPtr = Address::invalid(); + Address OpenMPLocalAddr = + getLangOpts().OpenMP + ? CGM.getOpenMPRuntime().getAddressOfLocalVariable(*this, &D) + : Address::invalid(); + if (OpenMPLocalAddr.isValid()) { + DeclPtr = OpenMPLocalAddr; + AllocaPtr = DeclPtr; + LValue Dst = MakeAddrLValue(DeclPtr, Ty); + if (Arg.isIndirect()) { + LValue Src = MakeAddrLValue(Arg.getIndirectAddress(), Ty); + callCStructCopyConstructor(Dst, Src); + PushCleanupIfNeeded(Arg.getIndirectAddress()); } else { - // Otherwise, create a temporary to hold the value. + EmitStoreOfScalar(Arg.getDirectValue(), Dst, /* isInitialization */ true); + } + } else { + bool DoStore = false; + bool IsScalar = hasScalarEvaluationKind(Ty); + // If we already have a pointer to the argument, reuse the input pointer. + if (Arg.isIndirect()) { + // If we have a prettier pointer type at this point, bitcast to that. + DeclPtr = Arg.getIndirectAddress(); + DeclPtr = Builder.CreateElementBitCast(DeclPtr, ConvertTypeForMem(Ty), + D.getName()); + // Indirect argument is in alloca address space, which may be different + // from the default address space. + auto AllocaAS = CGM.getASTAllocaAddressSpace(); + auto *V = DeclPtr.getPointer(); + AllocaPtr = DeclPtr; + auto SrcLangAS = getLangOpts().OpenCL ? LangAS::opencl_private : AllocaAS; + auto DestLangAS = + getLangOpts().OpenCL ? LangAS::opencl_private : LangAS::Default; + if (SrcLangAS != DestLangAS) { + assert(getContext().getTargetAddressSpace(SrcLangAS) == + CGM.getDataLayout().getAllocaAddrSpace()); + auto DestAS = getContext().getTargetAddressSpace(DestLangAS); + auto *T = DeclPtr.getElementType()->getPointerTo(DestAS); + DeclPtr = DeclPtr.withPointer(getTargetHooks().performAddrSpaceCast( + *this, V, SrcLangAS, DestLangAS, T, true)); + } + PushCleanupIfNeeded(DeclPtr); + } else { + // Create a temporary to hold the value. DeclPtr = CreateMemTemp(Ty, getContext().getDeclAlign(&D), D.getName() + ".addr", &AllocaPtr); + DoStore = true; } - DoStore = true; - } - llvm::Value *ArgVal = (DoStore ? Arg.getDirectValue() : nullptr); + llvm::Value *ArgVal = (DoStore ? Arg.getDirectValue() : nullptr); - LValue lv = MakeAddrLValue(DeclPtr, Ty); - if (IsScalar) { - Qualifiers qs = Ty.getQualifiers(); - if (Qualifiers::ObjCLifetime lt = qs.getObjCLifetime()) { - // We honor __attribute__((ns_consumed)) for types with lifetime. - // For __strong, it's handled by just skipping the initial retain; - // otherwise we have to balance out the initial +1 with an extra - // cleanup to do the release at the end of the function. - bool isConsumed = D.hasAttr(); + LValue lv = MakeAddrLValue(DeclPtr, Ty); + if (IsScalar) { + Qualifiers qs = Ty.getQualifiers(); + if (Qualifiers::ObjCLifetime lt = qs.getObjCLifetime()) { + // We honor __attribute__((ns_consumed)) for types with lifetime. + // For __strong, it's handled by just skipping the initial retain; + // otherwise we have to balance out the initial +1 with an extra + // cleanup to do the release at the end of the function. + bool isConsumed = D.hasAttr(); - // If a parameter is pseudo-strong then we can omit the implicit retain. - if (D.isARCPseudoStrong()) { - assert(lt == Qualifiers::OCL_Strong && - "pseudo-strong variable isn't strong?"); - assert(qs.hasConst() && "pseudo-strong variable should be const!"); - lt = Qualifiers::OCL_ExplicitNone; - } + // If a parameter is pseudo-strong then we can omit the implicit retain. + if (D.isARCPseudoStrong()) { + assert(lt == Qualifiers::OCL_Strong && + "pseudo-strong variable isn't strong?"); + assert(qs.hasConst() && "pseudo-strong variable should be const!"); + lt = Qualifiers::OCL_ExplicitNone; + } - // Load objects passed indirectly. - if (Arg.isIndirect() && !ArgVal) - ArgVal = Builder.CreateLoad(DeclPtr); + // Load objects passed indirectly. + if (Arg.isIndirect() && !ArgVal) + ArgVal = Builder.CreateLoad(DeclPtr); - if (lt == Qualifiers::OCL_Strong) { - if (!isConsumed) { - if (CGM.getCodeGenOpts().OptimizationLevel == 0) { - // use objc_storeStrong(&dest, value) for retaining the - // object. But first, store a null into 'dest' because - // objc_storeStrong attempts to release its old value. - llvm::Value *Null = CGM.EmitNullConstant(D.getType()); - EmitStoreOfScalar(Null, lv, /* isInitialization */ true); - EmitARCStoreStrongCall(lv.getAddress(*this), ArgVal, true); - DoStore = false; + if (lt == Qualifiers::OCL_Strong) { + if (!isConsumed) { + if (CGM.getCodeGenOpts().OptimizationLevel == 0) { + // use objc_storeStrong(&dest, value) for retaining the + // object. But first, store a null into 'dest' because + // objc_storeStrong attempts to release its old value. + llvm::Value *Null = CGM.EmitNullConstant(D.getType()); + EmitStoreOfScalar(Null, lv, /* isInitialization */ true); + EmitARCStoreStrongCall(lv.getAddress(*this), ArgVal, true); + DoStore = false; + } else + // Don't use objc_retainBlock for block pointers, because we + // don't want to Block_copy something just because we got it + // as a parameter. + ArgVal = EmitARCRetainNonBlock(ArgVal); + } + } else { + // Push the cleanup for a consumed parameter. + if (isConsumed) { + ARCPreciseLifetime_t precise = + (D.hasAttr() ? ARCPreciseLifetime + : ARCImpreciseLifetime); + EHStack.pushCleanup(getARCCleanupKind(), + ArgVal, precise); + } + + if (lt == Qualifiers::OCL_Weak) { + EmitARCInitWeak(DeclPtr, ArgVal); + DoStore = false; // The weak init is a store, no need to do two. } - else - // Don't use objc_retainBlock for block pointers, because we - // don't want to Block_copy something just because we got it - // as a parameter. - ArgVal = EmitARCRetainNonBlock(ArgVal); - } - } else { - // Push the cleanup for a consumed parameter. - if (isConsumed) { - ARCPreciseLifetime_t precise = (D.hasAttr() - ? ARCPreciseLifetime : ARCImpreciseLifetime); - EHStack.pushCleanup(getARCCleanupKind(), ArgVal, - precise); } - if (lt == Qualifiers::OCL_Weak) { - EmitARCInitWeak(DeclPtr, ArgVal); - DoStore = false; // The weak init is a store, no need to do two. - } + // Enter the cleanup scope. + EmitAutoVarWithLifetime(*this, D, DeclPtr, lt); } - - // Enter the cleanup scope. - EmitAutoVarWithLifetime(*this, D, DeclPtr, lt); } - } - // Store the initial value into the alloca. - if (DoStore) - EmitStoreOfScalar(ArgVal, lv, /* isInitialization */ true); + // Store the initial value into the alloca. + if (DoStore) + EmitStoreOfScalar(ArgVal, lv, /* isInitialization */ true); + } setAddrOfLocalVar(&D, DeclPtr); diff --git a/clang/test/OpenMP/globalization_byval_struct.c b/clang/test/OpenMP/globalization_byval_struct.c new file mode 100644 --- /dev/null +++ b/clang/test/OpenMP/globalization_byval_struct.c @@ -0,0 +1,33 @@ +// RUN: %clang_cc1 -no-opaque-pointers -verify -fopenmp -x c -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc +// RUN: %clang_cc1 -no-opaque-pointers -verify -fopenmp -x c -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s +// expected-no-diagnostics + +extern int printf(const char *, ...); + +struct S { + int a; + float b; +}; + +// CHECK: define{{.*}}void @test(%struct.S* noundef byval(%struct.S) align {{[0-9]+}} [[arg:%[0-9a-zA-Z]+]]) +// CHECK: [[g:%[0-9a-zA-Z]+]] = call align {{[0-9]+}} i8* @__kmpc_alloc_shared +// CHECK: bitcast i8* [[g]] to %struct.S* +// CHECK: bitcast %struct.S* [[arg]] to i8** +// CHECK: call void [[cc:@__copy_constructor[_0-9a-zA-Z]+]] +// CHECK: void [[cc]] +void test(struct S s) { +#pragma omp parallel for + for (int i = 0; i < s.a; ++i) { + printf("%i : %i : %f\n", i, s.a, s.b); + } +} + +void foo() { + #pragma omp target teams num_teams(1) + { + struct S s; + s.a = 7; + s.b = 11; + test(s); + } +}