clang 22.0.0git
CIRGenOpenACCRecipe.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// Helperes to emit OpenACC clause recipes as CIR code.
10//
11//===----------------------------------------------------------------------===//
12
13#include "CIRGenOpenACCRecipe.h"
14
15namespace clang::CIRGen {
16mlir::Block *OpenACCRecipeBuilderBase::createRecipeBlock(mlir::Region &region,
17 mlir::Type opTy,
18 mlir::Location loc,
19 size_t numBounds,
20 bool isInit) {
22 types.reserve(numBounds + 2);
23 types.push_back(opTy);
24 // The init section is the only one that doesn't have TWO copies of the
25 // operation-type. Copy has a to/from, and destroy has a
26 // 'reference'/'privatized' copy version.
27 if (!isInit)
28 types.push_back(opTy);
29
30 auto boundsTy = mlir::acc::DataBoundsType::get(&cgf.getMLIRContext());
31 for (size_t i = 0; i < numBounds; ++i)
32 types.push_back(boundsTy);
33
34 llvm::SmallVector<mlir::Location> locs{types.size(), loc};
35 return builder.createBlock(&region, region.end(), types, locs);
36}
37
38mlir::Value
39OpenACCRecipeBuilderBase::createBoundsLoop(mlir::Value subscriptedValue,
40 mlir::Value bound,
41 mlir::Location loc, bool inverse) {
42 mlir::Operation *bodyInsertLoc;
43
44 mlir::Type itrTy = cgf.cgm.convertType(cgf.getContext().UnsignedLongLongTy);
45 auto itrPtrTy = cir::PointerType::get(itrTy);
46 mlir::IntegerAttr itrAlign =
47 cgf.cgm.getSize(cgf.getContext().getTypeAlignInChars(
48 cgf.getContext().UnsignedLongLongTy));
49 auto idxType = mlir::IndexType::get(&cgf.getMLIRContext());
50
51 auto doSubscriptOp = [&](mlir::Value subVal,
52 cir::LoadOp idxLoad) -> mlir::Value {
53 auto eltTy = cast<cir::PointerType>(subVal.getType()).getPointee();
54
55 if (auto arrayTy = dyn_cast<cir::ArrayType>(eltTy))
56 return builder.getArrayElement(loc, loc, subVal, arrayTy.getElementType(),
57 idxLoad.getResult(),
58 /*shouldDecay=*/true);
59
60 assert(isa<cir::PointerType>(eltTy));
61
62 auto eltLoad = cir::LoadOp::create(builder, loc, {subVal});
63
64 return cir::PtrStrideOp::create(builder, loc, eltLoad.getType(), eltLoad,
65 idxLoad.getResult())
66 .getResult();
67 };
68
69 auto forStmtBuilder = [&]() {
70 // get the lower and upper bound for iterating over.
71 auto lowerBoundVal =
72 mlir::acc::GetLowerboundOp::create(builder, loc, idxType, bound);
73 auto lbConversion = mlir::UnrealizedConversionCastOp::create(
74 builder, loc, itrTy, lowerBoundVal.getResult());
75 auto upperBoundVal =
76 mlir::acc::GetUpperboundOp::create(builder, loc, idxType, bound);
77 auto ubConversion = mlir::UnrealizedConversionCastOp::create(
78 builder, loc, itrTy, upperBoundVal.getResult());
79
80 // Create a memory location for the iterator.
81 auto itr =
82 cir::AllocaOp::create(builder, loc, itrPtrTy, itrTy, "iter", itrAlign);
83 // Store to the iterator: either lower bound, or if inverse loop, upper
84 // bound.
85 if (inverse) {
86 cir::ConstantOp constOne = builder.getConstInt(loc, itrTy, 1);
87
88 auto sub =
89 cir::BinOp::create(builder, loc, itrTy, cir::BinOpKind::Sub,
90 ubConversion.getResult(0), constOne.getResult());
91
92 // Upperbound is exclusive, so subtract 1.
93 builder.CIRBaseBuilderTy::createStore(loc, sub.getResult(), itr);
94 } else {
95 // Lowerbound is inclusive, so we can include it.
96 builder.CIRBaseBuilderTy::createStore(loc, lbConversion.getResult(0),
97 itr);
98 }
99 // Save the 'end' iterator based on whether we are inverted or not. This
100 // end iterator never changes, so we can just get it and convert it, so no
101 // need to store/load/etc.
102 auto endItr = inverse ? lbConversion : ubConversion;
103
104 builder.createFor(
105 loc,
106 /*condBuilder=*/
107 [&](mlir::OpBuilder &b, mlir::Location loc) {
108 auto loadCur = cir::LoadOp::create(builder, loc, {itr});
109 // Use 'not equal' since we are just doing an increment/decrement.
110 auto cmp = builder.createCompare(
111 loc, inverse ? cir::CmpOpKind::ge : cir::CmpOpKind::lt,
112 loadCur.getResult(), endItr.getResult(0));
113 builder.createCondition(cmp);
114 },
115 /*bodyBuilder=*/
116 [&](mlir::OpBuilder &b, mlir::Location loc) {
117 auto load = cir::LoadOp::create(builder, loc, {itr});
118
119 if (subscriptedValue)
120 subscriptedValue = doSubscriptOp(subscriptedValue, load);
121 bodyInsertLoc = builder.createYield(loc);
122 },
123 /*stepBuilder=*/
124 [&](mlir::OpBuilder &b, mlir::Location loc) {
125 auto load = cir::LoadOp::create(builder, loc, {itr});
126 auto unary = cir::UnaryOp::create(builder, loc, load.getType(),
127 inverse ? cir::UnaryOpKind::Dec
128 : cir::UnaryOpKind::Inc,
129 load.getResult());
130 builder.CIRBaseBuilderTy::createStore(loc, unary.getResult(), itr);
131 builder.createYield(loc);
132 });
133 };
134
135 cir::ScopeOp::create(builder, loc,
136 [&](mlir::OpBuilder &b, mlir::Location loc) {
137 forStmtBuilder();
138 builder.createYield(loc);
139 });
140
141 // Leave the insertion point to be inside the body, so we can loop over
142 // these things.
143 builder.setInsertionPoint(bodyInsertLoc);
144 return subscriptedValue;
145}
146
147mlir::acc::ReductionOperator
149 switch (op) {
151 return mlir::acc::ReductionOperator::AccAdd;
153 return mlir::acc::ReductionOperator::AccMul;
155 return mlir::acc::ReductionOperator::AccMax;
157 return mlir::acc::ReductionOperator::AccMin;
159 return mlir::acc::ReductionOperator::AccIand;
161 return mlir::acc::ReductionOperator::AccIor;
163 return mlir::acc::ReductionOperator::AccXor;
165 return mlir::acc::ReductionOperator::AccLand;
167 return mlir::acc::ReductionOperator::AccLor;
169 llvm_unreachable("invalid reduction operator");
170 }
171
172 llvm_unreachable("invalid reduction operator");
173}
174
175// This function generates the 'destroy' section for a recipe. Note
176// that this function is not 'insertion point' clean, in that it alters the
177// insertion point to be inside of the 'destroy' section of the recipe, but
178// doesn't restore it aftewards.
180 mlir::Location loc, mlir::Location locEnd, mlir::Value mainOp,
181 CharUnits alignment, QualType origType, size_t numBounds, QualType baseType,
182 mlir::Region &destroyRegion) {
183 mlir::Block *block = createRecipeBlock(destroyRegion, mainOp.getType(), loc,
184 numBounds, /*isInit=*/false);
185 builder.setInsertionPointToEnd(&destroyRegion.back());
186 CIRGenFunction::LexicalScope ls(cgf, loc, block);
187
188 mlir::Type elementTy =
189 mlir::cast<cir::PointerType>(mainOp.getType()).getPointee();
190 auto emitDestroy = [&](mlir::Value var, mlir::Type ty) {
191 Address addr{var, ty, alignment};
192 cgf.emitDestroy(addr, origType,
193 cgf.getDestroyer(QualType::DK_cxx_destructor));
194 };
195
196 if (numBounds) {
197 mlir::OpBuilder::InsertionGuard guardCase(builder);
198 // Get the range of bounds arguments, which are all but the 1st 2. 1st is
199 // a 'reference', 2nd is the 'private' variant we need to destroy from.
201 block->getArguments().drop_front(2);
202
203 mlir::Value subscriptedValue = block->getArgument(1);
204 for (mlir::BlockArgument boundArg : llvm::reverse(boundsRange))
205 subscriptedValue = createBoundsLoop(subscriptedValue, boundArg, loc,
206 /*inverse=*/true);
207
208 emitDestroy(subscriptedValue, cgf.cgm.convertType(origType));
209 } else {
210 // If we don't have any bounds, we can just destroy the variable directly.
211 // The destroy region has a signature of "original item, privatized item".
212 // So the 2nd item is the one that needs destroying, the former is just
213 // for reference and we don't really have a need for it at the moment.
214 emitDestroy(block->getArgument(1), elementTy);
215 }
216
217 mlir::acc::YieldOp::create(builder, locEnd);
218}
219
220// TODO: OpenACC: When we get this implemented for the reduction/firstprivate,
221// this might end up re-merging with createRecipeInitCopy. For now, keep it
222// separate until we're sure what everything looks like to keep this as clean
223// as possible.
225 mlir::Location loc, mlir::Location locEnd, SourceRange exprRange,
226 mlir::Value mainOp, mlir::acc::PrivateRecipeOp recipe, size_t numBounds,
227 llvm::ArrayRef<QualType> boundTypes, const VarDecl *allocaDecl,
228 QualType origType, const Expr *initExpr) {
229 assert(allocaDecl && "Required recipe variable not set?");
230 CIRGenFunction::DeclMapRevertingRAII declMapRAII{cgf, allocaDecl};
231
232 mlir::Block *block =
233 createRecipeBlock(recipe.getInitRegion(), mainOp.getType(), loc,
234 numBounds, /*isInit=*/true);
235 builder.setInsertionPointToEnd(&recipe.getInitRegion().back());
236 CIRGenFunction::LexicalScope ls(cgf, loc, block);
237
238 const Type *allocaPointeeType =
239 allocaDecl->getType()->getPointeeOrArrayElementType();
240 // We are OK with no init for builtins, arrays of builtins, or pointers,
241 // else we should NYI so we know to go look for these.
242 if (cgf.getContext().getLangOpts().CPlusPlus && !allocaDecl->getInit() &&
243 !allocaDecl->getType()->isPointerType() &&
244 !allocaPointeeType->isBuiltinType() &&
245 !allocaPointeeType->isPointerType()) {
246 // If we don't have any initialization recipe, we failed during Sema to
247 // initialize this correctly. If we disable the
248 // Sema::TentativeAnalysisScopes in SemaOpenACC::CreateInitRecipe, it'll
249 // emit an error to tell us. However, emitting those errors during
250 // production is a violation of the standard, so we cannot do them.
251 cgf.cgm.errorNYI(exprRange, "private default-init recipe");
252 }
253
254 if (!numBounds) {
255 // This is an 'easy' case, we just have to use the builtin init stuff to
256 // initialize this variable correctly.
257 CIRGenFunction::AutoVarEmission tempDeclEmission =
258 cgf.emitAutoVarAlloca(*allocaDecl, builder.saveInsertionPoint());
259 cgf.emitAutoVarInit(tempDeclEmission);
260 } else {
261 cgf.cgm.errorNYI(exprRange, "private-init with bounds");
262 }
263
264 mlir::acc::YieldOp::create(builder, locEnd);
265}
266
268 mlir::Location loc, mlir::Location locEnd, mlir::Value mainOp,
269 CIRGenFunction::AutoVarEmission tempDeclEmission,
270 mlir::acc::FirstprivateRecipeOp recipe, const VarDecl *varRecipe,
271 const VarDecl *temporary) {
272 mlir::Block *block =
273 createRecipeBlock(recipe.getCopyRegion(), mainOp.getType(), loc,
274 /*numBounds=*/0, /*isInit=*/false);
275 builder.setInsertionPointToEnd(&recipe.getCopyRegion().back());
276 CIRGenFunction::LexicalScope ls(cgf, loc, block);
277
278 mlir::BlockArgument fromArg = block->getArgument(0);
279 mlir::BlockArgument toArg = block->getArgument(1);
280
281 mlir::Type elementTy =
282 mlir::cast<cir::PointerType>(mainOp.getType()).getPointee();
283
284 // Set the address of the emission to be the argument, so that we initialize
285 // that instead of the variable in the other block.
286 tempDeclEmission.setAllocatedAddress(
287 Address{toArg, elementTy, cgf.getContext().getDeclAlign(varRecipe)});
288 tempDeclEmission.EmittedAsOffload = true;
289
290 CIRGenFunction::DeclMapRevertingRAII declMapRAII{cgf, temporary};
291 cgf.setAddrOfLocalVar(
292 temporary,
293 Address{fromArg, elementTy, cgf.getContext().getDeclAlign(varRecipe)});
294
295 cgf.emitAutoVarInit(tempDeclEmission);
296 mlir::acc::YieldOp::create(builder, locEnd);
297}
298// This function generates the 'combiner' section for a reduction recipe. Note
299// that this function is not 'insertion point' clean, in that it alters the
300// insertion point to be inside of the 'combiner' section of the recipe, but
301// doesn't restore it aftewards.
303 mlir::Location loc, mlir::Location locEnd, mlir::Value mainOp,
304 mlir::acc::ReductionRecipeOp recipe) {
305 mlir::Block *block = builder.createBlock(
306 &recipe.getCombinerRegion(), recipe.getCombinerRegion().end(),
307 {mainOp.getType(), mainOp.getType()}, {loc, loc});
308 builder.setInsertionPointToEnd(&recipe.getCombinerRegion().back());
309 CIRGenFunction::LexicalScope ls(cgf, loc, block);
310
311 mlir::BlockArgument lhsArg = block->getArgument(0);
312
313 mlir::acc::YieldOp::create(builder, locEnd, lhsArg);
314}
315
316} // namespace clang::CIRGen
__device__ __2f16 b
void createFirstprivateRecipeCopy(mlir::Location loc, mlir::Location locEnd, mlir::Value mainOp, CIRGenFunction::AutoVarEmission tempDeclEmission, mlir::acc::FirstprivateRecipeOp recipe, const VarDecl *varRecipe, const VarDecl *temporary)
void createPrivateInitRecipe(mlir::Location loc, mlir::Location locEnd, SourceRange exprRange, mlir::Value mainOp, mlir::acc::PrivateRecipeOp recipe, size_t numBounds, llvm::ArrayRef< QualType > boundTypes, const VarDecl *allocaDecl, QualType origType, const Expr *initExpr)
mlir::Value createBoundsLoop(mlir::Value subscriptedValue, mlir::Value bound, mlir::Location loc, bool inverse)
mlir::acc::ReductionOperator convertReductionOp(OpenACCReductionOperator op)
void createReductionRecipeCombiner(mlir::Location loc, mlir::Location locEnd, mlir::Value mainOp, mlir::acc::ReductionRecipeOp recipe)
void createRecipeDestroySection(mlir::Location loc, mlir::Location locEnd, mlir::Value mainOp, CharUnits alignment, QualType origType, size_t numBounds, QualType baseType, mlir::Region &destroyRegion)
mlir::Block * createRecipeBlock(mlir::Region &region, mlir::Type opTy, mlir::Location loc, size_t numBounds, bool isInit)
CharUnits - This is an opaque type for sizes expressed in character units.
Definition CharUnits.h:38
This represents one expression.
Definition Expr.h:112
A (possibly-)qualified type.
Definition TypeBase.h:937
A trivial tuple used to represent a source range.
const Type * getPointeeOrArrayElementType() const
If this is a pointer type, return the pointee type.
Definition TypeBase.h:9058
bool isPointerType() const
Definition TypeBase.h:8522
bool isBuiltinType() const
Helper methods to distinguish type categories.
Definition TypeBase.h:8645
QualType getType() const
Definition Decl.h:722
Represents a variable declaration or definition.
Definition Decl.h:925
const Expr * getInit() const
Definition Decl.h:1367
@ Type
The l-value was considered opaque, so the alignment was determined from a type.
OpenACCReductionOperator
@ Invalid
Invalid Reduction Clause Kind.
bool isa(CodeGen::Address addr)
Definition Address.h:330
U cast(CodeGen::Address addr)
Definition Address.h:327
bool EmittedAsOffload
True if the variable was emitted as an offload recipe, and thus doesn't have the same sort of alloca ...
Represents a scope, including function bodies, compound statements, and the substatements of if/while...