[clang] 0eabf59 - Enable constexpr class members that are device-mapped to not be optimized out.

[Doru Bercea via cfe-commits]

Author: Doru Bercea
Date: 2023-03-23T10:17:25-04:00
New Revision: 0eabf59528f3c3f64923900cae740d9f26c45ae8

URL: https://github.com/llvm/llvm-project/commit/0eabf59528f3c3f64923900cae740d9f26c45ae8
DIFF: https://github.com/llvm/llvm-project/commit/0eabf59528f3c3f64923900cae740d9f26c45ae8.diff

LOG: Enable constexpr class members that are device-mapped to not be optimized out.

This patch fixes an issue whereby a constexpr class member which is
mapped to the device is being optimized out thus leading to a runtime
error.

Patch: https://reviews.llvm.org/D146552

Added: 
    clang/test/OpenMP/declare_target_constexpr_codegen.cpp
    openmp/libomptarget/test/offloading/target_constexpr_mapping.cpp

Modified: 
    clang/lib/CodeGen/CGOpenMPRuntime.cpp

Removed: 
    


################################################################################
diff  --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 5f21cfca66bb8..58a95d64ac50e 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -10387,7 +10387,9 @@ void CGOpenMPRuntime::registerTargetGlobalVariable(const VarDecl *VD,
     }
     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
index 0000000000000..27161feef05e0
--- /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]+anotherPi" --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
index 0000000000000..14cf92a7cc26e
--- /dev/null
+++ b/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;
+}