diff --git a/clang/docs/ReleaseNotes.rst b/clang/docs/ReleaseNotes.rst index 7238386231e1a..2f149f45bf52b 100644 --- a/clang/docs/ReleaseNotes.rst +++ b/clang/docs/ReleaseNotes.rst @@ -211,6 +211,8 @@ New Compiler Flags the preprocessed text to the output. This can greatly reduce the size of the preprocessed output, which can be helpful when trying to reduce a test case. +* ``-fopenacc`` was added as a part of the effort to support OpenACC in clang. + Deprecated Compiler Flags ------------------------- @@ -672,6 +674,17 @@ Miscellaneous Clang Crashes Fixed - Fixed a crash when an ObjC ivar has an invalid type. See (`#68001 `_) +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 318eb0d6f8890..5dc1b5b718ad5 100644 --- a/clang/include/clang/Basic/DiagnosticGroups.td +++ b/clang/include/clang/Basic/DiagnosticGroups.td @@ -1311,6 +1311,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 de180344fcc5c..c3d06053caa5e 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, 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; + // OpenMP support. def warn_pragma_omp_ignored : Warning< "unexpected '#pragma omp ...' in program">, InGroup, DefaultIgnore; diff --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def index c0ea4ecb9806a..872d693cc3ebb 100644 --- a/clang/include/clang/Basic/LangOptions.def +++ b/clang/include/clang/Basic/LangOptions.def @@ -283,6 +283,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 20a8ada60e0fe..f6482f5ee6697 100644 --- a/clang/include/clang/Basic/LangOptions.h +++ b/clang/include/clang/Basic/LangOptions.h @@ -502,6 +502,11 @@ class LangOptions : public LangOptionsBase { // 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 3ce317d318f9b..bbfb5cfccac2f 100644 --- a/clang/include/clang/Basic/TokenKinds.def +++ b/clang/include/clang/Basic/TokenKinds.def @@ -943,6 +943,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 2eb86caa6e6d4..4149a9d40d26c 100644 --- a/clang/include/clang/Driver/Options.td +++ b/clang/include/clang/Driver/Options.td @@ -1349,6 +1349,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, + HelpText<"Overrides the _OPENACC macro value for experimental testing " + "during OpenACC support development">; +def openacc_macro_override_EQ + : Joined<["-"], "fexperimental-openacc-macro-override=">, + Alias; + +// End Clang specific/exclusive options for OpenACC. + def libomptarget_amdgpu_bc_path_EQ : Joined<["--"], "libomptarget-amdgpu-bc-path=">, Group, HelpText<"Path to libomptarget-amdgcn bitcode library">; def libomptarget_amdgcn_bc_path_EQ : Joined<["--"], "libomptarget-amdgcn-bc-path=">, Group, @@ -3316,6 +3329,14 @@ def fno_openmp_target_debug : Flag<["-"], "fno-openmp-target-debug">; } // let Visibility = [ClangOption, CC1Option, FC1Option] } // let Flags = [NoArgumentUnused] +//===----------------------------------------------------------------------===// +// FlangOption + FC1 + ClangOption + CC1Option +//===----------------------------------------------------------------------===// +let Visibility = [FC1Option, FlangOption, CC1Option, ClangOption] in { +def fopenacc : Flag<["-"], "fopenacc">, Group, + HelpText<"Enable OpenACC">; +} // let Visibility = [FC1Option, FlangOption, CC1Option, ClangOption] + //===----------------------------------------------------------------------===// // Optimisation remark options //===----------------------------------------------------------------------===// @@ -6224,8 +6245,6 @@ file}]>; def ffixed_line_length_VALUE : Joined<["-"], "ffixed-line-length-">, Group, Alias; def fconvert_EQ : Joined<["-"], "fconvert=">, Group, HelpText<"Set endian conversion of data for unformatted files">; -def fopenacc : Flag<["-"], "fopenacc">, Group, - HelpText<"Enable OpenACC">; def fdefault_double_8 : Flag<["-"],"fdefault-double-8">, Group, HelpText<"Set the default double precision kind to an 8 byte wide type">; def fdefault_integer_8 : Flag<["-"],"fdefault-integer-8">, Group, diff --git a/clang/include/clang/Parse/Parser.h b/clang/include/clang/Parse/Parser.h index 79ac622fd03e2..28be840bdc34a 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 FPContractHandler; std::unique_ptr OpenCLExtensionHandler; std::unique_ptr OpenMPHandler; + std::unique_ptr OpenACCHandler; std::unique_ptr PCSectionHandler; std::unique_ptr MSCommentHandler; std::unique_ptr MSDetectMismatchHandler; @@ -3524,6 +3525,15 @@ class Parser : public CodeCompletionHandler { /// 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 601bbfb927746..b1f6362c58f9d 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; @@ -6597,6 +6614,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 4e6d7bb16f51b..8db78528bbe56 100644 --- a/clang/lib/Frontend/CompilerInvocation.cpp +++ b/clang/lib/Frontend/CompilerInvocation.cpp @@ -3532,6 +3532,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. @@ -4001,6 +4008,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 846e5fce6de7b..17948dcebd7e5 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 5a20e9da974fa..22e902f7e1bc5 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 78c3ab72979a0..f265faa8a73f7 100644 --- a/clang/lib/Parse/ParseDecl.cpp +++ b/clang/lib/Parse/ParseDecl.cpp @@ -4748,6 +4748,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 35b1a93a54a6a..e12215d74bcc8 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 0000000000000..2fba6cd2805cf --- /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 b3178aef64d72..d3fb7fcc8dfac 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 +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 { + PragmaNoOpenMPHandler() : PragmaNoSupportHandler("omp") {} +}; + +struct PragmaNoOpenACCHandler + : public PragmaNoSupportHandler { + 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 +struct PragmaSupportHandler : public PragmaHandler { + PragmaSupportHandler(StringRef Name) : PragmaHandler(Name) {} void HandlePragma(Preprocessor &PP, PragmaIntroducer Introducer, Token &FirstToken) override; }; +struct PragmaOpenMPHandler + : public PragmaSupportHandler { + PragmaOpenMPHandler() : PragmaSupportHandler("omp") {} +}; + +struct PragmaOpenACCHandler + : public PragmaSupportHandler { + PragmaOpenACCHandler() : PragmaSupportHandler("acc") {} +}; + /// PragmaCommentHandler - "\#pragma comment ...". struct PragmaCommentHandler : public PragmaHandler { PragmaCommentHandler(Sema &Actions) @@ -423,6 +456,12 @@ void Parser::initializePragmaHandlers() { OpenMPHandler = std::make_unique(); PP.AddPragmaHandler(OpenMPHandler.get()); + if (getLangOpts().OpenACC) + OpenACCHandler = std::make_unique(); + else + OpenACCHandler = std::make_unique(); + PP.AddPragmaHandler(OpenACCHandler.get()); + if (getLangOpts().MicrosoftExt || getTargetInfo().getTriple().isOSBinFormatELF()) { MSCommentHandler = std::make_unique(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 +void PragmaNoSupportHandler::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 +void PragmaSupportHandler::HandlePragma( + Preprocessor &PP, PragmaIntroducer Introducer, Token &FirstTok) { SmallVector 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 2531147c23196..924f27da8b52c 100644 --- a/clang/lib/Parse/ParseStmt.cpp +++ b/clang/lib/Parse/ParseStmt.cpp @@ -475,6 +475,9 @@ StmtResult Parser::ParseStatementOrDeclarationAfterAttributes( // 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 0f930248e7717..f9418db708b4d 100644 --- a/clang/lib/Parse/Parser.cpp +++ b/clang/lib/Parse/Parser.cpp @@ -317,6 +317,12 @@ bool Parser::SkipUntil(ArrayRef 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: @@ -850,6 +856,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 42b48d230af7a..dbc2d0b56c46d 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 739344b9a128d..02ac27fce2784 100644 --- a/clang/lib/Serialization/ASTWriter.cpp +++ b/clang/lib/Serialization/ASTWriter.cpp @@ -4499,6 +4499,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 0000000000000..c7f1d2545bd03 --- /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 0000000000000..a25192ff3467c --- /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 0000000000000..c1228c8f2b97f --- /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 0000000000000..095cbf570a41a --- /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 0000000000000..be7052f00e0ce --- /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 + +