[Clang][OpenMP] Fix the issue that globalization doesn't work with byval struct function argument

[tianshilei1992]

Summary

This patch fixes the issue that the globalized variable is not properly
initialized when it is a byval struct function argument.

Fixes #56218.

[tianshilei1992]

fix unused variable

[tianshilei1992]

rebase and fix test error

[jdoerfert]

cleanup/destructor test missing.

@ABataev WDYT?

[ABataev]

In D129008#3634073, @jdoerfert wrote:

cleanup/destructor test missing.

@ABataev WDYT?

Looks good in general, extra tests would be good

[tianshilei1992]

callCStructCopyConstructor is actually for Objective-C…Cannot use it here.

[jdoerfert]

In D129008#3640194, @tianshilei1992 wrote:

callCStructCopyConstructor is actually for Objective-C…Cannot use it here.

Don't we generate copies of things elsewhere already?

[tianshilei1992]

In D129008#3640233, @jdoerfert wrote:

In D129008#3640194, @tianshilei1992 wrote:

callCStructCopyConstructor is actually for Objective-C…Cannot use it here.

Don't we generate copies of things elsewhere already?

No we don't. I think the best place to emit copy is where the globalized variable is generated.

[tianshilei1992]

clang/lib/CodeGen/CGDecl.cpp

2501

Hi @aaron.ballman, is it possible to invoke the copy constructor of a struct/class here, either it is user defined or default one?

[aaron.ballman]

In D129008#3640194, @tianshilei1992 wrote:

callCStructCopyConstructor is actually for Objective-C…Cannot use it here.

Is it for ObjC? Looking at the comments, it looks like it's for C:

// These functions emit calls to the special functions of non-trivial C
// structs.

clang/lib/CodeGen/CGDecl.cpp

2501

CC @rjmccall and @efriedma for codegen owner opinions.

I'm honestly not certain; I would imagine that OpenMP defines what should happen, but this is outside of my area of expertise.

[tianshilei1992]

In D129008#4079660, @aaron.ballman wrote:

In D129008#3640194, @tianshilei1992 wrote:

callCStructCopyConstructor is actually for Objective-C…Cannot use it here.

Is it for ObjC? Looking at the comments, it looks like it's for C:

// These functions emit calls to the special functions of non-trivial C
// structs.

That's not ObjC. It can be C or C++.

My apology I should have added more context here. In OpenMP, we need to "globalize" certain captured local variables by allocating another buffers and then copy the memory. If the local variables are structs or even C++ classes, the copy has to be done by invoking the corresponding copy constructors. However, I don't know how to invoke the copy constructor here. callCStructCopyConstructor only does plain copy. It's not gonna work for C++ classes.

[ABataev]

In D129008#4079872, @tianshilei1992 wrote:

In D129008#4079660, @aaron.ballman wrote:

In D129008#3640194, @tianshilei1992 wrote:

callCStructCopyConstructor is actually for Objective-C…Cannot use it here.

Is it for ObjC? Looking at the comments, it looks like it's for C:

// These functions emit calls to the special functions of non-trivial C
// structs.

That's not ObjC. It can be C or C++.

My apology I should have added more context here. In OpenMP, we need to "globalize" certain captured local variables by allocating another buffers and then copy the memory. If the local variables are structs or even C++ classes, the copy has to be done by invoking the corresponding copy constructors. However, I don't know how to invoke the copy constructor here. callCStructCopyConstructor only does plain copy. It's not gonna work for C++ classes.

You need to do it manually. Usually we built the required copy constructor call in Sema (it requires correct lookup) and then emitted it in codegen.

[tianshilei1992]

In D129008#4079892, @ABataev wrote:

In D129008#4079872, @tianshilei1992 wrote:

In D129008#4079660, @aaron.ballman wrote:

In D129008#3640194, @tianshilei1992 wrote:

callCStructCopyConstructor is actually for Objective-C…Cannot use it here.

Is it for ObjC? Looking at the comments, it looks like it's for C:

// These functions emit calls to the special functions of non-trivial C
// structs.

That's not ObjC. It can be C or C++.

My apology I should have added more context here. In OpenMP, we need to "globalize" certain captured local variables by allocating another buffers and then copy the memory. If the local variables are structs or even C++ classes, the copy has to be done by invoking the corresponding copy constructors. However, I don't know how to invoke the copy constructor here. callCStructCopyConstructor only does plain copy. It's not gonna work for C++ classes.

You need to do it manually. Usually we built the required copy constructor call in Sema (it requires correct lookup) and then emitted it in codegen.

But we don't do globalization in Sema right?

[ABataev]

In D129008#4079914, @tianshilei1992 wrote:

In D129008#4079892, @ABataev wrote:

In D129008#4079872, @tianshilei1992 wrote:

In D129008#4079660, @aaron.ballman wrote:

In D129008#3640194, @tianshilei1992 wrote:

callCStructCopyConstructor is actually for Objective-C…Cannot use it here.

Is it for ObjC? Looking at the comments, it looks like it's for C:

// These functions emit calls to the special functions of non-trivial C
// structs.

That's not ObjC. It can be C or C++.

My apology I should have added more context here. In OpenMP, we need to "globalize" certain captured local variables by allocating another buffers and then copy the memory. If the local variables are structs or even C++ classes, the copy has to be done by invoking the corresponding copy constructors. However, I don't know how to invoke the copy constructor here. callCStructCopyConstructor only does plain copy. It's not gonna work for C++ classes.

