[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.

3062

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

[jhuber6]

clang/lib/CodeGen/CodeGenModule.cpp

3062

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

3062

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/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 (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))
         EmitOMPDeclareMapper(DMD);
       return;
     }
   }
 
   // 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;
 
       StringRef MangledName = getMangledName(GD);
 
       // Compute the function info and LLVM type.
       const CGFunctionInfo &FI = getTypes().arrangeGlobalDeclaration(GD);

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