[OpenACC][CIR] Basic infrastructure for OpenACC lowering (#134717)

This is the first of a few patches that will do infrastructure work to
enable the OpenACC lowering via the OpenACC dialect.

At the moment this just gets the various function calls that will end up
generating OpenACC, plus some tests to validate that we're doing the
diagnostics in OpenACC specific locations.

Additionally, this adds Stmt and Decl files for CIRGen.
diff --git a/clang/include/clang/AST/DeclOpenACC.h b/clang/include/clang/AST/DeclOpenACC.h
index 26cf721..8c612fb 100644
--- a/clang/include/clang/AST/DeclOpenACC.h
+++ b/clang/include/clang/AST/DeclOpenACC.h
@@ -59,6 +59,8 @@
   }
 
   ArrayRef<const OpenACCClause *> clauses() const { return Clauses; }
+  static bool classof(const Decl *D) { return classofKind(D->getKind()); }
+  static bool classofKind(Kind K);
 };
 
 class OpenACCDeclareDecl final
diff --git a/clang/include/clang/AST/GlobalDecl.h b/clang/include/clang/AST/GlobalDecl.h
index df11a79..baf5371d 100644
--- a/clang/include/clang/AST/GlobalDecl.h
+++ b/clang/include/clang/AST/GlobalDecl.h
@@ -17,6 +17,7 @@
 #include "clang/AST/Attr.h"
 #include "clang/AST/DeclCXX.h"
 #include "clang/AST/DeclObjC.h"
+#include "clang/AST/DeclOpenACC.h"
 #include "clang/AST/DeclOpenMP.h"
 #include "clang/AST/DeclTemplate.h"
 #include "clang/Basic/ABI.h"
@@ -86,6 +87,8 @@
   GlobalDecl(const ObjCMethodDecl *D) { Init(D); }
   GlobalDecl(const OMPDeclareReductionDecl *D) { Init(D); }
   GlobalDecl(const OMPDeclareMapperDecl *D) { Init(D); }
+  GlobalDecl(const OpenACCRoutineDecl *D) { Init(D); }
+  GlobalDecl(const OpenACCDeclareDecl *D) { Init(D); }
   GlobalDecl(const CXXConstructorDecl *D, CXXCtorType Type) : Value(D, Type) {}
   GlobalDecl(const CXXDestructorDecl *D, CXXDtorType Type) : Value(D, Type) {}
   GlobalDecl(const VarDecl *D, DynamicInitKind StubKind)
diff --git a/clang/include/clang/Basic/DiagnosticDriverKinds.td b/clang/include/clang/Basic/DiagnosticDriverKinds.td
index a96b4fb..c69ad3a 100644
--- a/clang/include/clang/Basic/DiagnosticDriverKinds.td
+++ b/clang/include/clang/Basic/DiagnosticDriverKinds.td
@@ -850,4 +850,9 @@
 
 def err_drv_malformed_warning_suppression_mapping : Error<
   "failed to process suppression mapping file '%0': %1">;
+
+def warn_drv_openacc_without_cir
+    : Warning<"OpenACC directives will result in no runtime behavior; use "
+              "-fclangir to enable runtime effect">,
+      InGroup<SourceUsesOpenACC>;
 }
diff --git a/clang/lib/AST/DeclOpenACC.cpp b/clang/lib/AST/DeclOpenACC.cpp
index 760c08d..e0fe7be 100644
--- a/clang/lib/AST/DeclOpenACC.cpp
+++ b/clang/lib/AST/DeclOpenACC.cpp
@@ -17,6 +17,11 @@
 
 using namespace clang;
 
+bool OpenACCConstructDecl::classofKind(Kind K) {
+  return OpenACCDeclareDecl::classofKind(K) ||
+         OpenACCRoutineDecl::classofKind(K);
+}
+
 OpenACCDeclareDecl *
 OpenACCDeclareDecl::Create(ASTContext &Ctx, DeclContext *DC,
                            SourceLocation StartLoc, SourceLocation DirLoc,
diff --git a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
index 5b832b4..d0eb648 100644
--- a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
@@ -15,6 +15,7 @@
 #include "mlir/IR/Location.h"
 #include "clang/AST/Attr.h"
 #include "clang/AST/Decl.h"
+#include "clang/AST/DeclOpenACC.h"
 #include "clang/AST/Expr.h"
 #include "clang/AST/ExprCXX.h"
 #include "clang/CIR/MissingFeatures.h"
@@ -266,6 +267,12 @@
     emitVarDecl(vd);
     return;
   }
