Replace separator in OpenMP variant name mangling.

[LukasSommerTu]

Summary

Nvidia PTX does not allow . to appear in identifiers, so OpenMP variant mangling now uses $ to separate segments of the mangled name for variants of functions declared via declare variant.

[jdoerfert]

Changes looks good to me.

We don't have tests that check the mangling, interesting (and my fault), should we add one?

[LukasSommerTu]

Added a small test checking for functions with the correctly mangled name.

[LukasSommerTu]

@jdoerfert: I've added a test. I kept it focused on the variants for NVPTX architecture (mangling is currently the same for all archs), let me know if that was what you had in mind.

[jdoerfert]

LGTM. Thanks a lot!

[LukasSommerTu]

@jdoerfert: Can you please commit for me, I do not have commit access. Thanks!

clang/include/clang/AST/Decl.h

 inline bool IsEnumDeclScoped(EnumDecl *ED) {
   return ED->isScoped();
 }
 
 /// OpenMP variants are mangled early based on their OpenMP context selector.
 /// The new name looks likes this:
 ///  <name> + OpenMPVariantManglingSeparatorStr + <mangled OpenMP context>
 static constexpr StringRef getOpenMPVariantManglingSeparatorStr() {
-  return ".ompvariant";
+  return "$ompvariant";
 }
 
 } // namespace clang
 
 #endif // LLVM_CLANG_AST_DECL_H

clang/lib/AST/OpenMPClause.cpp

for (const OMPTraitSet &Set : Sets) {
     OS << "}";
   }
 }
 
 std::string OMPTraitInfo::getMangledName() const {
   std::string MangledName;
   llvm::raw_string_ostream OS(MangledName);
   for (const OMPTraitSet &Set : Sets) {
-    OS << '.' << 'S' << unsigned(Set.Kind);
+    OS << '$' << 'S' << unsigned(Set.Kind);
     for (const OMPTraitSelector &Selector : Set.Selectors) {
 
       bool AllowsTraitScore = false;
       bool RequiresProperty = false;
       isValidTraitSelectorForTraitSet(
           Selector.Kind, Set.Kind, AllowsTraitScore, RequiresProperty);
-      OS << '.' << 's' << unsigned(Selector.Kind);
+      OS << '$' << 's' << unsigned(Selector.Kind);
 
       if (!RequiresProperty ||
           Selector.Kind == TraitSelector::user_condition)
         continue;
 
       for (const OMPTraitProperty &Property : Selector.Properties)
-        OS << '.' << 'P'
-           << getOpenMPContextTraitPropertyName(Property.Kind);
+        OS << '$' << 'P' << getOpenMPContextTraitPropertyName(Property.Kind);
     }
   }
   return OS.str();
 }
 
 OMPTraitInfo::OMPTraitInfo(StringRef MangledName) {
   unsigned long U;
   do {
-    if (!MangledName.consume_front(".S"))
+    if (!MangledName.consume_front("$S"))
       break;
     if (MangledName.consumeInteger(10, U))
       break;
     Sets.push_back(OMPTraitSet());
     OMPTraitSet &Set = Sets.back();
     Set.Kind = TraitSet(U);
     do {
-      if (!MangledName.consume_front(".s"))
+      if (!MangledName.consume_front("$s"))
         break;
       if (MangledName.consumeInteger(10, U))
         break;
       Set.Selectors.push_back(OMPTraitSelector());
       OMPTraitSelector &Selector = Set.Selectors.back();
       Selector.Kind = TraitSelector(U);
       do {
-        if (!MangledName.consume_front(".P"))
+        if (!MangledName.consume_front("$P"))
           break;
         Selector.Properties.push_back(OMPTraitProperty());
         OMPTraitProperty &Property = Selector.Properties.back();
-        std::pair<StringRef, StringRef> PropRestPair = MangledName.split('.');
+        std::pair<StringRef, StringRef> PropRestPair = MangledName.split('$');
         Property.Kind =
             getOpenMPContextTraitPropertyKind(Set.Kind, PropRestPair.first);
         MangledName = PropRestPair.second;
       } while (true);
     } while (true);
   } while (true);
 }
 

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 -fopenmp-version=50
+// 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 - -fopenmp-version=50 | 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 -emit-pch -o %t -fopenmp-version=50
+// 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 - -fopenmp-version=50 | FileCheck %s --implicit-check-not='call i32 {@_Z3bazv|@_Z3barv}'
+// expected-no-diagnostics
+
+// CHECK-DAG: @_Z3barv
+// CHECK-DAG: @_Z3bazv
+// CHECK-DAG: @"_Z54bar$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_anyv"
+// CHECK-DAG: @"_Z54baz$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_anyv"
+// CHECK-DAG: call i32 @"_Z54bar$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_anyv"()
+// CHECK-DAG: call i32 @"_Z54baz$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_anyv"()
+
+#ifndef HEADER
+#define HEADER
+
+#pragma omp declare target
+
+int bar() { return 1; }
+
+int baz() { return 5; }
+
+#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 end declare target
+
+int main() {
+  int res;
+#pragma omp target map(from \
+                       : res)
+  res = bar() + baz();
+  return res;
+}
+
+#endif
 No newline at end of file