https://github.com/jdoerfert updated https://github.com/llvm/llvm-project/pull/200301
>From 03048caa5d966a85198c1db271940165f41d19e2 Mon Sep 17 00:00:00 2001 From: Johannes Doerfert <[email protected]> Date: Fri, 29 May 2026 15:17:30 -0700 Subject: [PATCH] [OpenMP] Introduce the ompx_name clause for kernel naming This adds support for the ompx_name clause that allows users to specify custom kernel names for OpenMP target offloading regions. The clause accepts a string literal and overrides the default compiler-generated kernel names. Example usage: #pragma omp target ompx_name("my_kernel") { ... } In the process, the linkage of the offload entries was changed from weak to external (for kernels), or whatever the global variable linkage is. This makes sure we can link weak globals together but clash on kernels with the same name. Co-Authored-By: Claude (claude-sonnet-4.5) <[email protected]> --- clang/include/clang/AST/OpenMPClause.h | 33 ++++++++ clang/include/clang/AST/RecursiveASTVisitor.h | 9 ++- .../clang/Basic/DiagnosticSemaKinds.td | 7 ++ clang/include/clang/Sema/SemaOpenMP.h | 15 ++++ clang/lib/AST/OpenMPClause.cpp | 6 ++ clang/lib/AST/StmtProfile.cpp | 5 ++ clang/lib/CodeGen/CGOpenMPRuntime.cpp | 3 + clang/lib/Parse/ParseOpenMP.cpp | 1 + clang/lib/Sema/SemaOpenMP.cpp | 37 +++++++++ clang/lib/Sema/TreeTransform.h | 20 +++++ clang/lib/Serialization/ASTReader.cpp | 8 ++ clang/lib/Serialization/ASTWriter.cpp | 5 ++ clang/test/OpenMP/amdgcn_weak_alias.c | 4 +- clang/test/OpenMP/declare_target_codegen.cpp | 2 +- clang/test/OpenMP/ompx_name_codegen.cpp | 53 +++++++++++++ .../test/OpenMP/ompx_name_messages_errors.cpp | 48 ++++++++++++ clang/test/OpenMP/target_codegen.cpp | 22 +++--- clang/test/OpenMP/target_depend_codegen.cpp | 4 +- clang/test/OpenMP/target_indirect_codegen.cpp | 8 +- .../OpenMP/target_parallel_depend_codegen.cpp | 4 +- .../target_parallel_for_depend_codegen.cpp | 4 +- ...arget_parallel_for_simd_depend_codegen.cpp | 4 +- clang/test/OpenMP/target_simd_codegen.cpp | 16 ++-- .../OpenMP/target_simd_depend_codegen.cpp | 4 +- .../OpenMP/target_teams_depend_codegen.cpp | 4 +- ...target_teams_distribute_depend_codegen.cpp | 4 +- ...distribute_parallel_for_depend_codegen.cpp | 4 +- ...ibute_parallel_for_simd_depend_codegen.cpp | 4 +- ...t_teams_distribute_simd_depend_codegen.cpp | 4 +- clang/tools/libclang/CIndex.cpp | 1 + .../llvm/Frontend/Offloading/Utility.h | 3 +- llvm/include/llvm/Frontend/OpenMP/OMP.td | 18 +++++ .../llvm/Frontend/OpenMP/OMPIRBuilder.h | 13 ++-- llvm/lib/Frontend/Offloading/Utility.cpp | 5 +- llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 16 ++-- offload/test/offloading/ompx_name.c | 78 +++++++++++++++++++ .../offloading/ompx_name_duplicate_link.c | 35 +++++++++ 37 files changed, 451 insertions(+), 60 deletions(-) create mode 100644 clang/test/OpenMP/ompx_name_codegen.cpp create mode 100644 clang/test/OpenMP/ompx_name_messages_errors.cpp create mode 100644 offload/test/offloading/ompx_name.c create mode 100644 offload/test/offloading/ompx_name_duplicate_link.c diff --git a/clang/include/clang/AST/OpenMPClause.h b/clang/include/clang/AST/OpenMPClause.h index ccf2c40bc5efa..2772d56262997 100644 --- a/clang/include/clang/AST/OpenMPClause.h +++ b/clang/include/clang/AST/OpenMPClause.h @@ -10413,6 +10413,39 @@ class OMPXBareClause : public OMPNoChildClause<llvm::omp::OMPC_ompx_bare> { OMPXBareClause() = default; }; +/// This represents the 'ompx_name' clause in the '#pragma omp target' +/// directive. +/// +/// \code +/// #pragma omp target ompx_name("foo") +/// \endcode +/// In this example directive '#pragma omp target' has simple 'ompx_name' +/// clause with the name "foo". +class OMPXNameClause final + : public OMPOneStmtClause<llvm::omp::OMPC_ompx_name, OMPClause> { + friend class OMPClauseReader; + + /// Set name. + void setName(Expr *A) { setStmt(A); } + +public: + /// Build 'ompx_name' clause with the given name. + /// + /// \param A Name. + /// \param StartLoc Starting location of the clause. + /// \param LParenLoc Location of '('. + /// \param EndLoc Ending location of the clause. + OMPXNameClause(Expr *A, SourceLocation StartLoc, SourceLocation LParenLoc, + SourceLocation EndLoc) + : OMPOneStmtClause(A, StartLoc, LParenLoc, EndLoc) {} + + /// Build an empty clause. + OMPXNameClause() : OMPOneStmtClause() {} + + /// Returns name. + Expr *getName() const { return getStmtAs<Expr>(); } +}; + } // namespace clang #endif // LLVM_CLANG_AST_OPENMPCLAUSE_H diff --git a/clang/include/clang/AST/RecursiveASTVisitor.h b/clang/include/clang/AST/RecursiveASTVisitor.h index febdf715698d9..bf2a6b5bdf628 100644 --- a/clang/include/clang/AST/RecursiveASTVisitor.h +++ b/clang/include/clang/AST/RecursiveASTVisitor.h @@ -3483,7 +3483,14 @@ bool RecursiveASTVisitor<Derived>::VisitOMPAllocatorClause( } template <typename Derived> -bool RecursiveASTVisitor<Derived>::VisitOMPAllocateClause(OMPAllocateClause *C) { +bool RecursiveASTVisitor<Derived>::VisitOMPXNameClause(OMPXNameClause *C) { + TRY_TO(TraverseStmt(C->getName())); + return true; +} + +template <typename Derived> +bool RecursiveASTVisitor<Derived>::VisitOMPAllocateClause( + OMPAllocateClause *C) { TRY_TO(TraverseStmt(C->getAllocator())); TRY_TO(VisitOMPClauseList(C)); return true; diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 077aace321264..71fb51adc677e 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12024,6 +12024,13 @@ let CategoryName = "OpenMP Issue" in { // OpenMP support. def err_omp_expected_var_arg : Error< "%0 is not a global variable, static local variable or static data member">; +def err_ompx_name_argument_not_string : Error< + "argument to 'ompx_name' clause must be a string literal">; +def warn_ompx_name_duplicate : Warning< + "OpenMP target kernel name '%0' is used more than once in this translation unit">, + InGroup<OpenMPTarget>; +def note_ompx_name_previous : Note< + "previous use of this kernel name is here">; def err_omp_expected_var_arg_suggest : Error< "%0 is not a global variable, static local variable or static data member; " "did you mean %1">; diff --git a/clang/include/clang/Sema/SemaOpenMP.h b/clang/include/clang/Sema/SemaOpenMP.h index 3621ce96b8724..1009001c557b9 100644 --- a/clang/include/clang/Sema/SemaOpenMP.h +++ b/clang/include/clang/Sema/SemaOpenMP.h @@ -29,6 +29,8 @@ #include "clang/Sema/Ownership.h" #include "clang/Sema/SemaBase.h" #include "llvm/ADT/DenseMap.h" +#include "llvm/ADT/DenseSet.h" +#include "llvm/ADT/StringMap.h" #include "llvm/Frontend/OpenMP/OMP.h.inc" #include "llvm/Frontend/OpenMP/OMPConstants.h" #include <optional> @@ -1180,6 +1182,11 @@ class SemaOpenMP : public SemaBase { SourceLocation LParenLoc, SourceLocation EndLoc); + /// Called on well-formed 'ompx_name' clause. + OMPClause *ActOnOpenMPOmpxNameClause(Expr *Name, SourceLocation StartLoc, + SourceLocation LParenLoc, + SourceLocation EndLoc); + /// Data used for processing a list of variables in OpenMP clauses. struct OpenMPVarListDataTy final { Expr *DepModOrTailExpr = nullptr; @@ -1501,6 +1508,14 @@ class SemaOpenMP : public SemaBase { private: void *VarDataSharingAttributesStack; + /// User-provided target kernel names from 'ompx_name' clauses in this + /// translation unit, keyed to their first source location. + llvm::StringMap<SourceLocation> OMPKernelNames; + + /// Source locations for duplicate kernel names that have already been + /// diagnosed. This prevents repeated diagnostics during template transforms. + llvm::DenseSet<unsigned> DiagnosedOMPKernelNameLocs; + /// Number of nested '#pragma omp declare target' directives. SmallVector<DeclareTargetContextInfo, 4> DeclareTargetNesting; diff --git a/clang/lib/AST/OpenMPClause.cpp b/clang/lib/AST/OpenMPClause.cpp index 3a35e17aff40b..33be79e78015e 100644 --- a/clang/lib/AST/OpenMPClause.cpp +++ b/clang/lib/AST/OpenMPClause.cpp @@ -2060,6 +2060,12 @@ void OMPClausePrinter::VisitOMPAllocatorClause(OMPAllocatorClause *Node) { OS << ")"; } +void OMPClausePrinter::VisitOMPXNameClause(OMPXNameClause *Node) { + OS << "ompx_name("; + Node->getName()->printPretty(OS, nullptr, Policy, 0); + OS << ")"; +} + void OMPClausePrinter::VisitOMPCollapseClause(OMPCollapseClause *Node) { OS << "collapse("; Node->getNumForLoops()->printPretty(OS, nullptr, Policy, 0); diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp index eb25e5260fd1a..39d038a3e68a4 100644 --- a/clang/lib/AST/StmtProfile.cpp +++ b/clang/lib/AST/StmtProfile.cpp @@ -530,6 +530,11 @@ void OMPClauseProfiler::VisitOMPAllocatorClause(const OMPAllocatorClause *C) { Profiler->VisitStmt(C->getAllocator()); } +void OMPClauseProfiler::VisitOMPXNameClause(const OMPXNameClause *C) { + if (C->getName()) + Profiler->VisitStmt(C->getName()); +} + void OMPClauseProfiler::VisitOMPCollapseClause(const OMPCollapseClause *C) { if (C->getNumForLoops()) Profiler->VisitStmt(C->getNumForLoops()); diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index ec059f9dfef82..0ae5d4a96578d 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -6371,6 +6371,9 @@ void CGOpenMPRuntime::emitTargetOutlinedFunctionHelper( llvm::TargetRegionEntryInfo EntryInfo = getEntryInfoFromPresumedLoc(CGM, OMPBuilder, D.getBeginLoc(), ParentName); + if (auto *C = D.getSingleClause<OMPXNameClause>()) + if (auto *S = dyn_cast<StringLiteral>(C->getName()->IgnoreParenImpCasts())) + EntryInfo.UserProvidedName = S->getString().str(); CodeGenFunction CGF(CGM, true); llvm::OpenMPIRBuilder::FunctionGenCallback &&GenerateOutlinedFunction = diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp index c4177c1622521..c0a81a079d962 100644 --- a/clang/lib/Parse/ParseOpenMP.cpp +++ b/clang/lib/Parse/ParseOpenMP.cpp @@ -3251,6 +3251,7 @@ OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind, case OMPC_align: case OMPC_message: case OMPC_ompx_dyn_cgroup_mem: + case OMPC_ompx_name: case OMPC_dyn_groupprivate: case OMPC_transparent: // OpenMP [2.5, Restrictions] diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp index 76b40a5039180..f3b99c3631896 100644 --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -16877,6 +16877,9 @@ OMPClause *SemaOpenMP::ActOnOpenMPSingleExprClause(OpenMPClauseKind Kind, case OMPC_message: Res = ActOnOpenMPMessageClause(Expr, StartLoc, LParenLoc, EndLoc); break; + case OMPC_ompx_name: + Res = ActOnOpenMPOmpxNameClause(Expr, StartLoc, LParenLoc, EndLoc); + break; case OMPC_align: Res = ActOnOpenMPAlignClause(Expr, StartLoc, LParenLoc, EndLoc); break; @@ -17984,6 +17987,40 @@ OMPClause *SemaOpenMP::ActOnOpenMPMessageClause(Expr *ME, ME, HelperValStmt, CaptureRegion, StartLoc, LParenLoc, EndLoc); } +OMPClause *SemaOpenMP::ActOnOpenMPOmpxNameClause(Expr *Name, + SourceLocation StartLoc, + SourceLocation LParenLoc, + SourceLocation EndLoc) { + if (!Name) { + Diag(StartLoc, diag::err_ompx_name_argument_not_string); + return nullptr; + } + + if (!Name->isTypeDependent() && !Name->isValueDependent()) { + if (auto *PE = dyn_cast<PredefinedExpr>(Name->IgnoreParenCasts())) + Name = PE->getFunctionName(); + auto *SL = dyn_cast<StringLiteral>(Name->IgnoreParenCasts()); + if (!SL) { + Diag(Name->getExprLoc(), diag::err_ompx_name_argument_not_string); + return nullptr; + } + + StringRef KernelName = SL->getString(); + SourceLocation NameLoc = SL->getBeginLoc(); + auto It = OMPKernelNames.find(KernelName); + if (It == OMPKernelNames.end()) { + OMPKernelNames[KernelName] = NameLoc; + } else if (It->second != NameLoc && + DiagnosedOMPKernelNameLocs.insert(NameLoc.getRawEncoding()) + .second) { + Diag(NameLoc, diag::warn_ompx_name_duplicate) << KernelName; + Diag(It->second, diag::note_ompx_name_previous); + } + } + return new (getASTContext()) + OMPXNameClause(Name, StartLoc, LParenLoc, EndLoc); +} + OMPClause *SemaOpenMP::ActOnOpenMPOrderClause( OpenMPOrderClauseModifier Modifier, OpenMPOrderClauseKind Kind, SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation MLoc, diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h index c3bf71dbf10df..4d683e8d5833e 100644 --- a/clang/lib/Sema/TreeTransform.h +++ b/clang/lib/Sema/TreeTransform.h @@ -1822,6 +1822,17 @@ class TreeTransform { EndLoc); } + /// Build a new OpenMP 'ompx_name' clause. + /// + /// By default, performs semantic analysis to build the new OpenMP clause. + /// Subclasses may override this routine to provide different behavior. + OMPClause *RebuildOMPXNameClause(Expr *Name, SourceLocation StartLoc, + SourceLocation LParenLoc, + SourceLocation EndLoc) { + return getSema().OpenMP().ActOnOpenMPOmpxNameClause(Name, StartLoc, + LParenLoc, EndLoc); + } + /// Build a new OpenMP 'collapse' clause. /// /// By default, performs semantic analysis to build the new OpenMP clause. @@ -10602,6 +10613,15 @@ TreeTransform<Derived>::TransformOMPAllocatorClause(OMPAllocatorClause *C) { E.get(), C->getBeginLoc(), C->getLParenLoc(), C->getEndLoc()); } +template <typename Derived> +OMPClause *TreeTransform<Derived>::TransformOMPXNameClause(OMPXNameClause *C) { + ExprResult E = getDerived().TransformExpr(C->getName()); + if (E.isInvalid()) + return nullptr; + return getDerived().RebuildOMPXNameClause(E.get(), C->getBeginLoc(), + C->getLParenLoc(), C->getEndLoc()); +} + template <typename Derived> OMPClause * TreeTransform<Derived>::TransformOMPSimdlenClause(OMPSimdlenClause *C) { diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp index 74a7b51368c28..0e999e33fc281 100644 --- a/clang/lib/Serialization/ASTReader.cpp +++ b/clang/lib/Serialization/ASTReader.cpp @@ -11501,6 +11501,9 @@ OMPClause *OMPClauseReader::readClause() { case llvm::omp::OMPC_allocator: C = new (Context) OMPAllocatorClause(); break; + case llvm::omp::OMPC_ompx_name: + C = new (Context) OMPXNameClause(); + break; case llvm::omp::OMPC_collapse: C = new (Context) OMPCollapseClause(); break; @@ -11933,6 +11936,11 @@ void OMPClauseReader::VisitOMPAllocatorClause(OMPAllocatorClause *C) { C->setLParenLoc(Record.readSourceLocation()); } +void OMPClauseReader::VisitOMPXNameClause(OMPXNameClause *C) { + C->setName(Record.readExpr()); + C->setLParenLoc(Record.readSourceLocation()); +} + void OMPClauseReader::VisitOMPCollapseClause(OMPCollapseClause *C) { C->setNumForLoops(Record.readSubExpr()); C->setLParenLoc(Record.readSourceLocation()); diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp index 074b0fccdb65d..4d2b21a13bb9a 100644 --- a/clang/lib/Serialization/ASTWriter.cpp +++ b/clang/lib/Serialization/ASTWriter.cpp @@ -8141,6 +8141,11 @@ void OMPClauseWriter::VisitOMPAllocatorClause(OMPAllocatorClause *C) { Record.AddSourceLocation(C->getLParenLoc()); } +void OMPClauseWriter::VisitOMPXNameClause(OMPXNameClause *C) { + Record.AddStmt(C->getName()); + Record.AddSourceLocation(C->getLParenLoc()); +} + void OMPClauseWriter::VisitOMPCollapseClause(OMPCollapseClause *C) { Record.AddStmt(C->getNumForLoops()); Record.AddSourceLocation(C->getLParenLoc()); diff --git a/clang/test/OpenMP/amdgcn_weak_alias.c b/clang/test/OpenMP/amdgcn_weak_alias.c index 4cc54b9f15b43..6292bb5640a79 100644 --- a/clang/test/OpenMP/amdgcn_weak_alias.c +++ b/clang/test/OpenMP/amdgcn_weak_alias.c @@ -10,9 +10,9 @@ // HOST: @__Two_var = global i32 2, align 4 // HOST: @__Three_var = global i32 3, align 4 // HOST: @.offloading.entry_name = internal unnamed_addr constant [10 x i8] c"__Two_var\00", section ".llvm.rodata.offloading", align 1 -// HOST: @.offloading.entry.__Two_var = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @__Two_var, ptr @.offloading.entry_name, i64 4, i64 0, ptr null }, section "llvm_offload_entries", align 8 +// HOST: @.offloading.entry.__Two_var = constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @__Two_var, ptr @.offloading.entry_name, i64 4, i64 0, ptr null }, section "llvm_offload_entries", align 8 // HOST: @.offloading.entry_name.1 = internal unnamed_addr constant [12 x i8] c"__Three_var\00", section ".llvm.rodata.offloading", align 1 -// HOST: @.offloading.entry.__Three_var = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @__Three_var, ptr @.offloading.entry_name.1, i64 4, i64 0, ptr null }, section "llvm_offload_entries", align 8 +// HOST: @.offloading.entry.__Three_var = constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @__Three_var, ptr @.offloading.entry_name.1, i64 4, i64 0, ptr null }, section "llvm_offload_entries", align 8 // HOST: @One = weak alias i32 (), ptr @__One // HOST: @One_ = alias i32 (), ptr @__One // HOST: @One_var = weak alias i32, ptr @__One_var diff --git a/clang/test/OpenMP/declare_target_codegen.cpp b/clang/test/OpenMP/declare_target_codegen.cpp index 6add3eed1f226..89899a1cb4f30 100644 --- a/clang/test/OpenMP/declare_target_codegen.cpp +++ b/clang/test/OpenMP/declare_target_codegen.cpp @@ -30,7 +30,7 @@ // CHECK-DAG: @dx = {{protected | }}global i32 0, // CHECK-DAG: @dy = {{protected | }}global i32 0, // CHECK-DAG: @bbb = {{protected | }}global i32 0, -// CHECK-DAG: weak constant %struct.__tgt_offload_entry { +// CHECK-DAG: constant %struct.__tgt_offload_entry { // CHECK-DAG: @ccc = external global i32, // CHECK-DAG: @ddd = {{protected | }}global i32 0, // CHECK-DAG: @hhh_decl_tgt_ref_ptr = weak global ptr null diff --git a/clang/test/OpenMP/ompx_name_codegen.cpp b/clang/test/OpenMP/ompx_name_codegen.cpp new file mode 100644 index 0000000000000..60bc31c84df02 --- /dev/null +++ b/clang/test/OpenMP/ompx_name_codegen.cpp @@ -0,0 +1,53 @@ +// Test for ompx_name clause code generation +// +// This test verifies that the ompx_name clause correctly sets the kernel name. +// +// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc +// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s + +// expected-no-diagnostics + +#define TO_STR(x) #x + +// CHECK: define {{.*}} @my_custom_kernel( +void test_target() { + #pragma omp target ompx_name("my_custom_kernel") + { + } +} + +// CHECK: define {{.*}} @another_kernel( +void test_target_parallel() { + #pragma omp target parallel ompx_name("another_kernel") + { + } +} + +// CHECK: define {{.*}} @teams_kernel_name( +void test_target_teams() { + #pragma omp target teams ompx_name("teams_kernel_" "name") + { + } +} + +// CHECK: define {{.*}} @simd_kernel_name( +void test_target_simd() { + #pragma omp target simd ompx_name("simd_kernel_name") + for (int i = 0; i < 10; i++) + ; +} + +// CHECK: define {{.*}} @parallel_for_kernel_3( +void test_target_parallel_for() { + #pragma omp target parallel for ompx_name("parallel_for_kernel_" TO_STR(3)) + for (int i = 0; i < 10; i++) + ; +} + +// Verify default kernel name generation without ompx_name +// CHECK: define {{.*}} @__omp_offloading_{{[0-9a-f]+}}_{{[0-9a-f]+}}_{{.*}}_l{{[0-9]+}}( +void test_default_name() { + #pragma omp target + { + } +} diff --git a/clang/test/OpenMP/ompx_name_messages_errors.cpp b/clang/test/OpenMP/ompx_name_messages_errors.cpp new file mode 100644 index 0000000000000..6861eee365713 --- /dev/null +++ b/clang/test/OpenMP/ompx_name_messages_errors.cpp @@ -0,0 +1,48 @@ +// Test for ompx_name clause error checking +// RUN: %clang_cc1 -std=c++20 -verify -fopenmp %s + +static void foo() { +} + +void bar() { + int x = 5; + + // expected-error@+1 {{argument to 'ompx_name' clause must be a string literal}} + #pragma omp target ompx_name(x) + { + } + + // expected-error@+1 {{argument to 'ompx_name' clause must be a string literal}} + #pragma omp target ompx_name(123) + { + } + + // This should work - string literal + #pragma omp target ompx_name("valid_name") + { + } + +// expected-note@+1 {{previous use of this kernel name is here}} +#pragma omp target ompx_name("baz") + foo(); + +// expected-warning@+1 {{OpenMP target kernel name 'baz' is used more than once in this translation unit}} +#pragma omp target ompx_name("baz") + foo(); + +#pragma omp target ompx_name(foo) // expected-error {{argument to 'ompx_name' clause must be a string literal}} + foo(); + +#pragma omp target ompx_name("foo", "bar") // expected-error {{expected ')'}} expected-note {{to match this '('}} + foo(); +} + +consteval const char* getStr() { + return "foobar3"; +} + +void foobar() { +// CHECK: define {{.*}} @foobar3( + #pragma omp target ompx_name(getStr()) // expected-error {{argument to 'ompx_name' clause must be a string literal}} + {} +} diff --git a/clang/test/OpenMP/target_codegen.cpp b/clang/test/OpenMP/target_codegen.cpp index 431b09e81714b..34a02d85858ac 100644 --- a/clang/test/OpenMP/target_codegen.cpp +++ b/clang/test/OpenMP/target_codegen.cpp @@ -102,17 +102,17 @@ // CHECK-DAG: @{{.*}} = weak constant i8 0 // CHECK-DAG: @{{.*}} = weak constant i8 0 -// TCHECK: @{{.+}} = weak constant [[ENTTY]] -// TCHECK: @{{.+}} = weak constant [[ENTTY]] -// TCHECK: @{{.+}} = weak constant [[ENTTY]] -// TCHECK: @{{.+}} = weak constant [[ENTTY]] -// TCHECK: @{{.+}} = weak constant [[ENTTY]] -// TCHECK: @{{.+}} = weak constant [[ENTTY]] -// TCHECK: @{{.+}} = weak constant [[ENTTY]] -// TCHECK: @{{.+}} = weak constant [[ENTTY]] -// TCHECK: @{{.+}} = weak constant [[ENTTY]] -// TCHECK: @{{.+}} = weak constant [[ENTTY]] -// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] +// TCHECK-NOT: @{{.+}} = constant [[ENTTY]] template<typename tx, typename ty> struct TT{ diff --git a/clang/test/OpenMP/target_depend_codegen.cpp b/clang/test/OpenMP/target_depend_codegen.cpp index 86b70dd73680d..cf221386eaee9 100644 --- a/clang/test/OpenMP/target_depend_codegen.cpp +++ b/clang/test/OpenMP/target_depend_codegen.cpp @@ -47,9 +47,9 @@ // CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [4 x i64] [i64 544, i64 800, i64 3, i64 288] // CHECK-DAG: @{{.*}} = weak constant i8 0 -// TCHECK: @{{.+}} = weak constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] // TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]] -// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]] +// TCHECK-NOT: @{{.+}} = constant [[ENTTY]] template<typename tx, typename ty> struct TT{ diff --git a/clang/test/OpenMP/target_indirect_codegen.cpp b/clang/test/OpenMP/target_indirect_codegen.cpp index fd8b6c76d0881..ba161ff8cf94d 100644 --- a/clang/test/OpenMP/target_indirect_codegen.cpp +++ b/clang/test/OpenMP/target_indirect_codegen.cpp @@ -23,13 +23,13 @@ // HOST: @indirect_foo = global ptr @_Z3foov, align 8 // HOST: @indirect_array = global [3 x ptr] [ptr @_Z3foov, ptr @_ZL3barv, ptr @_Z3bazv], align 8 // HOST: @[[FOO_ENTRY_NAME:.+]] = internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[FOO_NAME:__omp_offloading_[0-9a-z]+_[0-9a-z]+_foo_l[0-9]+]]\00" -// HOST: @.offloading.entry.[[FOO_NAME]] = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 8, ptr @_Z3foov, ptr @[[FOO_ENTRY_NAME]], i64 8, i64 0, ptr null } +// HOST: @.offloading.entry.[[FOO_NAME]] = weak_odr constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 8, ptr @_Z3foov, ptr @[[FOO_ENTRY_NAME]], i64 8, i64 0, ptr null } // HOST: @[[BAZ_ENTRY_NAME:.+]] = internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[BAZ_NAME:__omp_offloading_[0-9a-z]+_[0-9a-z]+_baz_l[0-9]+]]\00" -// HOST: @.offloading.entry.[[BAZ_NAME]] = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 8, ptr @_Z3bazv, ptr @[[BAZ_ENTRY_NAME]], i64 8, i64 0, ptr null } +// HOST: @.offloading.entry.[[BAZ_NAME]] = weak_odr constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 8, ptr @_Z3bazv, ptr @[[BAZ_ENTRY_NAME]], i64 8, i64 0, ptr null } // HOST: @[[VAR_ENTRY_NAME:.+]] = internal unnamed_addr constant [4 x i8] c"var\00" -// HOST: @.offloading.entry.var = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @[[VAR]], ptr @[[VAR_ENTRY_NAME]], i64 1, i64 0, ptr null } +// HOST: @.offloading.entry.var = constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @[[VAR]], ptr @[[VAR_ENTRY_NAME]], i64 1, i64 0, ptr null } // HOST: @[[BAR_ENTRY_NAME:.+]] = internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[BAR_NAME:__omp_offloading_[0-9a-z]+_[0-9a-z]+_bar_l[0-9]+]]\00" -// HOST: @.offloading.entry.[[BAR_NAME]] = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 8, ptr @_ZL3barv, ptr @[[BAR_ENTRY_NAME]], i64 8, i64 0, ptr null } +// HOST: @.offloading.entry.[[BAR_NAME]] = weak_odr constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 8, ptr @_ZL3barv, ptr @[[BAR_ENTRY_NAME]], i64 8, i64 0, ptr null } //. // DEVICE: @[[FOO_NAME:__omp_offloading_[0-9a-z]+_[0-9a-z]+_foo_l[0-9]+]] = protected addrspace(1) constant {{ptr|ptr addrspace\(9\)}} @_Z3foov // DEVICE: @[[BAZ_NAME:__omp_offloading_[0-9a-z]+_[0-9a-z]+_baz_l[0-9]+]] = protected addrspace(1) constant {{ptr|ptr addrspace\(9\)}} @_Z3bazv diff --git a/clang/test/OpenMP/target_parallel_depend_codegen.cpp b/clang/test/OpenMP/target_parallel_depend_codegen.cpp index ae41454a5c376..6966f1bbc5a7b 100644 --- a/clang/test/OpenMP/target_parallel_depend_codegen.cpp +++ b/clang/test/OpenMP/target_parallel_depend_codegen.cpp @@ -47,9 +47,9 @@ // CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [3 x i64] [i64 544, i64 800, i64 288] // CHECK-DAG: @{{.*}} = weak constant i8 0 -// TCHECK: @{{.+}} = weak constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] // TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]] -// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]] +// TCHECK-NOT: @{{.+}} = constant [[ENTTY]] template<typename tx, typename ty> struct TT{ diff --git a/clang/test/OpenMP/target_parallel_for_depend_codegen.cpp b/clang/test/OpenMP/target_parallel_for_depend_codegen.cpp index 7cccb2549c2c8..8d79b37ea46c9 100644 --- a/clang/test/OpenMP/target_parallel_for_depend_codegen.cpp +++ b/clang/test/OpenMP/target_parallel_for_depend_codegen.cpp @@ -47,9 +47,9 @@ // CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [3 x i64] [i64 544, i64 800, i64 288] // CHECK-DAG: @{{.*}} = weak constant i8 0 -// TCHECK: @{{.+}} = weak constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] // TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]] -// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]] +// TCHECK-NOT: @{{.+}} = constant [[ENTTY]] template<typename tx, typename ty> struct TT{ diff --git a/clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp b/clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp index f1391cd26e2d4..cacde85ca6e82 100644 --- a/clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp +++ b/clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp @@ -47,9 +47,9 @@ // CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [3 x i64] [i64 544, i64 800, i64 288] // CHECK-DAG: @{{.*}} = weak constant i8 0 -// TCHECK: @{{.+}} = weak constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] // TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]] -// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]] +// TCHECK-NOT: @{{.+}} = constant [[ENTTY]] template<typename tx, typename ty> struct TT{ diff --git a/clang/test/OpenMP/target_simd_codegen.cpp b/clang/test/OpenMP/target_simd_codegen.cpp index 0c2dde23f6c46..141fa6ffe385b 100644 --- a/clang/test/OpenMP/target_simd_codegen.cpp +++ b/clang/test/OpenMP/target_simd_codegen.cpp @@ -101,14 +101,14 @@ // CHECK-DAG: @{{.*}} = weak constant i8 0 // CHECK-DAG: @{{.*}} = weak constant i8 0 -// TCHECK: @{{.+}} = weak constant [[ENTTY]] -// TCHECK: @{{.+}} = weak constant [[ENTTY]] -// TCHECK: @{{.+}} = weak constant [[ENTTY]] -// TCHECK: @{{.+}} = weak constant [[ENTTY]] -// TCHECK: @{{.+}} = weak constant [[ENTTY]] -// TCHECK: @{{.+}} = weak constant [[ENTTY]] -// TCHECK: @{{.+}} = weak constant [[ENTTY]] -// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] +// TCHECK-NOT: @{{.+}} = constant [[ENTTY]] template<typename tx, typename ty> struct TT{ diff --git a/clang/test/OpenMP/target_simd_depend_codegen.cpp b/clang/test/OpenMP/target_simd_depend_codegen.cpp index e399998869cf5..53a4f6ce9897b 100644 --- a/clang/test/OpenMP/target_simd_depend_codegen.cpp +++ b/clang/test/OpenMP/target_simd_depend_codegen.cpp @@ -47,9 +47,9 @@ // CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [3 x i64] [i64 544, i64 800, i64 288] // CHECK-DAG: @{{.*}} = weak constant i8 0 -// TCHECK: @{{.+}} = weak constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] // TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]] -// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]] +// TCHECK-NOT: @{{.+}} = constant [[ENTTY]] template<typename tx, typename ty> struct TT{ diff --git a/clang/test/OpenMP/target_teams_depend_codegen.cpp b/clang/test/OpenMP/target_teams_depend_codegen.cpp index 1b7e25ee7e936..3bc16dc41c610 100644 --- a/clang/test/OpenMP/target_teams_depend_codegen.cpp +++ b/clang/test/OpenMP/target_teams_depend_codegen.cpp @@ -47,9 +47,9 @@ // CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [3 x i64] [i64 544, i64 800, i64 288] // CHECK-DAG: @{{.*}} = weak constant i8 0 -// TCHECK: @{{.+}} = weak constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] // TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]] -// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]] +// TCHECK-NOT: @{{.+}} = constant [[ENTTY]] template<typename tx, typename ty> struct TT{ diff --git a/clang/test/OpenMP/target_teams_distribute_depend_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_depend_codegen.cpp index 5bf4615fe7b70..c146a36ec9b90 100644 --- a/clang/test/OpenMP/target_teams_distribute_depend_codegen.cpp +++ b/clang/test/OpenMP/target_teams_distribute_depend_codegen.cpp @@ -47,9 +47,9 @@ // CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [3 x i64] [i64 544, i64 800, i64 288] // CHECK-DAG: @{{.*}} = weak constant i8 0 -// TCHECK: @{{.+}} = weak constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] // TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]] -// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]] +// TCHECK-NOT: @{{.+}} = constant [[ENTTY]] template<typename tx, typename ty> struct TT{ diff --git a/clang/test/OpenMP/target_teams_distribute_parallel_for_depend_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_parallel_for_depend_codegen.cpp index 9fd3ca822a38b..f4d6c005d7d54 100644 --- a/clang/test/OpenMP/target_teams_distribute_parallel_for_depend_codegen.cpp +++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_depend_codegen.cpp @@ -47,9 +47,9 @@ // CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [3 x i64] [i64 544, i64 800, i64 288] // CHECK-DAG: @{{.*}} = weak constant i8 0 -// TCHECK: @{{.+}} = weak constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] // TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]] -// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]] +// TCHECK-NOT: @{{.+}} = constant [[ENTTY]] template<typename tx, typename ty> struct TT{ diff --git a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_depend_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_depend_codegen.cpp index 9393d9d0474bd..fc8114ed70f7f 100644 --- a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_depend_codegen.cpp +++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_depend_codegen.cpp @@ -47,9 +47,9 @@ // CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [3 x i64] [i64 544, i64 800, i64 288] // CHECK-DAG: @{{.*}} = weak constant i8 0 -// TCHECK: @{{.+}} = weak constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] // TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]] -// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]] +// TCHECK-NOT: @{{.+}} = constant [[ENTTY]] template<typename tx, typename ty> struct TT{ diff --git a/clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp index fd5cea7ebd9a0..47cef10da1b4e 100644 --- a/clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp +++ b/clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp @@ -61,9 +61,9 @@ // OMP50-DAG: @{{.*}} = weak constant i8 0 -// TCHECK: @{{.+}} = weak constant [[ENTTY]] +// TCHECK: @{{.+}} = constant [[ENTTY]] // TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]] -// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]] +// TCHECK-NOT: @{{.+}} = constant [[ENTTY]] template<typename tx, typename ty> struct TT{ diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp index 85760968cdbde..14484eb67d42a 100644 --- a/clang/tools/libclang/CIndex.cpp +++ b/clang/tools/libclang/CIndex.cpp @@ -2790,6 +2790,7 @@ void OMPClauseEnqueue::VisitOMPDoacrossClause(const OMPDoacrossClause *C) { void OMPClauseEnqueue::VisitOMPXAttributeClause(const OMPXAttributeClause *C) { } void OMPClauseEnqueue::VisitOMPXBareClause(const OMPXBareClause *C) {} +void OMPClauseEnqueue::VisitOMPXNameClause(const OMPXNameClause *C) {} } // namespace diff --git a/llvm/include/llvm/Frontend/Offloading/Utility.h b/llvm/include/llvm/Frontend/Offloading/Utility.h index 4c0bc87786dfb..6e8002a2cf9af 100644 --- a/llvm/include/llvm/Frontend/Offloading/Utility.h +++ b/llvm/include/llvm/Frontend/Offloading/Utility.h @@ -94,7 +94,8 @@ LLVM_ABI StringRef getOffloadEntrySection(Module &M); LLVM_ABI GlobalVariable * emitOffloadingEntry(Module &M, object::OffloadKind Kind, Constant *Addr, StringRef Name, uint64_t Size, uint32_t Flags, - uint64_t Data, Constant *AuxAddr = nullptr); + uint64_t Data, Constant *AuxAddr = nullptr, + GlobalValue::LinkageTypes Linkage = GlobalValue::WeakAnyLinkage); /// Create a constant struct initializer used to register this global at /// runtime. diff --git a/llvm/include/llvm/Frontend/OpenMP/OMP.td b/llvm/include/llvm/Frontend/OpenMP/OMP.td index e1e66df72dfc5..1e5c9379ff90b 100644 --- a/llvm/include/llvm/Frontend/OpenMP/OMP.td +++ b/llvm/include/llvm/Frontend/OpenMP/OMP.td @@ -426,6 +426,10 @@ def OMPC_OMPX_DynCGroupMem : Clause<[Spelling<"ompx_dyn_cgroup_mem">]> { let clangClass = "OMPXDynCGroupMemClause"; let flangClass = "ScalarIntExpr"; } +def OMPC_OMPX_Name : Clause<[Spelling<"ompx_name">]> { + let clangClass = "OMPXNameClause"; + let flangClass = "OmpXNameClause"; +} def OMP_ORDER_concurrent : EnumVal<"concurrent",1,1> {} def OMP_ORDER_unknown : EnumVal<"unknown",2,0> { let isDefault = 1; } def OMPC_Order : Clause<[Spelling<"order">]> { @@ -1224,6 +1228,7 @@ def OMP_Target : Directive<[Spelling<"target">]> { VersionedClause<OMPC_If>, VersionedClause<OMPC_NoWait>, VersionedClause<OMPC_OMPX_Bare>, + VersionedClause<OMPC_OMPX_Name>, VersionedClause<OMPC_OMPX_DynCGroupMem>, VersionedClause<OMPC_Replayable, 60>, VersionedClause<OMPC_ThreadLimit, 51>, @@ -1707,6 +1712,7 @@ def OMP_target_loop : Directive<[Spelling<"target loop">]> { VersionedClause<OMPC_Device>, VersionedClause<OMPC_DefaultMap>, VersionedClause<OMPC_NoWait>, + VersionedClause<OMPC_OMPX_Name>, ]; let leafConstructs = [OMP_Target, OMP_loop]; let category = CA_Executable; @@ -2195,6 +2201,7 @@ def OMP_TargetParallel : Directive<[Spelling<"target parallel">]> { VersionedClause<OMPC_ProcBind>, VersionedClause<OMPC_Severity, 60>, VersionedClause<OMPC_ThreadLimit, 51>, + VersionedClause<OMPC_OMPX_Name>, ]; let leafConstructs = [OMP_Target, OMP_Parallel]; let category = CA_Executable; @@ -2305,6 +2312,7 @@ def OMP_TargetParallelFor : Directive<[Spelling<"target parallel for">]> { let allowedOnceClauses = [ VersionedClause<OMPC_OMPX_DynCGroupMem>, VersionedClause<OMPC_ThreadLimit, 51>, + VersionedClause<OMPC_OMPX_Name>, ]; let leafConstructs = [OMP_Target, OMP_Parallel, OMP_For]; let category = CA_Executable; @@ -2348,6 +2356,7 @@ def OMP_TargetParallelForSimd let allowedOnceClauses = [ VersionedClause<OMPC_OMPX_DynCGroupMem>, VersionedClause<OMPC_ThreadLimit, 51>, + VersionedClause<OMPC_OMPX_Name>, ]; let leafConstructs = [OMP_Target, OMP_Parallel, OMP_For, OMP_Simd]; let category = CA_Executable; @@ -2384,6 +2393,7 @@ def OMP_target_parallel_loop : Directive<[Spelling<"target parallel loop">]> { VersionedClause<OMPC_ProcBind>, VersionedClause<OMPC_Severity, 60>, VersionedClause<OMPC_ThreadLimit, 51>, + VersionedClause<OMPC_OMPX_Name>, ]; let leafConstructs = [OMP_Target, OMP_Parallel, OMP_loop]; let category = CA_Executable; @@ -2423,6 +2433,7 @@ def OMP_TargetSimd : Directive<[Spelling<"target simd">]> { VersionedClause<OMPC_Severity, 60>, VersionedClause<OMPC_SimdLen>, VersionedClause<OMPC_ThreadLimit, 51>, + VersionedClause<OMPC_OMPX_Name>, ]; let leafConstructs = [OMP_Target, OMP_Simd]; let category = CA_Executable; @@ -2451,6 +2462,7 @@ def OMP_TargetTeams : Directive<[Spelling<"target teams">]> { VersionedClause<OMPC_NumTeams>, VersionedClause<OMPC_OMPX_DynCGroupMem>, VersionedClause<OMPC_OMPX_Bare>, + VersionedClause<OMPC_OMPX_Name>, VersionedClause<OMPC_ThreadLimit>, ]; let leafConstructs = [OMP_Target, OMP_Teams]; @@ -2485,6 +2497,7 @@ def OMP_TargetTeamsDistribute VersionedClause<OMPC_OMPX_DynCGroupMem>, VersionedClause<OMPC_Order, 50>, VersionedClause<OMPC_ThreadLimit>, + VersionedClause<OMPC_OMPX_Name>, ]; let leafConstructs = [OMP_Target, OMP_Teams, OMP_Distribute]; let category = CA_Executable; @@ -2605,6 +2618,7 @@ def OMP_TargetTeamsDistributeParallelFor ]; let allowedOnceClauses = [ VersionedClause<OMPC_OMPX_DynCGroupMem>, + VersionedClause<OMPC_OMPX_Name>, ]; let leafConstructs = [OMP_Target, OMP_Teams, OMP_Distribute, OMP_Parallel, OMP_For]; @@ -2650,6 +2664,7 @@ def OMP_TargetTeamsDistributeParallelForSimd ]; let allowedOnceClauses = [ VersionedClause<OMPC_OMPX_DynCGroupMem>, + VersionedClause<OMPC_OMPX_Name>, ]; let leafConstructs = [OMP_Target, OMP_Teams, OMP_Distribute, OMP_Parallel, OMP_For, OMP_Simd]; @@ -2690,6 +2705,7 @@ def OMP_TargetTeamsDistributeSimd VersionedClause<OMPC_SafeLen>, VersionedClause<OMPC_SimdLen>, VersionedClause<OMPC_ThreadLimit>, + VersionedClause<OMPC_OMPX_Name>, ]; let leafConstructs = [OMP_Target, OMP_Teams, OMP_Distribute, OMP_Simd]; let category = CA_Executable; @@ -2717,6 +2733,7 @@ def OMP_TargetTeamsWorkdistribute : Directive<[Spelling<"target teams workdistri VersionedClause<OMPC_NumTeams>, VersionedClause<OMPC_OMPX_DynCGroupMem>, VersionedClause<OMPC_OMPX_Bare>, + VersionedClause<OMPC_OMPX_Name>, VersionedClause<OMPC_ThreadLimit>, ]; let leafConstructs = [OMP_Target, OMP_Teams, OMP_Workdistribute]; @@ -2751,6 +2768,7 @@ def OMP_target_teams_loop : Directive<[Spelling<"target teams loop">]> { VersionedClause<OMPC_OMPX_DynCGroupMem>, VersionedClause<OMPC_Order>, VersionedClause<OMPC_ThreadLimit>, + VersionedClause<OMPC_OMPX_Name>, ]; let leafConstructs = [OMP_Target, OMP_Teams, OMP_loop]; let category = CA_Executable; diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h index 2b790458f3c32..7de40e6ae3d77 100644 --- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h +++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h @@ -239,22 +239,25 @@ struct TargetRegionEntryInfo { unsigned FileID; unsigned Line; unsigned Count; + std::string UserProvidedName; TargetRegionEntryInfo() : DeviceID(0), FileID(0), Line(0), Count(0) {} TargetRegionEntryInfo(StringRef ParentName, unsigned DeviceID, - unsigned FileID, unsigned Line, unsigned Count = 0) + unsigned FileID, unsigned Line, unsigned Count = 0, + StringRef UserProvidedName = "") : ParentName(ParentName), DeviceID(DeviceID), FileID(FileID), Line(Line), - Count(Count) {} + Count(Count), UserProvidedName(UserProvidedName) {} LLVM_ABI static void getTargetRegionEntryFnName(SmallVectorImpl<char> &Name, StringRef ParentName, unsigned DeviceID, unsigned FileID, unsigned Line, - unsigned Count); + unsigned Count, StringRef UserProvidedName = ""); bool operator<(const TargetRegionEntryInfo &RHS) const { - return std::make_tuple(ParentName, DeviceID, FileID, Line, Count) < + return std::make_tuple(ParentName, DeviceID, FileID, Line, Count, + UserProvidedName) < std::make_tuple(RHS.ParentName, RHS.DeviceID, RHS.FileID, RHS.Line, - RHS.Count); + RHS.Count, RHS.UserProvidedName); } }; diff --git a/llvm/lib/Frontend/Offloading/Utility.cpp b/llvm/lib/Frontend/Offloading/Utility.cpp index c07d276244ee1..3787f0cdae766 100644 --- a/llvm/lib/Frontend/Offloading/Utility.cpp +++ b/llvm/lib/Frontend/Offloading/Utility.cpp @@ -104,7 +104,8 @@ getOffloadEntryBoundarySymbols(const Triple &T, StringRef SectionName) { GlobalVariable *offloading::emitOffloadingEntry( Module &M, object::OffloadKind Kind, Constant *Addr, StringRef Name, - uint64_t Size, uint32_t Flags, uint64_t Data, Constant *AuxAddr) { + uint64_t Size, uint32_t Flags, uint64_t Data, Constant *AuxAddr, + GlobalValue::LinkageTypes Linkage) { const llvm::Triple &Triple = M.getTargetTriple(); StringRef SectionName = getOffloadEntrySection(M); @@ -115,7 +116,7 @@ GlobalVariable *offloading::emitOffloadingEntry( Triple.isNVPTX() ? "$offloading$entry$" : ".offloading.entry."; auto *Entry = new GlobalVariable( M, getEntryTy(M), - /*isConstant=*/true, GlobalValue::WeakAnyLinkage, EntryInitializer, + /*isConstant=*/true, Linkage, EntryInitializer, Prefix + Name, nullptr, GlobalValue::NotThreadLocal, M.getDataLayout().getDefaultGlobalsAddressSpace()); diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp index ecff0c9b0aac4..4c81e640b3776 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -11769,12 +11769,12 @@ std::unique_ptr<CodeExtractor> DeviceSharedMemOutlineInfo::createCodeExtractor( void OpenMPIRBuilder::createOffloadEntry(Constant *ID, Constant *Addr, uint64_t Size, int32_t Flags, - GlobalValue::LinkageTypes, + GlobalValue::LinkageTypes Linkage, StringRef Name) { if (!Config.isGPU()) { llvm::offloading::emitOffloadingEntry( M, object::OffloadKind::OFK_OpenMP, ID, - Name.empty() ? Addr->getName() : Name, Size, Flags, /*Data=*/0); + Name.empty() ? Addr->getName() : Name, Size, Flags, /*Data=*/0, /*AuxAddr*/nullptr, Linkage); return; } // TODO: Add support for global variables on the device after declare target @@ -11885,7 +11885,8 @@ void OpenMPIRBuilder::createOffloadEntriesAndInfoMetadata( } createOffloadEntry(CE->getID(), CE->getAddress(), /*Size=*/0, CE->getFlags(), - GlobalValue::WeakAnyLinkage); + GlobalValue::ExternalLinkage, + E.second.UserProvidedName); } else if (const auto *CE = dyn_cast< OffloadEntriesInfoManager::OffloadEntryInfoDeviceGlobalVar>( E.first)) { @@ -11969,7 +11970,12 @@ void OpenMPIRBuilder::createOffloadEntriesAndInfoMetadata( void TargetRegionEntryInfo::getTargetRegionEntryFnName( SmallVectorImpl<char> &Name, StringRef ParentName, unsigned DeviceID, - unsigned FileID, unsigned Line, unsigned Count) { + unsigned FileID, unsigned Line, unsigned Count, + StringRef UserProvidedName) { + if (!UserProvidedName.empty()) { + Name.append(UserProvidedName.begin(), UserProvidedName.end()); + return; + } raw_svector_ostream OS(Name); OS << KernelNamePrefix << llvm::format("%x", DeviceID) << llvm::format("_%x_", FileID) << ParentName << "_l" << Line; @@ -11982,7 +11988,7 @@ void OffloadEntriesInfoManager::getTargetRegionEntryFnName( unsigned NewCount = getTargetRegionEntryInfoCount(EntryInfo); TargetRegionEntryInfo::getTargetRegionEntryFnName( Name, EntryInfo.ParentName, EntryInfo.DeviceID, EntryInfo.FileID, - EntryInfo.Line, NewCount); + EntryInfo.Line, NewCount, EntryInfo.UserProvidedName); } TargetRegionEntryInfo diff --git a/offload/test/offloading/ompx_name.c b/offload/test/offloading/ompx_name.c new file mode 100644 index 0000000000000..d734b65f17f09 --- /dev/null +++ b/offload/test/offloading/ompx_name.c @@ -0,0 +1,78 @@ +// RUN: %libomptarget-compile-generic +// RUN: env LIBOMPTARGET_INFO=63 %libomptarget-run-generic 2>&1 | \ +// RUN: %fcheck-generic +// +// REQUIRES: gpu + +#include <stdio.h> + +int main() { + int result = 0; + +// CHECK: PluginInterface device {{[0-9]}} info: Launching kernel +// CHECK-SAME: my_custom_kernel +#pragma omp target ompx_name("my_custom_kernel") map(from : result) + { + result = 42; + } + + if (result != 42) { + printf("FAIL: result = %d\n", result); + return 1; + } + + result = 0; + +// CHECK: PluginInterface device {{[0-9]}} info: Launching kernel +// CHECK-SAME: parallel_kernel_name +#pragma omp target parallel ompx_name("parallel_kernel_name") \ + map(tofrom : result) + { +#pragma omp atomic + result++; + } + + if (result == 0) { + printf("FAIL: parallel result = %d\n", result); + return 1; + } + + result = 0; + +// CHECK: PluginInterface device {{[0-9]}} info: Launching kernel +// CHECK-SAME: teams_kernel_name +#pragma omp target teams ompx_name("teams_kernel_name") map(tofrom : result) + { +#pragma omp atomic + result++; + } + + if (result == 0) { + printf("FAIL: teams result = %d\n", result); + return 1; + } + + int data[100]; + for (int i = 0; i < 100; i++) + data[i] = 0; + +// CHECK: PluginInterface device {{[0-9]}} info: Launching kernel +// CHECK-SAME: parallel_for_kernel +#pragma omp target parallel for ompx_name("parallel_for_kernel") \ + map(tofrom : data[0 : 100]) + for (int i = 0; i < 100; i++) { + data[i] = i; + } + + for (int i = 0; i < 100; i++) { + if (data[i] != i) { + printf("FAIL: data[%d] = %d\n", i, data[i]); + return 1; + } + } + + // CHECK: PASS + printf("PASS\n"); + + return 0; +} diff --git a/offload/test/offloading/ompx_name_duplicate_link.c b/offload/test/offloading/ompx_name_duplicate_link.c new file mode 100644 index 0000000000000..ae80ca3aba305 --- /dev/null +++ b/offload/test/offloading/ompx_name_duplicate_link.c @@ -0,0 +1,35 @@ +// RUN: %libomptarget-compile-generic -DFIRST -c -o %t.first.o +// RUN: %libomptarget-compile-generic -DSECOND -c -o %t.second.o +// RUN: not %clang-generic %t.second.o %t.first.o -o %t 2>&1 | %fcheck-plain-generic %s +// +// REQUIRES: gpu +// +// CHECK: multiple definition + +#include <stdio.h> + +#ifdef FIRST +void first(void) { + int x = 0; +#pragma omp target ompx_name("duplicate_link_kernel") map(tofrom : x) + { x = 1; } + printf("x: %i\n", x); +} +#endif + +#ifdef SECOND +void second(void) { + int x = 0; +#pragma omp target ompx_name("duplicate_link_kernel") map(tofrom : x) + { x = 2; } + printf("x: %i\n", x); +} + +void first(void); + +int main(void) { + first(); + second(); + return 0; +} +#endif _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
