clang 22.0.0git
CIRGenStmtOpenACC.cpp
Go to the documentation of this file.
1//===----------------------------------------------------------------------===//
2//
3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4// See https://llvm.org/LICENSE.txt for license information.
5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6//
7//===----------------------------------------------------------------------===//
8//
9// Emit OpenACC Stmt nodes as CIR code.
10//
11//===----------------------------------------------------------------------===//
12
13#include "CIRGenBuilder.h"
14#include "CIRGenFunction.h"
15#include "mlir/Dialect/OpenACC/OpenACC.h"
18
19using namespace clang;
20using namespace clang::CIRGen;
21using namespace cir;
22using namespace mlir::acc;
23
24template <typename Op, typename TermOp>
25mlir::LogicalResult CIRGenFunction::emitOpenACCOpAssociatedStmt(
26 mlir::Location start, mlir::Location end, OpenACCDirectiveKind dirKind,
28 const Stmt *associatedStmt) {
29 mlir::LogicalResult res = mlir::success();
30
31 llvm::SmallVector<mlir::Type> retTy;
32 llvm::SmallVector<mlir::Value> operands;
33 auto op = Op::create(builder, start, retTy, operands);
34
35 emitOpenACCClauses(op, dirKind, dirLoc, clauses);
36
37 {
38 mlir::Block &block = op.getRegion().emplaceBlock();
39 mlir::OpBuilder::InsertionGuard guardCase(builder);
40 builder.setInsertionPointToEnd(&block);
41
42 LexicalScope ls{*this, start, builder.getInsertionBlock()};
43 res = emitStmt(associatedStmt, /*useCurrentScope=*/true);
44
45 TermOp::create(builder, end);
46 }
47 return res;
48}
49
50namespace {
51template <typename Op> struct CombinedType;
52template <> struct CombinedType<ParallelOp> {
53 static constexpr mlir::acc::CombinedConstructsType value =
54 mlir::acc::CombinedConstructsType::ParallelLoop;
55};
56template <> struct CombinedType<SerialOp> {
57 static constexpr mlir::acc::CombinedConstructsType value =
58 mlir::acc::CombinedConstructsType::SerialLoop;
59};
60template <> struct CombinedType<KernelsOp> {
61 static constexpr mlir::acc::CombinedConstructsType value =
62 mlir::acc::CombinedConstructsType::KernelsLoop;
63};
64} // namespace
65
66template <typename Op, typename TermOp>
67mlir::LogicalResult CIRGenFunction::emitOpenACCOpCombinedConstruct(
68 mlir::Location start, mlir::Location end, OpenACCDirectiveKind dirKind,
69 SourceLocation dirLoc, llvm::ArrayRef<const OpenACCClause *> clauses,
70 const Stmt *loopStmt) {
71 mlir::LogicalResult res = mlir::success();
72
73 llvm::SmallVector<mlir::Type> retTy;
74 llvm::SmallVector<mlir::Value> operands;
75
76 auto computeOp = Op::create(builder, start, retTy, operands);
77 computeOp.setCombinedAttr(builder.getUnitAttr());
78 mlir::acc::LoopOp loopOp;
79
80 // First, emit the bodies of both operations, with the loop inside the body of
81 // the combined construct.
82 {
83 mlir::Block &block = computeOp.getRegion().emplaceBlock();
84 mlir::OpBuilder::InsertionGuard guardCase(builder);
85 builder.setInsertionPointToEnd(&block);
86
87 LexicalScope ls{*this, start, builder.getInsertionBlock()};
88 auto loopOp = LoopOp::create(builder, start, retTy, operands);
89 loopOp.setCombinedAttr(mlir::acc::CombinedConstructsTypeAttr::get(
90 builder.getContext(), CombinedType<Op>::value));
91
92 {
93 mlir::Block &innerBlock = loopOp.getRegion().emplaceBlock();
94 mlir::OpBuilder::InsertionGuard guardCase(builder);
95 builder.setInsertionPointToEnd(&innerBlock);
96
97 LexicalScope ls{*this, start, builder.getInsertionBlock()};
98 ActiveOpenACCLoopRAII activeLoop{*this, &loopOp};
99
100 res = emitStmt(loopStmt, /*useCurrentScope=*/true);
101
102 mlir::acc::YieldOp::create(builder, end);
103 }
104
105 emitOpenACCClauses(computeOp, loopOp, dirKind, dirLoc, clauses);
106
107 updateLoopOpParallelism(loopOp, /*isOrphan=*/false, dirKind);
108
109 TermOp::create(builder, end);
110 }
111
112 return res;
113}
114
115template <typename Op>
116Op CIRGenFunction::emitOpenACCOp(
117 mlir::Location start, OpenACCDirectiveKind dirKind, SourceLocation dirLoc,
118 llvm::ArrayRef<const OpenACCClause *> clauses) {
119 llvm::SmallVector<mlir::Type> retTy;
120 llvm::SmallVector<mlir::Value> operands;
121 auto op = Op::create(builder, start, retTy, operands);
122
123 emitOpenACCClauses(op, dirKind, dirLoc, clauses);
124 return op;
125}
126
127mlir::LogicalResult
129 mlir::Location start = getLoc(s.getSourceRange().getBegin());
130 mlir::Location end = getLoc(s.getSourceRange().getEnd());
131
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());
145 default:
146 llvm_unreachable("invalid compute construct kind");
147 }
148}
149
150mlir::LogicalResult
152 mlir::Location start = getLoc(s.getSourceRange().getBegin());
153 mlir::Location end = getLoc(s.getSourceRange().getEnd());
154
155 return emitOpenACCOpAssociatedStmt<DataOp, mlir::acc::TerminatorOp>(
156 start, end, s.getDirectiveKind(), s.getDirectiveLoc(), s.clauses(),
157 s.getStructuredBlock());
158}
159
160mlir::LogicalResult
162 mlir::Location start = getLoc(s.getSourceRange().getBegin());
163 emitOpenACCOp<InitOp>(start, s.getDirectiveKind(), s.getDirectiveLoc(),
164 s.clauses());
165 return mlir::success();
166}
167
168mlir::LogicalResult
170 mlir::Location start = getLoc(s.getSourceRange().getBegin());
171 emitOpenACCOp<SetOp>(start, s.getDirectiveKind(), s.getDirectiveLoc(),
172 s.clauses());
173 return mlir::success();
174}
175
178 mlir::Location start = getLoc(s.getSourceRange().getBegin());
179 emitOpenACCOp<ShutdownOp>(start, s.getDirectiveKind(),
180 s.getDirectiveLoc(), s.clauses());
181 return mlir::success();
182}
183
184mlir::LogicalResult
186 mlir::Location start = getLoc(s.getSourceRange().getBegin());
187 auto waitOp = emitOpenACCOp<WaitOp>(start, s.getDirectiveKind(),
188 s.getDirectiveLoc(), s.clauses());
189
190 auto createIntExpr = [this](const Expr *intExpr) {
191 mlir::Value expr = emitScalarExpr(intExpr);
192 mlir::Location exprLoc = cgm.getLoc(intExpr->getBeginLoc());
193
194 mlir::IntegerType targetType = mlir::IntegerType::get(
195 &getMLIRContext(), getContext().getIntWidth(intExpr->getType()),
196 intExpr->getType()->isSignedIntegerOrEnumerationType()
197 ? mlir::IntegerType::SignednessSemantics::Signed
198 : mlir::IntegerType::SignednessSemantics::Unsigned);
199
200 auto conversionOp = mlir::UnrealizedConversionCastOp::create(
201 builder, exprLoc, targetType, expr);
202 return conversionOp.getResult(0);
203 };
204
205 // Emit the correct 'wait' clauses.
206 {
207 mlir::OpBuilder::InsertionGuard guardCase(builder);
208 builder.setInsertionPoint(waitOp);
209
210 if (s.hasDevNumExpr())
211 waitOp.getWaitDevnumMutable().append(createIntExpr(s.getDevNumExpr()));
212
213 for (Expr *QueueExpr : s.getQueueIdExprs())
214 waitOp.getWaitOperandsMutable().append(createIntExpr(QueueExpr));
215 }
216
217 return mlir::success();
218}
219
222 mlir::Location start = getLoc(s.getSourceRange().getBegin());
223 mlir::Location end = getLoc(s.getSourceRange().getEnd());
224
225 switch (s.getDirectiveKind()) {
227 return emitOpenACCOpCombinedConstruct<ParallelOp, mlir::acc::YieldOp>(
228 start, end, s.getDirectiveKind(), s.getDirectiveLoc(), s.clauses(),
229 s.getLoop());
231 return emitOpenACCOpCombinedConstruct<SerialOp, mlir::acc::YieldOp>(
232 start, end, s.getDirectiveKind(), s.getDirectiveLoc(), s.clauses(),
233 s.getLoop());
235 return emitOpenACCOpCombinedConstruct<KernelsOp, mlir::acc::TerminatorOp>(
236 start, end, s.getDirectiveKind(), s.getDirectiveLoc(), s.clauses(),
237 s.getLoop());
238 default:
239 llvm_unreachable("invalid compute construct kind");
240 }
241}
242
245 mlir::Location start = getLoc(s.getSourceRange().getBegin());
246 mlir::Location end = getLoc(s.getSourceRange().getEnd());
247
248 return emitOpenACCOpAssociatedStmt<HostDataOp, mlir::acc::TerminatorOp>(
249 start, end, s.getDirectiveKind(), s.getDirectiveLoc(), s.clauses(),
250 s.getStructuredBlock());
251}
252
255 mlir::Location start = getLoc(s.getSourceRange().getBegin());
256 emitOpenACCOp<EnterDataOp>(start, s.getDirectiveKind(), s.getDirectiveLoc(),
257 s.clauses());
258 return mlir::success();
259}
260
263 mlir::Location start = getLoc(s.getSourceRange().getBegin());
264 emitOpenACCOp<ExitDataOp>(start, s.getDirectiveKind(), s.getDirectiveLoc(),
265 s.clauses());
266 return mlir::success();
267}
268
269mlir::LogicalResult
271 mlir::Location start = getLoc(s.getSourceRange().getBegin());
272 emitOpenACCOp<UpdateOp>(start, s.getDirectiveKind(), s.getDirectiveLoc(),
273 s.clauses());
274 return mlir::success();
275}
276
277mlir::LogicalResult
279 // The 'cache' directive 'may' be at the top of a loop by standard, but
280 // doesn't have to be. Additionally, there is nothing that requires this be a
281 // loop affected by an OpenACC pragma. Sema doesn't do any level of
282 // enforcement here, since it isn't particularly valuable to do so thanks to
283 // that. Instead, we treat cache as a 'noop' if there is no acc.loop to apply
284 // it to.
285 if (!activeLoopOp)
286 return mlir::success();
287
288 mlir::acc::LoopOp loopOp = *activeLoopOp;
289
290 mlir::OpBuilder::InsertionGuard guard(builder);
291 builder.setInsertionPoint(loopOp);
292
293 for (const Expr *var : s.getVarList()) {
296
297 auto cacheOp = CacheOp::create(builder, opInfo.beginLoc, opInfo.varValue,
298 /*structured=*/false, /*implicit=*/false,
299 opInfo.name, opInfo.bounds);
300
301 loopOp.getCacheOperandsMutable().append(cacheOp.getResult());
302 }
303
304 return mlir::success();
305}
306
307const VarDecl *getLValueDecl(const Expr *e) {
308 // We are going to assume that after stripping implicit casts, that the LValue
309 // is just a DRE around the var-decl.
310
311 e = e->IgnoreImpCasts();
312
313 const auto *dre = cast<DeclRefExpr>(e);
314 return cast<VarDecl>(dre->getDecl());
315}
316
317mlir::LogicalResult
319 // For now, we are only support 'read'/'write'/'update', so diagnose. We can
320 // switch on the kind later once we implement the 'capture' form.
321 if (s.getAtomicKind() == OpenACCAtomicKind::Capture) {
322 cgm.errorNYI(s.getSourceRange(), "OpenACC Atomic Construct");
323 return mlir::failure();
324 }
325
326 // While Atomic is an 'associated statement' construct, it 'steals' the
327 // expression it is associated with rather than emitting it inside of it. So
328 // it has custom emit logic.
329 mlir::Location start = getLoc(s.getSourceRange().getBegin());
330 mlir::Location end = getLoc(s.getSourceRange().getEnd());
331 OpenACCAtomicConstruct::StmtInfo inf = s.getAssociatedStmtInfo();
332
333 switch (s.getAtomicKind()) {
335 llvm_unreachable("Unimplemented atomic construct type, should have "
336 "diagnosed/returned above");
337 return mlir::failure();
339
340 // Atomic 'read' only permits 'v = x', where v and x are both scalar L
341 // values. The getAssociatedStmtInfo strips off implicit casts, which
342 // includes implicit conversions and L-to-R-Value conversions, so we can
343 // just emit it as an L value. The Flang implementation has no problem with
344 // different types, so it appears that the dialect can handle the
345 // conversions.
346 mlir::Value v = emitLValue(inf.V).getPointer();
347 mlir::Value x = emitLValue(inf.X).getPointer();
348 mlir::Type resTy = convertType(inf.V->getType());
349 auto op = mlir::acc::AtomicReadOp::create(builder, start, x, v, resTy,
350 /*ifCond=*/{});
351 emitOpenACCClauses(op, s.getDirectiveKind(), s.getDirectiveLoc(),
352 s.clauses());
353 return mlir::success();
354 }
356 mlir::Value x = emitLValue(inf.X).getPointer();
357 mlir::Value expr = emitAnyExpr(inf.RefExpr).getValue();
358 auto op = mlir::acc::AtomicWriteOp::create(builder, start, x, expr,
359 /*ifCond=*/{});
360 emitOpenACCClauses(op, s.getDirectiveKind(), s.getDirectiveLoc(),
361 s.clauses());
362 return mlir::success();
363 }
366 mlir::Value x = emitLValue(inf.X).getPointer();
367 auto op =
368 mlir::acc::AtomicUpdateOp::create(builder, start, x, /*ifCond=*/{});
369 emitOpenACCClauses(op, s.getDirectiveKind(), s.getDirectiveLoc(),
370 s.clauses());
371 mlir::LogicalResult res = mlir::success();
372 {
373 mlir::OpBuilder::InsertionGuard guardCase(builder);
374 mlir::Type argTy = cast<cir::PointerType>(x.getType()).getPointee();
375 std::array<mlir::Type, 1> recipeType{argTy};
376 std::array<mlir::Location, 1> recipeLoc{start};
377 mlir::Block *recipeBlock = builder.createBlock(
378 &op.getRegion(), op.getRegion().end(), recipeType, recipeLoc);
379 builder.setInsertionPointToEnd(recipeBlock);
380
381 // Since we have an initial value that we know is a scalar type, we can
382 // just emit the entire statement here after sneaking-in our 'alloca' in
383 // the right place, then loading out of it. Flang does a lot less work
384 // (probably does its own emitting!), but we have more complicated AST
385 // nodes to worry about, so we can just count on opt to remove the extra
386 // alloca/load/store set.
387 auto alloca = cir::AllocaOp::create(
388 builder, start, x.getType(), argTy, "x_var",
389 cgm.getSize(getContext().getTypeAlignInChars(inf.X->getType())));
390
391 alloca.setInitAttr(mlir::UnitAttr::get(&getMLIRContext()));
392 builder.CIRBaseBuilderTy::createStore(start, recipeBlock->getArgument(0),
393 alloca);
394
395 const VarDecl *xval = getLValueDecl(inf.X);
396 CIRGenFunction::DeclMapRevertingRAII declMapRAII{*this, xval};
398 xval, Address{alloca, argTy, getContext().getDeclAlign(xval)});
399
400 res = emitStmt(s.getAssociatedStmt(), /*useCurrentScope=*/true);
401
402 auto load = cir::LoadOp::create(builder, start, {alloca});
403 mlir::acc::YieldOp::create(builder, end, {load});
404 }
405
406 return res;
407 }
408 }
409
410 llvm_unreachable("unknown OpenACC atomic kind");
411}
const VarDecl * getLValueDecl(const Expr *e)
This file defines OpenACC AST classes for statement-level contructs.
__device__ __2f16 float __ockl_bool s
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::Value getPointer() const
mlir::Value getValue() const
Return the value of this scalar value.
Definition CIRGenValue.h:56
This represents one expression.
Definition Expr.h:112
Expr * IgnoreImpCasts() LLVM_READONLY
Skip past any implicit casts which might surround this expression until reaching a fixed point.
Definition Expr.cpp:3065
QualType getType() const
Definition Expr.h:144
Encodes a location in the source.
Stmt - This represents one statement.
Definition Stmt.h:85
Represents a variable declaration or definition.
Definition Decl.h:926
const internal::VariadicDynCastAllOfMatcher< Stmt, Expr > expr
Matches expressions.
The JSON file list parser is used to communicate input to InstallAPI.
OpenACCDirectiveKind
OpenACCComputeConstruct(OpenACCDirectiveKind K, SourceLocation Start, SourceLocation DirectiveLoc, SourceLocation End, ArrayRef< const OpenACCClause * > Clauses, Stmt *StructuredBlock)
U cast(CodeGen::Address addr)
Definition Address.h:327
Represents a scope, including function bodies, compound statements, and the substatements of if/while...