tianshilei1992 updated this revision to Diff 441908.
tianshilei1992 added a comment.
rebase and fix test error
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D129008/new/
https://reviews.llvm.org/D129008
Files:
clang/lib/CodeGen/CGDecl.cpp
clang/test/OpenMP/globalization_byval_struct.c
Index: clang/test/OpenMP/globalization_byval_struct.c
===================================================================
--- /dev/null
+++ 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);
+ }
+}
Index: clang/lib/CodeGen/CGDecl.cpp
===================================================================
--- clang/lib/CodeGen/CGDecl.cpp
+++ 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<NSConsumedAttr>();
+ 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<NSConsumedAttr>();
- // 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<ObjCPreciseLifetimeAttr>() ? ARCPreciseLifetime
+ : ARCImpreciseLifetime);
+ EHStack.pushCleanup<ConsumeARCParameter>(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<ObjCPreciseLifetimeAttr>()
- ? ARCPreciseLifetime : ARCImpreciseLifetime);
- EHStack.pushCleanup<ConsumeARCParameter>(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);
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
