clang 22.0.0git
CIRGenOpenACCClause.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 clause nodes as CIR code.
10//
11//===----------------------------------------------------------------------===//
12
13#include <type_traits>
14
15#include "CIRGenCXXABI.h"
16#include "CIRGenFunction.h"
17
18#include "clang/AST/ExprCXX.h"
19
20#include "mlir/Dialect/Arith/IR/Arith.h"
21#include "mlir/Dialect/OpenACC/OpenACC.h"
22#include "llvm/ADT/TypeSwitch.h"
23
24using namespace clang;
25using namespace clang::CIRGen;
26
27namespace {
28// Simple type-trait to see if the first template arg is one of the list, so we
29// can tell whether to `if-constexpr` a bunch of stuff.
30template <typename ToTest, typename T, typename... Tys>
31constexpr bool isOneOfTypes =
32 std::is_same_v<ToTest, T> || isOneOfTypes<ToTest, Tys...>;
33template <typename ToTest, typename T>
34constexpr bool isOneOfTypes<ToTest, T> = std::is_same_v<ToTest, T>;
35
36// Holds information for emitting clauses for a combined construct. We
37// instantiate the clause emitter with this type so that it can use
38// if-constexpr to specially handle these.
39template <typename CompOpTy> struct CombinedConstructClauseInfo {
40 using ComputeOpTy = CompOpTy;
41 ComputeOpTy computeOp;
42 mlir::acc::LoopOp loopOp;
43};
44template <typename ToTest> constexpr bool isCombinedType = false;
45template <typename T>
46constexpr bool isCombinedType<CombinedConstructClauseInfo<T>> = true;
47
48template <typename OpTy>
49class OpenACCClauseCIREmitter final
50 : public OpenACCClauseVisitor<OpenACCClauseCIREmitter<OpTy>> {
51 // Necessary for combined constructs.
52 template <typename FriendOpTy> friend class OpenACCClauseCIREmitter;
53
54 OpTy &operation;
57
58 // This is necessary since a few of the clauses emit differently based on the
59 // directive kind they are attached to.
61 // TODO(cir): This source location should be able to go away once the NYI
62 // diagnostics are gone.
63 SourceLocation dirLoc;
64
66 // Keep track of the async-clause so that we can shortcut updating the data
67 // operands async clauses.
68 bool hasAsyncClause = false;
69 // Keep track of the data operands so that we can update their async clauses.
71
72 void clauseNotImplemented(const OpenACCClause &c) {
73 cgf.cgm.errorNYI(c.getSourceRange(), "OpenACC Clause", c.getClauseKind());
74 }
75
76 void setLastDeviceTypeClause(const OpenACCDeviceTypeClause &clause) {
77 lastDeviceTypeValues.clear();
78
79 for (const DeviceTypeArgument &arg : clause.getArchitectures())
80 lastDeviceTypeValues.push_back(decodeDeviceType(arg.getIdentifierInfo()));
81 }
82
83 mlir::Value emitIntExpr(const Expr *intExpr) {
84 return cgf.emitOpenACCIntExpr(intExpr);
85 }
86
87 // 'condition' as an OpenACC grammar production is used for 'if' and (some
88 // variants of) 'self'. It needs to be emitted as a signless-1-bit value, so
89 // this function emits the expression, then sets the unrealized conversion
90 // cast correctly, and returns the completed value.
91 mlir::Value createCondition(const Expr *condExpr) {
92 mlir::Value condition = cgf.evaluateExprAsBool(condExpr);
93 mlir::Location exprLoc = cgf.cgm.getLoc(condExpr->getBeginLoc());
94 mlir::IntegerType targetType = mlir::IntegerType::get(
95 &cgf.getMLIRContext(), /*width=*/1,
96 mlir::IntegerType::SignednessSemantics::Signless);
97 auto conversionOp = builder.create<mlir::UnrealizedConversionCastOp>(
98 exprLoc, targetType, condition);
99 return conversionOp.getResult(0);
100 }
101
102 mlir::Value createConstantInt(mlir::Location loc, unsigned width,
103 int64_t value) {
104 return cgf.createOpenACCConstantInt(loc, width, value);
105 mlir::IntegerType ty = mlir::IntegerType::get(
106 &cgf.getMLIRContext(), width,
107 mlir::IntegerType::SignednessSemantics::Signless);
108 auto constOp = builder.create<mlir::arith::ConstantOp>(
109 loc, builder.getIntegerAttr(ty, value));
110
111 return constOp.getResult();
112 }
113
114 mlir::Value createConstantInt(SourceLocation loc, unsigned width,
115 int64_t value) {
116 return createConstantInt(cgf.cgm.getLoc(loc), width, value);
117 }
118
119 mlir::acc::DeviceType decodeDeviceType(const IdentifierInfo *ii) {
120 // '*' case leaves no identifier-info, just a nullptr.
121 if (!ii)
122 return mlir::acc::DeviceType::Star;
123 return llvm::StringSwitch<mlir::acc::DeviceType>(ii->getName())
124 .CaseLower("default", mlir::acc::DeviceType::Default)
125 .CaseLower("host", mlir::acc::DeviceType::Host)
126 .CaseLower("multicore", mlir::acc::DeviceType::Multicore)
127 .CasesLower("nvidia", "acc_device_nvidia",
128 mlir::acc::DeviceType::Nvidia)
129 .CaseLower("radeon", mlir::acc::DeviceType::Radeon);
130 }
131
132 mlir::acc::GangArgType decodeGangType(OpenACCGangKind gk) {
133 switch (gk) {
135 return mlir::acc::GangArgType::Num;
137 return mlir::acc::GangArgType::Dim;
139 return mlir::acc::GangArgType::Static;
140 }
141 llvm_unreachable("unknown gang kind");
142 }
143
144 template <typename U = void,
145 typename = std::enable_if_t<isCombinedType<OpTy>, U>>
146 void applyToLoopOp(const OpenACCClause &c) {
147 mlir::OpBuilder::InsertionGuard guardCase(builder);
148 builder.setInsertionPoint(operation.loopOp);
149 OpenACCClauseCIREmitter<mlir::acc::LoopOp> loopEmitter{
150 operation.loopOp, cgf, builder, dirKind, dirLoc};
151 loopEmitter.lastDeviceTypeValues = lastDeviceTypeValues;
152 loopEmitter.Visit(&c);
153 }
154
155 template <typename U = void,
156 typename = std::enable_if_t<isCombinedType<OpTy>, U>>
157 void applyToComputeOp(const OpenACCClause &c) {
158 mlir::OpBuilder::InsertionGuard guardCase(builder);
159 builder.setInsertionPoint(operation.computeOp);
160 OpenACCClauseCIREmitter<typename OpTy::ComputeOpTy> computeEmitter{
161 operation.computeOp, cgf, builder, dirKind, dirLoc};
162
163 computeEmitter.lastDeviceTypeValues = lastDeviceTypeValues;
164
165 // Async handler uses the first data operand to figure out where to insert
166 // its information if it is present. This ensures that the new handler will
167 // correctly set the insertion point for async.
168 if (!dataOperands.empty())
169 computeEmitter.dataOperands.push_back(dataOperands.front());
170 computeEmitter.Visit(&c);
171
172 // Make sure all of the new data operands are kept track of here. The
173 // combined constructs always apply 'async' to only the compute component,
174 // so we need to collect these.
175 dataOperands.append(computeEmitter.dataOperands);
176 }
177
178 mlir::acc::DataClauseModifier
179 convertModifiers(OpenACCModifierKind modifiers) {
180 using namespace mlir::acc;
181 static_assert(static_cast<int>(OpenACCModifierKind::Zero) ==
182 static_cast<int>(DataClauseModifier::zero) &&
183 static_cast<int>(OpenACCModifierKind::Readonly) ==
184 static_cast<int>(DataClauseModifier::readonly) &&
185 static_cast<int>(OpenACCModifierKind::AlwaysIn) ==
186 static_cast<int>(DataClauseModifier::alwaysin) &&
187 static_cast<int>(OpenACCModifierKind::AlwaysOut) ==
188 static_cast<int>(DataClauseModifier::alwaysout) &&
189 static_cast<int>(OpenACCModifierKind::Capture) ==
190 static_cast<int>(DataClauseModifier::capture));
191
192 DataClauseModifier mlirModifiers{};
193
194 // The MLIR representation of this represents `always` as `alwaysin` +
195 // `alwaysout`. So do a small fixup here.
196 if (isOpenACCModifierBitSet(modifiers, OpenACCModifierKind::Always)) {
197 mlirModifiers = mlirModifiers | DataClauseModifier::always;
198 modifiers &= ~OpenACCModifierKind::Always;
199 }
200
201 mlirModifiers = mlirModifiers | static_cast<DataClauseModifier>(modifiers);
202 return mlirModifiers;
203 }
204
205 template <typename BeforeOpTy, typename AfterOpTy>
206 void addDataOperand(const Expr *varOperand, mlir::acc::DataClause dataClause,
207 OpenACCModifierKind modifiers, bool structured,
208 bool implicit) {
210 cgf.getOpenACCDataOperandInfo(varOperand);
211
212 auto beforeOp =
213 builder.create<BeforeOpTy>(opInfo.beginLoc, opInfo.varValue, structured,
214 implicit, opInfo.name, opInfo.bounds);
215 operation.getDataClauseOperandsMutable().append(beforeOp.getResult());
216
217 AfterOpTy afterOp;
218 {
219 mlir::OpBuilder::InsertionGuard guardCase(builder);
220 builder.setInsertionPointAfter(operation);
221
222 if constexpr (std::is_same_v<AfterOpTy, mlir::acc::DeleteOp> ||
223 std::is_same_v<AfterOpTy, mlir::acc::DetachOp>) {
224 // Detach/Delete ops don't have the variable reference here, so they
225 // take 1 fewer argument to their build function.
226 afterOp = builder.create<AfterOpTy>(
227 opInfo.beginLoc, beforeOp.getResult(), structured, implicit,
228 opInfo.name, opInfo.bounds);
229 } else {
230 afterOp = builder.create<AfterOpTy>(
231 opInfo.beginLoc, beforeOp.getResult(), opInfo.varValue, structured,
232 implicit, opInfo.name, opInfo.bounds);
233 }
234 }
235
236 // Set the 'rest' of the info for both operations.
237 beforeOp.setDataClause(dataClause);
238 afterOp.setDataClause(dataClause);
239 beforeOp.setModifiers(convertModifiers(modifiers));
240 afterOp.setModifiers(convertModifiers(modifiers));
241
242 // Make sure we record these, so 'async' values can be updated later.
243 dataOperands.push_back(beforeOp.getOperation());
244 dataOperands.push_back(afterOp.getOperation());
245 }
246
247 template <typename BeforeOpTy>
248 void addDataOperand(const Expr *varOperand, mlir::acc::DataClause dataClause,
249 OpenACCModifierKind modifiers, bool structured,
250 bool implicit) {
252 cgf.getOpenACCDataOperandInfo(varOperand);
253 auto beforeOp =
254 builder.create<BeforeOpTy>(opInfo.beginLoc, opInfo.varValue, structured,
255 implicit, opInfo.name, opInfo.bounds);
256 operation.getDataClauseOperandsMutable().append(beforeOp.getResult());
257
258 // Set the 'rest' of the info for the operation.
259 beforeOp.setDataClause(dataClause);
260 beforeOp.setModifiers(convertModifiers(modifiers));
261
262 // Make sure we record these, so 'async' values can be updated later.
263 dataOperands.push_back(beforeOp.getOperation());
264 }
265
266 // Helper function that covers for the fact that we don't have this function
267 // on all operation types.
268 mlir::ArrayAttr getAsyncOnlyAttr() {
269 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
270 mlir::acc::KernelsOp, mlir::acc::DataOp,
271 mlir::acc::UpdateOp>) {
272 return operation.getAsyncOnlyAttr();
273 } else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp,
274 mlir::acc::ExitDataOp>) {
275 if (!operation.getAsyncAttr())
276 return mlir::ArrayAttr{};
277
279 devTysTemp.push_back(mlir::acc::DeviceTypeAttr::get(
280 builder.getContext(), mlir::acc::DeviceType::None));
281 return mlir::ArrayAttr::get(builder.getContext(), devTysTemp);
282 } else if constexpr (isCombinedType<OpTy>) {
283 return operation.computeOp.getAsyncOnlyAttr();
284 }
285
286 // Note: 'wait' has async as well, but it cannot have data clauses, so we
287 // don't have to handle them here.
288
289 llvm_unreachable("getting asyncOnly when clause not valid on operation?");
290 }
291
292 // Helper function that covers for the fact that we don't have this function
293 // on all operation types.
294 mlir::ArrayAttr getAsyncOperandsDeviceTypeAttr() {
295 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
296 mlir::acc::KernelsOp, mlir::acc::DataOp,
297 mlir::acc::UpdateOp>) {
298 return operation.getAsyncOperandsDeviceTypeAttr();
299 } else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp,
300 mlir::acc::ExitDataOp>) {
301 if (!operation.getAsyncOperand())
302 return mlir::ArrayAttr{};
303
305 devTysTemp.push_back(mlir::acc::DeviceTypeAttr::get(
306 builder.getContext(), mlir::acc::DeviceType::None));
307 return mlir::ArrayAttr::get(builder.getContext(), devTysTemp);
308 } else if constexpr (isCombinedType<OpTy>) {
309 return operation.computeOp.getAsyncOperandsDeviceTypeAttr();
310 }
311
312 // Note: 'wait' has async as well, but it cannot have data clauses, so we
313 // don't have to handle them here.
314
315 llvm_unreachable(
316 "getting asyncOperandsDeviceType when clause not valid on operation?");
317 }
318
319 // Helper function that covers for the fact that we don't have this function
320 // on all operation types.
321 mlir::OperandRange getAsyncOperands() {
322 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
323 mlir::acc::KernelsOp, mlir::acc::DataOp,
324 mlir::acc::UpdateOp>)
325 return operation.getAsyncOperands();
326 else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp,
327 mlir::acc::ExitDataOp>)
328 return operation.getAsyncOperandMutable();
329 else if constexpr (isCombinedType<OpTy>)
330 return operation.computeOp.getAsyncOperands();
331
332 // Note: 'wait' has async as well, but it cannot have data clauses, so we
333 // don't have to handle them here.
334
335 llvm_unreachable(
336 "getting asyncOperandsDeviceType when clause not valid on operation?");
337 }
338
339 // The 'data' clauses all require that we add the 'async' values from the
340 // operation to them. We've collected the data operands along the way, so use
341 // that list to get the current 'async' values.
342 void updateDataOperandAsyncValues() {
343 if (!hasAsyncClause || dataOperands.empty())
344 return;
345
346 for (mlir::Operation *dataOp : dataOperands) {
347 llvm::TypeSwitch<mlir::Operation *, void>(dataOp)
348 .Case<ACC_DATA_ENTRY_OPS, ACC_DATA_EXIT_OPS>([&](auto op) {
349 op.setAsyncOnlyAttr(getAsyncOnlyAttr());
350 op.setAsyncOperandsDeviceTypeAttr(getAsyncOperandsDeviceTypeAttr());
351 op.getAsyncOperandsMutable().assign(getAsyncOperands());
352 })
353 .Default([&](mlir::Operation *) {
354 llvm_unreachable("Not a data operation?");
355 });
356 }
357 }
358
359 template <typename RecipeTy>
360 std::string getRecipeName(SourceRange loc, QualType baseType,
361 OpenACCReductionOperator reductionOp) {
362 std::string recipeName;
363 {
364 llvm::raw_string_ostream stream(recipeName);
365
366 if constexpr (std::is_same_v<RecipeTy, mlir::acc::PrivateRecipeOp>) {
367 stream << "privatization_";
368 } else if constexpr (std::is_same_v<RecipeTy,
369 mlir::acc::FirstprivateRecipeOp>) {
370 stream << "firstprivatization_";
371
372 } else if constexpr (std::is_same_v<RecipeTy,
373 mlir::acc::ReductionRecipeOp>) {
374 stream << "reduction_";
375 // Values here are a little weird (for bitwise and/or is 'i' prefix, and
376 // logical ops with 'l'), but are chosen to be the same as the MLIR
377 // dialect names as well as to match the Flang versions of these.
378 switch (reductionOp) {
380 stream << "add_";
381 break;
383 stream << "mul_";
384 break;
386 stream << "max_";
387 break;
389 stream << "min_";
390 break;
392 stream << "iand_";
393 break;
395 stream << "ior_";
396 break;
398 stream << "xor_";
399 break;
401 stream << "land_";
402 break;
404 stream << "lor_";
405 break;
407 llvm_unreachable("invalid reduction operator");
408 }
409 } else {
410 static_assert(!sizeof(RecipeTy), "Unknown Recipe op kind");
411 }
412
414 mc.mangleCanonicalTypeName(baseType, stream);
415 }
416 return recipeName;
417 }
418
419 void createFirstprivateRecipeCopy(
420 mlir::Location loc, mlir::Location locEnd, mlir::Value mainOp,
421 CIRGenFunction::AutoVarEmission tempDeclEmission,
422 mlir::acc::FirstprivateRecipeOp recipe, const VarDecl *varRecipe,
423 const VarDecl *temporary) {
424 mlir::Block *block = builder.createBlock(
425 &recipe.getCopyRegion(), recipe.getCopyRegion().end(),
426 {mainOp.getType(), mainOp.getType()}, {loc, loc});
427 builder.setInsertionPointToEnd(&recipe.getCopyRegion().back());
428 CIRGenFunction::LexicalScope ls(cgf, loc, block);
429
430 mlir::BlockArgument fromArg = block->getArgument(0);
431 mlir::BlockArgument toArg = block->getArgument(1);
432
433 mlir::Type elementTy =
434 mlir::cast<cir::PointerType>(mainOp.getType()).getPointee();
435
436 // Set the address of the emission to be the argument, so that we initialize
437 // that instead of the variable in the other block.
438 tempDeclEmission.setAllocatedAddress(
439 Address{toArg, elementTy, cgf.getContext().getDeclAlign(varRecipe)});
440 tempDeclEmission.EmittedAsOffload = true;
441
442 CIRGenFunction::DeclMapRevertingRAII declMapRAII{cgf, temporary};
444 temporary,
445 Address{fromArg, elementTy, cgf.getContext().getDeclAlign(varRecipe)});
446
447 cgf.emitAutoVarInit(tempDeclEmission);
448 mlir::acc::YieldOp::create(builder, locEnd);
449 }
450
451 // Create the 'init' section of the recipe, including the 'copy' section for
452 // 'firstprivate'. Note that this function is not 'insertion point' clean, in
453 // that it alters the insertion point to be inside of the 'destroy' section of
454 // the recipe, but doesn't restore it aftewards.
455 template <typename RecipeTy>
456 void createRecipeInitCopy(mlir::Location loc, mlir::Location locEnd,
457 SourceRange exprRange, mlir::Value mainOp,
458 RecipeTy recipe, const VarDecl *varRecipe,
459 const VarDecl *temporary) {
460 assert(varRecipe && "Required recipe variable not set?");
461
462 CIRGenFunction::AutoVarEmission tempDeclEmission{
464 CIRGenFunction::DeclMapRevertingRAII declMapRAII{cgf, varRecipe};
465
466 // Do the 'init' section of the recipe IR, which does an alloca, then the
467 // initialization (except for firstprivate).
468 mlir::Block *block = builder.createBlock(&recipe.getInitRegion(),
469 recipe.getInitRegion().end(),
470 {mainOp.getType()}, {loc});
471 builder.setInsertionPointToEnd(&recipe.getInitRegion().back());
472 CIRGenFunction::LexicalScope ls(cgf, loc, block);
473
474 tempDeclEmission =
475 cgf.emitAutoVarAlloca(*varRecipe, builder.saveInsertionPoint());
476
477 // 'firstprivate' doesn't do its initialization in the 'init' section,
478 // instead does it in the 'copy' section. SO only do init here.
479 // 'reduction' appears to use it too (rather than a 'copy' section), so
480 // we probably have to do it here too, but we can do that when we get to
481 // reduction implementation.
482 if constexpr (std::is_same_v<RecipeTy, mlir::acc::PrivateRecipeOp>) {
483 // We are OK with no init for builtins, arrays of builtins, or pointers,
484 // else we should NYI so we know to go look for these.
485 if (cgf.getContext().getLangOpts().CPlusPlus &&
486 !varRecipe->getType()
488 ->isBuiltinType() &&
489 !varRecipe->getType()->isPointerType() && !varRecipe->getInit()) {
490 // If we don't have any initialization recipe, we failed during Sema to
491 // initialize this correctly. If we disable the
492 // Sema::TentativeAnalysisScopes in SemaOpenACC::CreateInitRecipe, it'll
493 // emit an error to tell us. However, emitting those errors during
494 // production is a violation of the standard, so we cannot do them.
495 cgf.cgm.errorNYI(exprRange, "private default-init recipe");
496 }
497 cgf.emitAutoVarInit(tempDeclEmission);
498 } else if constexpr (std::is_same_v<RecipeTy,
499 mlir::acc::ReductionRecipeOp>) {
500 // Unlike Private, the recipe here is always required as it has to do
501 // init, not just 'default' init.
502 if (!varRecipe->getInit())
503 cgf.cgm.errorNYI(exprRange, "reduction init recipe");
504 cgf.emitAutoVarInit(tempDeclEmission);
505 }
506
507 mlir::acc::YieldOp::create(builder, locEnd);
508
509 if constexpr (std::is_same_v<RecipeTy, mlir::acc::FirstprivateRecipeOp>) {
510 if (!varRecipe->getInit()) {
511 // If we don't have any initialization recipe, we failed during Sema to
512 // initialize this correctly. If we disable the
513 // Sema::TentativeAnalysisScopes in SemaOpenACC::CreateInitRecipe, it'll
514 // emit an error to tell us. However, emitting those errors during
515 // production is a violation of the standard, so we cannot do them.
516 cgf.cgm.errorNYI(
517 exprRange, "firstprivate copy-init recipe not properly generated");
518 }
519
520 createFirstprivateRecipeCopy(loc, locEnd, mainOp, tempDeclEmission,
521 recipe, varRecipe, temporary);
522 }
523 }
524
525 // This function generates the 'combiner' section for a reduction recipe. Note
526 // that this function is not 'insertion point' clean, in that it alters the
527 // insertion point to be inside of the 'combiner' section of the recipe, but
528 // doesn't restore it aftewards.
529 void createReductionRecipeCombiner(mlir::Location loc, mlir::Location locEnd,
530 mlir::Value mainOp,
531 mlir::acc::ReductionRecipeOp recipe) {
532 mlir::Block *block = builder.createBlock(
533 &recipe.getCombinerRegion(), recipe.getCombinerRegion().end(),
534 {mainOp.getType(), mainOp.getType()}, {loc, loc});
535 builder.setInsertionPointToEnd(&recipe.getCombinerRegion().back());
536 CIRGenFunction::LexicalScope ls(cgf, loc, block);
537
538 mlir::BlockArgument lhsArg = block->getArgument(0);
539
540 mlir::acc::YieldOp::create(builder, locEnd, lhsArg);
541 }
542
543 // This function generates the 'destroy' section for a recipe. Note
544 // that this function is not 'insertion point' clean, in that it alters the
545 // insertion point to be inside of the 'destroy' section of the recipe, but
546 // doesn't restore it aftewards.
547 void createRecipeDestroySection(mlir::Location loc, mlir::Location locEnd,
548 mlir::Value mainOp, CharUnits alignment,
549 QualType baseType,
550 mlir::Region &destroyRegion) {
551 mlir::Block *block = builder.createBlock(
552 &destroyRegion, destroyRegion.end(), {mainOp.getType()}, {loc});
553 builder.setInsertionPointToEnd(&destroyRegion.back());
554 CIRGenFunction::LexicalScope ls(cgf, loc, block);
555
556 mlir::Type elementTy =
557 mlir::cast<cir::PointerType>(mainOp.getType()).getPointee();
558 Address addr{block->getArgument(0), elementTy, alignment};
559 cgf.emitDestroy(addr, baseType,
561
562 mlir::acc::YieldOp::create(builder, locEnd);
563 }
564
565 mlir::acc::ReductionOperator convertReductionOp(OpenACCReductionOperator op) {
566 switch (op) {
568 return mlir::acc::ReductionOperator::AccAdd;
570 return mlir::acc::ReductionOperator::AccMul;
572 return mlir::acc::ReductionOperator::AccMax;
574 return mlir::acc::ReductionOperator::AccMin;
576 return mlir::acc::ReductionOperator::AccIand;
578 return mlir::acc::ReductionOperator::AccIor;
580 return mlir::acc::ReductionOperator::AccXor;
582 return mlir::acc::ReductionOperator::AccLand;
584 return mlir::acc::ReductionOperator::AccLor;
586 llvm_unreachable("invalid reduction operator");
587 }
588
589 llvm_unreachable("invalid reduction operator");
590 }
591
592 template <typename RecipeTy>
593 RecipeTy getOrCreateRecipe(ASTContext &astCtx, const Expr *varRef,
594 const VarDecl *varRecipe, const VarDecl *temporary,
595 OpenACCReductionOperator reductionOp,
596 DeclContext *dc, QualType baseType,
597 mlir::Value mainOp) {
598 mlir::ModuleOp mod = builder.getBlock()
599 ->getParent()
600 ->template getParentOfType<mlir::ModuleOp>();
601
602 std::string recipeName = getRecipeName<RecipeTy>(varRef->getSourceRange(),
603 baseType, reductionOp);
604 if (auto recipe = mod.lookupSymbol<RecipeTy>(recipeName))
605 return recipe;
606
607 mlir::Location loc = cgf.cgm.getLoc(varRef->getBeginLoc());
608 mlir::Location locEnd = cgf.cgm.getLoc(varRef->getEndLoc());
609
610 mlir::OpBuilder modBuilder(mod.getBodyRegion());
611 RecipeTy recipe;
612
613 if constexpr (std::is_same_v<RecipeTy, mlir::acc::ReductionRecipeOp>) {
614 recipe = RecipeTy::create(modBuilder, loc, recipeName, mainOp.getType(),
615 convertReductionOp(reductionOp));
616 } else {
617 recipe = RecipeTy::create(modBuilder, loc, recipeName, mainOp.getType());
618 }
619
620 createRecipeInitCopy(loc, locEnd, varRef->getSourceRange(), mainOp, recipe,
621 varRecipe, temporary);
622
623 if constexpr (std::is_same_v<RecipeTy, mlir::acc::ReductionRecipeOp>) {
624 createReductionRecipeCombiner(loc, locEnd, mainOp, recipe);
625 }
626
627 if (varRecipe && varRecipe->needsDestruction(cgf.getContext()))
628 createRecipeDestroySection(loc, locEnd, mainOp,
629 cgf.getContext().getDeclAlign(varRecipe),
630 baseType, recipe.getDestroyRegion());
631 return recipe;
632 }
633
634public:
635 OpenACCClauseCIREmitter(OpTy &operation, CIRGen::CIRGenFunction &cgf,
637 OpenACCDirectiveKind dirKind, SourceLocation dirLoc)
638 : operation(operation), cgf(cgf), builder(builder), dirKind(dirKind),
639 dirLoc(dirLoc) {}
640
641 void VisitClause(const OpenACCClause &clause) {
642 clauseNotImplemented(clause);
643 }
644
645 // The entry point for the CIR emitter. All users should use this rather than
646 // 'visitClauseList', as this also handles the things that have to happen
647 // 'after' the clauses are all visited.
648 void emitClauses(ArrayRef<const OpenACCClause *> clauses) {
649 this->VisitClauseList(clauses);
650 updateDataOperandAsyncValues();
651 }
652
653 void VisitDefaultClause(const OpenACCDefaultClause &clause) {
654 // This type-trait checks if 'op'(the first arg) is one of the mlir::acc
655 // operations listed in the rest of the arguments.
656 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
657 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
658 switch (clause.getDefaultClauseKind()) {
660 operation.setDefaultAttr(mlir::acc::ClauseDefaultValue::None);
661 break;
663 operation.setDefaultAttr(mlir::acc::ClauseDefaultValue::Present);
664 break;
666 break;
667 }
668 } else if constexpr (isCombinedType<OpTy>) {
669 applyToComputeOp(clause);
670 } else {
671 llvm_unreachable("Unknown construct kind in VisitDefaultClause");
672 }
673 }
674
675 void VisitDeviceTypeClause(const OpenACCDeviceTypeClause &clause) {
676 setLastDeviceTypeClause(clause);
677
678 if constexpr (isOneOfTypes<OpTy, mlir::acc::InitOp,
679 mlir::acc::ShutdownOp>) {
680 for (const DeviceTypeArgument &arg : clause.getArchitectures())
681 operation.addDeviceType(builder.getContext(),
682 decodeDeviceType(arg.getIdentifierInfo()));
683 } else if constexpr (isOneOfTypes<OpTy, mlir::acc::SetOp>) {
684 assert(!operation.getDeviceTypeAttr() && "already have device-type?");
685 assert(clause.getArchitectures().size() <= 1);
686
687 if (!clause.getArchitectures().empty())
688 operation.setDeviceType(
689 decodeDeviceType(clause.getArchitectures()[0].getIdentifierInfo()));
690 } else if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
691 mlir::acc::SerialOp, mlir::acc::KernelsOp,
692 mlir::acc::DataOp, mlir::acc::LoopOp,
693 mlir::acc::UpdateOp>) {
694 // Nothing to do here, these constructs don't have any IR for these, as
695 // they just modify the other clauses IR. So setting of
696 // `lastDeviceTypeValues` (done above) is all we need.
697 } else if constexpr (isCombinedType<OpTy>) {
698 // Nothing to do here either, combined constructs are just going to use
699 // 'lastDeviceTypeValues' to set the value for the child visitor.
700 } else {
701 // TODO: When we've implemented this for everything, switch this to an
702 // unreachable. routine construct remains.
703 return clauseNotImplemented(clause);
704 }
705 }
706
707 void VisitNumWorkersClause(const OpenACCNumWorkersClause &clause) {
708 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
709 mlir::acc::KernelsOp>) {
710 operation.addNumWorkersOperand(builder.getContext(),
711 emitIntExpr(clause.getIntExpr()),
712 lastDeviceTypeValues);
713 } else if constexpr (isCombinedType<OpTy>) {
714 applyToComputeOp(clause);
715 } else {
716 llvm_unreachable("Unknown construct kind in VisitNumGangsClause");
717 }
718 }
719
720 void VisitVectorLengthClause(const OpenACCVectorLengthClause &clause) {
721 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
722 mlir::acc::KernelsOp>) {
723 operation.addVectorLengthOperand(builder.getContext(),
724 emitIntExpr(clause.getIntExpr()),
725 lastDeviceTypeValues);
726 } else if constexpr (isCombinedType<OpTy>) {
727 applyToComputeOp(clause);
728 } else {
729 llvm_unreachable("Unknown construct kind in VisitVectorLengthClause");
730 }
731 }
732
733 void VisitAsyncClause(const OpenACCAsyncClause &clause) {
734 hasAsyncClause = true;
735 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
736 mlir::acc::KernelsOp, mlir::acc::DataOp,
737 mlir::acc::EnterDataOp, mlir::acc::ExitDataOp,
738 mlir::acc::UpdateOp>) {
739 if (!clause.hasIntExpr()) {
740 operation.addAsyncOnly(builder.getContext(), lastDeviceTypeValues);
741 } else {
742
743 mlir::Value intExpr;
744 {
745 // Async int exprs can be referenced by the data operands, which means
746 // that the int-exprs have to appear before them. IF there is a data
747 // operand already, set the insertion point to 'before' it.
748 mlir::OpBuilder::InsertionGuard guardCase(builder);
749 if (!dataOperands.empty())
750 builder.setInsertionPoint(dataOperands.front());
751 intExpr = emitIntExpr(clause.getIntExpr());
752 }
753 operation.addAsyncOperand(builder.getContext(), intExpr,
754 lastDeviceTypeValues);
755 }
756 } else if constexpr (isOneOfTypes<OpTy, mlir::acc::WaitOp>) {
757 // Wait doesn't have a device_type, so its handling here is slightly
758 // different.
759 if (!clause.hasIntExpr())
760 operation.setAsync(true);
761 else
762 operation.getAsyncOperandMutable().append(
763 emitIntExpr(clause.getIntExpr()));
764 } else if constexpr (isCombinedType<OpTy>) {
765 applyToComputeOp(clause);
766 } else {
767 // TODO: When we've implemented this for everything, switch this to an
768 // unreachable. Combined constructs remain. update construct remains.
769 return clauseNotImplemented(clause);
770 }
771 }
772
773 void VisitSelfClause(const OpenACCSelfClause &clause) {
774 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
775 mlir::acc::KernelsOp>) {
776 if (clause.isEmptySelfClause()) {
777 operation.setSelfAttr(true);
778 } else if (clause.isConditionExprClause()) {
779 assert(clause.hasConditionExpr());
780 operation.getSelfCondMutable().append(
781 createCondition(clause.getConditionExpr()));
782 } else {
783 llvm_unreachable("var-list version of self shouldn't get here");
784 }
785 } else if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
786 assert(!clause.isEmptySelfClause() && !clause.isConditionExprClause() &&
787 "var-list version of self required for update");
788 for (const Expr *var : clause.getVarList())
789 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::UpdateHostOp>(
790 var, mlir::acc::DataClause::acc_update_self, {},
791 /*structured=*/false, /*implicit=*/false);
792 } else if constexpr (isCombinedType<OpTy>) {
793 applyToComputeOp(clause);
794 } else {
795 llvm_unreachable("Unknown construct kind in VisitSelfClause");
796 }
797 }
798
799 void VisitHostClause(const OpenACCHostClause &clause) {
800 if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
801 for (const Expr *var : clause.getVarList())
802 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::UpdateHostOp>(
803 var, mlir::acc::DataClause::acc_update_host, {},
804 /*structured=*/false, /*implicit=*/false);
805 } else {
806 llvm_unreachable("Unknown construct kind in VisitHostClause");
807 }
808 }
809
810 void VisitDeviceClause(const OpenACCDeviceClause &clause) {
811 if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
812 for (const Expr *var : clause.getVarList())
813 addDataOperand<mlir::acc::UpdateDeviceOp>(
814 var, mlir::acc::DataClause::acc_update_device, {},
815 /*structured=*/false, /*implicit=*/false);
816 } else {
817 llvm_unreachable("Unknown construct kind in VisitDeviceClause");
818 }
819 }
820
821 void VisitIfClause(const OpenACCIfClause &clause) {
822 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
823 mlir::acc::KernelsOp, mlir::acc::InitOp,
824 mlir::acc::ShutdownOp, mlir::acc::SetOp,
825 mlir::acc::DataOp, mlir::acc::WaitOp,
826 mlir::acc::HostDataOp, mlir::acc::EnterDataOp,
827 mlir::acc::ExitDataOp, mlir::acc::UpdateOp>) {
828 operation.getIfCondMutable().append(
829 createCondition(clause.getConditionExpr()));
830 } else if constexpr (isCombinedType<OpTy>) {
831 applyToComputeOp(clause);
832 } else {
833 llvm_unreachable("Unknown construct kind in VisitIfClause");
834 }
835 }
836
837 void VisitIfPresentClause(const OpenACCIfPresentClause &clause) {
838 if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp,
839 mlir::acc::UpdateOp>) {
840 operation.setIfPresent(true);
841 } else {
842 llvm_unreachable("unknown construct kind in VisitIfPresentClause");
843 }
844 }
845
846 void VisitDeviceNumClause(const OpenACCDeviceNumClause &clause) {
847 if constexpr (isOneOfTypes<OpTy, mlir::acc::InitOp, mlir::acc::ShutdownOp,
848 mlir::acc::SetOp>) {
849 operation.getDeviceNumMutable().append(emitIntExpr(clause.getIntExpr()));
850 } else {
851 llvm_unreachable(
852 "init, shutdown, set, are only valid device_num constructs");
853 }
854 }
855
856 void VisitNumGangsClause(const OpenACCNumGangsClause &clause) {
857 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
858 mlir::acc::KernelsOp>) {
860 for (const Expr *E : clause.getIntExprs())
861 values.push_back(emitIntExpr(E));
862
863 operation.addNumGangsOperands(builder.getContext(), values,
864 lastDeviceTypeValues);
865 } else if constexpr (isCombinedType<OpTy>) {
866 applyToComputeOp(clause);
867 } else {
868 llvm_unreachable("Unknown construct kind in VisitNumGangsClause");
869 }
870 }
871
872 void VisitWaitClause(const OpenACCWaitClause &clause) {
873 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
874 mlir::acc::KernelsOp, mlir::acc::DataOp,
875 mlir::acc::EnterDataOp, mlir::acc::ExitDataOp,
876 mlir::acc::UpdateOp>) {
877 if (!clause.hasExprs()) {
878 operation.addWaitOnly(builder.getContext(), lastDeviceTypeValues);
879 } else {
881 if (clause.hasDevNumExpr())
882 values.push_back(emitIntExpr(clause.getDevNumExpr()));
883 for (const Expr *E : clause.getQueueIdExprs())
884 values.push_back(emitIntExpr(E));
885 operation.addWaitOperands(builder.getContext(), clause.hasDevNumExpr(),
886 values, lastDeviceTypeValues);
887 }
888 } else if constexpr (isCombinedType<OpTy>) {
889 applyToComputeOp(clause);
890 } else {
891 // TODO: When we've implemented this for everything, switch this to an
892 // unreachable. update construct remains.
893 return clauseNotImplemented(clause);
894 }
895 }
896
897 void VisitDefaultAsyncClause(const OpenACCDefaultAsyncClause &clause) {
898 if constexpr (isOneOfTypes<OpTy, mlir::acc::SetOp>) {
899 operation.getDefaultAsyncMutable().append(
900 emitIntExpr(clause.getIntExpr()));
901 } else {
902 llvm_unreachable("set, is only valid device_num constructs");
903 }
904 }
905
906 void VisitSeqClause(const OpenACCSeqClause &clause) {
907 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
908 operation.addSeq(builder.getContext(), lastDeviceTypeValues);
909 } else if constexpr (isCombinedType<OpTy>) {
910 applyToLoopOp(clause);
911 } else {
912 // TODO: When we've implemented this for everything, switch this to an
913 // unreachable. Routine construct remains.
914 return clauseNotImplemented(clause);
915 }
916 }
917
918 void VisitAutoClause(const OpenACCAutoClause &clause) {
919 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
920 operation.addAuto(builder.getContext(), lastDeviceTypeValues);
921 } else if constexpr (isCombinedType<OpTy>) {
922 applyToLoopOp(clause);
923 } else {
924 // TODO: When we've implemented this for everything, switch this to an
925 // unreachable. Routine, construct remains.
926 return clauseNotImplemented(clause);
927 }
928 }
929
930 void VisitIndependentClause(const OpenACCIndependentClause &clause) {
931 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
932 operation.addIndependent(builder.getContext(), lastDeviceTypeValues);
933 } else if constexpr (isCombinedType<OpTy>) {
934 applyToLoopOp(clause);
935 } else {
936 // TODO: When we've implemented this for everything, switch this to an
937 // unreachable. Routine construct remains.
938 return clauseNotImplemented(clause);
939 }
940 }
941
942 void VisitCollapseClause(const OpenACCCollapseClause &clause) {
943 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
944 llvm::APInt value =
946
947 value = value.sextOrTrunc(64);
948 operation.setCollapseForDeviceTypes(builder.getContext(),
949 lastDeviceTypeValues, value);
950 } else if constexpr (isCombinedType<OpTy>) {
951 applyToLoopOp(clause);
952 } else {
953 llvm_unreachable("Unknown construct kind in VisitCollapseClause");
954 }
955 }
956
957 void VisitTileClause(const OpenACCTileClause &clause) {
958 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
960
961 for (const Expr *e : clause.getSizeExprs()) {
962 mlir::Location exprLoc = cgf.cgm.getLoc(e->getBeginLoc());
963
964 // We represent the * as -1. Additionally, this is a constant, so we
965 // can always just emit it as 64 bits to avoid having to do any more
966 // work to determine signedness or size.
967 if (isa<OpenACCAsteriskSizeExpr>(e)) {
968 values.push_back(createConstantInt(exprLoc, 64, -1));
969 } else {
970 llvm::APInt curValue =
971 e->EvaluateKnownConstInt(cgf.cgm.getASTContext());
972 values.push_back(createConstantInt(
973 exprLoc, 64, curValue.sextOrTrunc(64).getSExtValue()));
974 }
975 }
976
977 operation.setTileForDeviceTypes(builder.getContext(),
978 lastDeviceTypeValues, values);
979 } else if constexpr (isCombinedType<OpTy>) {
980 applyToLoopOp(clause);
981 } else {
982 llvm_unreachable("Unknown construct kind in VisitTileClause");
983 }
984 }
985
986 void VisitWorkerClause(const OpenACCWorkerClause &clause) {
987 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
988 if (clause.hasIntExpr())
989 operation.addWorkerNumOperand(builder.getContext(),
990 emitIntExpr(clause.getIntExpr()),
991 lastDeviceTypeValues);
992 else
993 operation.addEmptyWorker(builder.getContext(), lastDeviceTypeValues);
994
995 } else if constexpr (isCombinedType<OpTy>) {
996 applyToLoopOp(clause);
997 } else {
998 // TODO: When we've implemented this for everything, switch this to an
999 // unreachable. Combined constructs remain.
1000 return clauseNotImplemented(clause);
1001 }
1002 }
1003
1004 void VisitVectorClause(const OpenACCVectorClause &clause) {
1005 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
1006 if (clause.hasIntExpr())
1007 operation.addVectorOperand(builder.getContext(),
1008 emitIntExpr(clause.getIntExpr()),
1009 lastDeviceTypeValues);
1010 else
1011 operation.addEmptyVector(builder.getContext(), lastDeviceTypeValues);
1012
1013 } else if constexpr (isCombinedType<OpTy>) {
1014 applyToLoopOp(clause);
1015 } else {
1016 // TODO: When we've implemented this for everything, switch this to an
1017 // unreachable. Combined constructs remain.
1018 return clauseNotImplemented(clause);
1019 }
1020 }
1021
1022 void VisitGangClause(const OpenACCGangClause &clause) {
1023 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
1024 if (clause.getNumExprs() == 0) {
1025 operation.addEmptyGang(builder.getContext(), lastDeviceTypeValues);
1026 } else {
1029 for (unsigned i : llvm::index_range(0u, clause.getNumExprs())) {
1030 auto [kind, expr] = clause.getExpr(i);
1031 mlir::Location exprLoc = cgf.cgm.getLoc(expr->getBeginLoc());
1032 argTypes.push_back(decodeGangType(kind));
1033 if (kind == OpenACCGangKind::Dim) {
1034 llvm::APInt curValue =
1035 expr->EvaluateKnownConstInt(cgf.cgm.getASTContext());
1036 // The value is 1, 2, or 3, but the type isn't necessarily smaller
1037 // than 64.
1038 curValue = curValue.sextOrTrunc(64);
1039 values.push_back(
1040 createConstantInt(exprLoc, 64, curValue.getSExtValue()));
1041 } else if (isa<OpenACCAsteriskSizeExpr>(expr)) {
1042 values.push_back(createConstantInt(exprLoc, 64, -1));
1043 } else {
1044 values.push_back(emitIntExpr(expr));
1045 }
1046 }
1047
1048 operation.addGangOperands(builder.getContext(), lastDeviceTypeValues,
1049 argTypes, values);
1050 }
1051 } else if constexpr (isCombinedType<OpTy>) {
1052 applyToLoopOp(clause);
1053 } else {
1054 llvm_unreachable("Unknown construct kind in VisitGangClause");
1055 }
1056 }
1057
1058 void VisitCopyClause(const OpenACCCopyClause &clause) {
1059 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
1060 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
1061 for (const Expr *var : clause.getVarList())
1062 addDataOperand<mlir::acc::CopyinOp, mlir::acc::CopyoutOp>(
1063 var, mlir::acc::DataClause::acc_copy, clause.getModifierList(),
1064 /*structured=*/true,
1065 /*implicit=*/false);
1066 } else if constexpr (isCombinedType<OpTy>) {
1067 applyToComputeOp(clause);
1068 } else {
1069 // TODO: When we've implemented this for everything, switch this to an
1070 // unreachable. declare construct remains.
1071 return clauseNotImplemented(clause);
1072 }
1073 }
1074
1075 void VisitCopyInClause(const OpenACCCopyInClause &clause) {
1076 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
1077 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
1078 for (const Expr *var : clause.getVarList())
1079 addDataOperand<mlir::acc::CopyinOp, mlir::acc::DeleteOp>(
1080 var, mlir::acc::DataClause::acc_copyin, clause.getModifierList(),
1081 /*structured=*/true,
1082 /*implicit=*/false);
1083 } else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
1084 for (const Expr *var : clause.getVarList())
1085 addDataOperand<mlir::acc::CopyinOp>(
1086 var, mlir::acc::DataClause::acc_copyin, clause.getModifierList(),
1087 /*structured=*/false, /*implicit=*/false);
1088 } else if constexpr (isCombinedType<OpTy>) {
1089 applyToComputeOp(clause);
1090 } else {
1091 // TODO: When we've implemented this for everything, switch this to an
1092 // unreachable. declare construct remains.
1093 return clauseNotImplemented(clause);
1094 }
1095 }
1096
1097 void VisitCopyOutClause(const OpenACCCopyOutClause &clause) {
1098 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
1099 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
1100 for (const Expr *var : clause.getVarList())
1101 addDataOperand<mlir::acc::CreateOp, mlir::acc::CopyoutOp>(
1102 var, mlir::acc::DataClause::acc_copyout, clause.getModifierList(),
1103 /*structured=*/true,
1104 /*implicit=*/false);
1105 } else if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
1106 for (const Expr *var : clause.getVarList())
1107 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::CopyoutOp>(
1108 var, mlir::acc::DataClause::acc_copyout, clause.getModifierList(),
1109 /*structured=*/false,
1110 /*implicit=*/false);
1111 } else if constexpr (isCombinedType<OpTy>) {
1112 applyToComputeOp(clause);
1113 } else {
1114 // TODO: When we've implemented this for everything, switch this to an
1115 // unreachable. declare construct remains.
1116 return clauseNotImplemented(clause);
1117 }
1118 }
1119
1120 void VisitCreateClause(const OpenACCCreateClause &clause) {
1121 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
1122 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
1123 for (const Expr *var : clause.getVarList())
1124 addDataOperand<mlir::acc::CreateOp, mlir::acc::DeleteOp>(
1125 var, mlir::acc::DataClause::acc_create, clause.getModifierList(),
1126 /*structured=*/true,
1127 /*implicit=*/false);
1128 } else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
1129 for (const Expr *var : clause.getVarList())
1130 addDataOperand<mlir::acc::CreateOp>(
1131 var, mlir::acc::DataClause::acc_create, clause.getModifierList(),
1132 /*structured=*/false, /*implicit=*/false);
1133 } else if constexpr (isCombinedType<OpTy>) {
1134 applyToComputeOp(clause);
1135 } else {
1136 // TODO: When we've implemented this for everything, switch this to an
1137 // unreachable. declare construct remains.
1138 return clauseNotImplemented(clause);
1139 }
1140 }
1141
1142 void VisitDeleteClause(const OpenACCDeleteClause &clause) {
1143 if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
1144 for (const Expr *var : clause.getVarList())
1145 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::DeleteOp>(
1146 var, mlir::acc::DataClause::acc_delete, {},
1147 /*structured=*/false,
1148 /*implicit=*/false);
1149 } else {
1150 llvm_unreachable("Unknown construct kind in VisitDeleteClause");
1151 }
1152 }
1153
1154 void VisitDetachClause(const OpenACCDetachClause &clause) {
1155 if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
1156 for (const Expr *var : clause.getVarList())
1157 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::DetachOp>(
1158 var, mlir::acc::DataClause::acc_detach, {},
1159 /*structured=*/false,
1160 /*implicit=*/false);
1161 } else {
1162 llvm_unreachable("Unknown construct kind in VisitDetachClause");
1163 }
1164 }
1165
1166 void VisitFinalizeClause(const OpenACCFinalizeClause &clause) {
1167 if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
1168 operation.setFinalize(true);
1169 } else {
1170 llvm_unreachable("Unknown construct kind in VisitFinalizeClause");
1171 }
1172 }
1173
1174 void VisitUseDeviceClause(const OpenACCUseDeviceClause &clause) {
1175 if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp>) {
1176 for (const Expr *var : clause.getVarList())
1177 addDataOperand<mlir::acc::UseDeviceOp>(
1178 var, mlir::acc::DataClause::acc_use_device, {}, /*structured=*/true,
1179 /*implicit=*/false);
1180 } else {
1181 llvm_unreachable("Unknown construct kind in VisitUseDeviceClause");
1182 }
1183 }
1184
1185 void VisitDevicePtrClause(const OpenACCDevicePtrClause &clause) {
1186 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
1187 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
1188 for (const Expr *var : clause.getVarList())
1189 addDataOperand<mlir::acc::DevicePtrOp>(
1190 var, mlir::acc::DataClause::acc_deviceptr, {},
1191 /*structured=*/true,
1192 /*implicit=*/false);
1193 } else if constexpr (isCombinedType<OpTy>) {
1194 applyToComputeOp(clause);
1195 } else {
1196 // TODO: When we've implemented this for everything, switch this to an
1197 // unreachable. declare remains.
1198 return clauseNotImplemented(clause);
1199 }
1200 }
1201
1202 void VisitNoCreateClause(const OpenACCNoCreateClause &clause) {
1203 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
1204 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
1205 for (const Expr *var : clause.getVarList())
1206 addDataOperand<mlir::acc::NoCreateOp, mlir::acc::DeleteOp>(
1207 var, mlir::acc::DataClause::acc_no_create, {}, /*structured=*/true,
1208 /*implicit=*/false);
1209 } else if constexpr (isCombinedType<OpTy>) {
1210 applyToComputeOp(clause);
1211 } else {
1212 llvm_unreachable("Unknown construct kind in VisitNoCreateClause");
1213 }
1214 }
1215
1216 void VisitPresentClause(const OpenACCPresentClause &clause) {
1217 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
1218 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
1219 for (const Expr *var : clause.getVarList())
1220 addDataOperand<mlir::acc::PresentOp, mlir::acc::DeleteOp>(
1221 var, mlir::acc::DataClause::acc_present, {}, /*structured=*/true,
1222 /*implicit=*/false);
1223 } else if constexpr (isCombinedType<OpTy>) {
1224 applyToComputeOp(clause);
1225 } else {
1226 // TODO: When we've implemented this for everything, switch this to an
1227 // unreachable. declare remains.
1228 return clauseNotImplemented(clause);
1229 }
1230 }
1231
1232 void VisitAttachClause(const OpenACCAttachClause &clause) {
1233 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
1234 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
1235 for (const Expr *var : clause.getVarList())
1236 addDataOperand<mlir::acc::AttachOp, mlir::acc::DetachOp>(
1237 var, mlir::acc::DataClause::acc_attach, {}, /*structured=*/true,
1238 /*implicit=*/false);
1239 } else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
1240 for (const Expr *var : clause.getVarList())
1241 addDataOperand<mlir::acc::AttachOp>(
1242 var, mlir::acc::DataClause::acc_attach, {},
1243 /*structured=*/false, /*implicit=*/false);
1244 } else if constexpr (isCombinedType<OpTy>) {
1245 applyToComputeOp(clause);
1246 } else {
1247 llvm_unreachable("Unknown construct kind in VisitAttachClause");
1248 }
1249 }
1250
1251 void VisitPrivateClause(const OpenACCPrivateClause &clause) {
1252 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
1253 mlir::acc::LoopOp>) {
1254 for (const auto [varExpr, varRecipe] :
1255 llvm::zip_equal(clause.getVarList(), clause.getInitRecipes())) {
1257 cgf.getOpenACCDataOperandInfo(varExpr);
1258 auto privateOp = mlir::acc::PrivateOp::create(
1259 builder, opInfo.beginLoc, opInfo.varValue, /*structured=*/true,
1260 /*implicit=*/false, opInfo.name, opInfo.bounds);
1261 privateOp.setDataClause(mlir::acc::DataClause::acc_private);
1262
1263 {
1264 mlir::OpBuilder::InsertionGuard guardCase(builder);
1265 auto recipe = getOrCreateRecipe<mlir::acc::PrivateRecipeOp>(
1266 cgf.getContext(), varExpr, varRecipe, /*temporary=*/nullptr,
1267 OpenACCReductionOperator::Invalid,
1268
1270 privateOp.getResult());
1271 // TODO: OpenACC: The dialect is going to change in the near future to
1272 // have these be on a different operation, so when that changes, we
1273 // probably need to change these here.
1274 operation.addPrivatization(builder.getContext(), privateOp, recipe);
1275 }
1276 }
1277 } else if constexpr (isCombinedType<OpTy>) {
1278 // Despite this being valid on ParallelOp or SerialOp, combined type
1279 // applies to the 'loop'.
1280 applyToLoopOp(clause);
1281 } else {
1282 llvm_unreachable("Unknown construct kind in VisitPrivateClause");
1283 }
1284 }
1285
1286 void VisitFirstPrivateClause(const OpenACCFirstPrivateClause &clause) {
1287 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
1288 mlir::acc::SerialOp>) {
1289 for (const auto [varExpr, varRecipe] :
1290 llvm::zip_equal(clause.getVarList(), clause.getInitRecipes())) {
1292 cgf.getOpenACCDataOperandInfo(varExpr);
1293 auto firstPrivateOp = mlir::acc::FirstprivateOp::create(
1294 builder, opInfo.beginLoc, opInfo.varValue, /*structured=*/true,
1295 /*implicit=*/false, opInfo.name, opInfo.bounds);
1296
1297 firstPrivateOp.setDataClause(mlir::acc::DataClause::acc_firstprivate);
1298
1299 {
1300 mlir::OpBuilder::InsertionGuard guardCase(builder);
1301 auto recipe = getOrCreateRecipe<mlir::acc::FirstprivateRecipeOp>(
1302 cgf.getContext(), varExpr, varRecipe.RecipeDecl,
1303 varRecipe.InitFromTemporary, OpenACCReductionOperator::Invalid,
1305 firstPrivateOp.getResult());
1306
1307 // TODO: OpenACC: The dialect is going to change in the near future to
1308 // have these be on a different operation, so when that changes, we
1309 // probably need to change these here.
1310 operation.addFirstPrivatization(builder.getContext(), firstPrivateOp,
1311 recipe);
1312 }
1313 }
1314 } else if constexpr (isCombinedType<OpTy>) {
1315 // Unlike 'private', 'firstprivate' applies to the compute op, not the
1316 // loop op.
1317 applyToComputeOp(clause);
1318 } else {
1319 llvm_unreachable("Unknown construct kind in VisitFirstPrivateClause");
1320 }
1321 }
1322
1323 void VisitReductionClause(const OpenACCReductionClause &clause) {
1324 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
1325 mlir::acc::LoopOp>) {
1326 for (const auto [varExpr, varRecipe] :
1327 llvm::zip_equal(clause.getVarList(), clause.getRecipes())) {
1329 cgf.getOpenACCDataOperandInfo(varExpr);
1330
1331 auto reductionOp = mlir::acc::ReductionOp::create(
1332 builder, opInfo.beginLoc, opInfo.varValue, /*structured=*/true,
1333 /*implicit=*/false, opInfo.name, opInfo.bounds);
1334 reductionOp.setDataClause(mlir::acc::DataClause::acc_reduction);
1335
1336 {
1337 mlir::OpBuilder::InsertionGuard guardCase(builder);
1338
1339 auto recipe = getOrCreateRecipe<mlir::acc::ReductionRecipeOp>(
1340 cgf.getContext(), varExpr, varRecipe.RecipeDecl,
1341 /*temporary=*/nullptr, clause.getReductionOp(),
1343 reductionOp.getResult());
1344
1345 operation.addReduction(builder.getContext(), reductionOp, recipe);
1346 }
1347 }
1348 } else if constexpr (isCombinedType<OpTy>) {
1349 // Despite this being valid on ParallelOp or SerialOp, combined type
1350 // applies to the 'loop'.
1351 applyToLoopOp(clause);
1352 } else {
1353 llvm_unreachable("Unknown construct kind in VisitReductionClause");
1354 }
1355 }
1356};
1357
1358template <typename OpTy>
1359auto makeClauseEmitter(OpTy &op, CIRGen::CIRGenFunction &cgf,
1360 CIRGen::CIRGenBuilderTy &builder,
1361 OpenACCDirectiveKind dirKind, SourceLocation dirLoc) {
1362 return OpenACCClauseCIREmitter<OpTy>(op, cgf, builder, dirKind, dirLoc);
1363}
1364} // namespace
1365
1366template <typename Op>
1367void CIRGenFunction::emitOpenACCClauses(
1368 Op &op, OpenACCDirectiveKind dirKind, SourceLocation dirLoc,
1370 mlir::OpBuilder::InsertionGuard guardCase(builder);
1371
1372 // Sets insertion point before the 'op', since every new expression needs to
1373 // be before the operation.
1374 builder.setInsertionPoint(op);
1375 makeClauseEmitter(op, *this, builder, dirKind, dirLoc).emitClauses(clauses);
1376}
1377
1378#define EXPL_SPEC(N) \
1379 template void CIRGenFunction::emitOpenACCClauses<N>( \
1380 N &, OpenACCDirectiveKind, SourceLocation, \
1381 ArrayRef<const OpenACCClause *>);
1382EXPL_SPEC(mlir::acc::ParallelOp)
1383EXPL_SPEC(mlir::acc::SerialOp)
1384EXPL_SPEC(mlir::acc::KernelsOp)
1385EXPL_SPEC(mlir::acc::LoopOp)
1386EXPL_SPEC(mlir::acc::DataOp)
1387EXPL_SPEC(mlir::acc::InitOp)
1388EXPL_SPEC(mlir::acc::ShutdownOp)
1389EXPL_SPEC(mlir::acc::SetOp)
1390EXPL_SPEC(mlir::acc::WaitOp)
1391EXPL_SPEC(mlir::acc::HostDataOp)
1392EXPL_SPEC(mlir::acc::EnterDataOp)
1393EXPL_SPEC(mlir::acc::ExitDataOp)
1394EXPL_SPEC(mlir::acc::UpdateOp)
1395#undef EXPL_SPEC
1396
1397template <typename ComputeOp, typename LoopOp>
1398void CIRGenFunction::emitOpenACCClauses(
1399 ComputeOp &op, LoopOp &loopOp, OpenACCDirectiveKind dirKind,
1401 static_assert(std::is_same_v<mlir::acc::LoopOp, LoopOp>);
1402
1403 CombinedConstructClauseInfo<ComputeOp> inf{op, loopOp};
1404 // We cannot set the insertion point here and do so in the emitter, but make
1405 // sure we reset it with the 'guard' anyway.
1406 mlir::OpBuilder::InsertionGuard guardCase(builder);
1407 makeClauseEmitter(inf, *this, builder, dirKind, dirLoc).emitClauses(clauses);
1408}
1409
1410#define EXPL_SPEC(N) \
1411 template void CIRGenFunction::emitOpenACCClauses<N, mlir::acc::LoopOp>( \
1412 N &, mlir::acc::LoopOp &, OpenACCDirectiveKind, SourceLocation, \
1413 ArrayRef<const OpenACCClause *>);
1414
1415EXPL_SPEC(mlir::acc::ParallelOp)
1416EXPL_SPEC(mlir::acc::SerialOp)
1417EXPL_SPEC(mlir::acc::KernelsOp)
1418#undef EXPL_SPEC
#define EXPL_SPEC(N)
Expr * E
Defines the clang::Expr interface and subclasses for C++ expressions.
__device__ __2f16 float c
Holds long-lived AST nodes (such as types and decls) that can be referred to throughout the semantic ...
Definition: ASTContext.h:188
const LangOptions & getLangOpts() const
Definition: ASTContext.h:894
CharUnits getDeclAlign(const Decl *D, bool ForAlignof=false) const
Return a conservative estimate of the alignment of the specified decl D.
clang::MangleContext & getMangleContext()
Gets the mangle context.
Definition: CIRGenCXXABI.h:160
AutoVarEmission emitAutoVarAlloca(const clang::VarDecl &d, mlir::OpBuilder::InsertPoint ip={})
Definition: CIRGenDecl.cpp:27
const clang::Decl * curFuncDecl
mlir::Value evaluateExprAsBool(const clang::Expr *e)
Perform the usual unary conversions on the specified expression and compare the result against zero,...
Definition: CIRGenExpr.cpp:688
OpenACCDataOperandInfo getOpenACCDataOperandInfo(const Expr *e)
void emitAutoVarInit(const AutoVarEmission &emission)
Emit the initializer for an allocated variable.
Definition: CIRGenDecl.cpp:79
mlir::Value emitOpenACCIntExpr(const Expr *intExpr)
void emitDestroy(Address addr, QualType type, Destroyer *destroyer)
Immediately perform the destruction of the given object.
Definition: CIRGenDecl.cpp:727
mlir::MLIRContext & getMLIRContext()
Destroyer * getDestroyer(clang::QualType::DestructionKind kind)
Definition: CIRGenDecl.cpp:760
mlir::Value createOpenACCConstantInt(mlir::Location loc, unsigned width, int64_t value)
clang::ASTContext & getContext() const
void setAddrOfLocalVar(const clang::VarDecl *vd, Address addr)
Set the address of a local variable.
DiagnosticBuilder errorNYI(SourceLocation, llvm::StringRef)
Helpers to emit "not yet implemented" error diagnostics.
clang::ASTContext & getASTContext() const
Definition: CIRGenModule.h:102
mlir::Location getLoc(clang::SourceLocation cLoc)
Helpers to convert the presumed location of Clang's SourceLocation to an MLIR Location.
CIRGenCXXABI & getCXXABI() const
Definition: CIRGenModule.h:109
CharUnits - This is an opaque type for sizes expressed in character units.
Definition: CharUnits.h:38
DeclContext - This is used only as base class of specific decl types that can act as declaration cont...
Definition: DeclBase.h:1449
static DeclContext * castToDeclContext(const Decl *)
Definition: DeclBase.cpp:1063
This represents one expression.
Definition: Expr.h:112
llvm::APSInt EvaluateKnownConstInt(const ASTContext &Ctx, SmallVectorImpl< PartialDiagnosticAt > *Diag=nullptr) const
EvaluateKnownConstInt - Call EvaluateAsRValue and return the folded integer.
One of these records is kept for each identifier that is lexed.
StringRef getName() const
Return the actual identifier string.
A simple pair of identifier info and location.
MangleContext - Context for tracking state which persists across multiple calls to the C++ name mangl...
Definition: Mangle.h:52
virtual void mangleCanonicalTypeName(QualType T, raw_ostream &, bool NormalizeIntegers=false)=0
Generates a unique string for an externally visible type for use with TBAA or type uniquing.
void VisitClauseList(ArrayRef< const OpenACCClause * > List)
const Expr * getConditionExpr() const
ArrayRef< Expr * > getVarList()
This is the base type for all OpenACC Clauses.
Definition: OpenACCClause.h:27
Represents a 'collapse' clause on a 'loop' construct.
OpenACCModifierKind getModifierList() const
OpenACCModifierKind getModifierList() const
OpenACCModifierKind getModifierList() const
OpenACCModifierKind getModifierList() const
A 'default' clause, has the optional 'none' or 'present' argument.
OpenACCDefaultClauseKind getDefaultClauseKind() const
A 'device_type' or 'dtype' clause, takes a list of either an 'asterisk' or an identifier.
ArrayRef< DeviceTypeArgument > getArchitectures() const
ArrayRef< OpenACCFirstPrivateRecipe > getInitRecipes()
unsigned getNumExprs() const
std::pair< OpenACCGangKind, const Expr * > getExpr(unsigned I) const
An 'if' clause, which has a required condition expression.
ArrayRef< Expr * > getIntExprs()
ArrayRef< VarDecl * > getInitRecipes()
ArrayRef< OpenACCReductionRecipe > getRecipes()
OpenACCReductionOperator getReductionOp() const
A 'self' clause, which has an optional condition expression, or, in the event of an 'update' directiv...
const Expr * getConditionExpr() const
bool isConditionExprClause() const
ArrayRef< Expr * > getVarList()
bool hasConditionExpr() const
bool isEmptySelfClause() const
ArrayRef< Expr * > getSizeExprs()
ArrayRef< Expr * > getQueueIdExprs()
Expr * getDevNumExpr() const
A (possibly-)qualified type.
Definition: TypeBase.h:937
Encodes a location in the source.
A trivial tuple used to represent a source range.
SourceLocation getEndLoc() const LLVM_READONLY
Definition: Stmt.cpp:358
SourceRange getSourceRange() const LLVM_READONLY
SourceLocation tokens are not useful in isolation - they are low level value objects created/interpre...
Definition: Stmt.cpp:334
SourceLocation getBeginLoc() const LLVM_READONLY
Definition: Stmt.cpp:346
const Type * getPointeeOrArrayElementType() const
If this is a pointer type, return the pointee type.
Definition: TypeBase.h:9116
bool isPointerType() const
Definition: TypeBase.h:8580
bool isBuiltinType() const
Helper methods to distinguish type categories.
Definition: TypeBase.h:8703
QualType getType() const
Definition: Decl.h:722
Represents a variable declaration or definition.
Definition: Decl.h:925
QualType::DestructionKind needsDestruction(const ASTContext &Ctx) const
Would the destruction of this variable have any effect, and if so, what kind?
Definition: Decl.cpp:2851
const Expr * getInit() const
Definition: Decl.h:1367
const internal::VariadicDynCastAllOfMatcher< Stmt, Expr > expr
Matches expressions.
unsigned kind
All of the diagnostics that can be emitted by the frontend.
Definition: DiagnosticIDs.h:76
The JSON file list parser is used to communicate input to InstallAPI.
OpenACCDirectiveKind
Definition: OpenACCKinds.h:28
OpenACCReductionOperator
Definition: OpenACCKinds.h:547
@ Invalid
Invalid Reduction Clause Kind.
OpenACCModifierKind
Definition: OpenACCKinds.h:641
@ Invalid
Not a valid option.
bool isOpenACCModifierBitSet(OpenACCModifierKind List, OpenACCModifierKind Bit)
Definition: OpenACCKinds.h:652
const FunctionProtoType * T
OpenACCGangKind
Definition: OpenACCKinds.h:606
__DEVICE__ _Tp arg(const std::complex< _Tp > &__c)
Definition: complex_cmath.h:40
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...