aboutsummaryrefslogtreecommitdiff
path: root/clang
diff options
context:
space:
mode:
authorErich Keane <ekeane@nvidia.com>2023-11-17 06:29:02 -0800
committerGitHub <noreply@github.com>2023-11-17 06:29:02 -0800
commitff219ea9ca80f46ff85dbdb94622ffb319a0d237 (patch)
tree17213428b3e4b11cc4b516b42f02cb7f64f2cb60 /clang
parent9c0e64999b23046d0b8987a48ddc41a4c6129f9d (diff)
downloadllvm-ff219ea9ca80f46ff85dbdb94622ffb319a0d237.zip
llvm-ff219ea9ca80f46ff85dbdb94622ffb319a0d237.tar.gz
llvm-ff219ea9ca80f46ff85dbdb94622ffb319a0d237.tar.bz2
[OpenACC] Initial commits to support OpenACC (#70234)
Initial commits to support OpenACC. This patchset: adds a clang-command line argument '-fopenacc', and starts to define _OPENACC, albeit to '1' instead of the standardized value (since we don't properly implement OpenACC yet). The OpenACC spec defines `_OPENACC` to be equal to the latest standard implemented. However, since we're not done implementing any standard, we've defined this by default to be `1`. As it is useful to run our compiler against existing OpenACC workloads, we're providing a temporary override flag to change the `_OPENACC` value to be any entirely digit value, permitting testing against any existing OpenACC project. Exactly like the OpenMP parser, the OpenACC pragma parser needs to consume and reprocess the tokens. This patch sets up the infrastructure to do so by refactoring the OpenMP version of this into a more general version that works for OpenACC as well. Additionally, this adds a few diagnostics and token kinds to get us started.
Diffstat (limited to 'clang')
-rw-r--r--clang/docs/ReleaseNotes.rst13
-rw-r--r--clang/include/clang/Basic/DiagnosticGroups.td4
-rw-r--r--clang/include/clang/Basic/DiagnosticParseKinds.td9
-rw-r--r--clang/include/clang/Basic/LangOptions.def2
-rw-r--r--clang/include/clang/Basic/LangOptions.h5
-rw-r--r--clang/include/clang/Basic/TokenKinds.def6
-rw-r--r--clang/include/clang/Driver/Options.td23
-rw-r--r--clang/include/clang/Parse/Parser.h10
-rw-r--r--clang/lib/Driver/ToolChains/Clang.cpp20
-rw-r--r--clang/lib/Frontend/CompilerInvocation.cpp15
-rw-r--r--clang/lib/Frontend/InitPreprocessor.cpp11
-rw-r--r--clang/lib/Parse/CMakeLists.txt1
-rw-r--r--clang/lib/Parse/ParseDecl.cpp5
-rw-r--r--clang/lib/Parse/ParseDeclCXX.cpp2
-rw-r--r--clang/lib/Parse/ParseOpenACC.cpp27
-rw-r--r--clang/lib/Parse/ParsePragma.cpp92
-rw-r--r--clang/lib/Parse/ParseStmt.cpp3
-rw-r--r--clang/lib/Parse/Parser.cpp8
-rw-r--r--clang/lib/Serialization/ASTReader.cpp2
-rw-r--r--clang/lib/Serialization/ASTWriter.cpp2
-rw-r--r--clang/test/Driver/openacc.c14
-rw-r--r--clang/test/ParserOpenACC/disabled.c4
-rw-r--r--clang/test/ParserOpenACC/unimplemented.c20
-rw-r--r--clang/test/ParserOpenACC/unimplemented.cpp20
-rw-r--r--clang/test/Preprocessor/openacc.c13
25 files changed, 304 insertions, 27 deletions
diff --git a/clang/docs/ReleaseNotes.rst b/clang/docs/ReleaseNotes.rst
index 739831d..08ffb08 100644
--- a/clang/docs/ReleaseNotes.rst
+++ b/clang/docs/ReleaseNotes.rst
@@ -241,6 +241,8 @@ New Compiler Flags
handlers will be smaller. A throw expression of a type with a
potentially-throwing destructor will lead to an error.
+* ``-fopenacc`` was added as a part of the effort to support OpenACC in clang.
+
Deprecated Compiler Flags
-------------------------
@@ -731,6 +733,17 @@ Miscellaneous Clang Crashes Fixed
- Fixed a crash when ``-ast-dump=json`` was used for code using class
template deduction guides.
+OpenACC Specific Changes
+------------------------
+- OpenACC Implementation effort is beginning with semantic analysis and parsing
+ of OpenACC pragmas. The ``-fopenacc`` flag was added to enable these new,
+ albeit incomplete changes. The ``_OPENACC`` macro is currently defined to
+ ``1``, as support is too incomplete to update to a standards-required value.
+- Added ``-fexperimental-openacc-macro-override``, a command line option to
+ permit overriding the ``_OPENACC`` macro to be any digit-only value specified
+ by the user, which permits testing the compiler against existing OpenACC
+ workloads in order to evaluate implementation progress.
+
Target Specific Changes
-----------------------
diff --git a/clang/include/clang/Basic/DiagnosticGroups.td b/clang/include/clang/Basic/DiagnosticGroups.td
index 37559c7ff..ff028bb 100644
--- a/clang/include/clang/Basic/DiagnosticGroups.td
+++ b/clang/include/clang/Basic/DiagnosticGroups.td
@@ -1315,6 +1315,10 @@ def OpenMP : DiagGroup<"openmp", [
OpenMPMapping, OpenMP51Ext, OpenMPExtensions, OpenMPTargetException
]>;
+// OpenACC warnings.
+def SourceUsesOpenACC : DiagGroup<"source-uses-openacc">;
+def OpenACC : DiagGroup<"openacc", [SourceUsesOpenACC]>;
+
// Backend warnings.
def BackendInlineAsm : DiagGroup<"inline-asm">;
def BackendSourceMgr : DiagGroup<"source-mgr">;
diff --git a/clang/include/clang/Basic/DiagnosticParseKinds.td b/clang/include/clang/Basic/DiagnosticParseKinds.td
index de18034..c3d0605 100644
--- a/clang/include/clang/Basic/DiagnosticParseKinds.td
+++ b/clang/include/clang/Basic/DiagnosticParseKinds.td
@@ -1342,6 +1342,15 @@ def err_opencl_logical_exclusive_or : Error<
def err_openclcxx_virtual_function : Error<
"virtual functions are not supported in C++ for OpenCL">;
+// OpenACC Support.
+def warn_pragma_acc_ignored : Warning<
+ "unexpected '#pragma acc ...' in program">, InGroup<SourceUsesOpenACC>, DefaultIgnore;
+def err_acc_unexpected_directive : Error<
+ "unexpected OpenACC directive %select{|'#pragma acc %1'}0">;
+def warn_pragma_acc_unimplemented
+ : Warning<"OpenACC directives not yet implemented, pragma ignored">,
+ InGroup<SourceUsesOpenACC>;
+
// OpenMP support.
def warn_pragma_omp_ignored : Warning<
"unexpected '#pragma omp ...' in program">, InGroup<SourceUsesOpenMP>, DefaultIgnore;
diff --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def
index 412b116..cd77b22 100644
--- a/clang/include/clang/Basic/LangOptions.def
+++ b/clang/include/clang/Basic/LangOptions.def
@@ -285,6 +285,8 @@ LANGOPT(OffloadUniformBlock, 1, 0, "Assume that kernels are launched with unifor
LANGOPT(HIPStdPar, 1, 0, "Enable Standard Parallel Algorithm Acceleration for HIP (experimental)")
LANGOPT(HIPStdParInterposeAlloc, 1, 0, "Replace allocations / deallocations with HIP RT calls when Standard Parallel Algorithm Acceleration for HIP is enabled (Experimental)")
+LANGOPT(OpenACC , 1, 0, "OpenACC Enabled")
+
LANGOPT(SizedDeallocation , 1, 0, "sized deallocation")
LANGOPT(AlignedAllocation , 1, 0, "aligned allocation")
LANGOPT(AlignedAllocationUnavailable, 1, 0, "aligned allocation functions are unavailable")
diff --git a/clang/include/clang/Basic/LangOptions.h b/clang/include/clang/Basic/LangOptions.h
index ae99357..2d167dd 100644
--- a/clang/include/clang/Basic/LangOptions.h
+++ b/clang/include/clang/Basic/LangOptions.h
@@ -502,6 +502,11 @@ public:
// received as a result of a standard operator new (-fcheck-new)
bool CheckNew = false;
+ // In OpenACC mode, contains a user provided override for the _OPENACC macro.
+ // This exists so that we can override the macro value and test our incomplete
+ // implementation on real-world examples.
+ std::string OpenACCMacroOverride;
+
LangOptions();
/// Set language defaults for the given input language and
diff --git a/clang/include/clang/Basic/TokenKinds.def b/clang/include/clang/Basic/TokenKinds.def
index 82a503d..3ab4208 100644
--- a/clang/include/clang/Basic/TokenKinds.def
+++ b/clang/include/clang/Basic/TokenKinds.def
@@ -946,6 +946,12 @@ ANNOTATION(attr_openmp)
PRAGMA_ANNOTATION(pragma_openmp)
PRAGMA_ANNOTATION(pragma_openmp_end)
+// Annotations for OpenACC pragma directives - #pragma acc.
+// Like with OpenMP, these are produced by the lexer when it parses a
+// #pragma acc directive so it can be handled during parsing of the directives.
+PRAGMA_ANNOTATION(pragma_openacc)
+PRAGMA_ANNOTATION(pragma_openacc_end)
+
// Annotations for loop pragma directives #pragma clang loop ...
// The lexer produces these so that they only take effect when the parser
// handles #pragma loop ... directives.
diff --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td
index 8115504..df12ba8 100644
--- a/clang/include/clang/Driver/Options.td
+++ b/clang/include/clang/Driver/Options.td
@@ -1359,6 +1359,19 @@ def fno_hip_emit_relocatable : Flag<["-"], "fno-hip-emit-relocatable">,
HelpText<"Do not override toolchain to compile HIP source to relocatable">;
}
+// Clang specific/exclusive options for OpenACC.
+def openacc_macro_override
+ : Separate<["-"], "fexperimental-openacc-macro-override">,
+ Visibility<[ClangOption, CC1Option]>,
+ Group<f_Group>,
+ HelpText<"Overrides the _OPENACC macro value for experimental testing "
+ "during OpenACC support development">;
+def openacc_macro_override_EQ
+ : Joined<["-"], "fexperimental-openacc-macro-override=">,
+ Alias<openacc_macro_override>;
+
+// End Clang specific/exclusive options for OpenACC.
+
def libomptarget_amdgpu_bc_path_EQ : Joined<["--"], "libomptarget-amdgpu-bc-path=">, Group<i_Group>,
HelpText<"Path to libomptarget-amdgcn bitcode library">;
def libomptarget_amdgcn_bc_path_EQ : Joined<["--"], "libomptarget-amdgcn-bc-path=">, Group<i_Group>,
@@ -3343,6 +3356,14 @@ def fno_openmp_target_debug : Flag<["-"], "fno-openmp-target-debug">;
} // let Flags = [NoArgumentUnused]
//===----------------------------------------------------------------------===//
+// FlangOption + FC1 + ClangOption + CC1Option
+//===----------------------------------------------------------------------===//
+let Visibility = [FC1Option, FlangOption, CC1Option, ClangOption] in {
+def fopenacc : Flag<["-"], "fopenacc">, Group<f_Group>,
+ HelpText<"Enable OpenACC">;
+} // let Visibility = [FC1Option, FlangOption, CC1Option, ClangOption]
+
+//===----------------------------------------------------------------------===//
// Optimisation remark options
//===----------------------------------------------------------------------===//
@@ -6266,8 +6287,6 @@ file}]>;
def ffixed_line_length_VALUE : Joined<["-"], "ffixed-line-length-">, Group<f_Group>, Alias<ffixed_line_length_EQ>;
def fconvert_EQ : Joined<["-"], "fconvert=">, Group<f_Group>,
HelpText<"Set endian conversion of data for unformatted files">;
-def fopenacc : Flag<["-"], "fopenacc">, Group<f_Group>,
- HelpText<"Enable OpenACC">;
def fdefault_double_8 : Flag<["-"],"fdefault-double-8">, Group<f_Group>,
HelpText<"Set the default double precision kind to an 8 byte wide type">;
def fdefault_integer_8 : Flag<["-"],"fdefault-integer-8">, Group<f_Group>,
diff --git a/clang/include/clang/Parse/Parser.h b/clang/include/clang/Parse/Parser.h
index 30e0352..4631e9a 100644
--- a/clang/include/clang/Parse/Parser.h
+++ b/clang/include/clang/Parse/Parser.h
@@ -175,6 +175,7 @@ class Parser : public CodeCompletionHandler {
std::unique_ptr<PragmaHandler> FPContractHandler;
std::unique_ptr<PragmaHandler> OpenCLExtensionHandler;
std::unique_ptr<PragmaHandler> OpenMPHandler;
+ std::unique_ptr<PragmaHandler> OpenACCHandler;
std::unique_ptr<PragmaHandler> PCSectionHandler;
std::unique_ptr<PragmaHandler> MSCommentHandler;
std::unique_ptr<PragmaHandler> MSDetectMismatchHandler;
@@ -3524,6 +3525,15 @@ public:
/// where, map-type-modifier ::= always | close | mapper(mapper-identifier)
bool parseMapTypeModifiers(Sema::OpenMPVarListDataTy &Data);
+ //===--------------------------------------------------------------------===//
+ // OpenACC Parsing.
+
+ /// Placeholder for now, should just ignore the directives after emitting a
+ /// diagnostic. Eventually will be split into a few functions to parse
+ /// different situations.
+ DeclGroupPtrTy ParseOpenACCDirective();
+ StmtResult ParseOpenACCDirectiveStmt();
+
private:
//===--------------------------------------------------------------------===//
// C++ 14: Templates [temp]
diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index b462f5a..6dec117 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -3633,6 +3633,23 @@ static void RenderHLSLOptions(const ArgList &Args, ArgStringList &CmdArgs,
CmdArgs.push_back("-finclude-default-header");
}
+static void RenderOpenACCOptions(const Driver &D, const ArgList &Args,
+ ArgStringList &CmdArgs, types::ID InputType) {
+ if (!Args.hasArg(options::OPT_fopenacc))
+ return;
+
+ CmdArgs.push_back("-fopenacc");
+
+ if (Arg *A = Args.getLastArg(options::OPT_openacc_macro_override)) {
+ StringRef Value = A->getValue();
+ int Version;
+ if (!Value.getAsInteger(10, Version))
+ A->renderAsInput(Args, CmdArgs);
+ else
+ D.Diag(diag::err_drv_clang_unsupported) << Value;
+ }
+}
+
static void RenderARCMigrateToolOptions(const Driver &D, const ArgList &Args,
ArgStringList &CmdArgs) {
bool ARCMTEnabled = false;
@@ -6623,6 +6640,9 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
// Forward hlsl options to -cc1
RenderHLSLOptions(Args, CmdArgs, InputType);
+ // Forward OpenACC options to -cc1
+ RenderOpenACCOptions(D, Args, CmdArgs, InputType);
+
if (IsHIP) {
if (Args.hasFlag(options::OPT_fhip_new_launch_api,
options::OPT_fno_hip_new_launch_api, true))
diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp
index efcf073..3f4ca02 100644
--- a/clang/lib/Frontend/CompilerInvocation.cpp
+++ b/clang/lib/Frontend/CompilerInvocation.cpp
@@ -3549,6 +3549,13 @@ void CompilerInvocationBase::GenerateLangArgs(const LangOptions &Opts,
if (Opts.OpenMPCUDAMode)
GenerateArg(Consumer, OPT_fopenmp_cuda_mode);
+ if (Opts.OpenACC) {
+ GenerateArg(Consumer, OPT_fopenacc);
+ if (!Opts.OpenACCMacroOverride.empty())
+ GenerateArg(Consumer, OPT_openacc_macro_override,
+ Opts.OpenACCMacroOverride);
+ }
+
// The arguments used to set Optimize, OptimizeSize and NoInlineDefine are
// generated from CodeGenOptions.
@@ -4018,6 +4025,14 @@ bool CompilerInvocation::ParseLangArgs(LangOptions &Opts, ArgList &Args,
(T.isNVPTX() || T.isAMDGCN()) &&
Args.hasArg(options::OPT_fopenmp_cuda_mode);
+ // OpenACC Configuration.
+ if (Args.hasArg(options::OPT_fopenacc)) {
+ Opts.OpenACC = true;
+
+ if (Arg *A = Args.getLastArg(options::OPT_openacc_macro_override))
+ Opts.OpenACCMacroOverride = A->getValue();
+ }
+
// FIXME: Eliminate this dependency.
unsigned Opt = getOptimizationLevel(Args, IK, Diags),
OptSize = getOptimizationLevelSize(Args);
diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp
index 846e5fc..17948dc 100644
--- a/clang/lib/Frontend/InitPreprocessor.cpp
+++ b/clang/lib/Frontend/InitPreprocessor.cpp
@@ -605,6 +605,17 @@ static void InitializeStandardPredefinedMacros(const TargetInfo &TI,
Builder.defineMacro("HIP_API_PER_THREAD_DEFAULT_STREAM");
}
}
+
+ if (LangOpts.OpenACC) {
+ // FIXME: When we have full support for OpenACC, we should set this to the
+ // version we support. Until then, set as '1' by default, but provide a
+ // temporary mechanism for users to override this so real-world examples can
+ // be tested against.
+ if (!LangOpts.OpenACCMacroOverride.empty())
+ Builder.defineMacro("_OPENACC", LangOpts.OpenACCMacroOverride);
+ else
+ Builder.defineMacro("_OPENACC", "1");
+ }
}
/// Initialize the predefined C++ language feature test macros defined in
diff --git a/clang/lib/Parse/CMakeLists.txt b/clang/lib/Parse/CMakeLists.txt
index 5a20e9d..22e902f 100644
--- a/clang/lib/Parse/CMakeLists.txt
+++ b/clang/lib/Parse/CMakeLists.txt
@@ -23,6 +23,7 @@ add_clang_library(clangParse
ParseTemplate.cpp
ParseTentative.cpp
Parser.cpp
+ ParseOpenACC.cpp
LINK_LIBS
clangAST
diff --git a/clang/lib/Parse/ParseDecl.cpp b/clang/lib/Parse/ParseDecl.cpp
index 06f2c87..d86f477 100644
--- a/clang/lib/Parse/ParseDecl.cpp
+++ b/clang/lib/Parse/ParseDecl.cpp
@@ -4734,6 +4734,11 @@ void Parser::ParseStructUnionBody(SourceLocation RecordLoc,
continue;
}
+ if (Tok.is(tok::annot_pragma_openacc)) {
+ ParseOpenACCDirective();
+ continue;
+ }
+
if (tok::isPragmaAnnotation(Tok.getKind())) {
Diag(Tok.getLocation(), diag::err_pragma_misplaced_in_decl)
<< DeclSpec::getSpecifierName(
diff --git a/clang/lib/Parse/ParseDeclCXX.cpp b/clang/lib/Parse/ParseDeclCXX.cpp
index 35b1a93..e12215d 100644
--- a/clang/lib/Parse/ParseDeclCXX.cpp
+++ b/clang/lib/Parse/ParseDeclCXX.cpp
@@ -3429,6 +3429,8 @@ Parser::DeclGroupPtrTy Parser::ParseCXXClassMemberDeclarationWithPragmas(
case tok::annot_pragma_openmp:
return ParseOpenMPDeclarativeDirectiveWithExtDecl(
AS, AccessAttrs, /*Delayed=*/true, TagType, TagDecl);
+ case tok::annot_pragma_openacc:
+ return ParseOpenACCDirective();
default:
if (tok::isPragmaAnnotation(Tok.getKind())) {
diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
new file mode 100644
index 0000000..2fba6cd
--- /dev/null
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -0,0 +1,27 @@
+//===--- ParseOpenACC.cpp - OpenACC-specific parsing support --------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file implements the parsing logic for OpenACC language features.
+//
+//===----------------------------------------------------------------------===//
+
+#include "clang/Parse/ParseDiagnostic.h"
+#include "clang/Parse/Parser.h"
+
+using namespace clang;
+
+Parser::DeclGroupPtrTy Parser::ParseOpenACCDirective() {
+ Diag(Tok, diag::warn_pragma_acc_unimplemented);
+ SkipUntil(tok::annot_pragma_openacc_end);
+ return nullptr;
+}
+StmtResult Parser::ParseOpenACCDirectiveStmt() {
+ Diag(Tok, diag::warn_pragma_acc_unimplemented);
+ SkipUntil(tok::annot_pragma_openacc_end);
+ return StmtEmpty();
+}
diff --git a/clang/lib/Parse/ParsePragma.cpp b/clang/lib/Parse/ParsePragma.cpp
index b3178ae..d3fb7fc 100644
--- a/clang/lib/Parse/ParsePragma.cpp
+++ b/clang/lib/Parse/ParsePragma.cpp
@@ -166,18 +166,51 @@ struct PragmaFPHandler : public PragmaHandler {
Token &FirstToken) override;
};
-struct PragmaNoOpenMPHandler : public PragmaHandler {
- PragmaNoOpenMPHandler() : PragmaHandler("omp") { }
+// A pragma handler to be the base of the NoOpenMPHandler and NoOpenACCHandler,
+// which are identical other than the name given to them, and the diagnostic
+// emitted.
+template <diag::kind IgnoredDiag>
+struct PragmaNoSupportHandler : public PragmaHandler {
+ PragmaNoSupportHandler(StringRef Name) : PragmaHandler(Name) {}
void HandlePragma(Preprocessor &PP, PragmaIntroducer Introducer,
Token &FirstToken) override;
};
-struct PragmaOpenMPHandler : public PragmaHandler {
- PragmaOpenMPHandler() : PragmaHandler("omp") { }
+struct PragmaNoOpenMPHandler
+ : public PragmaNoSupportHandler<diag::warn_pragma_omp_ignored> {
+ PragmaNoOpenMPHandler() : PragmaNoSupportHandler("omp") {}
+};
+
+struct PragmaNoOpenACCHandler
+ : public PragmaNoSupportHandler<diag::warn_pragma_acc_ignored> {
+ PragmaNoOpenACCHandler() : PragmaNoSupportHandler("acc") {}
+};
+
+// A pragma handler to be the base for the OpenMPHandler and OpenACCHandler,
+// which are identical other than the tokens used for the start/end of a pragma
+// section, and some diagnostics.
+template <tok::TokenKind StartTok, tok::TokenKind EndTok,
+ diag::kind UnexpectedDiag>
+struct PragmaSupportHandler : public PragmaHandler {
+ PragmaSupportHandler(StringRef Name) : PragmaHandler(Name) {}
void HandlePragma(Preprocessor &PP, PragmaIntroducer Introducer,
Token &FirstToken) override;
};
+struct PragmaOpenMPHandler
+ : public PragmaSupportHandler<tok::annot_pragma_openmp,
+ tok::annot_pragma_openmp_end,
+ diag::err_omp_unexpected_directive> {
+ PragmaOpenMPHandler() : PragmaSupportHandler("omp") {}
+};
+
+struct PragmaOpenACCHandler
+ : public PragmaSupportHandler<tok::annot_pragma_openacc,
+ tok::annot_pragma_openacc_end,
+ diag::err_acc_unexpected_directive> {
+ PragmaOpenACCHandler() : PragmaSupportHandler("acc") {}
+};
+
/// PragmaCommentHandler - "\#pragma comment ...".
struct PragmaCommentHandler : public PragmaHandler {
PragmaCommentHandler(Sema &Actions)
@@ -423,6 +456,12 @@ void Parser::initializePragmaHandlers() {
OpenMPHandler = std::make_unique<PragmaNoOpenMPHandler>();
PP.AddPragmaHandler(OpenMPHandler.get());
+ if (getLangOpts().OpenACC)
+ OpenACCHandler = std::make_unique<PragmaOpenACCHandler>();
+ else
+ OpenACCHandler = std::make_unique<PragmaNoOpenACCHandler>();
+ PP.AddPragmaHandler(OpenACCHandler.get());
+
if (getLangOpts().MicrosoftExt ||
getTargetInfo().getTriple().isOSBinFormatELF()) {
MSCommentHandler = std::make_unique<PragmaCommentHandler>(Actions);
@@ -542,6 +581,9 @@ void Parser::resetPragmaHandlers() {
PP.RemovePragmaHandler(OpenMPHandler.get());
OpenMPHandler.reset();
+ PP.RemovePragmaHandler(OpenACCHandler.get());
+ OpenACCHandler.reset();
+
if (getLangOpts().MicrosoftExt ||
getTargetInfo().getTriple().isOSBinFormatELF()) {
PP.RemovePragmaHandler(MSCommentHandler.get());
@@ -2610,42 +2652,42 @@ void PragmaOpenCLExtensionHandler::HandlePragma(Preprocessor &PP,
StateLoc, State);
}
-/// Handle '#pragma omp ...' when OpenMP is disabled.
-///
-void PragmaNoOpenMPHandler::HandlePragma(Preprocessor &PP,
- PragmaIntroducer Introducer,
- Token &FirstTok) {
- if (!PP.getDiagnostics().isIgnored(diag::warn_pragma_omp_ignored,
- FirstTok.getLocation())) {
- PP.Diag(FirstTok, diag::warn_pragma_omp_ignored);
- PP.getDiagnostics().setSeverity(diag::warn_pragma_omp_ignored,
- diag::Severity::Ignored, SourceLocation());
+/// Handle '#pragma omp ...' when OpenMP is disabled and '#pragma acc ...' when
+/// OpenACC is disabled.
+template <diag::kind IgnoredDiag>
+void PragmaNoSupportHandler<IgnoredDiag>::HandlePragma(
+ Preprocessor &PP, PragmaIntroducer Introducer, Token &FirstTok) {
+ if (!PP.getDiagnostics().isIgnored(IgnoredDiag, FirstTok.getLocation())) {
+ PP.Diag(FirstTok, IgnoredDiag);
+ PP.getDiagnostics().setSeverity(IgnoredDiag, diag::Severity::Ignored,
+ SourceLocation());
}
PP.DiscardUntilEndOfDirective();
}
-/// Handle '#pragma omp ...' when OpenMP is enabled.
-///
-void PragmaOpenMPHandler::HandlePragma(Preprocessor &PP,
- PragmaIntroducer Introducer,
- Token &FirstTok) {
+/// Handle '#pragma omp ...' when OpenMP is enabled, and handle '#pragma acc...'
+/// when OpenACC is enabled.
+template <tok::TokenKind StartTok, tok::TokenKind EndTok,
+ diag::kind UnexpectedDiag>
+void PragmaSupportHandler<StartTok, EndTok, UnexpectedDiag>::HandlePragma(
+ Preprocessor &PP, PragmaIntroducer Introducer, Token &FirstTok) {
SmallVector<Token, 16> Pragma;
Token Tok;
Tok.startToken();
- Tok.setKind(tok::annot_pragma_openmp);
+ Tok.setKind(StartTok);
Tok.setLocation(Introducer.Loc);
while (Tok.isNot(tok::eod) && Tok.isNot(tok::eof)) {
Pragma.push_back(Tok);
PP.Lex(Tok);
- if (Tok.is(tok::annot_pragma_openmp)) {
- PP.Diag(Tok, diag::err_omp_unexpected_directive) << 0;
+ if (Tok.is(StartTok)) {
+ PP.Diag(Tok, UnexpectedDiag) << 0;
unsigned InnerPragmaCnt = 1;
while (InnerPragmaCnt != 0) {
PP.Lex(Tok);
- if (Tok.is(tok::annot_pragma_openmp))
+ if (Tok.is(StartTok))
++InnerPragmaCnt;
- else if (Tok.is(tok::annot_pragma_openmp_end))
+ else if (Tok.is(EndTok))
--InnerPragmaCnt;
}
PP.Lex(Tok);
@@ -2653,7 +2695,7 @@ void PragmaOpenMPHandler::HandlePragma(Preprocessor &PP,
}
SourceLocation EodLoc = Tok.getLocation();
Tok.startToken();
- Tok.setKind(tok::annot_pragma_openmp_end);
+ Tok.setKind(EndTok);
Tok.setLocation(EodLoc);
Pragma.push_back(Tok);
diff --git a/clang/lib/Parse/ParseStmt.cpp b/clang/lib/Parse/ParseStmt.cpp
index 2531147..924f27d 100644
--- a/clang/lib/Parse/ParseStmt.cpp
+++ b/clang/lib/Parse/ParseStmt.cpp
@@ -475,6 +475,9 @@ Retry:
// Do not prohibit attributes if they were OpenMP attributes.
return ParseOpenMPDeclarativeOrExecutableDirective(StmtCtx);
+ case tok::annot_pragma_openacc:
+ return ParseOpenACCDirectiveStmt();
+
case tok::annot_pragma_ms_pointers_to_members:
ProhibitAttributes(CXX11Attrs);
ProhibitAttributes(GNUAttrs);
diff --git a/clang/lib/Parse/Parser.cpp b/clang/lib/Parse/Parser.cpp
index 176d214..41b7462 100644
--- a/clang/lib/Parse/Parser.cpp
+++ b/clang/lib/Parse/Parser.cpp
@@ -318,6 +318,12 @@ bool Parser::SkipUntil(ArrayRef<tok::TokenKind> Toks, SkipUntilFlags Flags) {
return false;
ConsumeAnnotationToken();
break;
+ case tok::annot_pragma_openacc:
+ case tok::annot_pragma_openacc_end:
+ // FIXME: Like OpenMP above, we should not be doing this if we're parsing
+ // an OpenACC Directive.
+ ConsumeAnnotationToken();
+ break;
case tok::annot_module_begin:
case tok::annot_module_end:
case tok::annot_module_include:
@@ -851,6 +857,8 @@ Parser::ParseExternalDeclaration(ParsedAttributes &Attrs,
AccessSpecifier AS = AS_none;
return ParseOpenMPDeclarativeDirectiveWithExtDecl(AS, Attrs);
}
+ case tok::annot_pragma_openacc:
+ return ParseOpenACCDirective();
case tok::annot_pragma_ms_pointers_to_members:
HandlePragmaMSPointersToMembers();
return nullptr;
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index 42b48d2..dbc2d0b 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -1819,6 +1819,8 @@ Token ASTReader::ReadToken(ModuleFile &F, const RecordDataImpl &Record,
case tok::annot_pragma_openmp:
case tok::annot_pragma_openmp_end:
case tok::annot_pragma_unused:
+ case tok::annot_pragma_openacc:
+ case tok::annot_pragma_openacc_end:
break;
default:
llvm_unreachable("missing deserialization code for annotation token");
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index 0161ad1..5ebeb10 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -4501,6 +4501,8 @@ void ASTWriter::AddToken(const Token &Tok, RecordDataImpl &Record) {
case tok::annot_pragma_openmp:
case tok::annot_pragma_openmp_end:
case tok::annot_pragma_unused:
+ case tok::annot_pragma_openacc:
+ case tok::annot_pragma_openacc_end:
break;
default:
llvm_unreachable("missing serialization code for annotation token");
diff --git a/clang/test/Driver/openacc.c b/clang/test/Driver/openacc.c
new file mode 100644
index 0000000..c7f1d25
--- /dev/null
+++ b/clang/test/Driver/openacc.c
@@ -0,0 +1,14 @@
+// RUN: %clang -S -### -fopenacc %s 2>&1 | FileCheck %s --check-prefix=CHECK-DRIVER
+// CHECK-DRIVER: "-cc1" {{.*}} "-fopenacc"
+
+// RUN: %clang -S -### -fopenacc -fexperimental-openacc-macro-override=202211 %s 2>&1 | FileCheck %s --check-prefix=CHECK-MACRO-OVERRIDE
+// RUN: %clang -S -### -fopenacc -fexperimental-openacc-macro-override 202211 %s 2>&1 | FileCheck %s --check-prefix=CHECK-MACRO-OVERRIDE
+// CHECK-MACRO-OVERRIDE: "-cc1"{{.*}} "-fexperimental-openacc-macro-override" "202211"
+
+// RUN: not %clang -S -fopenacc -fexperimental-openacc-macro-override=202211L %s 2>&1 | FileCheck %s --check-prefix=INVALID
+// RUN: not %clang -S -fopenacc -fexperimental-openacc-macro-override 202211L %s 2>&1 | FileCheck %s --check-prefix=INVALID
+// RUN: not %clang -S -fopenacc -fexperimental-openacc-macro-override=L202211 %s 2>&1 | FileCheck %s --check-prefix=INVALID
+// RUN: not %clang -S -fopenacc -fexperimental-openacc-macro-override L202211 %s 2>&1 | FileCheck %s --check-prefix=INVALID
+// RUN: not %clang -S -fopenacc -fexperimental-openacc-macro-override=2022L11 %s 2>&1 | FileCheck %s --check-prefix=INVALID
+// RUN: not %clang -S -fopenacc -fexperimental-openacc-macro-override 2022L11 %s 2>&1 | FileCheck %s --check-prefix=INVALID
+// INVALID: error: the clang compiler does not support
diff --git a/clang/test/ParserOpenACC/disabled.c b/clang/test/ParserOpenACC/disabled.c
new file mode 100644
index 0000000..a25192f
--- /dev/null
+++ b/clang/test/ParserOpenACC/disabled.c
@@ -0,0 +1,4 @@
+// RUN: %clang_cc1 %s -verify -Wsource-uses-openacc
+// expected-warning@+1{{unexpected '#pragma acc ...' in program}}
+#pragma acc foo bar baz blitz.
+int foo;
diff --git a/clang/test/ParserOpenACC/unimplemented.c b/clang/test/ParserOpenACC/unimplemented.c
new file mode 100644
index 0000000..c1228c8
--- /dev/null
+++ b/clang/test/ParserOpenACC/unimplemented.c
@@ -0,0 +1,20 @@
+// RUN: %clang_cc1 %s -verify -fopenacc
+
+// Parser::ParseExternalDeclaration
+// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc not yet implemented
+int foo;
+
+struct S {
+// Parser::ParseStructUnionBody
+// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc not yet implemented
+ int foo;
+};
+
+void func() {
+// Parser::ParseStmtOrDeclarationAfterAttributes
+// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc not yet implemented
+ while(0) {}
+}
diff --git a/clang/test/ParserOpenACC/unimplemented.cpp b/clang/test/ParserOpenACC/unimplemented.cpp
new file mode 100644
index 0000000..095cbf5
--- /dev/null
+++ b/clang/test/ParserOpenACC/unimplemented.cpp
@@ -0,0 +1,20 @@
+// RUN: %clang_cc1 %s -verify -fopenacc
+
+// Parser::ParseExternalDeclaration
+// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc not yet implemented
+int foo;
+
+struct S {
+// Parser::ParseCXXClassMemberDeclarationWithPragmas
+// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc not yet implemented
+ int foo;
+};
+
+void func() {
+// Parser::ParseStmtOrDeclarationAfterAttributes
+// expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc not yet implemented
+ while(false) {}
+}
diff --git a/clang/test/Preprocessor/openacc.c b/clang/test/Preprocessor/openacc.c
new file mode 100644
index 0000000..be7052f0
--- /dev/null
+++ b/clang/test/Preprocessor/openacc.c
@@ -0,0 +1,13 @@
+// RUN: %clang_cc1 -E -fopenacc %s | FileCheck %s --check-prefix=DEFAULT
+// RUN: %clang_cc1 -E -fopenacc -fexperimental-openacc-macro-override 202211 %s | FileCheck %s --check-prefix=OVERRIDE
+
+// DEFAULT: OpenACC:1:
+// OVERRIDE: OpenACC:202211:
+OpenACC:_OPENACC:
+
+// RUN: %clang_cc1 -E -dM -fopenacc %s | FileCheck %s --check-prefix=MACRO_PRINT_DEF
+// RUN: %clang_cc1 -E -dM -fopenacc -fexperimental-openacc-macro-override 202211 %s | FileCheck %s --check-prefix=MACRO_PRINT_OVR
+// MACRO_PRINT_DEF: #define _OPENACC 1
+// MACRO_PRINT_OVR: #define _OPENACC 202211
+
+