https://github.com/abhinavgaba updated https://github.com/llvm/llvm-project/pull/170578
>From 5a840934ddc9fb897a25ea0f90322e3ef20c5482 Mon Sep 17 00:00:00 2001 From: Abhinav Gaba <[email protected]> Date: Mon, 1 Dec 2025 14:20:15 -0800 Subject: [PATCH 1/3] [OpenMP][Clang] Parsing/Sema support for `use_device_ptr(fb_preserve/fb_nullify)`. Depends on #169603. This is the `use_device_ptr` counterpart of #168905. With OpenMP 6.1, a `fallback` modifier can be specified on the `use_device_ptr` clause to control the behavior when a pointer lookup fails, i.e. there is no device pointer to translate into. The default is `fb_preserve` (i.e. retain the original pointer), while `fb_nullify` means: use `nullptr` as the translated pointer. --- clang/include/clang/AST/OpenMPClause.h | 39 ++++++++++++++++++++--- clang/include/clang/Basic/OpenMPKinds.def | 8 +++++ clang/include/clang/Basic/OpenMPKinds.h | 8 +++++ clang/include/clang/Sema/SemaOpenMP.h | 8 +++-- clang/lib/AST/OpenMPClause.cpp | 17 ++++++++-- clang/lib/Basic/OpenMPKinds.cpp | 22 +++++++++++-- clang/lib/Parse/ParseOpenMP.cpp | 18 +++++++++++ clang/lib/Sema/SemaOpenMP.cpp | 14 +++++--- clang/lib/Sema/TreeTransform.h | 12 ++++--- clang/lib/Serialization/ASTReader.cpp | 2 ++ clang/lib/Serialization/ASTWriter.cpp | 2 ++ 11 files changed, 130 insertions(+), 20 deletions(-) diff --git a/clang/include/clang/AST/OpenMPClause.h b/clang/include/clang/AST/OpenMPClause.h index d9c3cf239451e..2d63169021347 100644 --- a/clang/include/clang/AST/OpenMPClause.h +++ b/clang/include/clang/AST/OpenMPClause.h @@ -7989,6 +7989,13 @@ class OMPUseDevicePtrClause final friend OMPVarListClause; friend TrailingObjects; + /// Fallback modifier for the clause. + OpenMPUseDevicePtrFallbackModifier FallbackModifier = + OMPC_USE_DEVICE_PTR_FALLBACK_unknown; + + /// Location of the fallback modifier. + SourceLocation FallbackModifierLoc; + /// Build clause with number of variables \a NumVars. /// /// \param Locs Locations needed to build a mappable clause. It includes 1) @@ -7999,10 +8006,14 @@ class OMPUseDevicePtrClause final /// NumUniqueDeclarations: number of unique base declarations in this clause; /// 3) NumComponentLists: number of component lists in this clause; and 4) /// NumComponents: total number of expression components in the clause. + /// \param FallbackModifier The fallback modifier for the clause. + /// \param FallbackModifierLoc Location of the fallback modifier. explicit OMPUseDevicePtrClause(const OMPVarListLocTy &Locs, - const OMPMappableExprListSizeTy &Sizes) - : OMPMappableExprListClause(llvm::omp::OMPC_use_device_ptr, Locs, Sizes) { - } + const OMPMappableExprListSizeTy &Sizes, + OpenMPUseDevicePtrFallbackModifier FallbackModifier, + SourceLocation FallbackModifierLoc) + : OMPMappableExprListClause(llvm::omp::OMPC_use_device_ptr, Locs, Sizes), + FallbackModifier(FallbackModifier), FallbackModifierLoc(FallbackModifierLoc) {} /// Build an empty clause. /// @@ -8055,6 +8066,14 @@ class OMPUseDevicePtrClause final return {getPrivateCopies().end(), varlist_size()}; } + /// Set the fallback modifier for the clause. + void setFallbackModifier(OpenMPUseDevicePtrFallbackModifier M) { + FallbackModifier = M; + } + + /// Set the location of the fallback modifier. + void setFallbackModifierLoc(SourceLocation Loc) { FallbackModifierLoc = Loc; } + public: /// Creates clause with a list of variables \a Vars. /// @@ -8067,11 +8086,15 @@ class OMPUseDevicePtrClause final /// \param Inits Expressions referring to private copy initializers. /// \param Declarations Declarations used in the clause. /// \param ComponentLists Component lists used in the clause. + /// \param FallbackModifier The fallback modifier for the clause. + /// \param FallbackModifierLoc Location of the fallback modifier. static OMPUseDevicePtrClause * Create(const ASTContext &C, const OMPVarListLocTy &Locs, ArrayRef<Expr *> Vars, ArrayRef<Expr *> PrivateVars, ArrayRef<Expr *> Inits, ArrayRef<ValueDecl *> Declarations, - MappableExprComponentListsRef ComponentLists); + MappableExprComponentListsRef ComponentLists, + OpenMPUseDevicePtrFallbackModifier FallbackModifier, + SourceLocation FallbackModifierLoc); /// Creates an empty clause with the place for \a NumVars variables. /// @@ -8084,6 +8107,14 @@ class OMPUseDevicePtrClause final static OMPUseDevicePtrClause * CreateEmpty(const ASTContext &C, const OMPMappableExprListSizeTy &Sizes); + /// Get the fallback modifier for the clause. + OpenMPUseDevicePtrFallbackModifier getFallbackModifier() const { + return FallbackModifier; + } + + /// Get the location of the fallback modifier. + SourceLocation getFallbackModifierLoc() const { return FallbackModifierLoc; } + using private_copies_iterator = MutableArrayRef<Expr *>::iterator; using private_copies_const_iterator = ArrayRef<const Expr *>::iterator; using private_copies_range = llvm::iterator_range<private_copies_iterator>; diff --git a/clang/include/clang/Basic/OpenMPKinds.def b/clang/include/clang/Basic/OpenMPKinds.def index ceac89d3aba6d..e61ee0ddc08da 100644 --- a/clang/include/clang/Basic/OpenMPKinds.def +++ b/clang/include/clang/Basic/OpenMPKinds.def @@ -110,6 +110,9 @@ #ifndef OPENMP_NEED_DEVICE_PTR_KIND #define OPENMP_NEED_DEVICE_PTR_KIND(Name) #endif +#ifndef OPENMP_USE_DEVICE_PTR_FALLBACK_MODIFIER +#define OPENMP_USE_DEVICE_PTR_FALLBACK_MODIFIER(Name) +#endif // Static attributes for 'schedule' clause. OPENMP_SCHEDULE_KIND(static) @@ -282,6 +285,10 @@ OPENMP_THREADSET_KIND(omp_team) OPENMP_NEED_DEVICE_PTR_KIND(fb_nullify) OPENMP_NEED_DEVICE_PTR_KIND(fb_preserve) +// OpenMP 6.1 modifiers for 'use_device_ptr' clause. +OPENMP_USE_DEVICE_PTR_FALLBACK_MODIFIER(fb_nullify) +OPENMP_USE_DEVICE_PTR_FALLBACK_MODIFIER(fb_preserve) + #undef OPENMP_NUMTASKS_MODIFIER #undef OPENMP_NUMTHREADS_MODIFIER #undef OPENMP_DYN_GROUPPRIVATE_MODIFIER @@ -315,3 +322,4 @@ OPENMP_NEED_DEVICE_PTR_KIND(fb_preserve) #undef OPENMP_ALLOCATE_MODIFIER #undef OPENMP_THREADSET_KIND #undef OPENMP_NEED_DEVICE_PTR_KIND +#undef OPENMP_USE_DEVICE_PTR_FALLBACK_MODIFIER diff --git a/clang/include/clang/Basic/OpenMPKinds.h b/clang/include/clang/Basic/OpenMPKinds.h index 3b088b3efd998..4e83bfcd0128b 100644 --- a/clang/include/clang/Basic/OpenMPKinds.h +++ b/clang/include/clang/Basic/OpenMPKinds.h @@ -218,6 +218,14 @@ enum OpenMPNeedDevicePtrModifier { OMPC_NEED_DEVICE_PTR_unknown, }; +/// OpenMP 6.1 use_device_ptr fallback modifier +enum OpenMPUseDevicePtrFallbackModifier { +#define OPENMP_USE_DEVICE_PTR_FALLBACK_MODIFIER(Name) \ + OMPC_USE_DEVICE_PTR_FALLBACK_##Name, +#include "clang/Basic/OpenMPKinds.def" + OMPC_USE_DEVICE_PTR_FALLBACK_unknown, +}; + /// OpenMP bindings for the 'bind' clause. enum OpenMPBindClauseKind { #define OPENMP_BIND_KIND(Name) OMPC_BIND_##Name, diff --git a/clang/include/clang/Sema/SemaOpenMP.h b/clang/include/clang/Sema/SemaOpenMP.h index 2d05b4423140b..e4eb3345534a4 100644 --- a/clang/include/clang/Sema/SemaOpenMP.h +++ b/clang/include/clang/Sema/SemaOpenMP.h @@ -1176,6 +1176,9 @@ class SemaOpenMP : public SemaBase { int OriginalSharingModifier = 0; // Default is shared int NeedDevicePtrModifier = 0; SourceLocation NeedDevicePtrModifierLoc; + int UseDevicePtrFallbackModifier = + OMPC_USE_DEVICE_PTR_FALLBACK_unknown; ///< Fallback modifier for use_device_ptr clause. + SourceLocation UseDevicePtrFallbackModifierLoc; SmallVector<OpenMPMapModifierKind, NumberOfOMPMapClauseModifiers> MapTypeModifiers; SmallVector<SourceLocation, NumberOfOMPMapClauseModifiers> @@ -1364,8 +1367,9 @@ class SemaOpenMP : public SemaBase { ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs, ArrayRef<Expr *> UnresolvedMappers = {}); /// Called on well-formed 'use_device_ptr' clause. - OMPClause *ActOnOpenMPUseDevicePtrClause(ArrayRef<Expr *> VarList, - const OMPVarListLocTy &Locs); + OMPClause *ActOnOpenMPUseDevicePtrClause( + ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs, + OpenMPUseDevicePtrFallbackModifier FallbackModifier, SourceLocation FallbackModifierLoc); /// Called on well-formed 'use_device_addr' clause. OMPClause *ActOnOpenMPUseDeviceAddrClause(ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs); diff --git a/clang/lib/AST/OpenMPClause.cpp b/clang/lib/AST/OpenMPClause.cpp index 2183d77de8fa7..5a6a958595671 100644 --- a/clang/lib/AST/OpenMPClause.cpp +++ b/clang/lib/AST/OpenMPClause.cpp @@ -1441,7 +1441,9 @@ OMPUseDevicePtrClause *OMPUseDevicePtrClause::Create( const ASTContext &C, const OMPVarListLocTy &Locs, ArrayRef<Expr *> Vars, ArrayRef<Expr *> PrivateVars, ArrayRef<Expr *> Inits, ArrayRef<ValueDecl *> Declarations, - MappableExprComponentListsRef ComponentLists) { + MappableExprComponentListsRef ComponentLists, + OpenMPUseDevicePtrFallbackModifier FallbackModifier, + SourceLocation FallbackModifierLoc) { OMPMappableExprListSizeTy Sizes; Sizes.NumVars = Vars.size(); Sizes.NumUniqueDeclarations = getUniqueDeclarationsTotalNumber(Declarations); @@ -1465,7 +1467,8 @@ OMPUseDevicePtrClause *OMPUseDevicePtrClause::Create( Sizes.NumUniqueDeclarations + Sizes.NumComponentLists, Sizes.NumComponents)); - OMPUseDevicePtrClause *Clause = new (Mem) OMPUseDevicePtrClause(Locs, Sizes); + OMPUseDevicePtrClause *Clause = new (Mem) + OMPUseDevicePtrClause(Locs, Sizes, FallbackModifier, FallbackModifierLoc); Clause->setVarRefs(Vars); Clause->setPrivateCopies(PrivateVars); @@ -2753,7 +2756,15 @@ void OMPClausePrinter::VisitOMPDefaultmapClause(OMPDefaultmapClause *Node) { void OMPClausePrinter::VisitOMPUseDevicePtrClause(OMPUseDevicePtrClause *Node) { if (!Node->varlist_empty()) { OS << "use_device_ptr"; - VisitOMPClauseList(Node, '('); + if (Node->getFallbackModifier() != OMPC_USE_DEVICE_PTR_FALLBACK_unknown) { + OS << "(" + << getOpenMPSimpleClauseTypeName(OMPC_use_device_ptr, + Node->getFallbackModifier()) + << ":"; + VisitOMPClauseList(Node, ' '); + } else { + VisitOMPClauseList(Node, '('); + } OS << ")"; } } diff --git a/clang/lib/Basic/OpenMPKinds.cpp b/clang/lib/Basic/OpenMPKinds.cpp index 03485b7e81abc..7ba2c89638c05 100644 --- a/clang/lib/Basic/OpenMPKinds.cpp +++ b/clang/lib/Basic/OpenMPKinds.cpp @@ -238,6 +238,16 @@ unsigned clang::getOpenMPSimpleClauseType(OpenMPClauseKind Kind, StringRef Str, return OMPC_NUMTHREADS_unknown; return Type; } + case OMPC_use_device_ptr: { + unsigned Type = llvm::StringSwitch<unsigned>(Str) +#define OPENMP_USE_DEVICE_PTR_FALLBACK_MODIFIER(Name) \ + .Case(#Name, OMPC_USE_DEVICE_PTR_FALLBACK_##Name) +#include "clang/Basic/OpenMPKinds.def" + .Default(OMPC_USE_DEVICE_PTR_FALLBACK_unknown); + if (LangOpts.OpenMP < 61) + return OMPC_USE_DEVICE_PTR_FALLBACK_unknown; + return Type; + } case OMPC_unknown: case OMPC_threadprivate: case OMPC_groupprivate: @@ -280,7 +290,6 @@ unsigned clang::getOpenMPSimpleClauseType(OpenMPClauseKind Kind, StringRef Str, case OMPC_nogroup: case OMPC_hint: case OMPC_uniform: - case OMPC_use_device_ptr: case OMPC_use_device_addr: case OMPC_is_device_ptr: case OMPC_has_device_addr: @@ -608,6 +617,16 @@ const char *clang::getOpenMPSimpleClauseTypeName(OpenMPClauseKind Kind, #include "clang/Basic/OpenMPKinds.def" } llvm_unreachable("Invalid OpenMP 'threadset' clause modifier"); + case OMPC_use_device_ptr: + switch (Type) { + case OMPC_USE_DEVICE_PTR_FALLBACK_unknown: + return "unknown"; +#define OPENMP_USE_DEVICE_PTR_FALLBACK_MODIFIER(Name) \ + case OMPC_USE_DEVICE_PTR_FALLBACK_##Name: \ + return #Name; +#include "clang/Basic/OpenMPKinds.def" + } + llvm_unreachable("Invalid OpenMP 'use_device_ptr' clause modifier"); case OMPC_unknown: case OMPC_threadprivate: case OMPC_groupprivate: @@ -650,7 +669,6 @@ const char *clang::getOpenMPSimpleClauseTypeName(OpenMPClauseKind Kind, case OMPC_nogroup: case OMPC_hint: case OMPC_uniform: - case OMPC_use_device_ptr: case OMPC_use_device_addr: case OMPC_is_device_ptr: case OMPC_has_device_addr: diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp index 15c3f7594bf44..7862723a2a19e 100644 --- a/clang/lib/Parse/ParseOpenMP.cpp +++ b/clang/lib/Parse/ParseOpenMP.cpp @@ -5053,6 +5053,24 @@ bool Parser::ParseOpenMPVarList(OpenMPDirectiveKind DKind, ExpectAndConsume(tok::colon, diag::warn_pragma_expected_colon, "adjust-op"); } + } else if (Kind == OMPC_use_device_ptr && getLangOpts().OpenMP >= 61) { + // Handle optional fallback modifier for use_device_ptr clause. + // use_device_ptr([fb_preserve | fb_nullify :] list) + // Default is fb_preserve. + if (Tok.is(tok::identifier)) { + auto FallbackModifier = static_cast<OpenMPUseDevicePtrFallbackModifier>( + getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok), getLangOpts())); + if (FallbackModifier != OMPC_USE_DEVICE_PTR_FALLBACK_unknown) { + Data.UseDevicePtrFallbackModifier = FallbackModifier; + Data.UseDevicePtrFallbackModifierLoc = Tok.getLocation(); + ConsumeToken(); + if (Tok.is(tok::colon)) { + Data.ColonLoc = ConsumeToken(); + } else { + Diag(Tok, diag::err_modifier_expected_colon) << "fallback"; + } + } + } } bool IsComma = diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp index 431c545c07e47..f28829a4a2b7b 100644 --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -18724,7 +18724,11 @@ OMPClause *SemaOpenMP::ActOnOpenMPVarListClause(OpenMPClauseKind Kind, VarList, Locs); break; case OMPC_use_device_ptr: - Res = ActOnOpenMPUseDevicePtrClause(VarList, Locs); + Res = ActOnOpenMPUseDevicePtrClause( + VarList, Locs, + static_cast<OpenMPUseDevicePtrFallbackModifier>( + Data.UseDevicePtrFallbackModifier), + Data.UseDevicePtrFallbackModifierLoc); break; case OMPC_use_device_addr: Res = ActOnOpenMPUseDeviceAddrClause(VarList, Locs); @@ -24539,9 +24543,9 @@ OMPClause *SemaOpenMP::ActOnOpenMPFromClause( MapperId); } -OMPClause * -SemaOpenMP::ActOnOpenMPUseDevicePtrClause(ArrayRef<Expr *> VarList, - const OMPVarListLocTy &Locs) { +OMPClause *SemaOpenMP::ActOnOpenMPUseDevicePtrClause( + ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs, + OpenMPUseDevicePtrFallbackModifier FallbackModifier, SourceLocation FallbackModifierLoc) { MappableVarListInfo MVLI(VarList); SmallVector<Expr *, 8> PrivateCopies; SmallVector<Expr *, 8> Inits; @@ -24622,7 +24626,7 @@ SemaOpenMP::ActOnOpenMPUseDevicePtrClause(ArrayRef<Expr *> VarList, return OMPUseDevicePtrClause::Create( getASTContext(), Locs, MVLI.ProcessedVarList, PrivateCopies, Inits, - MVLI.VarBaseDeclarations, MVLI.VarComponents); + MVLI.VarBaseDeclarations, MVLI.VarComponents, FallbackModifier, FallbackModifierLoc); } OMPClause * diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h index 8e5dbeb792348..f77267b11b75c 100644 --- a/clang/lib/Sema/TreeTransform.h +++ b/clang/lib/Sema/TreeTransform.h @@ -2252,9 +2252,12 @@ class TreeTransform { /// /// By default, performs semantic analysis to build the new OpenMP clause. /// Subclasses may override this routine to provide different behavior. - OMPClause *RebuildOMPUseDevicePtrClause(ArrayRef<Expr *> VarList, - const OMPVarListLocTy &Locs) { - return getSema().OpenMP().ActOnOpenMPUseDevicePtrClause(VarList, Locs); + OMPClause *RebuildOMPUseDevicePtrClause( + ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs, + OpenMPUseDevicePtrFallbackModifier FallbackModifier, + SourceLocation FallbackModifierLoc) { + return getSema().OpenMP().ActOnOpenMPUseDevicePtrClause( + VarList, Locs, FallbackModifier, FallbackModifierLoc); } /// Build a new OpenMP 'use_device_addr' clause. @@ -11591,7 +11594,8 @@ OMPClause *TreeTransform<Derived>::TransformOMPUseDevicePtrClause( Vars.push_back(EVar.get()); } OMPVarListLocTy Locs(C->getBeginLoc(), C->getLParenLoc(), C->getEndLoc()); - return getDerived().RebuildOMPUseDevicePtrClause(Vars, Locs); + return getDerived().RebuildOMPUseDevicePtrClause( + Vars, Locs, C->getFallbackModifier(), C->getFallbackModifierLoc()); } template <typename Derived> diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp index aec61322fb8be..27281996bf8f4 100644 --- a/clang/lib/Serialization/ASTReader.cpp +++ b/clang/lib/Serialization/ASTReader.cpp @@ -12503,6 +12503,8 @@ void OMPClauseReader::VisitOMPFromClause(OMPFromClause *C) { void OMPClauseReader::VisitOMPUseDevicePtrClause(OMPUseDevicePtrClause *C) { C->setLParenLoc(Record.readSourceLocation()); + C->setFallbackModifier(Record.readEnum<OpenMPUseDevicePtrFallbackModifier>()); + C->setFallbackModifierLoc(Record.readSourceLocation()); auto NumVars = C->varlist_size(); auto UniqueDecls = C->getUniqueDeclarationsNum(); auto TotalLists = C->getTotalComponentListNum(); diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp index 667e04049dac8..df7e34f94cb79 100644 --- a/clang/lib/Serialization/ASTWriter.cpp +++ b/clang/lib/Serialization/ASTWriter.cpp @@ -8487,6 +8487,8 @@ void OMPClauseWriter::VisitOMPUseDevicePtrClause(OMPUseDevicePtrClause *C) { Record.push_back(C->getTotalComponentListNum()); Record.push_back(C->getTotalComponentsNum()); Record.AddSourceLocation(C->getLParenLoc()); + Record.writeEnum(C->getFallbackModifier()); + Record.AddSourceLocation(C->getFallbackModifierLoc()); for (auto *E : C->varlist()) Record.AddStmt(E); for (auto *VE : C->private_copies()) >From 25017e6054efe88b90b528cd8f1702cf810f5350 Mon Sep 17 00:00:00 2001 From: Abhinav Gaba <[email protected]> Date: Wed, 3 Dec 2025 16:06:54 -0800 Subject: [PATCH 2/3] Add tests. --- ...data_use_device_ptr_fallback_ast_print.cpp | 36 +++++++++++++++++++ ..._data_use_device_ptr_fallback_messages.cpp | 28 +++++++++++++++ 2 files changed, 64 insertions(+) create mode 100644 clang/test/OpenMP/target_data_use_device_ptr_fallback_ast_print.cpp create mode 100644 clang/test/OpenMP/target_data_use_device_ptr_fallback_messages.cpp diff --git a/clang/test/OpenMP/target_data_use_device_ptr_fallback_ast_print.cpp b/clang/test/OpenMP/target_data_use_device_ptr_fallback_ast_print.cpp new file mode 100644 index 0000000000000..060f64f6e86a8 --- /dev/null +++ b/clang/test/OpenMP/target_data_use_device_ptr_fallback_ast_print.cpp @@ -0,0 +1,36 @@ +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=61 -ast-print %s | FileCheck %s +// RUN: %clang_cc1 -fopenmp -fopenmp-version=61 -x c++ -std=c++11 -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fopenmp-version=61 -std=c++11 -include-pch %t -verify %s -ast-print | FileCheck %s + +// expected-no-diagnostics + +#ifndef HEADER +#define HEADER + +// CHECK-LABEL:void f1(int *p, int *q) +void f1(int *p, int *q) { + +// CHECK: #pragma omp target data use_device_ptr(fb_preserve: p) +#pragma omp target data use_device_ptr(fb_preserve: p) + {} + +// CHECK: #pragma omp target data use_device_ptr(fb_nullify: p) +#pragma omp target data use_device_ptr(fb_nullify: p) + {} + +// Without any fallback modifier +// CHECK: #pragma omp target data use_device_ptr(p) +#pragma omp target data use_device_ptr(p) + {} + +// Multiple variables with fb_preserve +// CHECK: #pragma omp target data use_device_ptr(fb_preserve: p,q) +#pragma omp target data use_device_ptr(fb_preserve: p, q) + {} + +// Multiple variables with fb_nullify +// CHECK: #pragma omp target data use_device_ptr(fb_nullify: p,q) +#pragma omp target data use_device_ptr(fb_nullify: p, q) + {} +} +#endif diff --git a/clang/test/OpenMP/target_data_use_device_ptr_fallback_messages.cpp b/clang/test/OpenMP/target_data_use_device_ptr_fallback_messages.cpp new file mode 100644 index 0000000000000..7a22e95e7fee6 --- /dev/null +++ b/clang/test/OpenMP/target_data_use_device_ptr_fallback_messages.cpp @@ -0,0 +1,28 @@ +// RUN: %clang_cc1 -std=c++11 -fopenmp -fopenmp-version=60 -verify=omp60,expected -ferror-limit 200 %s +// RUN: %clang_cc1 -std=c++11 -fopenmp -fopenmp-version=61 -verify=omp61,expected -ferror-limit 200 %s + +void f1(int x, int *p, int *q) { + + // Test that fallback modifier is only recognized in OpenMP 6.1+ +#pragma omp target data map(x) use_device_ptr(fb_preserve: p) // omp60-error {{use of undeclared identifier 'fb_preserve'}} + {} + +#pragma omp target data map(x) use_device_ptr(fb_nullify: p) // omp60-error {{use of undeclared identifier 'fb_nullify'}} + {} + + // Without modifier (should work in both versions) +#pragma omp target data map(x) use_device_ptr(p) + {} + + // Unknown modifier: should fail in both versions +#pragma omp target data map(x) use_device_ptr(fb_abc: p) // expected-error {{use of undeclared identifier 'fb_abc'}} + {} + + // Multiple modifiers: should fail in both versions +#pragma omp target data map(x) use_device_ptr(fb_nullify, fb_preserve: p, q) // omp61-error {{missing ':' after fallback modifier}} omp61-error {{expected expression}} omp61-error {{use of undeclared identifier 'fb_preserve'}} omp60-error {{use of undeclared identifier 'fb_nullify'}} omp60-error {{use of undeclared identifier 'fb_preserve'}} + {} + + // Test missing colon after modifier in OpenMP 6.1 - should error +#pragma omp target data map(x) use_device_ptr(fb_preserve p) // omp61-error {{missing ':' after fallback modifier}} omp60-error {{use of undeclared identifier 'fb_preserve'}} + {} +} \ No newline at end of file >From d2f526deccca51e2ac1a5841a41524b7b3e43cc4 Mon Sep 17 00:00:00 2001 From: Abhinav Gaba <[email protected]> Date: Wed, 3 Dec 2025 16:32:12 -0800 Subject: [PATCH 3/3] Minor formatting changes. --- clang/include/clang/AST/OpenMPClause.h | 11 ++++++----- clang/include/clang/Sema/SemaOpenMP.h | 6 ++++-- clang/lib/Parse/ParseOpenMP.cpp | 5 ++--- clang/lib/Sema/SemaOpenMP.cpp | 6 ++++-- .../target_data_use_device_ptr_fallback_messages.cpp | 2 +- 5 files changed, 17 insertions(+), 13 deletions(-) diff --git a/clang/include/clang/AST/OpenMPClause.h b/clang/include/clang/AST/OpenMPClause.h index 2d63169021347..49896a86c0686 100644 --- a/clang/include/clang/AST/OpenMPClause.h +++ b/clang/include/clang/AST/OpenMPClause.h @@ -8008,12 +8008,13 @@ class OMPUseDevicePtrClause final /// NumComponents: total number of expression components in the clause. /// \param FallbackModifier The fallback modifier for the clause. /// \param FallbackModifierLoc Location of the fallback modifier. - explicit OMPUseDevicePtrClause(const OMPVarListLocTy &Locs, - const OMPMappableExprListSizeTy &Sizes, - OpenMPUseDevicePtrFallbackModifier FallbackModifier, - SourceLocation FallbackModifierLoc) + explicit OMPUseDevicePtrClause( + const OMPVarListLocTy &Locs, const OMPMappableExprListSizeTy &Sizes, + OpenMPUseDevicePtrFallbackModifier FallbackModifier, + SourceLocation FallbackModifierLoc) : OMPMappableExprListClause(llvm::omp::OMPC_use_device_ptr, Locs, Sizes), - FallbackModifier(FallbackModifier), FallbackModifierLoc(FallbackModifierLoc) {} + FallbackModifier(FallbackModifier), + FallbackModifierLoc(FallbackModifierLoc) {} /// Build an empty clause. /// diff --git a/clang/include/clang/Sema/SemaOpenMP.h b/clang/include/clang/Sema/SemaOpenMP.h index e4eb3345534a4..1d4ea0f1cf3b0 100644 --- a/clang/include/clang/Sema/SemaOpenMP.h +++ b/clang/include/clang/Sema/SemaOpenMP.h @@ -1177,7 +1177,8 @@ class SemaOpenMP : public SemaBase { int NeedDevicePtrModifier = 0; SourceLocation NeedDevicePtrModifierLoc; int UseDevicePtrFallbackModifier = - OMPC_USE_DEVICE_PTR_FALLBACK_unknown; ///< Fallback modifier for use_device_ptr clause. + OMPC_USE_DEVICE_PTR_FALLBACK_unknown; ///< Fallback modifier for + ///< use_device_ptr clause. SourceLocation UseDevicePtrFallbackModifierLoc; SmallVector<OpenMPMapModifierKind, NumberOfOMPMapClauseModifiers> MapTypeModifiers; @@ -1369,7 +1370,8 @@ class SemaOpenMP : public SemaBase { /// Called on well-formed 'use_device_ptr' clause. OMPClause *ActOnOpenMPUseDevicePtrClause( ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs, - OpenMPUseDevicePtrFallbackModifier FallbackModifier, SourceLocation FallbackModifierLoc); + OpenMPUseDevicePtrFallbackModifier FallbackModifier, + SourceLocation FallbackModifierLoc); /// Called on well-formed 'use_device_addr' clause. OMPClause *ActOnOpenMPUseDeviceAddrClause(ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs); diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp index 7862723a2a19e..b8548f24de1e0 100644 --- a/clang/lib/Parse/ParseOpenMP.cpp +++ b/clang/lib/Parse/ParseOpenMP.cpp @@ -5064,11 +5064,10 @@ bool Parser::ParseOpenMPVarList(OpenMPDirectiveKind DKind, Data.UseDevicePtrFallbackModifier = FallbackModifier; Data.UseDevicePtrFallbackModifierLoc = Tok.getLocation(); ConsumeToken(); - if (Tok.is(tok::colon)) { + if (Tok.is(tok::colon)) Data.ColonLoc = ConsumeToken(); - } else { + else Diag(Tok, diag::err_modifier_expected_colon) << "fallback"; - } } } } diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp index f28829a4a2b7b..7a998546c6134 100644 --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -24545,7 +24545,8 @@ OMPClause *SemaOpenMP::ActOnOpenMPFromClause( OMPClause *SemaOpenMP::ActOnOpenMPUseDevicePtrClause( ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs, - OpenMPUseDevicePtrFallbackModifier FallbackModifier, SourceLocation FallbackModifierLoc) { + OpenMPUseDevicePtrFallbackModifier FallbackModifier, + SourceLocation FallbackModifierLoc) { MappableVarListInfo MVLI(VarList); SmallVector<Expr *, 8> PrivateCopies; SmallVector<Expr *, 8> Inits; @@ -24626,7 +24627,8 @@ OMPClause *SemaOpenMP::ActOnOpenMPUseDevicePtrClause( return OMPUseDevicePtrClause::Create( getASTContext(), Locs, MVLI.ProcessedVarList, PrivateCopies, Inits, - MVLI.VarBaseDeclarations, MVLI.VarComponents, FallbackModifier, FallbackModifierLoc); + MVLI.VarBaseDeclarations, MVLI.VarComponents, FallbackModifier, + FallbackModifierLoc); } OMPClause * diff --git a/clang/test/OpenMP/target_data_use_device_ptr_fallback_messages.cpp b/clang/test/OpenMP/target_data_use_device_ptr_fallback_messages.cpp index 7a22e95e7fee6..fff2dcf15e29e 100644 --- a/clang/test/OpenMP/target_data_use_device_ptr_fallback_messages.cpp +++ b/clang/test/OpenMP/target_data_use_device_ptr_fallback_messages.cpp @@ -25,4 +25,4 @@ void f1(int x, int *p, int *q) { // Test missing colon after modifier in OpenMP 6.1 - should error #pragma omp target data map(x) use_device_ptr(fb_preserve p) // omp61-error {{missing ':' after fallback modifier}} omp60-error {{use of undeclared identifier 'fb_preserve'}} {} -} \ No newline at end of file +} _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
