diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -3786,6 +3786,14 @@ }]; } +def OMPDeclareVariantNoMangle : InheritableAttr { + // This attribute has no spellings as it is only ever created implicitly. + let Spellings = []; + let Args = [DeclArgument]; + let Subjects = SubjectList<[Function]>; + let Documentation = [Undocumented]; +} + def Assumption : InheritableAttr { let Spellings = [Clang<"assume">]; let Subjects = SubjectList<[Function, ObjCMethod]>; diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -4293,6 +4293,7 @@ match_none disable_implicit_base allow_templates + keep_original_name The match extensions change when the *entire* context selector is considered a match for an OpenMP context. The default is ``all``, with ``none`` no trait in the @@ -4308,7 +4309,9 @@ applied to a definition. If ``allow_templates`` is given, template function definitions are considered as specializations of existing or assumed template declarations with the same name. The template parameters for the base functions -are used to instantiate the specialization. +are used to instantiate the specialization. if ``keep_original_name`` is given, +the variant functions will not be mangled. This allows the user to overload the +meaning of existing symbols as long as their definitions can be dropped. }]; } diff --git a/clang/lib/CodeGen/CodeGenModule.h b/clang/lib/CodeGen/CodeGenModule.h --- a/clang/lib/CodeGen/CodeGenModule.h +++ b/clang/lib/CodeGen/CodeGenModule.h @@ -370,6 +370,10 @@ /// for the same decl. llvm::DenseSet DiagnosedConflictingDefinitions; + /// Set of OpenMP declare variant globals actually emitted. Required to not + /// issue a mangling conflict on a variant function that is never emitted. + llvm::DenseSet VariantGlobalsEmitted; + /// A queue of (optional) vtables to consider emitting. std::vector DeferredVTables; diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1455,6 +1455,9 @@ // Keep the first result in the case of a mangling collision. const auto *ND = cast(GD.getDecl()); + if (LangOpts.OpenMP) + if (auto *A = ND->getAttr()) + ND = A->getFunction(); std::string MangledName = getMangledNameImpl(*this, GD, ND); // Ensure either we have different ABIs between host and device compilations, @@ -3053,6 +3056,11 @@ EmitOMPDeclareMapper(DMD); return; } + + // If this variant should keep its original name, mark the original + // declaration so we can make sure we don't conflict with it. + if (auto *A = Global->getAttr()) + VariantGlobalsEmitted.insert(A->getFunction()); } // Ignore declarations, they will be emitted on their first use. @@ -3794,7 +3802,9 @@ if (lookupRepresentativeDecl(MangledName, OtherGD) && (GD.getCanonicalDecl().getDecl() != OtherGD.getCanonicalDecl().getDecl()) && + VariantGlobalsEmitted.contains(OtherGD) && DiagnosedConflictingDefinitions.insert(GD).second) { + getDiags().Report(D->getLocation(), diag::err_duplicate_mangled_name) << MangledName; getDiags().Report(OtherGD.getDecl()->getLocation(), @@ -4112,7 +4122,7 @@ if (D && lookupRepresentativeDecl(MangledName, OtherGD) && (D->getCanonicalDecl() != OtherGD.getCanonicalDecl().getDecl()) && (OtherD = dyn_cast(OtherGD.getDecl())) && - OtherD->hasInit() && + OtherD->hasInit() && VariantGlobalsEmitted.contains(OtherGD) && DiagnosedConflictingDefinitions.insert(D).second) { getDiags().Report(D->getLocation(), diag::err_duplicate_mangled_name) << MangledName; diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp --- a/clang/lib/Parse/ParseOpenMP.cpp +++ b/clang/lib/Parse/ParseOpenMP.cpp @@ -957,6 +957,10 @@ TraitProperty::implementation_extension_allow_templates) return true; + if (TIProperty.Kind == + TraitProperty::implementation_extension_keep_original_name) + return true; + auto IsMatchExtension = [](OMPTraitProperty &TP) { return (TP.Kind == llvm::omp::TraitProperty::implementation_extension_match_all || diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -6910,6 +6910,11 @@ /*AppendArgs=*/nullptr, /*AppendArgsSize=*/0); for (FunctionDecl *BaseFD : Bases) BaseFD->addAttr(OMPDeclareVariantA); + if (DVScope.TI->isExtensionActive( + llvm::omp::TraitProperty:: + implementation_extension_keep_original_name)) + D->addAttr( + OMPDeclareVariantNoMangleAttr::CreateImplicit(Context, Bases.front())); } ExprResult Sema::ActOnOpenMPCall(ExprResult Call, Scope *Scope, diff --git a/clang/test/OpenMP/declare_variant_no_mangling.cpp b/clang/test/OpenMP/declare_variant_no_mangling.cpp new file mode 100644 --- /dev/null +++ b/clang/test/OpenMP/declare_variant_no_mangling.cpp @@ -0,0 +1,53 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs +// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s +// expected-no-diagnostics + +#ifndef HEADER +#define HEADER + +static int foo() { return 0; } +extern inline int bar() { return 0; } +int baz() { return 0; } + +#pragma omp begin declare variant match( \ + device = {arch(ppc64le, ppc64)}, \ + implementation = {extension(match_any, keep_original_name)}) + +static int foo() { return 1; } +extern inline int bar() { return 1; } + +#pragma omp end declare variant + +int main() { + return foo() + bar(); +} + +#endif +// CHECK-LABEL: define {{[^@]+}}@_Z3bazv +// CHECK-SAME: () #[[ATTR0:[0-9]+]] { +// CHECK-NEXT: entry: +// CHECK-NEXT: ret i32 0 +// +// +// CHECK-LABEL: define {{[^@]+}}@main +// CHECK-SAME: () #[[ATTR1:[0-9]+]] { +// CHECK-NEXT: entry: +// CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4 +// CHECK-NEXT: store i32 0, i32* [[RETVAL]], align 4 +// CHECK-NEXT: [[CALL:%.*]] = call noundef signext i32 @_ZL3foov() +// CHECK-NEXT: [[CALL1:%.*]] = call noundef signext i32 @_Z3barv() +// CHECK-NEXT: [[ADD:%.*]] = add nsw i32 [[CALL]], [[CALL1]] +// CHECK-NEXT: ret i32 [[ADD]] +// +// +// CHECK-LABEL: define {{[^@]+}}@_ZL3foov +// CHECK-SAME: () #[[ATTR0]] { +// CHECK-NEXT: entry: +// CHECK-NEXT: ret i32 1 +// +// +// CHECK-LABEL: define {{[^@]+}}@_Z3barv +// CHECK-SAME: () #[[ATTR0]] comdat { +// CHECK-NEXT: entry: +// CHECK-NEXT: ret i32 1 +// diff --git a/clang/test/OpenMP/declare_variant_no_mangling_messages.cpp b/clang/test/OpenMP/declare_variant_no_mangling_messages.cpp new file mode 100644 --- /dev/null +++ b/clang/test/OpenMP/declare_variant_no_mangling_messages.cpp @@ -0,0 +1,18 @@ +// RUN: not %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - 2>&1 | FileCheck %s + +// CHECK: definition with same mangled name '_Z3foov' as another definition + +#ifndef HEADER +#define HEADER + +int foo() { return 0; } + +#pragma omp begin declare variant match( \ + device = {arch(ppc64le, ppc64)}, \ + implementation = {extension(match_any, keep_original_name)}) + +int foo() { return 1; } + +#pragma omp end declare variant + +#endif diff --git a/clang/test/OpenMP/nvptx_declare_variant_name_mangling.cpp b/clang/test/OpenMP/nvptx_declare_variant_name_mangling.cpp --- a/clang/test/OpenMP/nvptx_declare_variant_name_mangling.cpp +++ b/clang/test/OpenMP/nvptx_declare_variant_name_mangling.cpp @@ -1,13 +1,15 @@ // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --implicit-check-not='call i32 {@_Z3bazv|@_Z3barv}' +// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --implicit-check-not='call i32 {@_Z3bazv|@_Z3barv|@_ZL3foov}' // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t -// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - | FileCheck %s --implicit-check-not='call i32 {@_Z3bazv|@_Z3barv}' +// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - | FileCheck %s --implicit-check-not='call i32 {@_Z3bazv|@_Z3barv|@_ZL3foov}' // expected-no-diagnostics // CHECK-DAG: @_Z3barv // CHECK-DAG: @_Z3bazv +// CHECK-DAG: call noundef i32 @_ZL3foov() // CHECK-DAG: define{{.*}} @"_Z53bar$ompvariant$S2$s7$Pnvptx$Pnvptx64$S3$s9$Pmatch_anyv" // CHECK-DAG: define{{.*}} @"_Z53baz$ompvariant$S2$s7$Pnvptx$Pnvptx64$S3$s9$Pmatch_anyv" +// CHECK-DAG: define{{.*}} @_ZL3foov // CHECK-DAG: call noundef i32 @"_Z53bar$ompvariant$S2$s7$Pnvptx$Pnvptx64$S3$s9$Pmatch_anyv"() // CHECK-DAG: call noundef i32 @"_Z53baz$ompvariant$S2$s7$Pnvptx$Pnvptx64$S3$s9$Pmatch_anyv"() @@ -20,6 +22,8 @@ int baz() { return 5; } +static int foo() { return 3; } + #pragma omp begin declare variant match(device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)}) int bar() { return 2; } @@ -28,13 +32,19 @@ #pragma omp end declare variant +#pragma omp begin declare variant match(device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any, keep_original_name)}) + +static int foo() { return 4; } + +#pragma omp end declare variant + #pragma omp end declare target int main() { int res; #pragma omp target map(from \ : res) - res = bar() + baz(); + res = bar() + baz() + foo(); return res; } diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def --- a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def +++ b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def @@ -1157,6 +1157,7 @@ __OMP_TRAIT_PROPERTY(implementation, extension, match_none) __OMP_TRAIT_PROPERTY(implementation, extension, disable_implicit_base) __OMP_TRAIT_PROPERTY(implementation, extension, allow_templates) +__OMP_TRAIT_PROPERTY(implementation, extension, keep_original_name) __OMP_TRAIT_SET(user)