[OpenMP] Add OpenMP variant extension to keep the unmangled name

[jhuber6]

Summary

The OpenMP variant clause current ensures unique function identifiers by
mangling the name with $ompvariant.... This patch adds an extension to
prevent this by stripping off the mangling during code generation. Doing
this will conflict with the other symbols in the module, but it will
allow us to overload the names of existing symbols if their definitions
are allowed to be dropped. This is most useful for overloading the
meaning of math functions without changing the symbol so they can still
be identified by LLVM.

[jhuber6]

The test currently fails because we don't do the stripping when we emit the precompiled header, it will simply be imported using the mangled version and it will not have in the original input info associated with it. I will need to figure out where we emit the PCH definitions and also demangle it there.

[jhuber6]

Updating to use an attribute instead to get around the loss of the identifier info when using PCH. Unfortunately this still causes problems when using PCH.

[jhuber6]

Updating to stop erroring if a function with the same mangled error exists, but is never emitted.

[jdoerfert]

There is no documentation of the extension and attribute. See also below.

clang/lib/CodeGen/CodeGenModule.cpp

1459

Check OpenMP in the lang opts.

3063

This looks like you now disable the diagnostic for pretty much everything, no?

[jhuber6]

clang/lib/CodeGen/CodeGenModule.cpp

3063

This should only get called if we plan to emit this global, if the global has the attribute stating that it should not be mangled we're basically just asserting that its associated non-variant declaration should not be found. So even if we hit a name mangling conflict, as long as we haven't tried to emit that global we should be fine. I can add a test for this.

[jhuber6]

Make suggested changes.

[jdoerfert]

clang/lib/CodeGen/CodeGenModule.cpp

3063

We only put things in this set for a special OpenMP case, correct?
The error condition below triggers only if something is in this set.

So all non-openmp codes will never see this diagnostic, or what am I missing?

clang/include/clang/Basic/Attr.td

def OMPDeclareVariant : InheritableAttr {
   ];
   let AdditionalMembers = [{
     OMPTraitInfo &getTraitInfo() { return *traitInfos; }
     void printPrettyPragma(raw_ostream & OS, const PrintingPolicy &Policy)
         const;
   }];
 }
 
+def OMPDeclareVariantNoMangle : InheritableAttr {
+  // This attribute has no spellings as it is only ever created implicitly.
+  let Spellings = [];
+  let Args = [DeclArgument<Function, "Function">];
+  let Subjects = SubjectList<[Function]>;
+  let Documentation = [Undocumented];
+}
+
 def Assumption : InheritableAttr {
   let Spellings = [Clang<"assume">];
   let Subjects = SubjectList<[Function, ObjCMethod]>;
   let InheritEvenIfAlreadyPresent = 1;
   let Documentation = [AssumptionDocs];
   let Args = [StringArgument<"Assumption">];
 }
 

clang/include/clang/Basic/AttrDocs.td

