[PATCH] D103995: [OpenMP] Add type to firstprivate symbol for const firstprivate values
Joseph Huber via Phabricator via cfe-commits
cfe-commits at lists.llvm.org
Wed Jun 9 16:41:00 PDT 2021
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;
-------------- next part --------------
A non-text attachment was scrubbed...
Name: D103995.351005.patch
Type: text/x-patch
Size: 2322 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20210609/72bc612e/attachment.bin>
More information about the cfe-commits
mailing list