+  case Decl::OpenACCDeclare:
+    emitOpenACCDeclare(cast<OpenACCDeclareDecl>(d));
+    return;
+  case Decl::OpenACCRoutine:
+    emitOpenACCRoutine(cast<OpenACCRoutineDecl>(d));
+    return;
   default:
     cgm.errorNYI(d.getSourceRange(), "emitDecl: unhandled decl type");
   }
diff --git a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
new file mode 100644
index 0000000..b588a50
--- /dev/null
+++ b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
@@ -0,0 +1,34 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 contains code to emit Decl nodes as CIR code.
+//
+//===----------------------------------------------------------------------===//
+
+#include "CIRGenFunction.h"
+#include "clang/AST/DeclOpenACC.h"
+
+using namespace clang;
+using namespace clang::CIRGen;
+
+void CIRGenFunction::emitOpenACCDeclare(const OpenACCDeclareDecl &d) {
+  getCIRGenModule().errorNYI(d.getSourceRange(), "OpenACC Declare Construct");
+}
+
+void CIRGenFunction::emitOpenACCRoutine(const OpenACCRoutineDecl &d) {
+  getCIRGenModule().errorNYI(d.getSourceRange(), "OpenACC Routine Construct");
+}
+
+void CIRGenModule::emitGlobalOpenACCDecl(const OpenACCConstructDecl *d) {
+  if (isa<OpenACCRoutineDecl>(d))
+    errorNYI(d->getSourceRange(), "OpenACC Routine Construct");
+  else if (isa<OpenACCDeclareDecl>(d))
+    errorNYI(d->getSourceRange(), "OpenACC Declare Construct");
+  else
+    llvm_unreachable("unknown OpenACC declaration kind?");
+}
diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h
index 1bedbe2..f505ed8 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunction.h
+++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h
@@ -509,6 +509,36 @@
 public:
   Address createTempAlloca(mlir::Type ty, CharUnits align, mlir::Location loc,
                            const Twine &name, bool insertIntoFnEntryBlock);
+
+  //===--------------------------------------------------------------------===//
+  //                         OpenACC Emission
+  //===--------------------------------------------------------------------===//
+public:
+  mlir::LogicalResult
+  emitOpenACCComputeConstruct(const OpenACCComputeConstruct &s);
+  mlir::LogicalResult emitOpenACCLoopConstruct(const OpenACCLoopConstruct &s);
+  mlir::LogicalResult
+  emitOpenACCCombinedConstruct(const OpenACCCombinedConstruct &s);
+  mlir::LogicalResult emitOpenACCDataConstruct(const OpenACCDataConstruct &s);
+  mlir::LogicalResult
+  emitOpenACCEnterDataConstruct(const OpenACCEnterDataConstruct &s);
+  mlir::LogicalResult
+  emitOpenACCExitDataConstruct(const OpenACCExitDataConstruct &s);
+  mlir::LogicalResult
+  emitOpenACCHostDataConstruct(const OpenACCHostDataConstruct &s);
+  mlir::LogicalResult emitOpenACCWaitConstruct(const OpenACCWaitConstruct &s);
+  mlir::LogicalResult emitOpenACCInitConstruct(const OpenACCInitConstruct &s);
+  mlir::LogicalResult
+  emitOpenACCShutdownConstruct(const OpenACCShutdownConstruct &s);
+  mlir::LogicalResult emitOpenACCSetConstruct(const OpenACCSetConstruct &s);
+  mlir::LogicalResult
+  emitOpenACCUpdateConstruct(const OpenACCUpdateConstruct &s);
+  mlir::LogicalResult
+  emitOpenACCAtomicConstruct(const OpenACCAtomicConstruct &s);
+  mlir::LogicalResult emitOpenACCCacheConstruct(const OpenACCCacheConstruct &s);
+
+  void emitOpenACCDeclare(const OpenACCDeclareDecl &d);
+  void emitOpenACCRoutine(const OpenACCRoutineDecl &d);
 };
 
 } // namespace clang::CIRGen
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index d3b3b06..f0e9b03 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -16,6 +16,7 @@
 
 #include "clang/AST/ASTContext.h"
 #include "clang/AST/DeclBase.h"
