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 @@ -10387,7 +10387,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)) diff --git a/clang/test/OpenMP/declare_target_constexpr_codegen.cpp b/clang/test/OpenMP/declare_target_constexpr_codegen.cpp new file mode 100644 --- /dev/null +++ b/clang/test/OpenMP/declare_target_constexpr_codegen.cpp @@ -0,0 +1,40 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals --prefix-filecheck-ir-name _ --global-value-regex "llvm.compiler.used" "_[0-9a-zA-Z]+A[0-9a-zA-Z]+pi[0-9a-zA-Z]+" "_[0-9a-zA-Z]+A[0-9a-zA-Z]+anotherPi[0-9a-zA-Z]+" --version 2 +// REQUIRES: amdgpu-registered-target + + +// Test target codegen - host bc file has to be created first. +// RUN: %clang_cc1 -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 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-target-debug -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK + +// expected-no-diagnostics + +#ifndef HEADER +#define HEADER + +#pragma omp declare target +class A { +public: + static constexpr double pi = 3.141592653589793116; +//. +// CHECK: @_ZN1A2piE = linkonce_odr constant double 0x400921FB54442D18, comdat, align 8 +// CHECK: @_ZL9anotherPi = internal constant double 3.140000e+00, align 8 +// CHECK: @llvm.compiler.used = appending global [2 x ptr] [ptr @"__ZN1A2piE$ref", ptr @"__ZL9anotherPi$ref"], section "llvm.metadata" +//. + A() { ; } + ~A() { ; } +}; +#pragma omp end declare target + +void F(const double &); +void Test() { F(A::pi); } + +#pragma omp declare target +constexpr static double anotherPi = 3.14; +#pragma omp end declare target + +#endif + + +// +//// NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line: +// CHECK: {{.*}} diff --git a/openmp/libomptarget/test/offloading/target_constexpr_mapping.cpp b/openmp/libomptarget/test/offloading/target_constexpr_mapping.cpp new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/offloading/target_constexpr_mapping.cpp @@ -0,0 +1,34 @@ +// RUN: %libomptarget-compileoptxx-run-and-check-generic + +#include +#include + +#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; +}