diff --git a/clang/include/clang/Basic/IdentifierTable.h b/clang/include/clang/Basic/IdentifierTable.h
--- a/clang/include/clang/Basic/IdentifierTable.h
+++ b/clang/include/clang/Basic/IdentifierTable.h
@@ -138,6 +138,9 @@
   // True if this is a mangled OpenMP variant name.
   unsigned IsMangledOpenMPVariantName : 1;
 
+  // True if we should keep the original name after mangling.
+  unsigned KeepOpenMPVariantName : 1;
+
   // True if this is a deprecated macro.
   unsigned IsDeprecatedMacro : 1;
 
@@ -161,7 +164,8 @@
         NeedsHandleIdentifier(false), IsFromAST(false), ChangedAfterLoad(false),
         FEChangedAfterLoad(false), RevertedTokenID(false), OutOfDate(false),
         IsModulesImport(false), IsMangledOpenMPVariantName(false),
-        IsDeprecatedMacro(false), IsRestrictExpansion(false), IsFinal(false) {}
+        KeepOpenMPVariantName(false), IsDeprecatedMacro(false),
+        IsRestrictExpansion(false), IsFinal(false) {}
 
 public:
   IdentifierInfo(const IdentifierInfo &) = delete;
@@ -441,6 +445,12 @@
   /// Set whether this is the mangled name of an OpenMP variant.
   void setMangledOpenMPVariantName(bool I) { IsMangledOpenMPVariantName = I; }
 
+  /// Determine if we should use the original name instead of the mangled one.
+  bool keepOpenMPVariantName() const { return KeepOpenMPVariantName; }
+
+  /// Set whether we should emit the unmangled name at code generation.
+  void setKeepOpenMPVariantName(bool I) { KeepOpenMPVariantName = I; }
+
   /// Return true if this identifier is an editor placeholder.
   ///
   /// Editor placeholders are produced by the code-completion engine and are
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
@@ -1477,6 +1477,14 @@
                  GD.getWithKernelReferenceKind(KernelReferenceKind::Kernel),
                  ND));
 
+  // If this is an OpenMP variant with the `keep_original_name` extension we
+  // remove the variant mangling.
+  const IdentifierInfo *II = ND->getIdentifier();
+  if (II && II->isMangledOpenMPVariantName() && II->keepOpenMPVariantName())
+    MangledName = StringRef(MangledName)
+                      .split(getOpenMPVariantManglingSeparatorStr())
+                      .first.str();
+
   auto Result = Manglings.insert(std::make_pair(MangledName, GD));
   return MangledDeclNames[CanonicalGD] = Result.first->first();
 }
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
@@ -6881,6 +6881,8 @@
   IdentifierInfo &VariantII = Context.Idents.get(MangledName);
 
   VariantII.setMangledOpenMPVariantName(true);
+  VariantII.setKeepOpenMPVariantName(DVScope.TI->isExtensionActive(
+      llvm::omp::TraitProperty::implementation_extension_keep_original_name));
   D.SetIdentifier(&VariantII, D.getBeginLoc());
 }
 
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
@@ -8,12 +8,16 @@
 // CHECK-DAG: @_Z3bazv
 // 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{{.*}} @_ZL73foo
 // 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"()
+// CHECK-DAG: call noundef i32 @_ZL73foo()
 
 #ifndef HEADER
 #define HEADER
 
+static int foo() { return 3; }
+
 #pragma omp declare target
 
 int bar() { return 1; }
@@ -28,13 +32,22 @@
 
 #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 2; }
+
+#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();
+#pragma omp target map(from \
+                       : res)
+  res += 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)