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))
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))
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits