[Openmp-commits] [PATCH] D146552: [Clang][OpenMP] Enable device-mapped constexpr class members to not be optimized out

[Gheorghe-Teodor Bercea via Phabricator via Openmp-commits]

doru1004 created this revision.
doru1004 added reviewers: ronl, carlo.bertolli, jhuber6, jdoerfert, gregrodgers, dhruvachak.
doru1004 added a project: OpenMP.
Herald added subscribers: sunshaoce, guansong, yaxunl.
Herald added a project: All.
doru1004 requested review of this revision.
Herald added subscribers: openmp-commits, cfe-commits, jplehr, sstefan1.
Herald added a project: clang.

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:

  Libomptarget error: Unable to generate entries table for device id 0.
  Libomptarget error: Failed to init globals on device 0

This is due to the optimized-out variable not being present when host entry table values are matched with their device counterparts.

A currently failing example is included in the runtime test included in this patch.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D146552

Files:
  clang/lib/CodeGen/CGOpenMPRuntime.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/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))


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D146552.507062.patch
Type: text/x-patch
Size: 1674 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20230321/2907052d/attachment-0001.bin>