15#include "mlir/Dialect/OpenACC/OpenACC.h"
24template <
typename Op,
typename TermOp>
25mlir::LogicalResult CIRGenFunction::emitOpenACCOpAssociatedStmt(
28 mlir::LogicalResult res = mlir::success();
30 llvm::SmallVector<mlir::Type> retTy;
31 llvm::SmallVector<mlir::Value> operands;
32 auto op = Op::create(builder, start, retTy, operands);
34 emitOpenACCClauses(op, dirKind, clauses);
37 mlir::Block &block = op.getRegion().emplaceBlock();
38 mlir::OpBuilder::InsertionGuard guardCase(builder);
39 builder.setInsertionPointToEnd(&block);
41 LexicalScope ls{*
this, start, builder.getInsertionBlock()};
42 res =
emitStmt(associatedStmt,
true);
44 TermOp::create(builder, end);
50template <
typename Op>
struct CombinedType;
51template <>
struct CombinedType<ParallelOp> {
52 static constexpr mlir::acc::CombinedConstructsType value =
53 mlir::acc::CombinedConstructsType::ParallelLoop;
55template <>
struct CombinedType<SerialOp> {
56 static constexpr mlir::acc::CombinedConstructsType value =
57 mlir::acc::CombinedConstructsType::SerialLoop;
59template <>
struct CombinedType<KernelsOp> {
60 static constexpr mlir::acc::CombinedConstructsType value =
61 mlir::acc::CombinedConstructsType::KernelsLoop;
65template <
typename Op,
typename TermOp>
66mlir::LogicalResult CIRGenFunction::emitOpenACCOpCombinedConstruct(
68 llvm::ArrayRef<const OpenACCClause *> clauses,
const Stmt *loopStmt) {
69 mlir::LogicalResult res = mlir::success();
71 llvm::SmallVector<mlir::Type> retTy;
72 llvm::SmallVector<mlir::Value> operands;
74 auto computeOp = Op::create(builder, start, retTy, operands);
75 computeOp.setCombinedAttr(builder.getUnitAttr());
76 mlir::acc::LoopOp loopOp;
81 mlir::Block &block = computeOp.getRegion().emplaceBlock();
82 mlir::OpBuilder::InsertionGuard guardCase(builder);
83 builder.setInsertionPointToEnd(&block);
85 LexicalScope ls{*
this, start, builder.getInsertionBlock()};
86 auto loopOp = LoopOp::create(builder, start, retTy, operands);
87 loopOp.setCombinedAttr(mlir::acc::CombinedConstructsTypeAttr::get(
88 builder.getContext(), CombinedType<Op>::value));
91 mlir::Block &innerBlock = loopOp.getRegion().emplaceBlock();
92 mlir::OpBuilder::InsertionGuard guardCase(builder);
93 builder.setInsertionPointToEnd(&innerBlock);
95 LexicalScope ls{*
this, start, builder.getInsertionBlock()};
96 ActiveOpenACCLoopRAII activeLoop{*
this, &loopOp};
100 mlir::acc::YieldOp::create(builder, end);
103 emitOpenACCClauses(computeOp, loopOp, dirKind, clauses);
105 updateLoopOpParallelism(loopOp,
false, dirKind);
107 TermOp::create(builder, end);
113template <
typename Op>
114Op CIRGenFunction::emitOpenACCOp(
116 llvm::ArrayRef<const OpenACCClause *> clauses) {
117 llvm::SmallVector<mlir::Type> retTy;
118 llvm::SmallVector<mlir::Value> operands;
119 auto op = Op::create(builder, start, retTy, operands);
121 emitOpenACCClauses(op, dirKind, clauses);
127 mlir::Location start =
getLoc(
s.getSourceRange().getBegin());
128 mlir::Location end =
getLoc(
s.getSourceRange().getEnd());
130 switch (
s.getDirectiveKind()) {
132 return emitOpenACCOpAssociatedStmt<ParallelOp, mlir::acc::YieldOp>(
133 start, end,
s.getDirectiveKind(),
s.clauses(),
s.getStructuredBlock());
135 return emitOpenACCOpAssociatedStmt<SerialOp, mlir::acc::YieldOp>(
136 start, end,
s.getDirectiveKind(),
s.clauses(),
s.getStructuredBlock());
138 return emitOpenACCOpAssociatedStmt<KernelsOp, mlir::acc::TerminatorOp>(
139 start, end,
s.getDirectiveKind(),
s.clauses(),
s.getStructuredBlock());
141 llvm_unreachable(
"invalid compute construct kind");
147 mlir::Location start =
getLoc(
s.getSourceRange().getBegin());
148 mlir::Location end =
getLoc(
s.getSourceRange().getEnd());
150 return emitOpenACCOpAssociatedStmt<DataOp, mlir::acc::TerminatorOp>(
151 start, end,
s.getDirectiveKind(),
s.clauses(),
s.getStructuredBlock());
156 mlir::Location start =
getLoc(
s.getSourceRange().getBegin());
157 emitOpenACCOp<InitOp>(start,
s.getDirectiveKind(),
s.clauses());
158 return mlir::success();
163 mlir::Location start =
getLoc(
s.getSourceRange().getBegin());
164 emitOpenACCOp<SetOp>(start,
s.getDirectiveKind(),
s.clauses());
165 return mlir::success();
170 mlir::Location start =
getLoc(
s.getSourceRange().getBegin());
171 emitOpenACCOp<ShutdownOp>(start,
s.getDirectiveKind(),
s.clauses());
172 return mlir::success();
177 mlir::Location start =
getLoc(
s.getSourceRange().getBegin());
178 auto waitOp = emitOpenACCOp<WaitOp>(start,
s.getDirectiveKind(),
s.clauses());
180 auto createIntExpr = [
this](
const Expr *intExpr) {
182 mlir::Location exprLoc =
cgm.getLoc(intExpr->getBeginLoc());
184 mlir::IntegerType targetType = mlir::IntegerType::get(
186 intExpr->getType()->isSignedIntegerOrEnumerationType()
187 ? mlir::IntegerType::SignednessSemantics::Signed
188 : mlir::IntegerType::SignednessSemantics::Unsigned);
190 auto conversionOp = mlir::UnrealizedConversionCastOp::create(
191 builder, exprLoc, targetType,
expr);
192 return conversionOp.getResult(0);
197 mlir::OpBuilder::InsertionGuard guardCase(builder);
198 builder.setInsertionPoint(waitOp);
200 if (
s.hasDevNumExpr())
201 waitOp.getWaitDevnumMutable().append(createIntExpr(
s.getDevNumExpr()));
203 for (
Expr *QueueExpr :
s.getQueueIdExprs())
204 waitOp.getWaitOperandsMutable().append(createIntExpr(QueueExpr));
207 return mlir::success();
212 mlir::Location start =
getLoc(
s.getSourceRange().getBegin());
213 mlir::Location end =
getLoc(
s.getSourceRange().getEnd());
215 switch (
s.getDirectiveKind()) {
217 return emitOpenACCOpCombinedConstruct<ParallelOp, mlir::acc::YieldOp>(
218 start, end,
s.getDirectiveKind(),
s.clauses(),
s.getLoop());
220 return emitOpenACCOpCombinedConstruct<SerialOp, mlir::acc::YieldOp>(
221 start, end,
s.getDirectiveKind(),
s.clauses(),
s.getLoop());
223 return emitOpenACCOpCombinedConstruct<KernelsOp, mlir::acc::TerminatorOp>(
224 start, end,
s.getDirectiveKind(),
s.clauses(),
s.getLoop());
226 llvm_unreachable(
"invalid compute construct kind");
232 mlir::Location start =
getLoc(
s.getSourceRange().getBegin());
233 mlir::Location end =
getLoc(
s.getSourceRange().getEnd());
235 return emitOpenACCOpAssociatedStmt<HostDataOp, mlir::acc::TerminatorOp>(
236 start, end,
s.getDirectiveKind(),
s.clauses(),
s.getStructuredBlock());
241 mlir::Location start =
getLoc(
s.getSourceRange().getBegin());
242 emitOpenACCOp<EnterDataOp>(start,
s.getDirectiveKind(),
s.clauses());
243 return mlir::success();
248 mlir::Location start =
getLoc(
s.getSourceRange().getBegin());
249 emitOpenACCOp<ExitDataOp>(start,
s.getDirectiveKind(),
s.clauses());
250 return mlir::success();
255 mlir::Location start =
getLoc(
s.getSourceRange().getBegin());
256 emitOpenACCOp<UpdateOp>(start,
s.getDirectiveKind(),
s.clauses());
257 return mlir::success();
269 return mlir::success();
271 mlir::acc::LoopOp loopOp = *activeLoopOp;
273 mlir::OpBuilder::InsertionGuard guard(builder);
274 builder.setInsertionPoint(loopOp);
276 for (
const Expr *var :
s.getVarList()) {
284 loopOp.getCacheOperandsMutable().append(cacheOp.getResult());
287 return mlir::success();
300static mlir::acc::AtomicReadOp
302 mlir::Location start,
313 return mlir::acc::AtomicReadOp::create(builder, start, x, v, resTy,
317static mlir::acc::AtomicWriteOp
319 mlir::Location start,
323 return mlir::acc::AtomicWriteOp::create(builder, start, x,
expr,
327static std::pair<mlir::LogicalResult, mlir::acc::AtomicUpdateOp>
329 mlir::Location start, mlir::Location end,
332 auto op = mlir::acc::AtomicUpdateOp::create(builder, start, x, {});
334 mlir::LogicalResult res = mlir::success();
336 mlir::OpBuilder::InsertionGuard guardCase(builder);
338 std::array<mlir::Type, 1> recipeType{argTy};
339 std::array<mlir::Location, 1> recipeLoc{start};
340 auto *recipeBlock = builder.createBlock(
341 &op.getRegion(), op.getRegion().end(), recipeType, recipeLoc);
342 builder.setInsertionPointToEnd(recipeBlock);
349 auto alloca = cir::AllocaOp::create(
350 builder, start, x.getType(), argTy,
"x_var",
354 alloca.setInitAttr(builder.getUnitAttr());
355 builder.CIRBaseBuilderTy::createStore(start, recipeBlock->getArgument(0),
365 auto load = cir::LoadOp::create(builder, start, {alloca});
366 mlir::acc::YieldOp::create(builder, end, {load});
377 mlir::Location start =
getLoc(
s.getSourceRange().getBegin());
378 mlir::Location end =
getLoc(
s.getSourceRange().getEnd());
381 switch (
s.getAtomicKind()) {
384 mlir::acc::AtomicReadOp op =
386 emitOpenACCClauses(op,
s.getDirectiveKind(),
s.clauses());
387 return mlir::success();
392 emitOpenACCClauses(op,
s.getDirectiveKind(),
s.clauses());
393 return mlir::success();
399 emitOpenACCClauses(op,
s.getDirectiveKind(),
s.clauses());
406 auto op = mlir::acc::AtomicCaptureOp::create(builder, start, {});
407 emitOpenACCClauses(op,
s.getDirectiveKind(),
s.clauses());
408 mlir::LogicalResult res = mlir::success();
410 mlir::OpBuilder::InsertionGuard guardCase(builder);
413 builder.createBlock(&op.getRegion(), op.getRegion().end(), {}, {});
415 builder.setInsertionPointToStart(block);
417 auto terminator = mlir::acc::TerminatorOp::create(builder, end);
422 builder.setInsertionPoint(op);
426 llvm_unreachable(
"invalid form for Capture");
428 mlir::acc::AtomicReadOp first =
430 mlir::acc::AtomicWriteOp second =
433 first->moveBefore(terminator);
434 second->moveBefore(terminator);
438 mlir::acc::AtomicReadOp first =
440 auto [this_res, second] =
444 first->moveBefore(terminator);
445 second->moveBefore(terminator);
449 auto [this_res, first] =
452 mlir::acc::AtomicReadOp second =
455 first->moveBefore(terminator);
456 second->moveBefore(terminator);
465 llvm_unreachable(
"unknown OpenACC atomic kind");
static mlir::acc::AtomicReadOp emitAtomicRead(CIRGenFunction &cgf, CIRGenBuilderTy &builder, mlir::Location start, const OpenACCAtomicConstruct::SingleStmtInfo &inf)
static std::pair< mlir::LogicalResult, mlir::acc::AtomicUpdateOp > emitAtomicUpdate(CIRGenFunction &cgf, CIRGenBuilderTy &builder, mlir::Location start, mlir::Location end, const OpenACCAtomicConstruct::SingleStmtInfo &inf)
static mlir::acc::AtomicWriteOp emitAtomicWrite(CIRGenFunction &cgf, CIRGenBuilderTy &builder, mlir::Location start, const OpenACCAtomicConstruct::SingleStmtInfo &inf)
const VarDecl * getLValueDecl(const Expr *e)
This file defines OpenACC AST classes for statement-level contructs.
__device__ __2f16 float __ockl_bool s
CharUnits getTypeAlignInChars(QualType T) const
Return the ABI-specified alignment of a (complete) type T, in characters.
CharUnits getDeclAlign(const Decl *D, bool ForAlignof=false) const
Return a conservative estimate of the alignment of the specified decl D.
mlir::Type convertType(clang::QualType t)
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)
void replaceAddrOfLocalVar(const clang::VarDecl *vd, Address addr)
LValue emitLValue(const clang::Expr *e)
Emit code to compute a designator that specifies the location of the expression.
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::LogicalResult emitOpenACCHostDataConstruct(const OpenACCHostDataConstruct &s)
mlir::Value emitScalarExpr(const clang::Expr *e, bool ignoreResultAssign=false)
Emit the computation of the specified expression of scalar type.
mlir::MLIRContext & getMLIRContext()
mlir::LogicalResult emitOpenACCEnterDataConstruct(const OpenACCEnterDataConstruct &s)
RValue emitAnyExpr(const clang::Expr *e, AggValueSlot aggSlot=AggValueSlot::ignored(), bool ignoreResult=false)
Emit code to compute the specified expression which can have any type.
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)
mlir::IntegerAttr getSize(CharUnits size)
mlir::Value getPointer() const
mlir::Value getValue() const
Return the value of this scalar value.
This represents one expression.
Expr * IgnoreImpCasts() LLVM_READONLY
Skip past any implicit casts which might surround this expression until reaching a fixed point.
Stmt - This represents one statement.
Represents a variable declaration or definition.
const internal::VariadicDynCastAllOfMatcher< Stmt, Expr > expr
Matches expressions.
The JSON file list parser is used to communicate input to InstallAPI.
OpenACCComputeConstruct(OpenACCDirectiveKind K, SourceLocation Start, SourceLocation DirectiveLoc, SourceLocation End, ArrayRef< const OpenACCClause * > Clauses, Stmt *StructuredBlock)
U cast(CodeGen::Address addr)
enum OpenACCAtomicConstruct::StmtInfo::StmtForm Form
Represents a scope, including function bodies, compound statements, and the substatements of if/while...
llvm::SmallVector< mlir::Value > bounds