https://github.com/abhinavgaba created
https://github.com/llvm/llvm-project/pull/170578
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.
>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/2] [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/2] 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
_______________________________________________
llvm-branch-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
