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 = builder.create<mlir::UnrealizedConversionCastOp>(
100 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 = builder.create<mlir::arith::ConstantOp>(
111 loc, builder.getIntegerAttr(ty, value));
112
113 return constOp.getResult();
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 builder.create<BeforeOpTy>(opInfo.beginLoc, opInfo.varValue, structured,
221 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 = builder.create<AfterOpTy>(
234 opInfo.beginLoc, beforeOp.getResult(), structured, implicit,
235 opInfo.name, opInfo.bounds);
236 } else {
237 afterOp = builder.create<AfterOpTy>(
238 opInfo.beginLoc, beforeOp.getResult(), opInfo.varValue, structured,
239 implicit, 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 builder.create<BeforeOpTy>(opInfo.beginLoc, opInfo.varValue, structured,
262 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<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
557 mlir::acc::KernelsOp, mlir::acc::InitOp,
558 mlir::acc::ShutdownOp, mlir::acc::SetOp,
559 mlir::acc::DataOp, mlir::acc::WaitOp,
560 mlir::acc::HostDataOp, mlir::acc::EnterDataOp,
561 mlir::acc::ExitDataOp, mlir::acc::UpdateOp>) {
562 operation.getIfCondMutable().append(
563 createCondition(clause.getConditionExpr()));
564 } else if constexpr (isCombinedType<OpTy>) {
565 applyToComputeOp(clause);
566 } else {
567 llvm_unreachable("Unknown construct kind in VisitIfClause");
568 }
569 }
570
571 void VisitIfPresentClause(const OpenACCIfPresentClause &clause) {
572 if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp,
573 mlir::acc::UpdateOp>) {
574 operation.setIfPresent(true);
575 } else {
576 llvm_unreachable("unknown construct kind in VisitIfPresentClause");
577 }
578 }
579
580 void VisitDeviceNumClause(const OpenACCDeviceNumClause &clause) {
581 if constexpr (isOneOfTypes<OpTy, mlir::acc::InitOp, mlir::acc::ShutdownOp,
582 mlir::acc::SetOp>) {
583 operation.getDeviceNumMutable().append(emitIntExpr(clause.getIntExpr()));
584 } else {
585 llvm_unreachable(
586 "init, shutdown, set, are only valid device_num constructs");
587 }
588 }
589
590 void VisitNumGangsClause(const OpenACCNumGangsClause &clause) {
591 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
592 mlir::acc::KernelsOp>) {
593 llvm::SmallVector<mlir::Value> values;
594 for (const Expr *E : clause.getIntExprs())
595 values.push_back(emitIntExpr(E));
596
597 operation.addNumGangsOperands(builder.getContext(), values,
598 lastDeviceTypeValues);
599 } else if constexpr (isCombinedType<OpTy>) {
600 applyToComputeOp(clause);
601 } else {
602 llvm_unreachable("Unknown construct kind in VisitNumGangsClause");
603 }
604 }
605
606 void VisitWaitClause(const OpenACCWaitClause &clause) {
607 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
608 mlir::acc::KernelsOp, mlir::acc::DataOp,
609 mlir::acc::EnterDataOp, mlir::acc::ExitDataOp,
610 mlir::acc::UpdateOp>) {
611 if (!clause.hasExprs()) {
612 operation.addWaitOnly(builder.getContext(), lastDeviceTypeValues);
613 } else {
614 llvm::SmallVector<mlir::Value> values;
615 if (clause.hasDevNumExpr())
616 values.push_back(emitIntExpr(clause.getDevNumExpr()));
617 for (const Expr *E : clause.getQueueIdExprs())
618 values.push_back(emitIntExpr(E));
619 operation.addWaitOperands(builder.getContext(), clause.hasDevNumExpr(),
620 values, lastDeviceTypeValues);
621 }
622 } else if constexpr (isCombinedType<OpTy>) {
623 applyToComputeOp(clause);
624 } else {
625 // TODO: When we've implemented this for everything, switch this to an
626 // unreachable. update construct remains.
627 return clauseNotImplemented(clause);
628 }
629 }
630
631 void VisitDefaultAsyncClause(const OpenACCDefaultAsyncClause &clause) {
632 if constexpr (isOneOfTypes<OpTy, mlir::acc::SetOp>) {
633 operation.getDefaultAsyncMutable().append(
634 emitIntExpr(clause.getIntExpr()));
635 } else {
636 llvm_unreachable("set, is only valid device_num constructs");
637 }
638 }
639
640 void VisitSeqClause(const OpenACCSeqClause &clause) {
641 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
642 operation.addSeq(builder.getContext(), lastDeviceTypeValues);
643 } else if constexpr (isCombinedType<OpTy>) {
644 applyToLoopOp(clause);
645 } else {
646 // TODO: When we've implemented this for everything, switch this to an
647 // unreachable. Routine construct remains.
648 return clauseNotImplemented(clause);
649 }
650 }
651
652 void VisitAutoClause(const OpenACCAutoClause &clause) {
653 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
654 operation.addAuto(builder.getContext(), lastDeviceTypeValues);
655 } else if constexpr (isCombinedType<OpTy>) {
656 applyToLoopOp(clause);
657 } else {
658 // TODO: When we've implemented this for everything, switch this to an
659 // unreachable. Routine, construct remains.
660 return clauseNotImplemented(clause);
661 }
662 }
663
664 void VisitIndependentClause(const OpenACCIndependentClause &clause) {
665 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
666 operation.addIndependent(builder.getContext(), lastDeviceTypeValues);
667 } else if constexpr (isCombinedType<OpTy>) {
668 applyToLoopOp(clause);
669 } else {
670 // TODO: When we've implemented this for everything, switch this to an
671 // unreachable. Routine construct remains.
672 return clauseNotImplemented(clause);
673 }
674 }
675
676 void VisitCollapseClause(const OpenACCCollapseClause &clause) {
677 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
678 llvm::APInt value =
679 clause.getIntExpr()->EvaluateKnownConstInt(cgf.cgm.getASTContext());
680
681 value = value.sextOrTrunc(64);
682 operation.setCollapseForDeviceTypes(builder.getContext(),
683 lastDeviceTypeValues, value);
684 } else if constexpr (isCombinedType<OpTy>) {
685 applyToLoopOp(clause);
686 } else {
687 llvm_unreachable("Unknown construct kind in VisitCollapseClause");
688 }
689 }
690
691 void VisitTileClause(const OpenACCTileClause &clause) {
692 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
693 llvm::SmallVector<mlir::Value> values;
694
695 for (const Expr *e : clause.getSizeExprs()) {
696 mlir::Location exprLoc = cgf.cgm.getLoc(e->getBeginLoc());
697
698 // We represent the * as -1. Additionally, this is a constant, so we
699 // can always just emit it as 64 bits to avoid having to do any more
700 // work to determine signedness or size.
702 values.push_back(createConstantInt(exprLoc, 64, -1));
703 } else {
704 llvm::APInt curValue =
705 e->EvaluateKnownConstInt(cgf.cgm.getASTContext());
706 values.push_back(createConstantInt(
707 exprLoc, 64, curValue.sextOrTrunc(64).getSExtValue()));
708 }
709 }
710
711 operation.setTileForDeviceTypes(builder.getContext(),
712 lastDeviceTypeValues, values);
713 } else if constexpr (isCombinedType<OpTy>) {
714 applyToLoopOp(clause);
715 } else {
716 llvm_unreachable("Unknown construct kind in VisitTileClause");
717 }
718 }
719
720 void VisitWorkerClause(const OpenACCWorkerClause &clause) {
721 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
722 if (clause.hasIntExpr())
723 operation.addWorkerNumOperand(builder.getContext(),
724 emitIntExpr(clause.getIntExpr()),
725 lastDeviceTypeValues);
726 else
727 operation.addEmptyWorker(builder.getContext(), lastDeviceTypeValues);
728
729 } else if constexpr (isCombinedType<OpTy>) {
730 applyToLoopOp(clause);
731 } else {
732 // TODO: When we've implemented this for everything, switch this to an
733 // unreachable. Combined constructs remain.
734 return clauseNotImplemented(clause);
735 }
736 }
737
738 void VisitVectorClause(const OpenACCVectorClause &clause) {
739 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
740 if (clause.hasIntExpr())
741 operation.addVectorOperand(builder.getContext(),
742 emitIntExpr(clause.getIntExpr()),
743 lastDeviceTypeValues);
744 else
745 operation.addEmptyVector(builder.getContext(), lastDeviceTypeValues);
746
747 } else if constexpr (isCombinedType<OpTy>) {
748 applyToLoopOp(clause);
749 } else {
750 // TODO: When we've implemented this for everything, switch this to an
751 // unreachable. Combined constructs remain.
752 return clauseNotImplemented(clause);
753 }
754 }
755
756 void VisitGangClause(const OpenACCGangClause &clause) {
757 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
758 if (clause.getNumExprs() == 0) {
759 operation.addEmptyGang(builder.getContext(), lastDeviceTypeValues);
760 } else {
761 llvm::SmallVector<mlir::Value> values;
762 llvm::SmallVector<mlir::acc::GangArgType> argTypes;
763 for (unsigned i : llvm::index_range(0u, clause.getNumExprs())) {
764 auto [kind, expr] = clause.getExpr(i);
765 mlir::Location exprLoc = cgf.cgm.getLoc(expr->getBeginLoc());
766 argTypes.push_back(decodeGangType(kind));
767 if (kind == OpenACCGangKind::Dim) {
768 llvm::APInt curValue =
769 expr->EvaluateKnownConstInt(cgf.cgm.getASTContext());
770 // The value is 1, 2, or 3, but the type isn't necessarily smaller
771 // than 64.
772 curValue = curValue.sextOrTrunc(64);
773 values.push_back(
774 createConstantInt(exprLoc, 64, curValue.getSExtValue()));
776 values.push_back(createConstantInt(exprLoc, 64, -1));
777 } else {
778 values.push_back(emitIntExpr(expr));
779 }
780 }
781
782 operation.addGangOperands(builder.getContext(), lastDeviceTypeValues,
783 argTypes, values);
784 }
785 } else if constexpr (isCombinedType<OpTy>) {
786 applyToLoopOp(clause);
787 } else {
788 llvm_unreachable("Unknown construct kind in VisitGangClause");
789 }
790 }
791
792 void VisitCopyClause(const OpenACCCopyClause &clause) {
793 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
794 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
795 for (const Expr *var : clause.getVarList())
796 addDataOperand<mlir::acc::CopyinOp, mlir::acc::CopyoutOp>(
797 var, mlir::acc::DataClause::acc_copy, clause.getModifierList(),
798 /*structured=*/true,
799 /*implicit=*/false);
800 } else if constexpr (isCombinedType<OpTy>) {
801 applyToComputeOp(clause);
802 } else {
803 // TODO: When we've implemented this for everything, switch this to an
804 // unreachable. declare construct remains.
805 return clauseNotImplemented(clause);
806 }
807 }
808
809 void VisitCopyInClause(const OpenACCCopyInClause &clause) {
810 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
811 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
812 for (const Expr *var : clause.getVarList())
813 addDataOperand<mlir::acc::CopyinOp, mlir::acc::DeleteOp>(
814 var, mlir::acc::DataClause::acc_copyin, clause.getModifierList(),
815 /*structured=*/true,
816 /*implicit=*/false);
817 } else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
818 for (const Expr *var : clause.getVarList())
819 addDataOperand<mlir::acc::CopyinOp>(
820 var, mlir::acc::DataClause::acc_copyin, clause.getModifierList(),
821 /*structured=*/false, /*implicit=*/false);
822 } else if constexpr (isCombinedType<OpTy>) {
823 applyToComputeOp(clause);
824 } else {
825 // TODO: When we've implemented this for everything, switch this to an
826 // unreachable. declare construct remains.
827 return clauseNotImplemented(clause);
828 }
829 }
830
831 void VisitCopyOutClause(const OpenACCCopyOutClause &clause) {
832 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
833 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
834 for (const Expr *var : clause.getVarList())
835 addDataOperand<mlir::acc::CreateOp, mlir::acc::CopyoutOp>(
836 var, mlir::acc::DataClause::acc_copyout, clause.getModifierList(),
837 /*structured=*/true,
838 /*implicit=*/false);
839 } else if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
840 for (const Expr *var : clause.getVarList())
841 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::CopyoutOp>(
842 var, mlir::acc::DataClause::acc_copyout, clause.getModifierList(),
843 /*structured=*/false,
844 /*implicit=*/false);
845 } else if constexpr (isCombinedType<OpTy>) {
846 applyToComputeOp(clause);
847 } else {
848 // TODO: When we've implemented this for everything, switch this to an
849 // unreachable. declare construct remains.
850 return clauseNotImplemented(clause);
851 }
852 }
853
854 void VisitCreateClause(const OpenACCCreateClause &clause) {
855 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
856 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
857 for (const Expr *var : clause.getVarList())
858 addDataOperand<mlir::acc::CreateOp, mlir::acc::DeleteOp>(
859 var, mlir::acc::DataClause::acc_create, clause.getModifierList(),
860 /*structured=*/true,
861 /*implicit=*/false);
862 } else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
863 for (const Expr *var : clause.getVarList())
864 addDataOperand<mlir::acc::CreateOp>(
865 var, mlir::acc::DataClause::acc_create, clause.getModifierList(),
866 /*structured=*/false, /*implicit=*/false);
867 } else if constexpr (isCombinedType<OpTy>) {
868 applyToComputeOp(clause);
869 } else {
870 // TODO: When we've implemented this for everything, switch this to an
871 // unreachable. declare construct remains.
872 return clauseNotImplemented(clause);
873 }
874 }
875
876 void VisitDeleteClause(const OpenACCDeleteClause &clause) {
877 if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
878 for (const Expr *var : clause.getVarList())
879 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::DeleteOp>(
880 var, mlir::acc::DataClause::acc_delete, {},
881 /*structured=*/false,
882 /*implicit=*/false);
883 } else {
884 llvm_unreachable("Unknown construct kind in VisitDeleteClause");
885 }
886 }
887
888 void VisitDetachClause(const OpenACCDetachClause &clause) {
889 if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
890 for (const Expr *var : clause.getVarList())
891 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::DetachOp>(
892 var, mlir::acc::DataClause::acc_detach, {},
893 /*structured=*/false,
894 /*implicit=*/false);
895 } else {
896 llvm_unreachable("Unknown construct kind in VisitDetachClause");
897 }
898 }
899
900 void VisitFinalizeClause(const OpenACCFinalizeClause &clause) {
901 if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
902 operation.setFinalize(true);
903 } else {
904 llvm_unreachable("Unknown construct kind in VisitFinalizeClause");
905 }
906 }
907
908 void VisitUseDeviceClause(const OpenACCUseDeviceClause &clause) {
909 if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp>) {
910 for (const Expr *var : clause.getVarList())
911 addDataOperand<mlir::acc::UseDeviceOp>(
912 var, mlir::acc::DataClause::acc_use_device, {}, /*structured=*/true,
913 /*implicit=*/false);
914 } else {
915 llvm_unreachable("Unknown construct kind in VisitUseDeviceClause");
916 }
917 }
918
919 void VisitDevicePtrClause(const OpenACCDevicePtrClause &clause) {
920 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
921 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
922 for (const Expr *var : clause.getVarList())
923 addDataOperand<mlir::acc::DevicePtrOp>(
924 var, mlir::acc::DataClause::acc_deviceptr, {},
925 /*structured=*/true,
926 /*implicit=*/false);
927 } else if constexpr (isCombinedType<OpTy>) {
928 applyToComputeOp(clause);
929 } else {
930 // TODO: When we've implemented this for everything, switch this to an
931 // unreachable. declare remains.
932 return clauseNotImplemented(clause);
933 }
934 }
935
936 void VisitNoCreateClause(const OpenACCNoCreateClause &clause) {
937 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
938 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
939 for (const Expr *var : clause.getVarList())
940 addDataOperand<mlir::acc::NoCreateOp, mlir::acc::DeleteOp>(
941 var, mlir::acc::DataClause::acc_no_create, {}, /*structured=*/true,
942 /*implicit=*/false);
943 } else if constexpr (isCombinedType<OpTy>) {
944 applyToComputeOp(clause);
945 } else {
946 llvm_unreachable("Unknown construct kind in VisitNoCreateClause");
947 }
948 }
949
950 void VisitPresentClause(const OpenACCPresentClause &clause) {
951 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
952 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
953 for (const Expr *var : clause.getVarList())
954 addDataOperand<mlir::acc::PresentOp, mlir::acc::DeleteOp>(
955 var, mlir::acc::DataClause::acc_present, {}, /*structured=*/true,
956 /*implicit=*/false);
957 } else if constexpr (isCombinedType<OpTy>) {
958 applyToComputeOp(clause);
959 } else {
960 // TODO: When we've implemented this for everything, switch this to an
961 // unreachable. declare remains.
962 return clauseNotImplemented(clause);
963 }
964 }
965
966 void VisitAttachClause(const OpenACCAttachClause &clause) {
967 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
968 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
969 for (const Expr *var : clause.getVarList())
970 addDataOperand<mlir::acc::AttachOp, mlir::acc::DetachOp>(
971 var, mlir::acc::DataClause::acc_attach, {}, /*structured=*/true,
972 /*implicit=*/false);
973 } else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
974 for (const Expr *var : clause.getVarList())
975 addDataOperand<mlir::acc::AttachOp>(
976 var, mlir::acc::DataClause::acc_attach, {},
977 /*structured=*/false, /*implicit=*/false);
978 } else if constexpr (isCombinedType<OpTy>) {
979 applyToComputeOp(clause);
980 } else {
981 llvm_unreachable("Unknown construct kind in VisitAttachClause");
982 }
983 }
984
985 void VisitPrivateClause(const OpenACCPrivateClause &clause) {
986 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
987 mlir::acc::LoopOp>) {
988 for (const auto [varExpr, varRecipe] :
989 llvm::zip_equal(clause.getVarList(), clause.getInitRecipes())) {
990 CIRGenFunction::OpenACCDataOperandInfo opInfo =
991 cgf.getOpenACCDataOperandInfo(varExpr);
992 auto privateOp = mlir::acc::PrivateOp::create(
993 builder, opInfo.beginLoc, opInfo.varValue, /*structured=*/true,
994 /*implicit=*/false, opInfo.name, opInfo.bounds);
995 privateOp.setDataClause(mlir::acc::DataClause::acc_private);
996
997 {
998 mlir::OpBuilder::InsertionGuard guardCase(builder);
999
1000 auto recipe =
1001 OpenACCRecipeBuilder<mlir::acc::PrivateRecipeOp>(cgf, builder)
1002 .getOrCreateRecipe(
1003 cgf.getContext(), recipeInsertLocation, varExpr,
1004 varRecipe.AllocaDecl, varRecipe.InitExpr,
1005 /*temporary=*/nullptr, OpenACCReductionOperator::Invalid,
1006 Decl::castToDeclContext(cgf.curFuncDecl), opInfo.origType,
1007 opInfo.bounds.size(), opInfo.boundTypes, opInfo.baseType,
1008 privateOp.getResult());
1009 // TODO: OpenACC: The dialect is going to change in the near future to
1010 // have these be on a different operation, so when that changes, we
1011 // probably need to change these here.
1012 operation.addPrivatization(builder.getContext(), privateOp, recipe);
1013 }
1014 }
1015 } else if constexpr (isCombinedType<OpTy>) {
1016 // Despite this being valid on ParallelOp or SerialOp, combined type
1017 // applies to the 'loop'.
1018 applyToLoopOp(clause);
1019 } else {
1020 llvm_unreachable("Unknown construct kind in VisitPrivateClause");
1021 }
1022 }
1023
1024 void VisitFirstPrivateClause(const OpenACCFirstPrivateClause &clause) {
1025 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
1026 mlir::acc::SerialOp>) {
1027 for (const auto [varExpr, varRecipe] :
1028 llvm::zip_equal(clause.getVarList(), clause.getInitRecipes())) {
1029 CIRGenFunction::OpenACCDataOperandInfo opInfo =
1030 cgf.getOpenACCDataOperandInfo(varExpr);
1031 auto firstPrivateOp = mlir::acc::FirstprivateOp::create(
1032 builder, opInfo.beginLoc, opInfo.varValue, /*structured=*/true,
1033 /*implicit=*/false, opInfo.name, opInfo.bounds);
1034
1035 firstPrivateOp.setDataClause(mlir::acc::DataClause::acc_firstprivate);
1036
1037 {
1038 mlir::OpBuilder::InsertionGuard guardCase(builder);
1039 // TODO: OpenACC: At the moment this is a bit of a hacky way of doing
1040 // this, and won't work when we get to bounds/etc. Do this for now to
1041 // limit the scope of this refactor.
1042 VarDecl *allocaDecl = varRecipe.AllocaDecl;
1043 allocaDecl->setInit(varRecipe.InitExpr);
1044 allocaDecl->setInitStyle(VarDecl::CallInit);
1045
1046 auto recipe =
1047 OpenACCRecipeBuilder<mlir::acc::FirstprivateRecipeOp>(cgf,
1048 builder)
1049 .getOrCreateRecipe(
1050 cgf.getContext(), recipeInsertLocation, varExpr,
1051 varRecipe.AllocaDecl, varRecipe.InitExpr,
1052 varRecipe.InitFromTemporary,
1053 OpenACCReductionOperator::Invalid,
1054 Decl::castToDeclContext(cgf.curFuncDecl), opInfo.origType,
1055 opInfo.bounds.size(), opInfo.boundTypes, opInfo.baseType,
1056 firstPrivateOp.getResult());
1057
1058 // TODO: OpenACC: The dialect is going to change in the near future to
1059 // have these be on a different operation, so when that changes, we
1060 // probably need to change these here.
1061 operation.addFirstPrivatization(builder.getContext(), firstPrivateOp,
1062 recipe);
1063 }
1064 }
1065 } else if constexpr (isCombinedType<OpTy>) {
1066 // Unlike 'private', 'firstprivate' applies to the compute op, not the
1067 // loop op.
1068 applyToComputeOp(clause);
1069 } else {
1070 llvm_unreachable("Unknown construct kind in VisitFirstPrivateClause");
1071 }
1072 }
1073
1074 void VisitReductionClause(const OpenACCReductionClause &clause) {
1075 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
1076 mlir::acc::LoopOp>) {
1077 for (const auto [varExpr, varRecipe] :
1078 llvm::zip_equal(clause.getVarList(), clause.getRecipes())) {
1079 CIRGenFunction::OpenACCDataOperandInfo opInfo =
1080 cgf.getOpenACCDataOperandInfo(varExpr);
1081
1082 auto reductionOp = mlir::acc::ReductionOp::create(
1083 builder, opInfo.beginLoc, opInfo.varValue, /*structured=*/true,
1084 /*implicit=*/false, opInfo.name, opInfo.bounds);
1085 reductionOp.setDataClause(mlir::acc::DataClause::acc_reduction);
1086
1087 {
1088 mlir::OpBuilder::InsertionGuard guardCase(builder);
1089 // TODO: OpenACC: At the moment this is a bit of a hacky way of doing
1090 // this, and won't work when we get to bounds/etc. Do this for now to
1091 // limit the scope of this refactor.
1092 VarDecl *allocaDecl = varRecipe.AllocaDecl;
1093 allocaDecl->setInit(varRecipe.InitExpr);
1094 allocaDecl->setInitStyle(VarDecl::CallInit);
1095
1096 auto recipe =
1097 OpenACCRecipeBuilder<mlir::acc::ReductionRecipeOp>(cgf, builder)
1098 .getOrCreateRecipe(
1099 cgf.getContext(), recipeInsertLocation, varExpr,
1100 varRecipe.AllocaDecl, varRecipe.InitExpr,
1101 /*temporary=*/nullptr, clause.getReductionOp(),
1102 Decl::castToDeclContext(cgf.curFuncDecl), opInfo.origType,
1103 opInfo.bounds.size(), opInfo.boundTypes, opInfo.baseType,
1104 reductionOp.getResult());
1105
1106 operation.addReduction(builder.getContext(), reductionOp, recipe);
1107 }
1108 }
1109 } else if constexpr (isCombinedType<OpTy>) {
1110 // Despite this being valid on ParallelOp or SerialOp, combined type
1111 // applies to the 'loop'.
1112 applyToLoopOp(clause);
1113 } else {
1114 llvm_unreachable("Unknown construct kind in VisitReductionClause");
1115 }
1116 }
1117};
1118
1119template <typename OpTy>
1120auto makeClauseEmitter(OpTy &op,
1121 mlir::OpBuilder::InsertPoint &recipeInsertLocation,
1123 CIRGen::CIRGenBuilderTy &builder,
1124 OpenACCDirectiveKind dirKind, SourceLocation dirLoc) {
1125 return OpenACCClauseCIREmitter<OpTy>(op, recipeInsertLocation, cgf, builder,
1126 dirKind, dirLoc);
1127}
1128} // namespace
1129
1130template <typename Op>
1131void CIRGenFunction::emitOpenACCClauses(
1132 Op &op, OpenACCDirectiveKind dirKind, SourceLocation dirLoc,
1134 mlir::OpBuilder::InsertionGuard guardCase(builder);
1135
1136 // Sets insertion point before the 'op', since every new expression needs to
1137 // be before the operation.
1138 builder.setInsertionPoint(op);
1139 makeClauseEmitter(op, lastRecipeLocation, *this, builder, dirKind, dirLoc)
1140 .emitClauses(clauses);
1141}
1142
1143#define EXPL_SPEC(N) \
1144 template void CIRGenFunction::emitOpenACCClauses<N>( \
1145 N &, OpenACCDirectiveKind, SourceLocation, \
1146 ArrayRef<const OpenACCClause *>);
1147EXPL_SPEC(mlir::acc::ParallelOp)
1148EXPL_SPEC(mlir::acc::SerialOp)
1149EXPL_SPEC(mlir::acc::KernelsOp)
1150EXPL_SPEC(mlir::acc::LoopOp)
1151EXPL_SPEC(mlir::acc::DataOp)
1152EXPL_SPEC(mlir::acc::InitOp)
1153EXPL_SPEC(mlir::acc::ShutdownOp)
1154EXPL_SPEC(mlir::acc::SetOp)
1155EXPL_SPEC(mlir::acc::WaitOp)
1156EXPL_SPEC(mlir::acc::HostDataOp)
1157EXPL_SPEC(mlir::acc::EnterDataOp)
1158EXPL_SPEC(mlir::acc::ExitDataOp)
1159EXPL_SPEC(mlir::acc::UpdateOp)
1160#undef EXPL_SPEC
1161
1162template <typename ComputeOp, typename LoopOp>
1163void CIRGenFunction::emitOpenACCClauses(
1164 ComputeOp &op, LoopOp &loopOp, OpenACCDirectiveKind dirKind,
1166 static_assert(std::is_same_v<mlir::acc::LoopOp, LoopOp>);
1167
1168 CombinedConstructClauseInfo<ComputeOp> inf{op, loopOp};
1169 // We cannot set the insertion point here and do so in the emitter, but make
1170 // sure we reset it with the 'guard' anyway.
1171 mlir::OpBuilder::InsertionGuard guardCase(builder);
1172 makeClauseEmitter(inf, lastRecipeLocation, *this, builder, dirKind, dirLoc)
1173 .emitClauses(clauses);
1174}
1175
1176#define EXPL_SPEC(N) \
1177 template void CIRGenFunction::emitOpenACCClauses<N, mlir::acc::LoopOp>( \
1178 N &, mlir::acc::LoopOp &, OpenACCDirectiveKind, SourceLocation, \
1179 ArrayRef<const OpenACCClause *>);
1180
1181EXPL_SPEC(mlir::acc::ParallelOp)
1182EXPL_SPEC(mlir::acc::SerialOp)
1183EXPL_SPEC(mlir::acc::KernelsOp)
1184#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:346
void setInitStyle(InitializationStyle Style)
Definition Decl.h:1451
@ CallInit
Call-style initialization (C++98)
Definition Decl.h:933
void setInit(Expr *I)
Definition Decl.cpp:2477
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)