You need to do it manually. Usually we built the required copy constructor call in Sema (it requires correct lookup) and then emitted it in codegen.

But we don't do globalization in Sema right?

Right. If we really needed the constructor call, we emitted special helper expressions/statements and stored them in the associated construct/clause, and then emitted it, if required.

[rjmccall]

Right. C structs can require more than just memcpy to copy in several different situations (all involving language extensions), but it doesn't require Sema to be involved: it doesn't introduce uses of user declarations, and the compiler can figure out the work it needs to do with a straightforward recursive inspection of the field types, so IRGen just synthesizes those operations automatically when requested. Neither condition holds in C++, so Sema has to synthesize a copy expression which resolves the correct copy constructor and sets up any extra arguments it might require, triggering appropriate diagnostics and/or tracking of used decls. There are a bunch of places we do that for various language extensions already, including a bunch in the OpenMP code, and we usually just store the expression in the associated AST node.

[tianshilei1992]

@rjmccall @ABataev Thanks for the help. So the basic idea is to build the expressions in Sema for those captured decls and then emit them accordingly in code gen, right?

There are a bunch of places we do that for various language extensions already, including a bunch in the OpenMP code, and we usually just store the expression in the associated AST node.

Could you please point me a rough location where we already have similar things such that I can learn and do the same thing for globalization? Thanks.

[ABataev]

In D129008#4083150, @tianshilei1992 wrote:

@rjmccall @ABataev Thanks for the help. So the basic idea is to build the expressions in Sema for those captured decls and then emit them accordingly in code gen, right?

Yes.

There are a bunch of places we do that for various language extensions already, including a bunch in the OpenMP code, and we usually just store the expression in the associated AST node.

Could you please point me a rough location where we already have similar things such that I can learn and do the same thing for globalization? Thanks.

You can check ActOnOpenMPFirstprivateClause, lines 18322-18353 (building of VDInitRefExpr)

[tianshilei1992]

In D129008#4083185, @ABataev wrote:

In D129008#4083150, @tianshilei1992 wrote:

There are a bunch of places we do that for various language extensions already, including a bunch in the OpenMP code, and we usually just store the expression in the associated AST node.

Could you please point me a rough location where we already have similar things such that I can learn and do the same thing for globalization? Thanks.

You can check ActOnOpenMPFirstprivateClause, lines 18322-18353 (building of VDInitRefExpr)

Really appreciate that.

clang/lib/CodeGen/CGDecl.cpp

if (BlockInfo) {
       return;
     }
     // Suppressing debug info for ThreadPrivateVar parameters, else it hides
     // debug info of TLS variables.
     NoDebugInfo =
         (IPD->getParameterKind() == ImplicitParamDecl::ThreadPrivateVar);
   }
 
+  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.
+    if (Ty->isRecordType() && !CurFuncIsThunk &&
+        Ty->castAs<RecordType>()->getDecl()->isParamDestroyedInCallee()) {
+      if (QualType::DestructionKind DtorKind =
+              D.needsDestruction(getContext())) {
+        assert((DtorKind == QualType::DK_cxx_destructor ||
+                DtorKind == QualType::DK_nontrivial_c_struct) &&
+               "unexpected destructor type");
+        pushDestroy(DtorKind, DeclPtr, Ty);
+        CalleeDestructedParamCleanups[cast<ParmVarDecl>(&D)] =
+            EHStack.stable_begin();
+      }
+    }
+  };
+
   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 {
+      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));
       }
-
-    // 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.
-    if (Ty->isRecordType() && !CurFuncIsThunk &&
-        Ty->castAs<RecordType>()->getDecl()->isParamDestroyedInCallee()) {
-      if (QualType::DestructionKind DtorKind =
-              D.needsDestruction(getContext())) {
-        assert((DtorKind == QualType::DK_cxx_destructor ||
-                DtorKind == QualType::DK_nontrivial_c_struct) &&
-               "unexpected destructor type");
-        pushDestroy(DtorKind, DeclPtr, Ty);
-        CalleeDestructedParamCleanups[cast<ParmVarDecl>(&D)] =
-            EHStack.stable_begin();
-      }
-    }
+      PushCleanupIfNeeded(DeclPtr);
     } 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;
-    } else {
-      // Otherwise, create a temporary to hold the value.
+      // Create a temporary to hold the value.
       DeclPtr = CreateMemTemp(Ty, getContext().getDeclAlign(&D),
                               D.getName() + ".addr", &AllocaPtr);
-    }
       DoStore = true;
     }
 
     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>();
 
         // 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);
 
         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
+            } 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);
+            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);
       }
     }
 
     // Store the initial value into the alloca.
     if (DoStore)
       EmitStoreOfScalar(ArgVal, lv, /* isInitialization */ true);
+  }
 
   setAddrOfLocalVar(&D, DeclPtr);
 
   // Emit debug info for param declarations in non-thunk functions.
   if (CGDebugInfo *DI = getDebugInfo()) {
     if (CGM.getCodeGenOpts().hasReducedDebugInfo() && !CurFuncIsThunk &&
         !NoDebugInfo) {
       llvm::DILocalVariable *DILocalVar = DI->EmitDeclareOfArgVariable(

clang/test/OpenMP/globalization_byval_struct.c

+// 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);
+  }
+}