diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -1928,6 +1928,8 @@
       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);
@@ -1972,6 +1974,8 @@
       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);
diff --git a/clang/test/OpenMP/amdgcn_target_global_constructor.cpp b/clang/test/OpenMP/amdgcn_target_global_constructor.cpp
--- a/clang/test/OpenMP/amdgcn_target_global_constructor.cpp
+++ b/clang/test/OpenMP/amdgcn_target_global_constructor.cpp
@@ -27,7 +27,7 @@
 // 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__fd02_6139fe30_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]+]]
@@ -45,7 +45,7 @@
 // CHECK-NEXT:    ret void
 //
 //
-// CHECK-LABEL: define {{[^@]+}}@__omp_offloading__fd02_613a0d56_A_l19_dtor
+// CHECK-LABEL: define {{[^@]+}}@__omp_offloading__fd02_6139fe30_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]+]]
@@ -92,11 +92,11 @@
 // 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 1631190576, !"__omp_offloading__fd02_6139fe30_A_l19_ctor", i32 19, i32 1}
+// CHECK: !1 = !{i32 0, i32 64770, i32 1631190576, !"__omp_offloading__fd02_6139fe30_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 ()* @__omp_offloading__fd02_6139fe30_A_l19_ctor, !"kernel", i32 1}
+// CHECK: !4 = !{void ()* @__omp_offloading__fd02_6139fe30_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}
diff --git a/openmp/libomptarget/test/offloading/global_constructor.cpp b/openmp/libomptarget/test/offloading/global_constructor.cpp
--- a/openmp/libomptarget/test/offloading/global_constructor.cpp
+++ b/openmp/libomptarget/test/offloading/global_constructor.cpp
@@ -1,23 +1,25 @@
 // 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");
 }