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