+#include "clang/AST/DeclOpenACC.h"
 #include "clang/AST/GlobalDecl.h"
 #include "clang/Basic/SourceManager.h"
 #include "clang/CIR/Dialect/IR/CIRDialect.h"
@@ -91,6 +92,11 @@
 }
 
 void CIRGenModule::emitGlobal(clang::GlobalDecl gd) {
+  if (const auto *cd = dyn_cast<clang::OpenACCConstructDecl>(gd.getDecl())) {
+    emitGlobalOpenACCDecl(cd);
+    return;
+  }
+
   const auto *global = cast<ValueDecl>(gd.getDecl());
 
   if (const auto *fd = dyn_cast<FunctionDecl>(global)) {
@@ -423,6 +429,12 @@
     emitGlobal(vd);
     break;
   }
+  case Decl::OpenACCRoutine:
+    emitGlobalOpenACCDecl(cast<OpenACCRoutineDecl>(decl));
+    break;
+  case Decl::OpenACCDeclare:
+    emitGlobalOpenACCDecl(cast<OpenACCDeclareDecl>(decl));
+    break;
   }
 }
 
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h
index 6ba1ccc..ab4545e 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.h
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.h
@@ -113,6 +113,8 @@
   void emitGlobalVarDefinition(const clang::VarDecl *vd,
                                bool isTentative = false);
 
+  void emitGlobalOpenACCDecl(const clang::OpenACCConstructDecl *cd);
+
   /// Return the result of value-initializing the given type, i.e. a null
   /// expression of the given type.
   mlir::Value emitNullConstant(QualType t, mlir::Location loc);
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmt.cpp b/clang/lib/CIR/CodeGen/CIRGenStmt.cpp
index 00d33e7..2551a67 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmt.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmt.cpp
@@ -16,6 +16,7 @@
 #include "mlir/IR/Builders.h"
 #include "clang/AST/ExprCXX.h"
 #include "clang/AST/Stmt.h"
+#include "clang/AST/StmtOpenACC.h"
 
 using namespace clang;
 using namespace clang::CIRGen;
@@ -85,7 +86,34 @@
     return emitWhileStmt(cast<WhileStmt>(*s));
   case Stmt::DoStmtClass:
     return emitDoStmt(cast<DoStmt>(*s));
-
+  case Stmt::OpenACCComputeConstructClass:
+    return emitOpenACCComputeConstruct(cast<OpenACCComputeConstruct>(*s));
+  case Stmt::OpenACCLoopConstructClass:
+    return emitOpenACCLoopConstruct(cast<OpenACCLoopConstruct>(*s));
+  case Stmt::OpenACCCombinedConstructClass:
+    return emitOpenACCCombinedConstruct(cast<OpenACCCombinedConstruct>(*s));
+  case Stmt::OpenACCDataConstructClass:
+    return emitOpenACCDataConstruct(cast<OpenACCDataConstruct>(*s));
+  case Stmt::OpenACCEnterDataConstructClass:
+    return emitOpenACCEnterDataConstruct(cast<OpenACCEnterDataConstruct>(*s));
+  case Stmt::OpenACCExitDataConstructClass:
+    return emitOpenACCExitDataConstruct(cast<OpenACCExitDataConstruct>(*s));
+  case Stmt::OpenACCHostDataConstructClass:
+    return emitOpenACCHostDataConstruct(cast<OpenACCHostDataConstruct>(*s));
+  case Stmt::OpenACCWaitConstructClass:
+    return emitOpenACCWaitConstruct(cast<OpenACCWaitConstruct>(*s));
+  case Stmt::OpenACCInitConstructClass:
+    return emitOpenACCInitConstruct(cast<OpenACCInitConstruct>(*s));
+  case Stmt::OpenACCShutdownConstructClass:
+    return emitOpenACCShutdownConstruct(cast<OpenACCShutdownConstruct>(*s));
+  case Stmt::OpenACCSetConstructClass:
+    return emitOpenACCSetConstruct(cast<OpenACCSetConstruct>(*s));
+  case Stmt::OpenACCUpdateConstructClass:
+    return emitOpenACCUpdateConstruct(cast<OpenACCUpdateConstruct>(*s));
+  case Stmt::OpenACCCacheConstructClass:
+    return emitOpenACCCacheConstruct(cast<OpenACCCacheConstruct>(*s));
+  case Stmt::OpenACCAtomicConstructClass:
+    return emitOpenACCAtomicConstruct(cast<OpenACCAtomicConstruct>(*s));
   case Stmt::OMPScopeDirectiveClass:
   case Stmt::OMPErrorDirectiveClass:
   case Stmt::NoStmtClass:
