doru1004 updated this revision to Diff 507190.
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D146552/new/
https://reviews.llvm.org/D146552
Files:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/test/OpenMP/declare_target_constexpr_codegen.cpp
openmp/libomptarget/test/offloading/target_constexpr_mapping.cpp
Index: openmp/libomptarget/test/offloading/target_constexpr_mapping.cpp
===================================================================
--- /dev/null
+++ openmp/libomptarget/test/offloading/target_constexpr_mapping.cpp
@@ -0,0 +1,34 @@
+// RUN: %libomptarget-compileoptxx-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+#pragma omp declare target
+class A {
+public:
+ constexpr static double pi = 3.141592653589793116;
+ A() { ; }
+ ~A() { ; }
+};
+#pragma omp end declare target
+
+#pragma omp declare target
+constexpr static double anotherPi = 3.14;
+#pragma omp end declare target
+
+int main() {
+ double a[2];
+#pragma omp target map(tofrom : a[:2])
+ {
+ a[0] = A::pi;
+ a[1] = anotherPi;
+ }
+
+ // CHECK: pi = 3.141592653589793116
+ printf("pi = %.18f\n", a[0]);
+
+ // CHECK: anotherPi = 3.14
+ printf("anotherPi = %.2f\n", a[1]);
+
+ return 0;
+}
Index: clang/test/OpenMP/declare_target_constexpr_codegen.cpp
===================================================================
--- /dev/null
+++ clang/test/OpenMP/declare_target_constexpr_codegen.cpp
@@ -0,0 +1,79 @@
+// 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]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _
+// REQUIRES: amdgpu-registered-target
+
+
+// RUN: %clang_cc1 -verify -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 -verify -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 - -disable-llvm-optzns | FileCheck %s --check-prefix=CHECK1
+// expected-no-diagnostics
+
+#pragma omp declare target
+class A {
+public:
+ constexpr static double pi = 3.141592653589793116;
+ A() { ; }
+ ~A() { ; }
+};
+#pragma omp end declare target
+
+#pragma omp declare target
+constexpr static double anotherPi = 3.14;
+#pragma omp end declare target
+
+int main() {
+ double a[2];
+#pragma omp target map(tofrom : a[:2])
+ {
+ a[0] = A::pi;
+ a[1] = anotherPi;
+ }
+
+ return (int)(a[0] + a[1]);
+}
+
+//.
+// CHECK1: @__omp_rtl_debug_kind = weak_odr hidden addrspace(1) constant i32 0
+// CHECK1: @__omp_rtl_assume_teams_oversubscription = weak_odr hidden addrspace(1) constant i32 0
+// CHECK1: @__omp_rtl_assume_threads_oversubscription = weak_odr hidden addrspace(1) constant i32 0
+// CHECK1: @__omp_rtl_assume_no_thread_state = weak_odr hidden addrspace(1) constant i32 0
+// CHECK1: @__omp_rtl_assume_no_nested_parallelism = weak_odr hidden addrspace(1) constant i32 0
+// CHECK1: @_ZN1A2piE = available_externally addrspace(4) constant double 0x400921FB54442D18, align 8
+// CHECK1: @0 = private unnamed_addr constant [23 x i8] c"
+// CHECK1: @1 = private unnamed_addr addrspace(1) constant %struct.ident_t { i32 0, i32 2, i32 0, i32 22, ptr @0 }, align 8
+// CHECK1: @__omp_offloading_fd00_240171e_main_l24_exec_mode = weak protected addrspace(1) constant i8 1
+// CHECK1: @_ZL9anotherPi = internal addrspace(4) constant double 3.140000e+00, align 8
+// CHECK1: @"__ZL9anotherPi$ref" = internal constant ptr addrspace(4) @_ZL9anotherPi
+// CHECK1: @llvm.compiler.used = appending addrspace(1) global [2 x ptr] [ptr addrspacecast (ptr addrspace(1) @__omp_offloading_fd00_240171e_main_l24_exec_mode to ptr), ptr @"__ZL9anotherPi$ref"], section "llvm.metadata"
+//.
+// CHECK1-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_main_l24
+// CHECK1-SAME: (ptr noundef nonnull align 8 dereferenceable(16) [[A:%.*]]) #[[ATTR0:[0-9]+]] {
+// CHECK1-NEXT: entry:
+// CHECK1-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
+// CHECK1-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
+// CHECK1-NEXT: store ptr [[A]], ptr [[A_ADDR_ASCAST]], align 8
+// CHECK1-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR_ASCAST]], align 8
+// CHECK1-NEXT: [[TMP1:%.*]] = call i32 @__kmpc_target_init(ptr addrspacecast (ptr addrspace(1) @[[GLOB1:[0-9]+]] to ptr), i8 1, i1 true)
+// CHECK1-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP1]], -1
+// CHECK1-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
+// CHECK1: user_code.entry:
+// CHECK1-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [2 x double], ptr [[TMP0]], i64 0, i64 0
+// CHECK1-NEXT: store double 0x400921FB54442D18, ptr [[ARRAYIDX]], align 8
+// CHECK1-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds [2 x double], ptr [[TMP0]], i64 0, i64 1
+// CHECK1-NEXT: store double 3.140000e+00, ptr [[ARRAYIDX1]], align 8
+// CHECK1-NEXT: call void @__kmpc_target_deinit(ptr addrspacecast (ptr addrspace(1) @[[GLOB1]] to ptr), i8 1)
+// CHECK1-NEXT: ret void
+// CHECK1: worker.exit:
+// CHECK1-NEXT: ret void
+//
+//.
+// CHECK1: attributes #0 = { convergent noinline norecurse nounwind optnone "kernel" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" }
+//.
+// CHECK1: !0 = !{i32 0, i32 64768, i32 37754654, !"main", i32 24, i32 0, i32 2}
+// CHECK1: !1 = !{i32 1, !"_ZL9anotherPi", i32 0, i32 1}
+// CHECK1: !2 = !{i32 1, !"_ZN1A2piE", i32 0, i32 0}
+// CHECK1: !3 = !{ptr @__omp_offloading_fd00_240171e_main_l24, !"kernel", i32 1}
+// CHECK1: !4 = !{i32 1, !"amdgpu_code_object_version", i32 400}
+// CHECK1: !5 = !{i32 1, !"wchar_size", i32 4}
+// CHECK1: !6 = !{i32 7, !"openmp", i32 50}
+// CHECK1: !7 = !{i32 7, !"openmp-device", i32 50}
+// CHECK1: !8 = !{!"AMD clang version 17.0.0 (g...@github.com:doru1004/llvm-project.git 7321ab818f82e60ceeb1f4e85b9c712b78d2646f)"}
+//.
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -10377,7 +10377,9 @@
}
Linkage = CGM.getLLVMLinkageVarDefinition(VD, /*IsConstant=*/false);
// Temp solution to prevent optimizations of the internal variables.
- if (CGM.getLangOpts().OpenMPIsDevice && !VD->isExternallyVisible()) {
+ if (CGM.getLangOpts().OpenMPIsDevice &&
+ (!VD->isExternallyVisible() ||
+ Linkage == llvm::GlobalValue::LinkOnceODRLinkage)) {
// Do not create a "ref-variable" if the original is not also available
// on the host.
if (!OffloadEntriesInfoManager.hasDeviceGlobalVarEntryInfo(VarName))
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
