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"
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.
__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.
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
__DEVICE__ _Tp arg(const std::complex< _Tp > &__c)