diff --git a/clang/include/clang/AST/DeclOpenACC.h b/clang/include/clang/AST/DeclOpenACC.h index 26cf721561fb1..8c612fbf1ec07 100644 --- a/clang/include/clang/AST/DeclOpenACC.h +++ b/clang/include/clang/AST/DeclOpenACC.h @@ -59,6 +59,8 @@ class OpenACCConstructDecl : public Decl { } ArrayRef 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 386693cabb1fb..8bf2eb2a5b78f 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 @@ class GlobalDecl { 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 df24cca49aaae..adf4d1f04af11 100644 --- a/clang/include/clang/Basic/DiagnosticDriverKinds.td +++ b/clang/include/clang/Basic/DiagnosticDriverKinds.td @@ -849,4 +849,9 @@ def warn_missing_include_dirs : Warning< 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; } diff --git a/clang/lib/AST/DeclOpenACC.cpp b/clang/lib/AST/DeclOpenACC.cpp index 760c08d21cccd..e0fe7be8fc1a3 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 5b832b463e752..d0eb648683e8c 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 @@ void CIRGenFunction::emitDecl(const Decl &d) { emitVarDecl(vd); return; } + case Decl::OpenACCDeclare: + emitOpenACCDeclare(cast(d)); + return; + case Decl::OpenACCRoutine: + emitOpenACCRoutine(cast(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 0000000000000..b588a50aa0404 --- /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(d)) + errorNYI(d->getSourceRange(), "OpenACC Routine Construct"); + else if (isa(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 1bedbe28ae625..f505ed8fff311 100644 --- a/clang/lib/CIR/CodeGen/CIRGenFunction.h +++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h @@ -509,6 +509,36 @@ class CIRGenFunction : public CIRGenTypeCache { 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 d3b3b0632c2f0..f0e9b0349f709 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 @@ mlir::Location CIRGenModule::getLoc(SourceRange cRange) { } void CIRGenModule::emitGlobal(clang::GlobalDecl gd) { + if (const auto *cd = dyn_cast(gd.getDecl())) { + emitGlobalOpenACCDecl(cd); + return; + } + const auto *global = cast(gd.getDecl()); if (const auto *fd = dyn_cast(global)) { @@ -423,6 +429,12 @@ void CIRGenModule::emitTopLevelDecl(Decl *decl) { emitGlobal(vd); break; } + case Decl::OpenACCRoutine: + emitGlobalOpenACCDecl(cast(decl)); + break; + case Decl::OpenACCDeclare: + emitGlobalOpenACCDecl(cast(decl)); + break; } } diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h index 6ba1ccc4ddd9f..ab4545effde45 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.h +++ b/clang/lib/CIR/CodeGen/CIRGenModule.h @@ -113,6 +113,8 @@ class CIRGenModule : public CIRGenTypeCache { 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 00d33e7feddff..2551a670b5325 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 @@ mlir::LogicalResult CIRGenFunction::emitStmt(const Stmt *s, return emitWhileStmt(cast(*s)); case Stmt::DoStmtClass: return emitDoStmt(cast(*s)); - + case Stmt::OpenACCComputeConstructClass: + return emitOpenACCComputeConstruct(cast(*s)); + case Stmt::OpenACCLoopConstructClass: + return emitOpenACCLoopConstruct(cast(*s)); + case Stmt::OpenACCCombinedConstructClass: + return emitOpenACCCombinedConstruct(cast(*s)); + case Stmt::OpenACCDataConstructClass: + return emitOpenACCDataConstruct(cast(*s)); + case Stmt::OpenACCEnterDataConstructClass: + return emitOpenACCEnterDataConstruct(cast(*s)); + case Stmt::OpenACCExitDataConstructClass: + return emitOpenACCExitDataConstruct(cast(*s)); + case Stmt::OpenACCHostDataConstructClass: + return emitOpenACCHostDataConstruct(cast(*s)); + case Stmt::OpenACCWaitConstructClass: + return emitOpenACCWaitConstruct(cast(*s)); + case Stmt::OpenACCInitConstructClass: + return emitOpenACCInitConstruct(cast(*s)); + case Stmt::OpenACCShutdownConstructClass: + return emitOpenACCShutdownConstruct(cast(*s)); + case Stmt::OpenACCSetConstructClass: + return emitOpenACCSetConstruct(cast(*s)); + case Stmt::OpenACCUpdateConstructClass: + return emitOpenACCUpdateConstruct(cast(*s)); + case Stmt::OpenACCCacheConstructClass: + return emitOpenACCCacheConstruct(cast(*s)); + case Stmt::OpenACCAtomicConstructClass: + return emitOpenACCAtomicConstruct(cast(*s)); case Stmt::OMPScopeDirectiveClass: case Stmt::OMPErrorDirectiveClass: case Stmt::NoStmtClass: @@ -192,20 +220,6 @@ mlir::LogicalResult CIRGenFunction::emitStmt(const Stmt *s, 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 0000000000000..cbae170162ffe --- /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 da8d63ca569af..11902c708c505 100644 --- a/clang/lib/CIR/CodeGen/CMakeLists.txt +++ b/clang/lib/CIR/CodeGen/CMakeLists.txt @@ -9,6 +9,7 @@ get_property(dialect_libs GLOBAL PROPERTY MLIR_DIALECT_LIBS) add_clang_library(clangCIR CIRGenerator.cpp CIRGenDecl.cpp + CIRGenDeclOpenACC.cpp CIRGenExpr.cpp CIRGenExprAggregate.cpp CIRGenExprConstant.cpp @@ -16,6 +17,7 @@ add_clang_library(clangCIR 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 572c71ef1001c..cfc5c069b0849 100644 --- a/clang/lib/Frontend/CompilerInvocation.cpp +++ b/clang/lib/Frontend/CompilerInvocation.cpp @@ -4674,6 +4674,51 @@ static bool isStrictlyPreprocessorAction(frontend::ActionKind Action) { 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 @@ bool CompilerInvocation::CreateFromArgsImpl( 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 0000000000000..2aa32b0484f2c --- /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 0000000000000..61bed79dc14ea --- /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 0000000000000..7b67df2b6b886 --- /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