``implementation={extension(EXTENSION)}``:
 
   .. code-block:: none
 
     match_all
     match_any
     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
 selector is allowed to be in the OpenMP context, with ``any`` a single trait in
 both the selector and OpenMP context is sufficient. Only a single match
 extension trait is allowed per context selector.
 The disable extensions remove default effects of the ``begin declare variant``
 applied to a definition. If ``disable_implicit_base`` is given, we will not
 introduce an implicit base function for a variant if no base function was
 found. The variant is still generated but will never be called, due to the
 absence of a base function and consequently calls to a base function.
 The allow extensions change when the ``begin declare variant`` effect is
 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.
 
   }];
 }
 
 def LeafDocs : Documentation {
   let Category = DocCatVariable;
   let Content = [{
 

clang/lib/CodeGen/CodeGenModule.h

private:
   /// values along with the corresponding globals, for opportunistic reuse.
   llvm::DenseMap<const VarDecl*, llvm::GlobalVariable*> InitializerConstants;
 
   /// Set of global decls for which we already diagnosed mangled name conflict.
   /// Required to not issue a warning (on a mangling conflict) multiple times
   /// for the same decl.
   llvm::DenseSet<GlobalDecl> 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<GlobalDecl> VariantGlobalsEmitted;
+
   /// A queue of (optional) vtables to consider emitting.
   std::vector<const CXXRecordDecl*> DeferredVTables;
 
   /// A queue of (optional) vtables that may be emitted opportunistically.
   std::vector<const CXXRecordDecl *> OpportunisticVTables;
 
   /// List of global values which are required to be present in the object file;
   /// bitcast to i8*. This is used for forcing visibility of symbols which may

clang/lib/CodeGen/CodeGenModule.cpp

if (!LangOpts.CUDAIsDevice ||
       !getContext().mayExternalizeStaticVar(GD.getDecl())) {
     auto FoundName = MangledDeclNames.find(CanonicalGD);
     if (FoundName != MangledDeclNames.end())
       return FoundName->second;
   }
 
   // Keep the first result in the case of a mangling collision.
   const auto *ND = cast<NamedDecl>(GD.getDecl());
+  if (LangOpts.OpenMP)
+    if (auto *A = ND->getAttr<OMPDeclareVariantNoMangleAttr>())
+      ND = A->getFunction();
   std::string MangledName = getMangledNameImpl(*this, GD, ND);
 
   // Ensure either we have different ABIs between host and device compilations,
   // says host compilation following MSVC ABI but device compilation follows
   // Itanium C++ ABI or, if they follow the same ABI, kernel names after
   // mangling should be the same after name stubbing. The later checking is
   // very important as the device kernel name being mangled in host-compilation
   // is used to resolve the device binaries to be executed. Inconsistent naming
if (auto *DRD = dyn_cast<OMPDeclareReductionDecl>(Global)) {
       if (MustBeEmitted(Global))
         EmitOMPDeclareReduction(DRD);
       return;
     } else if (auto *DMD = dyn_cast<OMPDeclareMapperDecl>(Global)) {
       if (MustBeEmitted(Global))
         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<OMPDeclareVariantNoMangleAttr>())
+      VariantGlobalsEmitted.insert(A->getFunction());
   }
 
   // Ignore declarations, they will be emitted on their first use.
   if (const auto *FD = dyn_cast<FunctionDecl>(Global)) {
     // Forward declarations are emitted lazily on first use.
     if (!FD->doesThisDeclarationHaveABody()) {
       if (!FD->doesDeclarationForceExternallyVisibleDefinition())
         return;
if (Entry) {
     // error.
     if (IsForDefinition && !Entry->isDeclaration()) {
       GlobalDecl OtherGD;
       // Check that GD is not yet in DiagnosedConflictingDefinitions is required
       // to make sure that we issue an error only once.
       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(),
                           diag::note_previous_definition);
       }
     }
 
     if ((isa<llvm::Function>(Entry) || isa<llvm::GlobalAlias>(Entry)) &&
if (IsForDefinition && !Entry->isDeclaration()) {
       GlobalDecl OtherGD;
       const VarDecl *OtherD;
 
       // Check that D is not yet in DiagnosedConflictingDefinitions is required
       // to make sure that we issue an error only once.
       if (D && lookupRepresentativeDecl(MangledName, OtherGD) &&
           (D->getCanonicalDecl() != OtherGD.getCanonicalDecl().getDecl()) &&
           (OtherD = dyn_cast<VarDecl>(OtherGD.getDecl())) &&
-          OtherD->hasInit() &&
+          OtherD->hasInit() && VariantGlobalsEmitted.contains(OtherGD) &&
           DiagnosedConflictingDefinitions.insert(D).second) {
         getDiags().Report(D->getLocation(), diag::err_duplicate_mangled_name)
             << MangledName;
         getDiags().Report(OtherGD.getDecl()->getLocation(),
                           diag::note_previous_definition);
       }
     }
 

clang/lib/Parse/ParseOpenMP.cpp

static bool checkExtensionProperty(Parser &P, SourceLocation Loc,
   if (TIProperty.Kind ==
       TraitProperty::implementation_extension_disable_implicit_base)
     return true;
 
   if (TIProperty.Kind ==
       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 ||
             TP.Kind ==
                 llvm::omp::TraitProperty::implementation_extension_match_any ||
             TP.Kind ==
                 llvm::omp::TraitProperty::implementation_extension_match_none);
   };

clang/lib/Sema/SemaOpenMP.cpp

void Sema::ActOnFinishedFunctionDefinitionInOpenMPDeclareVariantScope(
   OMPDeclareVariantScope &DVScope = OMPDeclareVariantScopes.back();
   auto *OMPDeclareVariantA = OMPDeclareVariantAttr::CreateImplicit(
       Context, VariantFuncRef, DVScope.TI,
       /*NothingArgs=*/nullptr, /*NothingArgsSize=*/0,
       /*NeedDevicePtrArgs=*/nullptr, /*NeedDevicePtrArgsSize=*/0,
       /*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,
                                  SourceLocation LParenLoc,
                                  MultiExprArg ArgExprs,
                                  SourceLocation RParenLoc, Expr *ExecConfig) {
   // The common case is a regular call we do not want to specialize at all. Try
   // to make that case fast by bailing early.

clang/test/OpenMP/declare_variant_no_mangling.cpp

+// 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
+//

clang/test/OpenMP/declare_variant_no_mangling_messages.cpp

+// 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

clang/test/OpenMP/nvptx_declare_variant_name_mangling.cpp

 // 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"()
 
 #ifndef HEADER
 #define HEADER
 
 #pragma omp declare target
 
 int bar() { return 1; }
 
 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; }
 
 int baz() { return 6; }
 
 #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;
 }
 
 #endif

llvm/include/llvm/Frontend/OpenMP/OMPKinds.def

 __OMP_TRAIT_PROPERTY(implementation, vendor, unknown)
 
 __OMP_TRAIT_SELECTOR(implementation, extension, true)
 __OMP_TRAIT_PROPERTY(implementation, extension, match_all)
 __OMP_TRAIT_PROPERTY(implementation, extension, match_any)
 __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)
 
 __OMP_TRAIT_SELECTOR(user, condition, true)
 
 __OMP_TRAIT_PROPERTY(user, condition, true)
 __OMP_TRAIT_PROPERTY(user, condition, false)
 __OMP_TRAIT_PROPERTY(user, condition, unknown)