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