[OpenACC] Initial commits to support OpenACC

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 <https://github.com/llvm/llvm-project/pull/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
 -----------------------
 

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">;

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;

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")

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

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.

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<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>,
@@ -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<f_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<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>,

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 @@ 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]

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))

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);

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

clang/lib/Parse/CMakeLists.txt

@@ -23,6 +23,7 @@ add_clang_library(clangParse
   ParseTemplate.cpp
   ParseTentative.cpp
   Parser.cpp
+  ParseOpenACC.cpp
 
   LINK_LIBS
   clangAST

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(

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())) {

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();
+}