Author: Joseph Huber Date: 2021-06-10T09:02:20-04:00 New Revision: 0c32ffceedca2a0d7026fc142bab8ac259131386
URL: https://github.com/llvm/llvm-project/commit/0c32ffceedca2a0d7026fc142bab8ac259131386 DIFF: https://github.com/llvm/llvm-project/commit/0c32ffceedca2a0d7026fc142bab8ac259131386.diff LOG: [OpenMP] Add type to firstprivate symbol for const firstprivate values 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. Reviewed By: ABataev Differential Revision: https://reviews.llvm.org/D103995 Added: Modified: clang/lib/CodeGen/CGOpenMPRuntime.cpp clang/test/OpenMP/nvptx_target_firstprivate_codegen.cpp Removed: ################################################################################ diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index e4bb6dcaa4f65..8f65f38747d87 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -10593,7 +10593,12 @@ CGOpenMPRuntime::registerTargetFirstprivateCopy(CodeGenFunction &CGF, 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; diff --git a/clang/test/OpenMP/nvptx_target_firstprivate_codegen.cpp b/clang/test/OpenMP/nvptx_target_firstprivate_codegen.cpp index 70f5a1f849be4..2c470ffcf3950 100644 --- a/clang/test/OpenMP/nvptx_target_firstprivate_codegen.cpp +++ b/clang/test/OpenMP/nvptx_target_firstprivate_codegen.cpp @@ -13,11 +13,14 @@ struct TT { 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 @@ static int fstatic(int n) { 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 @@ int bar(int n, double *ptr) { a += fstatic(n); a += ftemplate<int>(n); + fconst(TT<int, int>{0, 0}); + fconst(TT<char, char>{0, 0}); + return a; } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
