jhuber6 updated this revision to Diff 419567.
jhuber6 added a comment.
Herald added a reviewer: aaron.ballman.
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.
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D122760/new/
https://reviews.llvm.org/D122760
Files:
clang/include/clang/Basic/Attr.td
clang/lib/CodeGen/CodeGenModule.cpp
clang/lib/Parse/ParseOpenMP.cpp
clang/lib/Sema/SemaOpenMP.cpp
clang/test/OpenMP/nvptx_declare_variant_name_mangling.cpp
llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
Index: llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
===================================================================
--- llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
+++ 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)
Index: clang/test/OpenMP/nvptx_declare_variant_name_mangling.cpp
===================================================================
--- clang/test/OpenMP/nvptx_declare_variant_name_mangling.cpp
+++ 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;
}
Index: clang/lib/Sema/SemaOpenMP.cpp
===================================================================
--- clang/lib/Sema/SemaOpenMP.cpp
+++ 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,
Index: clang/lib/Parse/ParseOpenMP.cpp
===================================================================
--- clang/lib/Parse/ParseOpenMP.cpp
+++ 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 ||
Index: clang/lib/CodeGen/CodeGenModule.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenModule.cpp
+++ clang/lib/CodeGen/CodeGenModule.cpp
@@ -1455,6 +1455,8 @@
// 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,
Index: clang/include/clang/Basic/Attr.td
===================================================================
--- clang/include/clang/Basic/Attr.td
+++ clang/include/clang/Basic/Attr.td
@@ -3779,6 +3779,14 @@
}];
}
+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]>;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
