[OpenMP] Add AMDGPU calling convention to ctor / dtor functions

[jhuber6]

Summary

This patch adds the necessary AMDGPU calling convention to the ctor /
dtor kernels. These are fundamentally device kenels called by the host
on image load. Without this calling convention information the AMDGPU
plugin is unable to identify them.

Depends on D122504

Fixes #54091

[jhuber6]

Making test better.

[jdoerfert]

LG

[jhuber6]

Remove leftover code.

[jhuber6]

Fix test

clang/lib/CodeGen/CGOpenMPRuntime.cpp

if (CGM.getLangOpts().OpenMPIsDevice) {
       // the threadprivate copy of the variable VD
       CodeGenFunction CtorCGF(CGM);
 
       const CGFunctionInfo &FI = CGM.getTypes().arrangeNullaryFunction();
       llvm::FunctionType *FTy = CGM.getTypes().GetFunctionType(FI);
       llvm::Function *Fn = CGM.CreateGlobalInitOrCleanUpFunction(
           FTy, Twine(Buffer, "_ctor"), FI, Loc, false,
           llvm::GlobalValue::WeakODRLinkage);
+      if (CGM.getTriple().isAMDGCN())
+        Fn->setCallingConv(llvm::CallingConv::AMDGPU_KERNEL);
       auto NL = ApplyDebugLocation::CreateEmpty(CtorCGF);
       CtorCGF.StartFunction(GlobalDecl(), CGM.getContext().VoidTy, Fn, FI,
                             FunctionArgList(), Loc, Loc);
       auto AL = ApplyDebugLocation::CreateArtificial(CtorCGF);
       llvm::Constant *AddrInAS0 = Addr;
       if (Addr->getAddressSpace() != 0)
         AddrInAS0 = llvm::ConstantExpr::getAddrSpaceCast(
             Addr, llvm::PointerType::getWithSamePointeeType(
if (CGM.getLangOpts().OpenMPIsDevice) {
       // copy of the variable VD
       CodeGenFunction DtorCGF(CGM);
 
       const CGFunctionInfo &FI = CGM.getTypes().arrangeNullaryFunction();
       llvm::FunctionType *FTy = CGM.getTypes().GetFunctionType(FI);
       llvm::Function *Fn = CGM.CreateGlobalInitOrCleanUpFunction(
           FTy, Twine(Buffer, "_dtor"), FI, Loc, false,
           llvm::GlobalValue::WeakODRLinkage);
+      if (CGM.getTriple().isAMDGCN())
+        Fn->setCallingConv(llvm::CallingConv::AMDGPU_KERNEL);
       auto NL = ApplyDebugLocation::CreateEmpty(DtorCGF);
       DtorCGF.StartFunction(GlobalDecl(), CGM.getContext().VoidTy, Fn, FI,
                             FunctionArgList(), Loc, Loc);
       // Create a scope with an artificial location for the body of this
       // function.
       auto AL = ApplyDebugLocation::CreateArtificial(DtorCGF);
       llvm::Constant *AddrInAS0 = Addr;
       if (Addr->getAddressSpace() != 0)

clang/test/OpenMP/amdgcn_target_global_constructor.cpp

-// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --check-globals --include-generated-funcs
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --check-globals --include-generated-funcs --replace-value-regex "__omp_offloading__[0-9a-z]_[0-9a-z]_"
 // REQUIRES: amdgpu-registered-target
 
 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host.bc
 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s
 // expected-no-diagnostics
 #ifndef HEADER
 #define HEADER
 

 #endif
 //.
 // CHECK: @__omp_rtl_debug_kind = weak_odr hidden addrspace(1) constant i32 0
 // CHECK: @__omp_rtl_assume_teams_oversubscription = weak_odr hidden addrspace(1) constant i32 0
 // CHECK: @__omp_rtl_assume_threads_oversubscription = weak_odr hidden addrspace(1) constant i32 0
 // CHECK: @__omp_rtl_assume_no_thread_state = weak_odr hidden addrspace(1) constant i32 0
 // CHECK: @A = addrspace(1) global %struct.S zeroinitializer, align 4
 //.
-// CHECK-LABEL: define {{[^@]+}}@__omp_offloading__fd02_613a0d56_A_l19_ctor
+// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading__[0-9a-z]+_[0-9a-z]+}}_A_l19_ctor
 // CHECK-SAME: () #[[ATTR0:[0-9]+]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    call void @_ZN1SC1Ev(%struct.S* noundef nonnull align 4 dereferenceable(4) addrspacecast ([[STRUCT_S:%.*]] addrspace(1)* @A to %struct.S*)) #[[ATTR3:[0-9]+]]
 // CHECK-NEXT:    ret void
 //
 //
 // CHECK-LABEL: define {{[^@]+}}@_ZN1SC1Ev
 // CHECK-SAME: (%struct.S* noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1:[0-9]+]] comdat align 2 {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[THIS_ADDR:%.*]] = alloca %struct.S*, align 8, addrspace(5)
 // CHECK-NEXT:    [[THIS_ADDR_ASCAST:%.*]] = addrspacecast %struct.S* addrspace(5)* [[THIS_ADDR]] to %struct.S**
 // CHECK-NEXT:    store %struct.S* [[THIS]], %struct.S** [[THIS_ADDR_ASCAST]], align 8
 // CHECK-NEXT:    [[THIS1:%.*]] = load %struct.S*, %struct.S** [[THIS_ADDR_ASCAST]], align 8
 // CHECK-NEXT:    call void @_ZN1SC2Ev(%struct.S* noundef nonnull align 4 dereferenceable(4) [[THIS1]]) #[[ATTR3]]
 // CHECK-NEXT:    ret void
 //
 //
-// CHECK-LABEL: define {{[^@]+}}@__omp_offloading__fd02_613a0d56_A_l19_dtor
+// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading__[0-9a-z]+_[0-9a-z]+}}_A_l19_dtor
 // CHECK-SAME: () #[[ATTR0]] {
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    call void @_ZN1SD1Ev(%struct.S* noundef nonnull align 4 dereferenceable(4) addrspacecast ([[STRUCT_S:%.*]] addrspace(1)* @A to %struct.S*)) #[[ATTR4:[0-9]+]]
 // CHECK-NEXT:    ret void
 //
 //
 // CHECK-LABEL: define {{[^@]+}}@_ZN1SD1Ev
 // CHECK-SAME: (%struct.S* noundef nonnull align 4 dereferenceable(4) [[THIS:%.*]]) unnamed_addr #[[ATTR1]] comdat align 2 {

 //
 //.
 // CHECK: attributes #0 = { nounwind "kernel" "min-legal-vector-width"="0" }
 // CHECK: attributes #1 = { convergent noinline nounwind optnone "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
 // CHECK: attributes #2 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
 // CHECK: attributes #3 = { convergent }
 // CHECK: attributes #4 = { convergent nounwind }
 //.
-// CHECK: !0 = !{i32 0, i32 64770, i32 1631194454, !"__omp_offloading__fd02_613a0d56_A_l19_ctor", i32 19, i32 1}
-// CHECK: !1 = !{i32 0, i32 64770, i32 1631194454, !"__omp_offloading__fd02_613a0d56_A_l19_dtor", i32 19, i32 2}
+// CHECK: !0 = !{i32 0, i32 64770, i32 1680388141, !"{{__omp_offloading__[0-9a-z]+_[0-9a-z]+}}_A_l19_ctor", i32 19, i32 1}
+// CHECK: !1 = !{i32 0, i32 64770, i32 1680388141, !"{{__omp_offloading__[0-9a-z]+_[0-9a-z]+}}_A_l19_dtor", i32 19, i32 2}
 // CHECK: !2 = !{i32 1, !"A", i32 0, i32 0}
-// CHECK: !3 = !{void ()* @__omp_offloading__fd02_613a0d56_A_l19_ctor, !"kernel", i32 1}
-// CHECK: !4 = !{void ()* @__omp_offloading__fd02_613a0d56_A_l19_dtor, !"kernel", i32 1}
+// CHECK: !3 = !{void ()* @{{.*}}_A_l19_ctor, !"kernel", i32 1}
+// CHECK: !4 = !{void ()* @{{.*}}_A_l19_dtor, !"kernel", i32 1}
 // CHECK: !5 = !{i32 1, !"wchar_size", i32 4}
 // CHECK: !6 = !{i32 7, !"openmp", i32 50}
 // CHECK: !7 = !{i32 7, !"openmp-device", i32 50}
 // CHECK: !8 = !{!"clang version 15.0.0"}
 //.

openmp/libomptarget/test/offloading/global_constructor.cpp

 // RUN: %libomptarget-compilexx-generic && %libomptarget-run-generic | %fcheck-generic
 
-// Fails in DAGToDAG on an address space problem
-// UNSUPPORTED: amdgcn-amd-amdhsa
-// UNSUPPORTED: amdgcn-amd-amdhsa-newDriver
-
-#include <cmath>
 #include <cstdio>
 
-const double Host = log(2.0) / log(2.0);
-#pragma omp declare target
-const double Device = log(2.0) / log(2.0);
-#pragma omp end declare target
+int foo() { return 1; }
+
+class C {
+public:
+  C() : x(foo()) {}
+
+  int x;
+};
+
+C c;
+#pragma omp declare target(c)
 
 int main() {
-  double X;
-#pragma omp target map(from : X)
-  { X = Device; }
+  int x = 0;
+#pragma omp target map(from : x)
+  { x = c.x; }
 
   // CHECK: PASS
-  if (X == Host)
+  if (x == 1)
     printf("PASS\n");
 }