jhuber6 updated this revision to Diff 418354.
jhuber6 added a comment.
Making test better.
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D122515/new/
https://reviews.llvm.org/D122515
Files:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/test/OpenMP/amdgcn_target_global_constructor.cpp
openmp/libomptarget/test/offloading/global_constructor.cpp
Index: openmp/libomptarget/test/offloading/global_constructor.cpp
===================================================================
--- openmp/libomptarget/test/offloading/global_constructor.cpp
+++ 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");
}
Index: clang/test/OpenMP/amdgcn_target_global_constructor.cpp
===================================================================
--- clang/test/OpenMP/amdgcn_target_global_constructor.cpp
+++ 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}
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ 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);
@@ -1990,6 +1994,8 @@
DtorCGF.FinishFunction();
Dtor = Fn;
ID = llvm::ConstantExpr::getBitCast(Fn, CGM.Int8PtrTy);
+ if (CGM.getTriple().isAMDGCN())
+ Fn->setCallingConv(llvm::CallingConv::AMDGPU_KERNEL);
} else {
Dtor = new llvm::GlobalVariable(
CGM.getModule(), CGM.Int8Ty, /*isConstant=*/true,
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