@@ -192,20 +220,6 @@
   case Stmt::OMPAssumeDirectiveClass:
   case Stmt::OMPMaskedDirectiveClass:
   case Stmt::OMPStripeDirectiveClass:
-  case Stmt::OpenACCComputeConstructClass:
-  case Stmt::OpenACCLoopConstructClass:
-  case Stmt::OpenACCCombinedConstructClass:
-  case Stmt::OpenACCDataConstructClass:
-  case Stmt::OpenACCEnterDataConstructClass:
-  case Stmt::OpenACCExitDataConstructClass:
-  case Stmt::OpenACCHostDataConstructClass:
-  case Stmt::OpenACCWaitConstructClass:
-  case Stmt::OpenACCInitConstructClass:
-  case Stmt::OpenACCShutdownConstructClass:
-  case Stmt::OpenACCSetConstructClass:
-  case Stmt::OpenACCUpdateConstructClass:
-  case Stmt::OpenACCCacheConstructClass:
-  case Stmt::OpenACCAtomicConstructClass:
   case Stmt::ObjCAtCatchStmtClass:
   case Stmt::ObjCAtFinallyStmtClass:
     cgm.errorNYI(s->getSourceRange(),
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
new file mode 100644
index 0000000..cbae170
--- /dev/null
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
@@ -0,0 +1,91 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// Emit OpenACC Stmt nodes as CIR code.
+//
+//===----------------------------------------------------------------------===//
+
+#include "CIRGenBuilder.h"
+#include "CIRGenFunction.h"
+#include "clang/AST/StmtOpenACC.h"
+
+using namespace clang;
+using namespace clang::CIRGen;
+using namespace cir;
+
+mlir::LogicalResult
+CIRGenFunction::emitOpenACCComputeConstruct(const OpenACCComputeConstruct &s) {
+  getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Compute Construct");
+  return mlir::failure();
+}
+
+mlir::LogicalResult
+CIRGenFunction::emitOpenACCLoopConstruct(const OpenACCLoopConstruct &s) {
+  getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Loop Construct");
+  return mlir::failure();
+}
+mlir::LogicalResult CIRGenFunction::emitOpenACCCombinedConstruct(
+    const OpenACCCombinedConstruct &s) {
+  getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Combined Construct");
+  return mlir::failure();
+}
+mlir::LogicalResult
+CIRGenFunction::emitOpenACCDataConstruct(const OpenACCDataConstruct &s) {
+  getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Data Construct");
+  return mlir::failure();
+}
+mlir::LogicalResult CIRGenFunction::emitOpenACCEnterDataConstruct(
+    const OpenACCEnterDataConstruct &s) {
+  getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC EnterData Construct");
+  return mlir::failure();
+}
+mlir::LogicalResult CIRGenFunction::emitOpenACCExitDataConstruct(
+    const OpenACCExitDataConstruct &s) {
+  getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC ExitData Construct");
+  return mlir::failure();
+}
+mlir::LogicalResult CIRGenFunction::emitOpenACCHostDataConstruct(
+    const OpenACCHostDataConstruct &s) {
+  getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC HostData Construct");
+  return mlir::failure();
+}
+mlir::LogicalResult
+CIRGenFunction::emitOpenACCWaitConstruct(const OpenACCWaitConstruct &s) {
+  getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Wait Construct");
+  return mlir::failure();
+}
+mlir::LogicalResult
+CIRGenFunction::emitOpenACCInitConstruct(const OpenACCInitConstruct &s) {
+  getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Init Construct");
+  return mlir::failure();
+}
+mlir::LogicalResult CIRGenFunction::emitOpenACCShutdownConstruct(
+    const OpenACCShutdownConstruct &s) {
+  getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Shutdown Construct");
+  return mlir::failure();
+}
+mlir::LogicalResult
+CIRGenFunction::emitOpenACCSetConstruct(const OpenACCSetConstruct &s) {
+  getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Set Construct");
+  return mlir::failure();
+}
+mlir::LogicalResult
+CIRGenFunction::emitOpenACCUpdateConstruct(const OpenACCUpdateConstruct &s) {
+  getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Update Construct");
+  return mlir::failure();
+}
+mlir::LogicalResult
+CIRGenFunction::emitOpenACCAtomicConstruct(const OpenACCAtomicConstruct &s) {
+  getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Atomic Construct");
+  return mlir::failure();
+}
+mlir::LogicalResult
+CIRGenFunction::emitOpenACCCacheConstruct(const OpenACCCacheConstruct &s) {
+  getCIRGenModule().errorNYI(s.getSourceRange(), "OpenACC Cache Construct");
+  return mlir::failure();
+}
diff --git a/clang/lib/CIR/CodeGen/CMakeLists.txt b/clang/lib/CIR/CodeGen/CMakeLists.txt
index da8d63c..11902c7 100644
--- a/clang/lib/CIR/CodeGen/CMakeLists.txt
+++ b/clang/lib/CIR/CodeGen/CMakeLists.txt
@@ -9,6 +9,7 @@
 add_clang_library(clangCIR
   CIRGenerator.cpp
   CIRGenDecl.cpp
+  CIRGenDeclOpenACC.cpp
   CIRGenExpr.cpp
   CIRGenExprAggregate.cpp
   CIRGenExprConstant.cpp
@@ -16,6 +17,7 @@
   CIRGenFunction.cpp
   CIRGenModule.cpp
   CIRGenStmt.cpp
+  CIRGenStmtOpenACC.cpp
   CIRGenTypes.cpp
 
   DEPENDS
diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp
index 572c71e..cfc5c06 100644
--- a/clang/lib/Frontend/CompilerInvocation.cpp
+++ b/clang/lib/Frontend/CompilerInvocation.cpp
@@ -4674,6 +4674,51 @@
   llvm_unreachable("invalid frontend action");
 }
 
+static bool isCodeGenAction(frontend::ActionKind Action) {
+  switch (Action) {
+  case frontend::EmitAssembly:
+  case frontend::EmitBC:
+  case frontend::EmitCIR:
+  case frontend::EmitHTML:
+  case frontend::EmitLLVM:
+  case frontend::EmitLLVMOnly:
+  case frontend::EmitCodeGenOnly:
+  case frontend::EmitObj:
+  case frontend::GenerateModule:
+  case frontend::GenerateModuleInterface:
+  case frontend::GenerateReducedModuleInterface:
+  case frontend::GenerateHeaderUnit:
+  case frontend::GeneratePCH:
+  case frontend::GenerateInterfaceStubs:
+    return true;
+  case frontend::ASTDeclList:
+  case frontend::ASTDump:
+  case frontend::ASTPrint:
+  case frontend::ASTView:
+  case frontend::ExtractAPI:
+  case frontend::FixIt:
+  case frontend::ParseSyntaxOnly:
+  case frontend::ModuleFileInfo:
+  case frontend::VerifyPCH:
+  case frontend::PluginAction:
+  case frontend::RewriteObjC:
+  case frontend::RewriteTest:
+  case frontend::RunAnalysis:
+  case frontend::TemplightDump:
+  case frontend::DumpCompilerOptions:
+  case frontend::DumpRawTokens:
+  case frontend::DumpTokens:
+  case frontend::InitOnly:
+  case frontend::PrintPreamble:
+  case frontend::PrintPreprocessedInput:
+  case frontend::RewriteMacros:
+  case frontend::RunPreprocessorOnly:
+  case frontend::PrintDependencyDirectivesSourceMinimizerOutput:
+    return false;
+  }
+  llvm_unreachable("invalid frontend action");
+}
+
 static void GeneratePreprocessorArgs(const PreprocessorOptions &Opts,
                                      ArgumentConsumer Consumer,
                                      const LangOptions &LangOpts,
@@ -5001,6 +5046,10 @@
       Res.getTargetOpts().HostTriple = Res.getFrontendOpts().AuxTriple;
   }
 
+  if (LangOpts.OpenACC && !Res.getFrontendOpts().UseClangIRPipeline &&
+      isCodeGenAction(Res.getFrontendOpts().ProgramAction))
+    Diags.Report(diag::warn_drv_openacc_without_cir);
+
   // Set the triple of the host for OpenMP device compile.
   if (LangOpts.OpenMPIsTargetDevice)
     Res.getTargetOpts().HostTriple = Res.getFrontendOpts().AuxTriple;
diff --git a/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented-global.cpp b/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented-global.cpp
new file mode 100644
index 0000000..2aa32b0
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented-global.cpp
@@ -0,0 +1,6 @@
+// RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-linux-gnu -fopenacc -fclangir -emit-cir %s -o %t.cir -verify
+// RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-linux-gnu -fopenacc -fclangir -emit-llvm %s -o %t-cir.ll -verify
+
+int Global;
+// expected-error@+1{{ClangIR code gen Not Yet Implemented: OpenACC Declare Construct}}
+#pragma acc declare create(Global)
diff --git a/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp b/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp
new file mode 100644
index 0000000..61bed79
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp
@@ -0,0 +1,20 @@
+// RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-linux-gnu -fopenacc -fclangir -emit-cir %s -o %t.cir -verify
+// RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-linux-gnu -fopenacc -fclangir -emit-llvm %s -o %t-cir.ll -verify
+
+void HelloWorld(int *A, int *B, int *C, int N) {
+
+// expected-error@+2{{ClangIR code gen Not Yet Implemented: OpenACC Compute Construct}}
+// expected-error@+1{{ClangIR code gen Not Yet Implemented: statement}}
+#pragma acc parallel
+  for (unsigned I = 0; I < N; ++I)
+    A[I] = B[I] + C[I];
+
+// expected-error@+2{{ClangIR code gen Not Yet Implemented: OpenACC Loop Construct}}
+// expected-error@+1{{ClangIR code gen Not Yet Implemented: statement}}
+#pragma acc loop
+  for (unsigned I = 0; I < N; ++I)
+    A[I] = B[I] + C[I];
+
+// expected-error@+1{{ClangIR code gen Not Yet Implemented: OpenACC Declare Construct}}
+#pragma acc declare create(A)
+}
diff --git a/clang/test/Driver/openacc-no-cir.c b/clang/test/Driver/openacc-no-cir.c
new file mode 100644
index 0000000..7b67df2b
--- /dev/null
+++ b/clang/test/Driver/openacc-no-cir.c
@@ -0,0 +1,6 @@
+// RUN: %clang -fopenacc -S %s 2>&1 | FileCheck %s -check-prefix=ERROR
+// RUN: %clang -fclangir -fopenacc -S %s 2>&1 | FileCheck %s --allow-empty -check-prefix=NOERROR
+// RUN: %clang -fopenacc -fclangir -S %s 2>&1 | FileCheck %s --allow-empty -check-prefix=NOERROR
+
+// ERROR: OpenACC directives will result in no runtime behavior; use -fclangir to enable runtime effect
+// NOERROR-NOT: OpenACC directives