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#include "CIRGenOpenACCRecipe.h"
18
19#include "clang/AST/ExprCXX.h"
20
21#include "mlir/Dialect/Arith/IR/Arith.h"
22#include "mlir/Dialect/OpenACC/OpenACC.h"
23#include "llvm/ADT/TypeSwitch.h"
24
25using namespace clang;
26using namespace clang::CIRGen;
27
28namespace {
29// Simple type-trait to see if the first template arg is one of the list, so we
30// can tell whether to `if-constexpr` a bunch of stuff.
31template <typename ToTest, typename T, typename... Tys>
32constexpr bool isOneOfTypes =
33 std::is_same_v<ToTest, T> || isOneOfTypes<ToTest, Tys...>;
34template <typename ToTest, typename T>
35constexpr bool isOneOfTypes<ToTest, T> = std::is_same_v<ToTest, T>;
36
37// Holds information for emitting clauses for a combined construct. We
38// instantiate the clause emitter with this type so that it can use
39// if-constexpr to specially handle these.
40template <typename CompOpTy> struct CombinedConstructClauseInfo {
41 using ComputeOpTy = CompOpTy;
42 ComputeOpTy computeOp;
43 mlir::acc::LoopOp loopOp;
44};
45template <typename ToTest> constexpr bool isCombinedType = false;
46template <typename T>
47constexpr bool isCombinedType<CombinedConstructClauseInfo<T>> = true;
48
49template <typename OpTy>
50class OpenACCClauseCIREmitter final
51 : public OpenACCClauseVisitor<OpenACCClauseCIREmitter<OpTy>> {
52 // Necessary for combined constructs.
53 template <typename FriendOpTy> friend class OpenACCClauseCIREmitter;
54
55 OpTy &operation;
56 mlir::OpBuilder::InsertPoint &recipeInsertLocation;
57 CIRGen::CIRGenFunction &cgf;
58 CIRGen::CIRGenBuilderTy &builder;
59
60 // This is necessary since a few of the clauses emit differently based on the
61 // directive kind they are attached to.
63 // TODO(cir): This source location should be able to go away once the NYI
64 // diagnostics are gone.
65 SourceLocation dirLoc;
66
67 llvm::SmallVector<mlir::acc::DeviceType> lastDeviceTypeValues;
68 // Keep track of the async-clause so that we can shortcut updating the data
69 // operands async clauses.
70 bool hasAsyncClause = false;
71 // Keep track of the data operands so that we can update their async clauses.
72 llvm::SmallVector<mlir::Operation *> dataOperands;
73
74 void clauseNotImplemented(const OpenACCClause &c) {
75 cgf.cgm.errorNYI(c.getSourceRange(), "OpenACC Clause", c.getClauseKind());
76 }
77
78 void setLastDeviceTypeClause(const OpenACCDeviceTypeClause &clause) {
79 lastDeviceTypeValues.clear();
80
81 for (const DeviceTypeArgument &arg : clause.getArchitectures())
82 lastDeviceTypeValues.push_back(decodeDeviceType(arg.getIdentifierInfo()));
83 }
84
85 mlir::Value emitIntExpr(const Expr *intExpr) {
86 return cgf.emitOpenACCIntExpr(intExpr);
87 }
88
89 // 'condition' as an OpenACC grammar production is used for 'if' and (some
90 // variants of) 'self'. It needs to be emitted as a signless-1-bit value, so
91 // this function emits the expression, then sets the unrealized conversion
92 // cast correctly, and returns the completed value.
93 mlir::Value createCondition(const Expr *condExpr) {
94 mlir::Value condition = cgf.evaluateExprAsBool(condExpr);
95 mlir::Location exprLoc = cgf.cgm.getLoc(condExpr->getBeginLoc());
96 mlir::IntegerType targetType = mlir::IntegerType::get(
97 &cgf.getMLIRContext(), /*width=*/1,
98 mlir::IntegerType::SignednessSemantics::Signless);
99 auto conversionOp = mlir::UnrealizedConversionCastOp::create(
100 builder, exprLoc, targetType, condition);
101 return conversionOp.getResult(0);
102 }
103
104 mlir::Value createConstantInt(mlir::Location loc, unsigned width,
105 int64_t value) {
106 return cgf.createOpenACCConstantInt(loc, width, value);
107 mlir::IntegerType ty = mlir::IntegerType::get(
108 &cgf.getMLIRContext(), width,
109 mlir::IntegerType::SignednessSemantics::Signless);
110 auto constOp = mlir::arith::ConstantOp::create(
111 builder, loc, builder.getIntegerAttr(ty, value));
112
113 return constOp;
114 }
115
116 mlir::Value createConstantInt(SourceLocation loc, unsigned width,
117 int64_t value) {
118 return createConstantInt(cgf.cgm.getLoc(loc), width, value);
119 }
120
121 mlir::acc::DeviceType decodeDeviceType(const IdentifierInfo *ii) {
122 // '*' case leaves no identifier-info, just a nullptr.
123 if (!ii)
124 return mlir::acc::DeviceType::Star;
125 return llvm::StringSwitch<mlir::acc::DeviceType>(ii->getName())
126 .CaseLower("default", mlir::acc::DeviceType::Default)
127 .CaseLower("host", mlir::acc::DeviceType::Host)
128 .CaseLower("multicore", mlir::acc::DeviceType::Multicore)
129 .CasesLower({"nvidia", "acc_device_nvidia"},
130 mlir::acc::DeviceType::Nvidia)
131 .CaseLower("radeon", mlir::acc::DeviceType::Radeon);
132 }
133
134 mlir::acc::GangArgType decodeGangType(OpenACCGangKind gk) {
135 switch (gk) {
136 case OpenACCGangKind::Num:
137 return mlir::acc::GangArgType::Num;
138 case OpenACCGangKind::Dim:
139 return mlir::acc::GangArgType::Dim;
140 case OpenACCGangKind::Static:
141 return mlir::acc::GangArgType::Static;
142 }
143 llvm_unreachable("unknown gang kind");
144 }
145
146 template <typename U = void,
147 typename = std::enable_if_t<isCombinedType<OpTy>, U>>
148 void applyToLoopOp(const OpenACCClause &c) {
149 mlir::OpBuilder::InsertionGuard guardCase(builder);
150 builder.setInsertionPoint(operation.loopOp);
151 OpenACCClauseCIREmitter<mlir::acc::LoopOp> loopEmitter{
152 operation.loopOp, recipeInsertLocation, cgf, builder, dirKind, dirLoc};
153 loopEmitter.lastDeviceTypeValues = lastDeviceTypeValues;
154 loopEmitter.Visit(&c);
155 }
156
157 template <typename U = void,
158 typename = std::enable_if_t<isCombinedType<OpTy>, U>>
159 void applyToComputeOp(const OpenACCClause &c) {
160 mlir::OpBuilder::InsertionGuard guardCase(builder);
161 builder.setInsertionPoint(operation.computeOp);
162 OpenACCClauseCIREmitter<typename OpTy::ComputeOpTy> computeEmitter{
163 operation.computeOp,
164 recipeInsertLocation,
165 cgf,
166 builder,
167 dirKind,
168 dirLoc};
169
170 computeEmitter.lastDeviceTypeValues = lastDeviceTypeValues;
171
172 // Async handler uses the first data operand to figure out where to insert
173 // its information if it is present. This ensures that the new handler will
174 // correctly set the insertion point for async.
175 if (!dataOperands.empty())
176 computeEmitter.dataOperands.push_back(dataOperands.front());
177 computeEmitter.Visit(&c);
178
179 // Make sure all of the new data operands are kept track of here. The
180 // combined constructs always apply 'async' to only the compute component,
181 // so we need to collect these.
182 dataOperands.append(computeEmitter.dataOperands);
183 }
184
185 mlir::acc::DataClauseModifier
186 convertModifiers(OpenACCModifierKind modifiers) {
187 using namespace mlir::acc;
188 static_assert(static_cast<int>(OpenACCModifierKind::Zero) ==
189 static_cast<int>(DataClauseModifier::zero) &&
190 static_cast<int>(OpenACCModifierKind::Readonly) ==
191 static_cast<int>(DataClauseModifier::readonly) &&
192 static_cast<int>(OpenACCModifierKind::AlwaysIn) ==
193 static_cast<int>(DataClauseModifier::alwaysin) &&
194 static_cast<int>(OpenACCModifierKind::AlwaysOut) ==
195 static_cast<int>(DataClauseModifier::alwaysout) &&
196 static_cast<int>(OpenACCModifierKind::Capture) ==
197 static_cast<int>(DataClauseModifier::capture));
198
199 DataClauseModifier mlirModifiers{};
200
201 // The MLIR representation of this represents `always` as `alwaysin` +
202 // `alwaysout`. So do a small fixup here.
203 if (isOpenACCModifierBitSet(modifiers, OpenACCModifierKind::Always)) {
204 mlirModifiers = mlirModifiers | DataClauseModifier::always;
205 modifiers &= ~OpenACCModifierKind::Always;
206 }
207
208 mlirModifiers = mlirModifiers | static_cast<DataClauseModifier>(modifiers);
209 return mlirModifiers;
210 }
211
212 template <typename BeforeOpTy, typename AfterOpTy>
213 void addDataOperand(const Expr *varOperand, mlir::acc::DataClause dataClause,
214 OpenACCModifierKind modifiers, bool structured,
215 bool implicit) {
216 CIRGenFunction::OpenACCDataOperandInfo opInfo =
217 cgf.getOpenACCDataOperandInfo(varOperand);
218
219 auto beforeOp =
220 BeforeOpTy::create(builder, opInfo.beginLoc, opInfo.varValue,
221 structured, implicit, opInfo.name, opInfo.bounds);
222 operation.getDataClauseOperandsMutable().append(beforeOp.getResult());
223
224 AfterOpTy afterOp;
225 {
226 mlir::OpBuilder::InsertionGuard guardCase(builder);
227 builder.setInsertionPointAfter(operation);
228
229 if constexpr (std::is_same_v<AfterOpTy, mlir::acc::DeleteOp> ||
230 std::is_same_v<AfterOpTy, mlir::acc::DetachOp>) {
231 // Detach/Delete ops don't have the variable reference here, so they
232 // take 1 fewer argument to their build function.
233 afterOp =
234 AfterOpTy::create(builder, opInfo.beginLoc, beforeOp, structured,
235 implicit, opInfo.name, opInfo.bounds);
236 } else {
237 afterOp = AfterOpTy::create(builder, opInfo.beginLoc, beforeOp,
238 opInfo.varValue, structured, implicit,
239 opInfo.name, opInfo.bounds);
240 }
241 }
242
243 // Set the 'rest' of the info for both operations.
244 beforeOp.setDataClause(dataClause);
245 afterOp.setDataClause(dataClause);
246 beforeOp.setModifiers(convertModifiers(modifiers));
247 afterOp.setModifiers(convertModifiers(modifiers));
248
249 // Make sure we record these, so 'async' values can be updated later.
250 dataOperands.push_back(beforeOp.getOperation());
251 dataOperands.push_back(afterOp.getOperation());
252 }
253
254 template <typename BeforeOpTy>
255 void addDataOperand(const Expr *varOperand, mlir::acc::DataClause dataClause,
256 OpenACCModifierKind modifiers, bool structured,
257 bool implicit) {
258 CIRGenFunction::OpenACCDataOperandInfo opInfo =
259 cgf.getOpenACCDataOperandInfo(varOperand);
260 auto beforeOp =
261 BeforeOpTy::create(builder, opInfo.beginLoc, opInfo.varValue,
262 structured, implicit, opInfo.name, opInfo.bounds);
263 operation.getDataClauseOperandsMutable().append(beforeOp.getResult());
264
265 // Set the 'rest' of the info for the operation.
266 beforeOp.setDataClause(dataClause);
267 beforeOp.setModifiers(convertModifiers(modifiers));
268
269 // Make sure we record these, so 'async' values can be updated later.
270 dataOperands.push_back(beforeOp.getOperation());
271 }
272
273 // Helper function that covers for the fact that we don't have this function
274 // on all operation types.
275 mlir::ArrayAttr getAsyncOnlyAttr() {
276 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
277 mlir::acc::KernelsOp, mlir::acc::DataOp,
278 mlir::acc::UpdateOp>) {
279 return operation.getAsyncOnlyAttr();
280 } else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp,
281 mlir::acc::ExitDataOp>) {
282 if (!operation.getAsyncAttr())
283 return mlir::ArrayAttr{};
284
285 llvm::SmallVector<mlir::Attribute> devTysTemp;
286 devTysTemp.push_back(mlir::acc::DeviceTypeAttr::get(
287 builder.getContext(), mlir::acc::DeviceType::None));
288 return mlir::ArrayAttr::get(builder.getContext(), devTysTemp);
289 } else if constexpr (isCombinedType<OpTy>) {
290 return operation.computeOp.getAsyncOnlyAttr();
291 }
292
293 // Note: 'wait' has async as well, but it cannot have data clauses, so we
294 // don't have to handle them here.
295
296 llvm_unreachable("getting asyncOnly when clause not valid on operation?");
297 }
298
299 // Helper function that covers for the fact that we don't have this function
300 // on all operation types.
301 mlir::ArrayAttr getAsyncOperandsDeviceTypeAttr() {
302 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
303 mlir::acc::KernelsOp, mlir::acc::DataOp,
304 mlir::acc::UpdateOp>) {
305 return operation.getAsyncOperandsDeviceTypeAttr();
306 } else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp,
307 mlir::acc::ExitDataOp>) {
308 if (!operation.getAsyncOperand())
309 return mlir::ArrayAttr{};
310
311 llvm::SmallVector<mlir::Attribute> devTysTemp;
312 devTysTemp.push_back(mlir::acc::DeviceTypeAttr::get(
313 builder.getContext(), mlir::acc::DeviceType::None));
314 return mlir::ArrayAttr::get(builder.getContext(), devTysTemp);
315 } else if constexpr (isCombinedType<OpTy>) {
316 return operation.computeOp.getAsyncOperandsDeviceTypeAttr();
317 }
318
319 // Note: 'wait' has async as well, but it cannot have data clauses, so we
320 // don't have to handle them here.
321
322 llvm_unreachable(
323 "getting asyncOperandsDeviceType when clause not valid on operation?");
324 }
325
326 // Helper function that covers for the fact that we don't have this function
327 // on all operation types.
328 mlir::OperandRange getAsyncOperands() {
329 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
330 mlir::acc::KernelsOp, mlir::acc::DataOp,
331 mlir::acc::UpdateOp>)
332 return operation.getAsyncOperands();
333 else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp,
334 mlir::acc::ExitDataOp>)
335 return operation.getAsyncOperandMutable();
336 else if constexpr (isCombinedType<OpTy>)
337 return operation.computeOp.getAsyncOperands();
338
339 // Note: 'wait' has async as well, but it cannot have data clauses, so we
340 // don't have to handle them here.
341
342 llvm_unreachable(
343 "getting asyncOperandsDeviceType when clause not valid on operation?");
344 }
345
346 // The 'data' clauses all require that we add the 'async' values from the
347 // operation to them. We've collected the data operands along the way, so use
348 // that list to get the current 'async' values.
349 void updateDataOperandAsyncValues() {
350 if (!hasAsyncClause || dataOperands.empty())
351 return;
352
353 for (mlir::Operation *dataOp : dataOperands) {
354 llvm::TypeSwitch<mlir::Operation *, void>(dataOp)
355 .Case<ACC_DATA_ENTRY_OPS, ACC_DATA_EXIT_OPS>([&](auto op) {
356 op.setAsyncOnlyAttr(getAsyncOnlyAttr());
357 op.setAsyncOperandsDeviceTypeAttr(getAsyncOperandsDeviceTypeAttr());
358 op.getAsyncOperandsMutable().assign(getAsyncOperands());
359 })
360 .Default([&](mlir::Operation *) {
361 llvm_unreachable("Not a data operation?");
362 });
363 }
364 }
365
366public:
367 OpenACCClauseCIREmitter(OpTy &operation,
368 mlir::OpBuilder::InsertPoint &recipeInsertLocation,
369 CIRGen::CIRGenFunction &cgf,
370 CIRGen::CIRGenBuilderTy &builder,
371 OpenACCDirectiveKind dirKind, SourceLocation dirLoc)
372 : operation(operation), recipeInsertLocation(recipeInsertLocation),
373 cgf(cgf), builder(builder), dirKind(dirKind), dirLoc(dirLoc) {}
374
375 void VisitClause(const OpenACCClause &clause) {
376 clauseNotImplemented(clause);
377 }
378
379 // The entry point for the CIR emitter. All users should use this rather than
380 // 'visitClauseList', as this also handles the things that have to happen
381 // 'after' the clauses are all visited.
382 void emitClauses(ArrayRef<const OpenACCClause *> clauses) {
383 this->VisitClauseList(clauses);
384 updateDataOperandAsyncValues();
385 }
386
387 void VisitDefaultClause(const OpenACCDefaultClause &clause) {
388 // This type-trait checks if 'op'(the first arg) is one of the mlir::acc
389 // operations listed in the rest of the arguments.
390 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
391 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
392 switch (clause.getDefaultClauseKind()) {
393 case OpenACCDefaultClauseKind::None:
394 operation.setDefaultAttr(mlir::acc::ClauseDefaultValue::None);
395 break;
396 case OpenACCDefaultClauseKind::Present:
397 operation.setDefaultAttr(mlir::acc::ClauseDefaultValue::Present);
398 break;
399 case OpenACCDefaultClauseKind::Invalid:
400 break;
401 }
402 } else if constexpr (isCombinedType<OpTy>) {
403 applyToComputeOp(clause);
404 } else {
405 llvm_unreachable("Unknown construct kind in VisitDefaultClause");
406 }
407 }
408
409 void VisitDeviceTypeClause(const OpenACCDeviceTypeClause &clause) {
410 setLastDeviceTypeClause(clause);
411
412 if constexpr (isOneOfTypes<OpTy, mlir::acc::InitOp,
413 mlir::acc::ShutdownOp>) {
414 for (const DeviceTypeArgument &arg : clause.getArchitectures())
415 operation.addDeviceType(builder.getContext(),
416 decodeDeviceType(arg.getIdentifierInfo()));
417 } else if constexpr (isOneOfTypes<OpTy, mlir::acc::SetOp>) {
418 assert(!operation.getDeviceTypeAttr() && "already have device-type?");
419 assert(clause.getArchitectures().size() <= 1);
420
421 if (!clause.getArchitectures().empty())
422 operation.setDeviceType(
423 decodeDeviceType(clause.getArchitectures()[0].getIdentifierInfo()));
424 } else if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
425 mlir::acc::SerialOp, mlir::acc::KernelsOp,
426 mlir::acc::DataOp, mlir::acc::LoopOp,
427 mlir::acc::UpdateOp>) {
428 // Nothing to do here, these constructs don't have any IR for these, as
429 // they just modify the other clauses IR. So setting of
430 // `lastDeviceTypeValues` (done above) is all we need.
431 } else if constexpr (isCombinedType<OpTy>) {
432 // Nothing to do here either, combined constructs are just going to use
433 // 'lastDeviceTypeValues' to set the value for the child visitor.
434 } else {
435 // TODO: When we've implemented this for everything, switch this to an
436 // unreachable. routine construct remains.
437 return clauseNotImplemented(clause);
438 }
439 }
440
441 void VisitNumWorkersClause(const OpenACCNumWorkersClause &clause) {
442 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
443 mlir::acc::KernelsOp>) {
444 operation.addNumWorkersOperand(builder.getContext(),
445 emitIntExpr(clause.getIntExpr()),
446 lastDeviceTypeValues);
447 } else if constexpr (isCombinedType<OpTy>) {
448 applyToComputeOp(clause);
449 } else {
450 llvm_unreachable("Unknown construct kind in VisitNumGangsClause");
451 }
452 }
453
454 void VisitVectorLengthClause(const OpenACCVectorLengthClause &clause) {
455 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
456 mlir::acc::KernelsOp>) {
457 operation.addVectorLengthOperand(builder.getContext(),
458 emitIntExpr(clause.getIntExpr()),
459 lastDeviceTypeValues);
460 } else if constexpr (isCombinedType<OpTy>) {
461 applyToComputeOp(clause);
462 } else {
463 llvm_unreachable("Unknown construct kind in VisitVectorLengthClause");
464 }
465 }
466
467 void VisitAsyncClause(const OpenACCAsyncClause &clause) {
468 hasAsyncClause = true;
469 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
470 mlir::acc::KernelsOp, mlir::acc::DataOp,
471 mlir::acc::EnterDataOp, mlir::acc::ExitDataOp,
472 mlir::acc::UpdateOp>) {
473 if (!clause.hasIntExpr()) {
474 operation.addAsyncOnly(builder.getContext(), lastDeviceTypeValues);
475 } else {
476
477 mlir::Value intExpr;
478 {
479 // Async int exprs can be referenced by the data operands, which means
480 // that the int-exprs have to appear before them. IF there is a data
481 // operand already, set the insertion point to 'before' it.
482 mlir::OpBuilder::InsertionGuard guardCase(builder);
483 if (!dataOperands.empty())
484 builder.setInsertionPoint(dataOperands.front());
485 intExpr = emitIntExpr(clause.getIntExpr());
486 }
487 operation.addAsyncOperand(builder.getContext(), intExpr,
488 lastDeviceTypeValues);
489 }
490 } else if constexpr (isOneOfTypes<OpTy, mlir::acc::WaitOp>) {
491 // Wait doesn't have a device_type, so its handling here is slightly
492 // different.
493 if (!clause.hasIntExpr())
494 operation.setAsync(true);
495 else
496 operation.getAsyncOperandMutable().append(
497 emitIntExpr(clause.getIntExpr()));
498 } else if constexpr (isCombinedType<OpTy>) {
499 applyToComputeOp(clause);
500 } else {
501 // TODO: When we've implemented this for everything, switch this to an
502 // unreachable. Combined constructs remain. update construct remains.
503 return clauseNotImplemented(clause);
504 }
505 }
506
507 void VisitSelfClause(const OpenACCSelfClause &clause) {
508 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
509 mlir::acc::KernelsOp>) {
510 if (clause.isEmptySelfClause()) {
511 operation.setSelfAttr(true);
512 } else if (clause.isConditionExprClause()) {
513 assert(clause.hasConditionExpr());
514 operation.getSelfCondMutable().append(
515 createCondition(clause.getConditionExpr()));
516 } else {
517 llvm_unreachable("var-list version of self shouldn't get here");
518 }
519 } else if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
520 assert(!clause.isEmptySelfClause() && !clause.isConditionExprClause() &&
521 "var-list version of self required for update");
522 for (const Expr *var : clause.getVarList())
523 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::UpdateHostOp>(
524 var, mlir::acc::DataClause::acc_update_self, {},
525 /*structured=*/false, /*implicit=*/false);
526 } else if constexpr (isCombinedType<OpTy>) {
527 applyToComputeOp(clause);
528 } else {
529 llvm_unreachable("Unknown construct kind in VisitSelfClause");
530 }
531 }
532
533 void VisitHostClause(const OpenACCHostClause &clause) {
534 if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
535 for (const Expr *var : clause.getVarList())
536 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::UpdateHostOp>(
537 var, mlir::acc::DataClause::acc_update_host, {},
538 /*structured=*/false, /*implicit=*/false);
539 } else {
540 llvm_unreachable("Unknown construct kind in VisitHostClause");
541 }
542 }
543
544 void VisitDeviceClause(const OpenACCDeviceClause &clause) {
545 if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
546 for (const Expr *var : clause.getVarList())
547 addDataOperand<mlir::acc::UpdateDeviceOp>(
548 var, mlir::acc::DataClause::acc_update_device, {},
549 /*structured=*/false, /*implicit=*/false);
550 } else {
551 llvm_unreachable("Unknown construct kind in VisitDeviceClause");
552 }
553 }
554
555 void VisitIfClause(const OpenACCIfClause &clause) {
556 if constexpr (isOneOfTypes<
557 OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
558 mlir::acc::KernelsOp, mlir::acc::InitOp,
559 mlir::acc::ShutdownOp, mlir::acc::SetOp,
560 mlir::acc::DataOp, mlir::acc::WaitOp,
561 mlir::acc::HostDataOp, mlir::acc::EnterDataOp,
562 mlir::acc::ExitDataOp, mlir::acc::UpdateOp,
563 mlir::acc::AtomicReadOp, mlir::acc::AtomicWriteOp,
564 mlir::acc::AtomicUpdateOp, mlir::acc::AtomicCaptureOp>) {
565 operation.getIfCondMutable().append(
566 createCondition(clause.getConditionExpr()));
567 } else if constexpr (isCombinedType<OpTy>) {
568 applyToComputeOp(clause);
569 } else {
570 llvm_unreachable("Unknown construct kind in VisitIfClause");
571 }
572 }
573
574 void VisitIfPresentClause(const OpenACCIfPresentClause &clause) {
575 if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp,
576 mlir::acc::UpdateOp>) {
577 operation.setIfPresent(true);
578 } else {
579 llvm_unreachable("unknown construct kind in VisitIfPresentClause");
580 }
581 }
582
583 void VisitDeviceNumClause(const OpenACCDeviceNumClause &clause) {
584 if constexpr (isOneOfTypes<OpTy, mlir::acc::InitOp, mlir::acc::ShutdownOp,
585 mlir::acc::SetOp>) {
586 operation.getDeviceNumMutable().append(emitIntExpr(clause.getIntExpr()));
587 } else {
588 llvm_unreachable(
589 "init, shutdown, set, are only valid device_num constructs");
590 }
591 }
592
593 void VisitNumGangsClause(const OpenACCNumGangsClause &clause) {
594 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
595 mlir::acc::KernelsOp>) {
596 llvm::SmallVector<mlir::Value> values;
597 for (const Expr *E : clause.getIntExprs())
598 values.push_back(emitIntExpr(E));
599
600 operation.addNumGangsOperands(builder.getContext(), values,
601 lastDeviceTypeValues);
602 } else if constexpr (isCombinedType<OpTy>) {
603 applyToComputeOp(clause);
604 } else {
605 llvm_unreachable("Unknown construct kind in VisitNumGangsClause");
606 }
607 }
608
609 void VisitWaitClause(const OpenACCWaitClause &clause) {
610 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
611 mlir::acc::KernelsOp, mlir::acc::DataOp,
612 mlir::acc::EnterDataOp, mlir::acc::ExitDataOp,
613 mlir::acc::UpdateOp>) {
614 if (!clause.hasExprs()) {
615 operation.addWaitOnly(builder.getContext(), lastDeviceTypeValues);
616 } else {
617 llvm::SmallVector<mlir::Value> values;
618 if (clause.hasDevNumExpr())
619 values.push_back(emitIntExpr(clause.getDevNumExpr()));
620 for (const Expr *E : clause.getQueueIdExprs())
621 values.push_back(emitIntExpr(E));
622 operation.addWaitOperands(builder.getContext(), clause.hasDevNumExpr(),
623 values, lastDeviceTypeValues);
624 }
625 } else if constexpr (isCombinedType<OpTy>) {
626 applyToComputeOp(clause);
627 } else {
628 // TODO: When we've implemented this for everything, switch this to an
629 // unreachable. update construct remains.
630 return clauseNotImplemented(clause);
631 }
632 }
633
634 void VisitDefaultAsyncClause(const OpenACCDefaultAsyncClause &clause) {
635 if constexpr (isOneOfTypes<OpTy, mlir::acc::SetOp>) {
636 operation.getDefaultAsyncMutable().append(
637 emitIntExpr(clause.getIntExpr()));
638 } else {
639 llvm_unreachable("set, is only valid device_num constructs");
640 }
641 }
642
643 void VisitSeqClause(const OpenACCSeqClause &clause) {
644 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
645 operation.addSeq(builder.getContext(), lastDeviceTypeValues);
646 } else if constexpr (isCombinedType<OpTy>) {
647 applyToLoopOp(clause);
648 } else {
649 // TODO: When we've implemented this for everything, switch this to an
650 // unreachable. Routine construct remains.
651 return clauseNotImplemented(clause);
652 }
653 }
654
655 void VisitAutoClause(const OpenACCAutoClause &clause) {
656 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
657 operation.addAuto(builder.getContext(), lastDeviceTypeValues);
658 } else if constexpr (isCombinedType<OpTy>) {
659 applyToLoopOp(clause);
660 } else {
661 // TODO: When we've implemented this for everything, switch this to an
662 // unreachable. Routine, construct remains.
663 return clauseNotImplemented(clause);
664 }
665 }
666
667 void VisitIndependentClause(const OpenACCIndependentClause &clause) {
668 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
669 operation.addIndependent(builder.getContext(), lastDeviceTypeValues);
670 } else if constexpr (isCombinedType<OpTy>) {
671 applyToLoopOp(clause);
672 } else {
673 // TODO: When we've implemented this for everything, switch this to an
674 // unreachable. Routine construct remains.
675 return clauseNotImplemented(clause);
676 }
677 }
678
679 void VisitCollapseClause(const OpenACCCollapseClause &clause) {
680 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
681 llvm::APInt value =
682 clause.getIntExpr()->EvaluateKnownConstInt(cgf.cgm.getASTContext());
683
684 value = value.sextOrTrunc(64);
685 operation.setCollapseForDeviceTypes(builder.getContext(),
686 lastDeviceTypeValues, value);
687 } else if constexpr (isCombinedType<OpTy>) {
688 applyToLoopOp(clause);
689 } else {
690 llvm_unreachable("Unknown construct kind in VisitCollapseClause");
691 }
692 }
693
694 void VisitTileClause(const OpenACCTileClause &clause) {
695 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
696 llvm::SmallVector<mlir::Value> values;
697
698 for (const Expr *e : clause.getSizeExprs()) {
699 mlir::Location exprLoc = cgf.cgm.getLoc(e->getBeginLoc());
700
701 // We represent the * as -1. Additionally, this is a constant, so we
702 // can always just emit it as 64 bits to avoid having to do any more
703 // work to determine signedness or size.
705 values.push_back(createConstantInt(exprLoc, 64, -1));
706 } else {
707 llvm::APInt curValue =
708 e->EvaluateKnownConstInt(cgf.cgm.getASTContext());
709 values.push_back(createConstantInt(
710 exprLoc, 64, curValue.sextOrTrunc(64).getSExtValue()));
711 }
712 }
713
714 operation.setTileForDeviceTypes(builder.getContext(),
715 lastDeviceTypeValues, values);
716 } else if constexpr (isCombinedType<OpTy>) {
717 applyToLoopOp(clause);
718 } else {
719 llvm_unreachable("Unknown construct kind in VisitTileClause");
720 }
721 }
722
723 void VisitWorkerClause(const OpenACCWorkerClause &clause) {
724 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
725 if (clause.hasIntExpr())
726 operation.addWorkerNumOperand(builder.getContext(),
727 emitIntExpr(clause.getIntExpr()),
728 lastDeviceTypeValues);
729 else
730 operation.addEmptyWorker(builder.getContext(), lastDeviceTypeValues);
731
732 } else if constexpr (isCombinedType<OpTy>) {
733 applyToLoopOp(clause);
734 } else {
735 // TODO: When we've implemented this for everything, switch this to an
736 // unreachable. Combined constructs remain.
737 return clauseNotImplemented(clause);
738 }
739 }
740
741 void VisitVectorClause(const OpenACCVectorClause &clause) {
742 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
743 if (clause.hasIntExpr())
744 operation.addVectorOperand(builder.getContext(),
745 emitIntExpr(clause.getIntExpr()),
746 lastDeviceTypeValues);
747 else
748 operation.addEmptyVector(builder.getContext(), lastDeviceTypeValues);
749
750 } else if constexpr (isCombinedType<OpTy>) {
751 applyToLoopOp(clause);
752 } else {
753 // TODO: When we've implemented this for everything, switch this to an
754 // unreachable. Combined constructs remain.
755 return clauseNotImplemented(clause);
756 }
757 }
758
759 void VisitGangClause(const OpenACCGangClause &clause) {
760 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
761 if (clause.getNumExprs() == 0) {
762 operation.addEmptyGang(builder.getContext(), lastDeviceTypeValues);
763 } else {
764 llvm::SmallVector<mlir::Value> values;
765 llvm::SmallVector<mlir::acc::GangArgType> argTypes;
766 for (unsigned i : llvm::index_range(0u, clause.getNumExprs())) {
767 auto [kind, expr] = clause.getExpr(i);
768 mlir::Location exprLoc = cgf.cgm.getLoc(expr->getBeginLoc());
769 argTypes.push_back(decodeGangType(kind));
770 if (kind == OpenACCGangKind::Dim) {
771 llvm::APInt curValue =
772 expr->EvaluateKnownConstInt(cgf.cgm.getASTContext());
773 // The value is 1, 2, or 3, but the type isn't necessarily smaller
774 // than 64.
775 curValue = curValue.sextOrTrunc(64);
776 values.push_back(
777 createConstantInt(exprLoc, 64, curValue.getSExtValue()));
779 values.push_back(createConstantInt(exprLoc, 64, -1));
780 } else {
781 values.push_back(emitIntExpr(expr));
782 }
783 }
784
785 operation.addGangOperands(builder.getContext(), lastDeviceTypeValues,
786 argTypes, values);
787 }
788 } else if constexpr (isCombinedType<OpTy>) {
789 applyToLoopOp(clause);
790 } else {
791 llvm_unreachable("Unknown construct kind in VisitGangClause");
792 }
793 }
794
795 void VisitCopyClause(const OpenACCCopyClause &clause) {
796 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
797 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
798 for (const Expr *var : clause.getVarList())
799 addDataOperand<mlir::acc::CopyinOp, mlir::acc::CopyoutOp>(
800 var, mlir::acc::DataClause::acc_copy, clause.getModifierList(),
801 /*structured=*/true,
802 /*implicit=*/false);
803 } else if constexpr (isCombinedType<OpTy>) {
804 applyToComputeOp(clause);
805 } else {
806 // TODO: When we've implemented this for everything, switch this to an
807 // unreachable. declare construct remains.
808 return clauseNotImplemented(clause);
809 }
810 }
811
812 void VisitCopyInClause(const OpenACCCopyInClause &clause) {
813 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
814 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
815 for (const Expr *var : clause.getVarList())
816 addDataOperand<mlir::acc::CopyinOp, mlir::acc::DeleteOp>(
817 var, mlir::acc::DataClause::acc_copyin, clause.getModifierList(),
818 /*structured=*/true,
819 /*implicit=*/false);
820 } else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
821 for (const Expr *var : clause.getVarList())
822 addDataOperand<mlir::acc::CopyinOp>(
823 var, mlir::acc::DataClause::acc_copyin, clause.getModifierList(),
824 /*structured=*/false, /*implicit=*/false);
825 } else if constexpr (isCombinedType<OpTy>) {
826 applyToComputeOp(clause);
827 } else {
828 // TODO: When we've implemented this for everything, switch this to an
829 // unreachable. declare construct remains.
830 return clauseNotImplemented(clause);
831 }
832 }
833
834 void VisitCopyOutClause(const OpenACCCopyOutClause &clause) {
835 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
836 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
837 for (const Expr *var : clause.getVarList())
838 addDataOperand<mlir::acc::CreateOp, mlir::acc::CopyoutOp>(
839 var, mlir::acc::DataClause::acc_copyout, clause.getModifierList(),
840 /*structured=*/true,
841 /*implicit=*/false);
842 } else if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
843 for (const Expr *var : clause.getVarList())
844 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::CopyoutOp>(
845 var, mlir::acc::DataClause::acc_copyout, clause.getModifierList(),
846 /*structured=*/false,
847 /*implicit=*/false);
848 } else if constexpr (isCombinedType<OpTy>) {
849 applyToComputeOp(clause);
850 } else {
851 // TODO: When we've implemented this for everything, switch this to an
852 // unreachable. declare construct remains.
853 return clauseNotImplemented(clause);
854 }
855 }
856
857 void VisitCreateClause(const OpenACCCreateClause &clause) {
858 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
859 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
860 for (const Expr *var : clause.getVarList())
861 addDataOperand<mlir::acc::CreateOp, mlir::acc::DeleteOp>(
862 var, mlir::acc::DataClause::acc_create, clause.getModifierList(),
863 /*structured=*/true,
864 /*implicit=*/false);
865 } else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
866 for (const Expr *var : clause.getVarList())
867 addDataOperand<mlir::acc::CreateOp>(
868 var, mlir::acc::DataClause::acc_create, clause.getModifierList(),
869 /*structured=*/false, /*implicit=*/false);
870 } else if constexpr (isCombinedType<OpTy>) {
871 applyToComputeOp(clause);
872 } else {
873 // TODO: When we've implemented this for everything, switch this to an
874 // unreachable. declare construct remains.
875 return clauseNotImplemented(clause);
876 }
877 }
878
879 void VisitDeleteClause(const OpenACCDeleteClause &clause) {
880 if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
881 for (const Expr *var : clause.getVarList())
882 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::DeleteOp>(
883 var, mlir::acc::DataClause::acc_delete, {},
884 /*structured=*/false,
885 /*implicit=*/false);
886 } else {
887 llvm_unreachable("Unknown construct kind in VisitDeleteClause");
888 }
889 }
890
891 void VisitDetachClause(const OpenACCDetachClause &clause) {
892 if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
893 for (const Expr *var : clause.getVarList())
894 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::DetachOp>(
895 var, mlir::acc::DataClause::acc_detach, {},
896 /*structured=*/false,
897 /*implicit=*/false);
898 } else {
899 llvm_unreachable("Unknown construct kind in VisitDetachClause");
900 }
901 }
902
903 void VisitFinalizeClause(const OpenACCFinalizeClause &clause) {
904 if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
905 operation.setFinalize(true);
906 } else {
907 llvm_unreachable("Unknown construct kind in VisitFinalizeClause");
908 }
909 }
910
911 void VisitUseDeviceClause(const OpenACCUseDeviceClause &clause) {
912 if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp>) {
913 for (const Expr *var : clause.getVarList())
914 addDataOperand<mlir::acc::UseDeviceOp>(
915 var, mlir::acc::DataClause::acc_use_device, {}, /*structured=*/true,
916 /*implicit=*/false);
917 } else {
918 llvm_unreachable("Unknown construct kind in VisitUseDeviceClause");
919 }
920 }
921
922 void VisitDevicePtrClause(const OpenACCDevicePtrClause &clause) {
923 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
924 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
925 for (const Expr *var : clause.getVarList())
926 addDataOperand<mlir::acc::DevicePtrOp>(
927 var, mlir::acc::DataClause::acc_deviceptr, {},
928 /*structured=*/true,
929 /*implicit=*/false);
930 } else if constexpr (isCombinedType<OpTy>) {
931 applyToComputeOp(clause);
932 } else {
933 // TODO: When we've implemented this for everything, switch this to an
934 // unreachable. declare remains.
935 return clauseNotImplemented(clause);
936 }
937 }
938
939 void VisitNoCreateClause(const OpenACCNoCreateClause &clause) {
940 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
941 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
942 for (const Expr *var : clause.getVarList())
943 addDataOperand<mlir::acc::NoCreateOp, mlir::acc::DeleteOp>(
944 var, mlir::acc::DataClause::acc_no_create, {}, /*structured=*/true,
945 /*implicit=*/false);
946 } else if constexpr (isCombinedType<OpTy>) {
947 applyToComputeOp(clause);
948 } else {
949 llvm_unreachable("Unknown construct kind in VisitNoCreateClause");
950 }
951 }
952
953 void VisitPresentClause(const OpenACCPresentClause &clause) {
954 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
955 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
956 for (const Expr *var : clause.getVarList())
957 addDataOperand<mlir::acc::PresentOp, mlir::acc::DeleteOp>(
958 var, mlir::acc::DataClause::acc_present, {}, /*structured=*/true,
959 /*implicit=*/false);
960 } else if constexpr (isCombinedType<OpTy>) {
961 applyToComputeOp(clause);
962 } else {
963 // TODO: When we've implemented this for everything, switch this to an
964 // unreachable. declare remains.
965 return clauseNotImplemented(clause);
966 }
967 }
968
969 void VisitAttachClause(const OpenACCAttachClause &clause) {
970 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
971 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
972 for (const Expr *var : clause.getVarList())
973 addDataOperand<mlir::acc::AttachOp, mlir::acc::DetachOp>(
974 var, mlir::acc::DataClause::acc_attach, {}, /*structured=*/true,
975 /*implicit=*/false);
976 } else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
977 for (const Expr *var : clause.getVarList())
978 addDataOperand<mlir::acc::AttachOp>(
979 var, mlir::acc::DataClause::acc_attach, {},
980 /*structured=*/false, /*implicit=*/false);
981 } else if constexpr (isCombinedType<OpTy>) {
982 applyToComputeOp(clause);
983 } else {
984 llvm_unreachable("Unknown construct kind in VisitAttachClause");
985 }
986 }
987
988 void VisitPrivateClause(const OpenACCPrivateClause &clause) {
989 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
990 mlir::acc::LoopOp>) {
991 for (const auto [varExpr, varRecipe] :
992 llvm::zip_equal(clause.getVarList(), clause.getInitRecipes())) {
993 CIRGenFunction::OpenACCDataOperandInfo opInfo =
994 cgf.getOpenACCDataOperandInfo(varExpr);
995 auto privateOp = mlir::acc::PrivateOp::create(
996 builder, opInfo.beginLoc, opInfo.varValue, /*structured=*/true,
997 /*implicit=*/false, opInfo.name, opInfo.bounds);
998 privateOp.setDataClause(mlir::acc::DataClause::acc_private);
999
1000 {
1001 mlir::OpBuilder::InsertionGuard guardCase(builder);
1002
1003 auto recipe =
1004 OpenACCRecipeBuilder<mlir::acc::PrivateRecipeOp>(cgf, builder)
1005 .getOrCreateRecipe(
1006 cgf.getContext(), recipeInsertLocation, varExpr,
1007 varRecipe.AllocaDecl,
1008 /*temporary=*/nullptr, OpenACCReductionOperator::Invalid,
1009 Decl::castToDeclContext(cgf.curFuncDecl), opInfo.origType,
1010 opInfo.bounds.size(), opInfo.boundTypes, opInfo.baseType,
1011 privateOp, /*reductionCombinerRecipes=*/{});
1012 // TODO: OpenACC: The dialect is going to change in the near future to
1013 // have these be on a different operation, so when that changes, we
1014 // probably need to change these here.
1015 operation.addPrivatization(builder.getContext(), privateOp, recipe);
1016 }
1017 }
1018 } else if constexpr (isCombinedType<OpTy>) {
1019 // Despite this being valid on ParallelOp or SerialOp, combined type
1020 // applies to the 'loop'.
1021 applyToLoopOp(clause);
1022 } else {
1023 llvm_unreachable("Unknown construct kind in VisitPrivateClause");
1024 }
1025 }
1026
1027 void VisitFirstPrivateClause(const OpenACCFirstPrivateClause &clause) {
1028 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
1029 mlir::acc::SerialOp>) {
1030 for (const auto [varExpr, varRecipe] :
1031 llvm::zip_equal(clause.getVarList(), clause.getInitRecipes())) {
1032 CIRGenFunction::OpenACCDataOperandInfo opInfo =
1033 cgf.getOpenACCDataOperandInfo(varExpr);
1034 auto firstPrivateOp = mlir::acc::FirstprivateOp::create(
1035 builder, opInfo.beginLoc, opInfo.varValue, /*structured=*/true,
1036 /*implicit=*/false, opInfo.name, opInfo.bounds);
1037
1038 firstPrivateOp.setDataClause(mlir::acc::DataClause::acc_firstprivate);
1039
1040 {
1041 mlir::OpBuilder::InsertionGuard guardCase(builder);
1042
1043 auto recipe =
1044 OpenACCRecipeBuilder<mlir::acc::FirstprivateRecipeOp>(cgf,
1045 builder)
1046 .getOrCreateRecipe(
1047 cgf.getContext(), recipeInsertLocation, varExpr,
1048 varRecipe.AllocaDecl, varRecipe.InitFromTemporary,
1049 OpenACCReductionOperator::Invalid,
1050 Decl::castToDeclContext(cgf.curFuncDecl), opInfo.origType,
1051 opInfo.bounds.size(), opInfo.boundTypes, opInfo.baseType,
1052 firstPrivateOp, /*reductionCombinerRecipe=*/{});
1053
1054 // TODO: OpenACC: The dialect is going to change in the near future to
1055 // have these be on a different operation, so when that changes, we
1056 // probably need to change these here.
1057 operation.addFirstPrivatization(builder.getContext(), firstPrivateOp,
1058 recipe);
1059 }
1060 }
1061 } else if constexpr (isCombinedType<OpTy>) {
1062 // Unlike 'private', 'firstprivate' applies to the compute op, not the
1063 // loop op.
1064 applyToComputeOp(clause);
1065 } else {
1066 llvm_unreachable("Unknown construct kind in VisitFirstPrivateClause");
1067 }
1068 }
1069
1070 void VisitReductionClause(const OpenACCReductionClause &clause) {
1071 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
1072 mlir::acc::LoopOp>) {
1073 for (const auto [varExpr, varRecipe] :
1074 llvm::zip_equal(clause.getVarList(), clause.getRecipes())) {
1075 CIRGenFunction::OpenACCDataOperandInfo opInfo =
1076 cgf.getOpenACCDataOperandInfo(varExpr);
1077
1078 auto reductionOp = mlir::acc::ReductionOp::create(
1079 builder, opInfo.beginLoc, opInfo.varValue, /*structured=*/true,
1080 /*implicit=*/false, opInfo.name, opInfo.bounds);
1081 reductionOp.setDataClause(mlir::acc::DataClause::acc_reduction);
1082
1083 {
1084 mlir::OpBuilder::InsertionGuard guardCase(builder);
1085
1086 auto recipe =
1087 OpenACCRecipeBuilder<mlir::acc::ReductionRecipeOp>(cgf, builder)
1088 .getOrCreateRecipe(
1089 cgf.getContext(), recipeInsertLocation, varExpr,
1090 varRecipe.AllocaDecl,
1091 /*temporary=*/nullptr, clause.getReductionOp(),
1092 Decl::castToDeclContext(cgf.curFuncDecl), opInfo.origType,
1093 opInfo.bounds.size(), opInfo.boundTypes, opInfo.baseType,
1094 reductionOp, varRecipe.CombinerRecipes);
1095
1096 operation.addReduction(builder.getContext(), reductionOp, recipe);
1097 }
1098 }
1099 } else if constexpr (isCombinedType<OpTy>) {
1100 // Despite this being valid on ParallelOp or SerialOp, combined type
1101 // applies to the 'loop'.
1102 applyToLoopOp(clause);
1103 } else {
1104 llvm_unreachable("Unknown construct kind in VisitReductionClause");
1105 }
1106 }
1107};
1108
1109template <typename OpTy>
1110auto makeClauseEmitter(OpTy &op,
1111 mlir::OpBuilder::InsertPoint &recipeInsertLocation,
1113 CIRGen::CIRGenBuilderTy &builder,
1114 OpenACCDirectiveKind dirKind, SourceLocation dirLoc) {
1115 return OpenACCClauseCIREmitter<OpTy>(op, recipeInsertLocation, cgf, builder,
1116 dirKind, dirLoc);
1117}
1118} // namespace
1119
1120template <typename Op>
1121void CIRGenFunction::emitOpenACCClauses(
1122 Op &op, OpenACCDirectiveKind dirKind, SourceLocation dirLoc,
1124 mlir::OpBuilder::InsertionGuard guardCase(builder);
1125
1126 // Sets insertion point before the 'op', since every new expression needs to
1127 // be before the operation.
1128 builder.setInsertionPoint(op);
1129 makeClauseEmitter(op, lastRecipeLocation, *this, builder, dirKind, dirLoc)
1130 .emitClauses(clauses);
1131}
1132
1133#define EXPL_SPEC(N) \
1134 template void CIRGenFunction::emitOpenACCClauses<N>( \
1135 N &, OpenACCDirectiveKind, SourceLocation, \
1136 ArrayRef<const OpenACCClause *>);
1137EXPL_SPEC(mlir::acc::ParallelOp)
1138EXPL_SPEC(mlir::acc::SerialOp)
1139EXPL_SPEC(mlir::acc::KernelsOp)
1140EXPL_SPEC(mlir::acc::LoopOp)
1141EXPL_SPEC(mlir::acc::DataOp)
1142EXPL_SPEC(mlir::acc::InitOp)
1143EXPL_SPEC(mlir::acc::ShutdownOp)
1144EXPL_SPEC(mlir::acc::SetOp)
1145EXPL_SPEC(mlir::acc::WaitOp)
1146EXPL_SPEC(mlir::acc::HostDataOp)
1147EXPL_SPEC(mlir::acc::EnterDataOp)
1148EXPL_SPEC(mlir::acc::ExitDataOp)
1149EXPL_SPEC(mlir::acc::UpdateOp)
1150EXPL_SPEC(mlir::acc::AtomicReadOp)
1151EXPL_SPEC(mlir::acc::AtomicWriteOp)
1152EXPL_SPEC(mlir::acc::AtomicCaptureOp)
1153EXPL_SPEC(mlir::acc::AtomicUpdateOp)
1154#undef EXPL_SPEC
1155
1156template <typename ComputeOp, typename LoopOp>
1157void CIRGenFunction::emitOpenACCClauses(
1158 ComputeOp &op, LoopOp &loopOp, OpenACCDirectiveKind dirKind,
1160 static_assert(std::is_same_v<mlir::acc::LoopOp, LoopOp>);
1161
1162 CombinedConstructClauseInfo<ComputeOp> inf{op, loopOp};
1163 // We cannot set the insertion point here and do so in the emitter, but make
1164 // sure we reset it with the 'guard' anyway.
1165 mlir::OpBuilder::InsertionGuard guardCase(builder);
1166 makeClauseEmitter(inf, lastRecipeLocation, *this, builder, dirKind, dirLoc)
1167 .emitClauses(clauses);
1168}
1169
1170#define EXPL_SPEC(N) \
1171 template void CIRGenFunction::emitOpenACCClauses<N, mlir::acc::LoopOp>( \
1172 N &, mlir::acc::LoopOp &, OpenACCDirectiveKind, SourceLocation, \
1173 ArrayRef<const OpenACCClause *>);
1174
1175EXPL_SPEC(mlir::acc::ParallelOp)
1176EXPL_SPEC(mlir::acc::SerialOp)
1177EXPL_SPEC(mlir::acc::KernelsOp)
1178#undef EXPL_SPEC
#define EXPL_SPEC(N)
Defines the clang::Expr interface and subclasses for C++ expressions.
__device__ __2f16 float c
static DeclContext * castToDeclContext(const Decl *)
llvm::APSInt EvaluateKnownConstInt(const ASTContext &Ctx) const
EvaluateKnownConstInt - Call EvaluateAsRValue and return the folded integer.
StringRef getName() const
Return the actual identifier string.
const Expr * getConditionExpr() const
ArrayRef< Expr * > getVarList()
OpenACCModifierKind getModifierList() const
OpenACCModifierKind getModifierList() const
OpenACCModifierKind getModifierList() const
OpenACCModifierKind getModifierList() const
OpenACCDefaultClauseKind getDefaultClauseKind() const
ArrayRef< DeviceTypeArgument > getArchitectures() const
ArrayRef< OpenACCFirstPrivateRecipe > getInitRecipes()
unsigned getNumExprs() const
std::pair< OpenACCGangKind, const Expr * > getExpr(unsigned I) const
ArrayRef< Expr * > getIntExprs()
ArrayRef< OpenACCPrivateRecipe > getInitRecipes()
ArrayRef< OpenACCReductionRecipe > getRecipes()
OpenACCReductionOperator getReductionOp() const
const Expr * getConditionExpr() const
bool isConditionExprClause() const
ArrayRef< Expr * > getVarList()
ArrayRef< Expr * > getSizeExprs()
ArrayRef< Expr * > getQueueIdExprs()
Expr * getDevNumExpr() const
Encodes a location in the source.
SourceLocation getBeginLoc() const LLVM_READONLY
Definition Stmt.cpp:350
const internal::VariadicDynCastAllOfMatcher< Stmt, Expr > expr
Matches expressions.
unsigned kind
All of the diagnostics that can be emitted by the frontend.
The JSON file list parser is used to communicate input to InstallAPI.
OpenACCDirectiveKind
bool isa(CodeGen::Address addr)
Definition Address.h:330
OpenACCModifierKind
IdentifierLoc DeviceTypeArgument
const FunctionProtoType * T
bool isOpenACCModifierBitSet(OpenACCModifierKind List, OpenACCModifierKind Bit)
__DEVICE__ _Tp arg(const std::complex< _Tp > &__c)