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

[Joseph Huber via Phabricator via cfe-commits]

jhuber6 created this revision.
jhuber6 added reviewers: jdoerfert, ABataev.
Herald added subscribers: guansong, yaxunl.
jhuber6 requested review of this revision.
Herald added subscribers: cfe-commits, sstefan1.
Herald added a project: clang.

Clang will create a global value put in constant memory if an aggregate value
is declared firstprivate in the target device. The symbol name only uses the
name of the firstprivate variable, so symbol name conflicts will occur if the
variable is allowed to have different types through templates. An example of
this behvaiour is shown in https://godbolt.org/z/EsMjYh47n. This patch adds the
mangled type name to the symbol to avoid such naming conflicts. This fixes
https://bugs.llvm.org/show_bug.cgi?id=50642.


Repository:
  rG LLVM Github Monorepo

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,9 @@
                              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);
+    CGM.getCXXABI().getMangleContext().mangleTypeName(VD->getType(), OS);
+    OS << "_" << VD->getName() << "_l" << Line;
     VarName = OS.str();
   }
   Linkage = llvm::GlobalValue::InternalLinkage;



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,9 @@
                              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);
+    CGM.getCXXABI().getMangleContext().mangleTypeName(VD->getType(), OS);
+    OS << "_" << VD->getName() << "_l" << Line;
     VarName = OS.str();
   }
   Linkage = llvm::GlobalValue::InternalLinkage;

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits