[PATCH] D103995: [OpenMP] Add type to firstprivate symbol for const firstprivate values

[Joseph Huber via Phabricator via cfe-commits]

jhuber6 updated this revision to Diff 351012.
jhuber6 added a comment.

Changing to only mangle when using C++. It should be correct either way but is unnecessary with C.


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D103995/new/

https://reviews.llvm.org/D103995

Files:
  clang/lib/CodeGen/CGOpenMPRuntime.cpp
  clang/test/OpenMP/nvptx_target_firstprivate_codegen.cpp


Index: clang/test/OpenMP/nvptx_target_firstprivate_codegen.cpp
===================================================================
--- clang/test/OpenMP/nvptx_target_firstprivate_codegen.cpp
+++ clang/test/OpenMP/nvptx_target_firstprivate_codegen.cpp
@@ -13,11 +13,14 @@
   ty Y;
 };
 
-// TCHECK-DAG:  [[TT:%.+]] = type { i64, i8 }
 // TCHECK-DAG:  [[TTII:%.+]] = type { i32, i32 }
+// TCHECK-DAG:  [[TTIC:%.+]] = type { i8, i8 }
+// TCHECK-DAG:  [[TT:%.+]] = type { i64, i8 }
 // TCHECK-DAG:  [[S1:%.+]] = type { double }
 
-// TCHECK: @__omp_offloading_firstprivate__{{.+}}_e_l27 = internal addrspace(4) global [[TTII]] zeroinitializer
+// TCHECK: @__omp_offloading_firstprivate__{{.+}}_e_l30 = internal addrspace(4) global [[TTII]] zeroinitializer
+// TCHECK: @__omp_offloading_firstprivate__{{.+}}_ZTSK2TTIiiE_t_l143 = internal addrspace(4) global [[TTII]] zeroinitializer
+// TCHECK: @__omp_offloading_firstprivate__{{.+}}_ZTSK2TTIccE_t_l143 = internal addrspace(4) global [[TTIC]] zeroinitializer
 int foo(int n, double *ptr) {
   int a = 0;
   short aa = 0;
@@ -136,6 +139,12 @@
   return a;
 }
 
+template <typename tx>
+void fconst(const tx t) {
+#pragma omp target firstprivate(t)
+  { }
+}
+
 // TCHECK: define {{.*}}void @__omp_offloading_{{.+}}(i{{[0-9]+}}{{.*}} [[A_IN:%.+]], i{{[0-9]+}}{{.*}} [[A3_IN:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]])
 // TCHECK:  [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
 // TCHECK:  [[A3_ADDR:%.+]] = alloca i{{[0-9]+}},
@@ -199,6 +208,9 @@
   a += fstatic(n);
   a += ftemplate<int>(n);
 
+  fconst(TT<int, int>{0, 0});
+  fconst(TT<char, char>{0, 0});
+
   return a;
 }
 
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -10593,7 +10593,12 @@
                              FileID, Line);
     llvm::raw_svector_ostream OS(Buffer);
     OS << "__omp_offloading_firstprivate_" << llvm::format("_%x", DeviceID)
-       << llvm::format("_%x_", FileID) << VD->getName() << "_l" << Line;
+       << llvm::format("_%x_", FileID);
+    if (CGM.getLangOpts().CPlusPlus) {
+      CGM.getCXXABI().getMangleContext().mangleTypeName(VD->getType(), OS);
+      OS << "_";
+    }
+    OS << VD->getName() << "_l" << Line;
     VarName = OS.str();
   }
   Linkage = llvm::GlobalValue::InternalLinkage;


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D103995.351012.patch
Type: text/x-patch
Size: 2383 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20210610/8e0b7191/attachment.bin>