15#include "mlir/Dialect/OpenACC/OpenACC.h"
24template <
typename Op,
typename TermOp>
25mlir::LogicalResult CIRGenFunction::emitOpenACCOpAssociatedStmt(
28 const Stmt *associatedStmt) {
29 mlir::LogicalResult res = mlir::success();
33 auto op = builder.create<Op>(start, retTy, operands);
35 emitOpenACCClauses(op, dirKind, dirLoc, clauses);
38 mlir::Block &block = op.getRegion().emplaceBlock();
39 mlir::OpBuilder::InsertionGuard guardCase(builder);
40 builder.setInsertionPointToEnd(&block);
42 LexicalScope ls{*
this, start, builder.getInsertionBlock()};
43 res =
emitStmt(associatedStmt,
true);
45 builder.create<TermOp>(end);
51template <
typename Op>
struct CombinedType;
52template <>
struct CombinedType<ParallelOp> {
53 static constexpr mlir::acc::CombinedConstructsType value =
54 mlir::acc::CombinedConstructsType::ParallelLoop;
56template <>
struct CombinedType<SerialOp> {
57 static constexpr mlir::acc::CombinedConstructsType value =
58 mlir::acc::CombinedConstructsType::SerialLoop;
60template <>
struct CombinedType<KernelsOp> {
61 static constexpr mlir::acc::CombinedConstructsType value =
62 mlir::acc::CombinedConstructsType::KernelsLoop;
66template <
typename Op,
typename TermOp>
67mlir::LogicalResult CIRGenFunction::emitOpenACCOpCombinedConstruct(
70 const Stmt *loopStmt) {
71 mlir::LogicalResult res = mlir::success();
76 auto computeOp = builder.create<Op>(start, retTy, operands);
77 computeOp.setCombinedAttr(builder.getUnitAttr());
78 mlir::acc::LoopOp loopOp;
83 mlir::Block &block = computeOp.getRegion().emplaceBlock();
84 mlir::OpBuilder::InsertionGuard guardCase(builder);
85 builder.setInsertionPointToEnd(&block);
87 LexicalScope ls{*
this, start, builder.getInsertionBlock()};
88 auto loopOp = builder.create<LoopOp>(start, retTy, operands);
89 loopOp.setCombinedAttr(mlir::acc::CombinedConstructsTypeAttr::get(
90 builder.getContext(), CombinedType<Op>::value));
93 mlir::Block &innerBlock = loopOp.getRegion().emplaceBlock();
94 mlir::OpBuilder::InsertionGuard guardCase(builder);
95 builder.setInsertionPointToEnd(&innerBlock);
97 LexicalScope ls{*
this, start, builder.getInsertionBlock()};
98 ActiveOpenACCLoopRAII activeLoop{*
this, &loopOp};
102 builder.create<mlir::acc::YieldOp>(end);
105 emitOpenACCClauses(computeOp, loopOp, dirKind, dirLoc, clauses);
107 updateLoopOpParallelism(loopOp,
false, dirKind);
109 builder.create<TermOp>(end);
115template <
typename Op>
116Op CIRGenFunction::emitOpenACCOp(
121 auto op = builder.create<Op>(start, retTy, operands);
123 emitOpenACCClauses(op, dirKind, dirLoc, clauses);
129 mlir::Location start =
getLoc(
s.getSourceRange().getBegin());
130 mlir::Location end =
getLoc(
s.getSourceRange().getEnd());
132 switch (
s.getDirectiveKind()) {
134 return emitOpenACCOpAssociatedStmt<ParallelOp, mlir::acc::YieldOp>(
135 start, end,
s.getDirectiveKind(),
s.getDirectiveLoc(),
s.clauses(),
136 s.getStructuredBlock());
138 return emitOpenACCOpAssociatedStmt<SerialOp, mlir::acc::YieldOp>(
139 start, end,
s.getDirectiveKind(),
s.getDirectiveLoc(),
s.clauses(),
140 s.getStructuredBlock());
142 return emitOpenACCOpAssociatedStmt<KernelsOp, mlir::acc::TerminatorOp>(
143 start, end,
s.getDirectiveKind(),
s.getDirectiveLoc(),
s.clauses(),
144 s.getStructuredBlock());
146 llvm_unreachable(
"invalid compute construct kind");
152 mlir::Location start =
getLoc(
s.getSourceRange().getBegin());
153 mlir::Location end =
getLoc(
s.getSourceRange().getEnd());
155 return emitOpenACCOpAssociatedStmt<DataOp, mlir::acc::TerminatorOp>(
156 start, end,
s.getDirectiveKind(),
s.getDirectiveLoc(),
s.clauses(),
157 s.getStructuredBlock());
162 mlir::Location start =
getLoc(
s.getSourceRange().getBegin());
163 emitOpenACCOp<InitOp>(start,
s.getDirectiveKind(),
s.getDirectiveLoc(),
165 return mlir::success();
170 mlir::Location start =
getLoc(
s.getSourceRange().getBegin());
171 emitOpenACCOp<SetOp>(start,
s.getDirectiveKind(),
s.getDirectiveLoc(),
173 return mlir::success();
178 mlir::Location start =
getLoc(
s.getSourceRange().getBegin());
179 emitOpenACCOp<ShutdownOp>(start,
s.getDirectiveKind(),
180 s.getDirectiveLoc(),
s.clauses());
181 return mlir::success();
186 mlir::Location start =
getLoc(
s.getSourceRange().getBegin());
187 auto waitOp = emitOpenACCOp<WaitOp>(start,
s.getDirectiveKind(),
188 s.getDirectiveLoc(),
s.clauses());
190 auto createIntExpr = [
this](
const Expr *intExpr) {
192 mlir::Location exprLoc =
cgm.
getLoc(intExpr->getBeginLoc());
194 mlir::IntegerType targetType = mlir::IntegerType::get(
196 intExpr->getType()->isSignedIntegerOrEnumerationType()
197 ? mlir::IntegerType::SignednessSemantics::Signed
198 : mlir::IntegerType::SignednessSemantics::Unsigned);
200 auto conversionOp = builder.create<mlir::UnrealizedConversionCastOp>(
201 exprLoc, targetType,
expr);
202 return conversionOp.getResult(0);
207 mlir::OpBuilder::InsertionGuard guardCase(builder);
208 builder.setInsertionPoint(waitOp);
210 if (
s.hasDevNumExpr())
211 waitOp.getWaitDevnumMutable().append(createIntExpr(
s.getDevNumExpr()));
213 for (
Expr *QueueExpr :
s.getQueueIdExprs())
214 waitOp.getWaitOperandsMutable().append(createIntExpr(QueueExpr));
217 return mlir::success();
222 mlir::Location start =
getLoc(
s.getSourceRange().getBegin());
223 mlir::Location end =
getLoc(
s.getSourceRange().getEnd());
225 switch (
s.getDirectiveKind()) {
227 return emitOpenACCOpCombinedConstruct<ParallelOp, mlir::acc::YieldOp>(
228 start, end,
s.getDirectiveKind(),
s.getDirectiveLoc(),
s.clauses(),
231 return emitOpenACCOpCombinedConstruct<SerialOp, mlir::acc::YieldOp>(
232 start, end,
s.getDirectiveKind(),
s.getDirectiveLoc(),
s.clauses(),
235 return emitOpenACCOpCombinedConstruct<KernelsOp, mlir::acc::TerminatorOp>(
236 start, end,
s.getDirectiveKind(),
s.getDirectiveLoc(),
s.clauses(),
239 llvm_unreachable(
"invalid compute construct kind");
245 mlir::Location start =
getLoc(
s.getSourceRange().getBegin());
246 mlir::Location end =
getLoc(
s.getSourceRange().getEnd());
248 return emitOpenACCOpAssociatedStmt<HostDataOp, mlir::acc::TerminatorOp>(
249 start, end,
s.getDirectiveKind(),
s.getDirectiveLoc(),
s.clauses(),
250 s.getStructuredBlock());
255 mlir::Location start =
getLoc(
s.getSourceRange().getBegin());
256 emitOpenACCOp<EnterDataOp>(start,
s.getDirectiveKind(),
s.getDirectiveLoc(),
258 return mlir::success();
263 mlir::Location start =
getLoc(
s.getSourceRange().getBegin());
264 emitOpenACCOp<ExitDataOp>(start,
s.getDirectiveKind(),
s.getDirectiveLoc(),
266 return mlir::success();
271 mlir::Location start =
getLoc(
s.getSourceRange().getBegin());
272 emitOpenACCOp<UpdateOp>(start,
s.getDirectiveKind(),
s.getDirectiveLoc(),
274 return mlir::success();
286 return mlir::success();
288 mlir::acc::LoopOp loopOp = *activeLoopOp;
290 mlir::OpBuilder::InsertionGuard guard(builder);
291 builder.setInsertionPoint(loopOp);
293 for (
const Expr *var :
s.getVarList()) {
297 auto cacheOp = builder.create<CacheOp>(
301 loopOp.getCacheOperandsMutable().append(cacheOp.getResult());
304 return mlir::success();
309 cgm.
errorNYI(
s.getSourceRange(),
"OpenACC Atomic Construct");
310 return mlir::failure();
This file defines OpenACC AST classes for statement-level contructs.
__device__ __2f16 float __ockl_bool s
mlir::LogicalResult emitOpenACCDataConstruct(const OpenACCDataConstruct &s)
mlir::LogicalResult emitOpenACCCombinedConstruct(const OpenACCCombinedConstruct &s)
mlir::LogicalResult emitOpenACCWaitConstruct(const OpenACCWaitConstruct &s)
mlir::LogicalResult emitOpenACCUpdateConstruct(const OpenACCUpdateConstruct &s)
mlir::LogicalResult emitOpenACCCacheConstruct(const OpenACCCacheConstruct &s)
mlir::Location getLoc(clang::SourceLocation srcLoc)
Helpers to convert Clang's SourceLocation to a MLIR Location.
mlir::LogicalResult emitOpenACCInitConstruct(const OpenACCInitConstruct &s)
mlir::LogicalResult emitOpenACCSetConstruct(const OpenACCSetConstruct &s)
OpenACCDataOperandInfo getOpenACCDataOperandInfo(const Expr *e)
mlir::LogicalResult emitOpenACCComputeConstruct(const OpenACCComputeConstruct &s)
mlir::LogicalResult emitOpenACCShutdownConstruct(const OpenACCShutdownConstruct &s)
mlir::Value emitScalarExpr(const clang::Expr *e)
Emit the computation of the specified expression of scalar type.
mlir::LogicalResult emitOpenACCHostDataConstruct(const OpenACCHostDataConstruct &s)
mlir::MLIRContext & getMLIRContext()
mlir::LogicalResult emitOpenACCEnterDataConstruct(const OpenACCEnterDataConstruct &s)
clang::ASTContext & getContext() const
mlir::LogicalResult emitStmt(const clang::Stmt *s, bool useCurrentScope, llvm::ArrayRef< const Attr * > attrs={})
mlir::LogicalResult emitOpenACCExitDataConstruct(const OpenACCExitDataConstruct &s)
mlir::LogicalResult emitOpenACCAtomicConstruct(const OpenACCAtomicConstruct &s)
DiagnosticBuilder errorNYI(SourceLocation, llvm::StringRef)
Helpers to emit "not yet implemented" error diagnostics.
mlir::Location getLoc(clang::SourceLocation cLoc)
Helpers to convert the presumed location of Clang's SourceLocation to an MLIR Location.
This represents one expression.
This class represents a compute construct, representing a 'Kind' of ‘parallel’, 'serial',...
Encodes a location in the source.
Stmt - This represents one statement.
const internal::VariadicDynCastAllOfMatcher< Stmt, Expr > expr
Matches expressions.
The JSON file list parser is used to communicate input to InstallAPI.
llvm::SmallVector< mlir::Value > bounds