22#include "mlir/Dialect/Arith/IR/Arith.h"
23#include "mlir/Dialect/OpenACC/OpenACC.h"
24#include "llvm/ADT/TypeSwitch.h"
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>;
41template <
typename CompOpTy>
struct CombinedConstructClauseInfo {
42 using ComputeOpTy = CompOpTy;
43 ComputeOpTy computeOp;
44 mlir::acc::LoopOp loopOp;
46template <
typename ToTest>
constexpr bool isCombinedType =
false;
48constexpr bool isCombinedType<CombinedConstructClauseInfo<T>> =
true;
50template <
typename OpTy>
51class OpenACCClauseCIREmitter final
54 template <
typename FriendOpTy>
friend class OpenACCClauseCIREmitter;
57 mlir::OpBuilder::InsertPoint &recipeInsertLocation;
58 CIRGen::CIRGenFunction &cgf;
59 CIRGen::CIRGenBuilderTy &builder;
65 llvm::SmallVector<mlir::acc::DeviceType> lastDeviceTypeValues;
68 bool hasAsyncClause =
false;
70 llvm::SmallVector<mlir::Operation *> dataOperands;
72 void setLastDeviceTypeClause(
const OpenACCDeviceTypeClause &clause) {
73 lastDeviceTypeValues.clear();
79 mlir::Value emitIntExpr(
const Expr *intExpr) {
80 return cgf.emitOpenACCIntExpr(intExpr);
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(), 1,
92 mlir::IntegerType::SignednessSemantics::Signless);
93 auto conversionOp = mlir::UnrealizedConversionCastOp::create(
94 builder, exprLoc, targetType, condition);
95 return conversionOp.getResult(0);
98 mlir::Value createConstantInt(mlir::Location loc,
unsigned width,
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));
110 mlir::Value createConstantInt(SourceLocation loc,
unsigned width,
112 return createConstantInt(cgf.cgm.getLoc(loc), width, value);
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;
124 llvm_unreachable(
"unknown gang kind");
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);
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};
146 computeEmitter.lastDeviceTypeValues = lastDeviceTypeValues;
151 if (!dataOperands.empty())
152 computeEmitter.dataOperands.push_back(dataOperands.front());
153 computeEmitter.Visit(&
c);
158 dataOperands.append(computeEmitter.dataOperands);
161 template <
typename BeforeOpTy,
typename AfterOpTy>
162 void addDataOperand(
const Expr *varOperand, mlir::acc::DataClause dataClause,
165 CIRGenFunction::OpenACCDataOperandInfo opInfo =
166 cgf.getOpenACCDataOperandInfo(varOperand);
170 structured, implicit, opInfo.
name, opInfo.
bounds);
171 operation.getDataClauseOperandsMutable().append(beforeOp.getResult());
175 mlir::OpBuilder::InsertionGuard guardCase(builder);
176 builder.setInsertionPointAfter(operation);
178 if constexpr (std::is_same_v<AfterOpTy, mlir::acc::DeleteOp> ||
179 std::is_same_v<AfterOpTy, mlir::acc::DetachOp>) {
183 AfterOpTy::create(builder, opInfo.
beginLoc, beforeOp, structured,
186 afterOp = AfterOpTy::create(builder, opInfo.
beginLoc, beforeOp,
187 opInfo.
varValue, structured, implicit,
193 beforeOp.setDataClause(dataClause);
194 afterOp.setDataClause(dataClause);
199 dataOperands.push_back(beforeOp.getOperation());
200 dataOperands.push_back(afterOp.getOperation());
203 template <
typename BeforeOpTy>
204 void addDataOperand(
const Expr *varOperand, mlir::acc::DataClause dataClause,
207 CIRGenFunction::OpenACCDataOperandInfo opInfo =
208 cgf.getOpenACCDataOperandInfo(varOperand);
211 structured, implicit, opInfo.
name, opInfo.
bounds);
212 operation.getDataClauseOperandsMutable().append(beforeOp.getResult());
215 beforeOp.setDataClause(dataClause);
219 dataOperands.push_back(beforeOp.getOperation());
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{};
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();
245 llvm_unreachable(
"getting asyncOnly when clause not valid on operation?");
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{};
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();
272 "getting asyncOperandsDeviceType when clause not valid on operation?");
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();
292 "getting asyncOperandsDeviceType when clause not valid on operation?");
298 void updateDataOperandAsyncValues() {
299 if (!hasAsyncClause || dataOperands.empty())
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());
309 .
Default([&](mlir::Operation *) {
310 llvm_unreachable(
"Not a data operation?");
316 OpenACCClauseCIREmitter(OpTy &operation,
317 mlir::OpBuilder::InsertPoint &recipeInsertLocation,
318 CIRGen::CIRGenFunction &cgf,
319 CIRGen::CIRGenBuilderTy &builder,
321 : operation(operation), recipeInsertLocation(recipeInsertLocation),
322 cgf(cgf), builder(builder), dirKind(dirKind) {}
324 void VisitClause(
const OpenACCClause &clause) {
325 llvm_unreachable(
"Unknown/unhandled clause kind");
331 void emitClauses(ArrayRef<const OpenACCClause *> clauses) {
332 this->VisitClauseList(clauses);
333 updateDataOperandAsyncValues();
336 void VisitDefaultClause(
const OpenACCDefaultClause &clause) {
339 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
340 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
342 case OpenACCDefaultClauseKind::None:
343 operation.setDefaultAttr(mlir::acc::ClauseDefaultValue::None);
345 case OpenACCDefaultClauseKind::Present:
346 operation.setDefaultAttr(mlir::acc::ClauseDefaultValue::Present);
348 case OpenACCDefaultClauseKind::Invalid:
351 }
else if constexpr (isCombinedType<OpTy>) {
352 applyToComputeOp(clause);
354 llvm_unreachable(
"Unknown construct kind in VisitDefaultClause");
358 void VisitDeviceTypeClause(
const OpenACCDeviceTypeClause &clause) {
359 setLastDeviceTypeClause(clause);
361 if constexpr (isOneOfTypes<OpTy, mlir::acc::InitOp,
362 mlir::acc::ShutdownOp>) {
364 operation.addDeviceType(builder.getContext(),
366 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::SetOp>) {
367 assert(!operation.getDeviceTypeAttr() &&
"already have device-type?");
371 operation.setDeviceType(
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>) {
380 }
else if constexpr (isCombinedType<OpTy>) {
384 llvm_unreachable(
"Unknown construct kind in VisitDeviceTypeClause");
388 void VisitNumWorkersClause(
const OpenACCNumWorkersClause &clause) {
389 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
390 mlir::acc::KernelsOp>) {
391 operation.addNumWorkersOperand(builder.getContext(),
393 lastDeviceTypeValues);
394 }
else if constexpr (isCombinedType<OpTy>) {
395 applyToComputeOp(clause);
397 llvm_unreachable(
"Unknown construct kind in VisitNumGangsClause");
401 void VisitVectorLengthClause(
const OpenACCVectorLengthClause &clause) {
402 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
403 mlir::acc::KernelsOp>) {
404 operation.addVectorLengthOperand(builder.getContext(),
406 lastDeviceTypeValues);
407 }
else if constexpr (isCombinedType<OpTy>) {
408 applyToComputeOp(clause);
410 llvm_unreachable(
"Unknown construct kind in VisitVectorLengthClause");
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>) {
421 operation.addAsyncOnly(builder.getContext(), lastDeviceTypeValues);
429 mlir::OpBuilder::InsertionGuard guardCase(builder);
430 if (!dataOperands.empty())
431 builder.setInsertionPoint(dataOperands.front());
434 operation.addAsyncOperand(builder.getContext(), intExpr,
435 lastDeviceTypeValues);
437 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::WaitOp>) {
441 operation.setAsync(
true);
443 operation.getAsyncOperandMutable().append(
445 }
else if constexpr (isCombinedType<OpTy>) {
446 applyToComputeOp(clause);
448 llvm_unreachable(
"Unknown construct kind in VisitAsyncClause");
452 void VisitSelfClause(
const OpenACCSelfClause &clause) {
453 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
454 mlir::acc::KernelsOp>) {
456 operation.setSelfAttr(
true);
459 operation.getSelfCondMutable().append(
462 llvm_unreachable(
"var-list version of self shouldn't get here");
464 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
466 "var-list version of self required for update");
468 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::UpdateHostOp>(
469 var, mlir::acc::DataClause::acc_update_self, {},
471 }
else if constexpr (isCombinedType<OpTy>) {
472 applyToComputeOp(clause);
474 llvm_unreachable(
"Unknown construct kind in VisitSelfClause");
478 void VisitHostClause(
const OpenACCHostClause &clause) {
479 if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
481 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::UpdateHostOp>(
482 var, mlir::acc::DataClause::acc_update_host, {},
485 llvm_unreachable(
"Unknown construct kind in VisitHostClause");
489 void VisitDeviceClause(
const OpenACCDeviceClause &clause) {
490 if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
492 addDataOperand<mlir::acc::UpdateDeviceOp>(
493 var, mlir::acc::DataClause::acc_update_device, {},
496 llvm_unreachable(
"Unknown construct kind in VisitDeviceClause");
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(
512 }
else if constexpr (isCombinedType<OpTy>) {
513 applyToComputeOp(clause);
515 llvm_unreachable(
"Unknown construct kind in VisitIfClause");
519 void VisitIfPresentClause(
const OpenACCIfPresentClause &clause) {
520 if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp,
521 mlir::acc::UpdateOp>) {
522 operation.setIfPresent(
true);
524 llvm_unreachable(
"unknown construct kind in VisitIfPresentClause");
528 void VisitDeviceNumClause(
const OpenACCDeviceNumClause &clause) {
529 if constexpr (isOneOfTypes<OpTy, mlir::acc::InitOp, mlir::acc::ShutdownOp,
531 operation.getDeviceNumMutable().append(emitIntExpr(clause.
getIntExpr()));
534 "init, shutdown, set, are only valid device_num constructs");
538 void VisitNumGangsClause(
const OpenACCNumGangsClause &clause) {
539 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
540 mlir::acc::KernelsOp>) {
541 llvm::SmallVector<mlir::Value> values;
543 values.push_back(emitIntExpr(E));
545 operation.addNumGangsOperands(builder.getContext(), values,
546 lastDeviceTypeValues);
547 }
else if constexpr (isCombinedType<OpTy>) {
548 applyToComputeOp(clause);
550 llvm_unreachable(
"Unknown construct kind in VisitNumGangsClause");
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>) {
560 operation.addWaitOnly(builder.getContext(), lastDeviceTypeValues);
562 llvm::SmallVector<mlir::Value> values;
566 values.push_back(emitIntExpr(E));
567 operation.addWaitOperands(builder.getContext(), clause.
hasDevNumExpr(),
568 values, lastDeviceTypeValues);
570 }
else if constexpr (isCombinedType<OpTy>) {
571 applyToComputeOp(clause);
575 llvm_unreachable(
"Unknown construct kind in VisitWaitClause");
579 void VisitDefaultAsyncClause(
const OpenACCDefaultAsyncClause &clause) {
580 if constexpr (isOneOfTypes<OpTy, mlir::acc::SetOp>) {
581 operation.getDefaultAsyncMutable().append(
584 llvm_unreachable(
"set, is only valid device_num constructs");
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);
594 llvm_unreachable(
"Unknown construct kind in VisitSeqClause");
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);
604 llvm_unreachable(
"Unknown construct kind in VisitAutoClause");
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);
614 llvm_unreachable(
"Unknown construct kind in VisitIndependentClause");
618 void VisitCollapseClause(
const OpenACCCollapseClause &clause) {
619 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
623 value = value.sextOrTrunc(64);
624 operation.setCollapseForDeviceTypes(builder.getContext(),
625 lastDeviceTypeValues, value);
626 }
else if constexpr (isCombinedType<OpTy>) {
627 applyToLoopOp(clause);
629 llvm_unreachable(
"Unknown construct kind in VisitCollapseClause");
633 void VisitTileClause(
const OpenACCTileClause &clause) {
634 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
635 llvm::SmallVector<mlir::Value> values;
638 mlir::Location exprLoc = cgf.cgm.getLoc(e->getBeginLoc());
644 values.push_back(createConstantInt(exprLoc, 64, -1));
646 llvm::APInt curValue =
647 e->EvaluateKnownConstInt(cgf.cgm.getASTContext());
648 values.push_back(createConstantInt(
649 exprLoc, 64, curValue.sextOrTrunc(64).getSExtValue()));
653 operation.setTileForDeviceTypes(builder.getContext(),
654 lastDeviceTypeValues, values);
655 }
else if constexpr (isCombinedType<OpTy>) {
656 applyToLoopOp(clause);
658 llvm_unreachable(
"Unknown construct kind in VisitTileClause");
662 void VisitWorkerClause(
const OpenACCWorkerClause &clause) {
663 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
665 operation.addWorkerNumOperand(builder.getContext(),
667 lastDeviceTypeValues);
669 operation.addEmptyWorker(builder.getContext(), lastDeviceTypeValues);
671 }
else if constexpr (isCombinedType<OpTy>) {
672 applyToLoopOp(clause);
674 llvm_unreachable(
"Unknown construct kind in VisitWorkerClause");
678 void VisitVectorClause(
const OpenACCVectorClause &clause) {
679 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
681 operation.addVectorOperand(builder.getContext(),
683 lastDeviceTypeValues);
685 operation.addEmptyVector(builder.getContext(), lastDeviceTypeValues);
687 }
else if constexpr (isCombinedType<OpTy>) {
688 applyToLoopOp(clause);
690 llvm_unreachable(
"Unknown construct kind in VisitVectorClause");
694 void VisitGangClause(
const OpenACCGangClause &clause) {
695 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
697 operation.addEmptyGang(builder.getContext(), lastDeviceTypeValues);
699 llvm::SmallVector<mlir::Value> values;
700 llvm::SmallVector<mlir::acc::GangArgType> argTypes;
701 for (
unsigned i : llvm::index_range(0u, clause.
getNumExprs())) {
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());
710 curValue = curValue.sextOrTrunc(64);
712 createConstantInt(exprLoc, 64, curValue.getSExtValue()));
714 values.push_back(createConstantInt(exprLoc, 64, -1));
716 values.push_back(emitIntExpr(
expr));
720 operation.addGangOperands(builder.getContext(), lastDeviceTypeValues,
723 }
else if constexpr (isCombinedType<OpTy>) {
724 applyToLoopOp(clause);
726 llvm_unreachable(
"Unknown construct kind in VisitGangClause");
730 void VisitCopyClause(
const OpenACCCopyClause &clause) {
731 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
732 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
734 addDataOperand<mlir::acc::CopyinOp, mlir::acc::CopyoutOp>(
738 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::DeclareEnterOp>) {
740 addDataOperand<mlir::acc::CopyinOp>(
744 }
else if constexpr (isCombinedType<OpTy>) {
745 applyToComputeOp(clause);
747 llvm_unreachable(
"Unknown construct kind in VisitCopyClause");
751 void VisitCopyInClause(
const OpenACCCopyInClause &clause) {
752 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
753 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
755 addDataOperand<mlir::acc::CopyinOp, mlir::acc::DeleteOp>(
759 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
761 addDataOperand<mlir::acc::CopyinOp>(
764 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::DeclareEnterOp>) {
766 addDataOperand<mlir::acc::CopyinOp>(
770 }
else if constexpr (isCombinedType<OpTy>) {
771 applyToComputeOp(clause);
773 llvm_unreachable(
"Unknown construct kind in VisitCopyInClause");
777 void VisitCopyOutClause(
const OpenACCCopyOutClause &clause) {
778 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
779 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
781 addDataOperand<mlir::acc::CreateOp, mlir::acc::CopyoutOp>(
785 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
787 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::CopyoutOp>(
791 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::DeclareEnterOp>) {
793 addDataOperand<mlir::acc::CreateOp>(
797 }
else if constexpr (isCombinedType<OpTy>) {
798 applyToComputeOp(clause);
800 llvm_unreachable(
"Unknown construct kind in VisitCopyOutClause");
804 void VisitCreateClause(
const OpenACCCreateClause &clause) {
805 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
806 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
808 addDataOperand<mlir::acc::CreateOp, mlir::acc::DeleteOp>(
812 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
814 addDataOperand<mlir::acc::CreateOp>(
817 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::DeclareEnterOp>) {
819 addDataOperand<mlir::acc::CreateOp>(
823 }
else if constexpr (isCombinedType<OpTy>) {
824 applyToComputeOp(clause);
826 llvm_unreachable(
"Unknown construct kind in VisitCreateClause");
830 void VisitLinkClause(
const OpenACCLinkClause &clause) {
831 if constexpr (isOneOfTypes<OpTy, mlir::acc::DeclareEnterOp>) {
833 addDataOperand<mlir::acc::DeclareLinkOp>(
834 var, mlir::acc::DataClause::acc_declare_link, {},
838 llvm_unreachable(
"Unknown construct kind in VisitLinkClause");
842 void VisitDeleteClause(
const OpenACCDeleteClause &clause) {
843 if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
845 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::DeleteOp>(
846 var, mlir::acc::DataClause::acc_delete, {},
850 llvm_unreachable(
"Unknown construct kind in VisitDeleteClause");
854 void VisitDetachClause(
const OpenACCDetachClause &clause) {
855 if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
857 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::DetachOp>(
858 var, mlir::acc::DataClause::acc_detach, {},
862 llvm_unreachable(
"Unknown construct kind in VisitDetachClause");
866 void VisitFinalizeClause(
const OpenACCFinalizeClause &clause) {
867 if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
868 operation.setFinalize(
true);
870 llvm_unreachable(
"Unknown construct kind in VisitFinalizeClause");
874 void VisitUseDeviceClause(
const OpenACCUseDeviceClause &clause) {
875 if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp>) {
877 addDataOperand<mlir::acc::UseDeviceOp>(
878 var, mlir::acc::DataClause::acc_use_device, {},
true,
881 llvm_unreachable(
"Unknown construct kind in VisitUseDeviceClause");
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>) {
890 addDataOperand<mlir::acc::DevicePtrOp>(
891 var, mlir::acc::DataClause::acc_deviceptr, {},
894 }
else if constexpr (isCombinedType<OpTy>) {
895 applyToComputeOp(clause);
897 llvm_unreachable(
"Unknown construct kind in VisitDevicePtrClause");
901 void VisitNoCreateClause(
const OpenACCNoCreateClause &clause) {
902 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
903 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
905 addDataOperand<mlir::acc::NoCreateOp, mlir::acc::DeleteOp>(
906 var, mlir::acc::DataClause::acc_no_create, {},
true,
908 }
else if constexpr (isCombinedType<OpTy>) {
909 applyToComputeOp(clause);
911 llvm_unreachable(
"Unknown construct kind in VisitNoCreateClause");
915 void VisitPresentClause(
const OpenACCPresentClause &clause) {
916 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
917 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
919 addDataOperand<mlir::acc::PresentOp, mlir::acc::DeleteOp>(
920 var, mlir::acc::DataClause::acc_present, {},
true,
922 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::DeclareEnterOp>) {
924 addDataOperand<mlir::acc::PresentOp>(
925 var, mlir::acc::DataClause::acc_present, {},
928 }
else if constexpr (isCombinedType<OpTy>) {
929 applyToComputeOp(clause);
931 llvm_unreachable(
"Unknown construct kind in VisitPresentClause");
935 void VisitAttachClause(
const OpenACCAttachClause &clause) {
936 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
937 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
939 addDataOperand<mlir::acc::AttachOp, mlir::acc::DetachOp>(
940 var, mlir::acc::DataClause::acc_attach, {},
true,
942 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
944 addDataOperand<mlir::acc::AttachOp>(
945 var, mlir::acc::DataClause::acc_attach, {},
947 }
else if constexpr (isCombinedType<OpTy>) {
948 applyToComputeOp(clause);
950 llvm_unreachable(
"Unknown construct kind in VisitAttachClause");
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] :
959 CIRGenFunction::OpenACCDataOperandInfo opInfo =
960 cgf.getOpenACCDataOperandInfo(varExpr);
961 auto privateOp = mlir::acc::PrivateOp::create(
964 privateOp.setDataClause(mlir::acc::DataClause::acc_private);
967 mlir::OpBuilder::InsertionGuard guardCase(builder);
970 OpenACCRecipeBuilder<mlir::acc::PrivateRecipeOp>(cgf, builder)
972 cgf.getContext(), recipeInsertLocation, varExpr,
973 varRecipe.AllocaDecl,
974 nullptr, OpenACCReductionOperator::Invalid,
981 operation.addPrivatization(builder.getContext(), privateOp, recipe);
984 }
else if constexpr (isCombinedType<OpTy>) {
987 applyToLoopOp(clause);
989 llvm_unreachable(
"Unknown construct kind in VisitPrivateClause");
993 void VisitFirstPrivateClause(
const OpenACCFirstPrivateClause &clause) {
994 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
995 mlir::acc::SerialOp>) {
996 for (
const auto [varExpr, varRecipe] :
998 CIRGenFunction::OpenACCDataOperandInfo opInfo =
999 cgf.getOpenACCDataOperandInfo(varExpr);
1000 auto firstPrivateOp = mlir::acc::FirstprivateOp::create(
1004 firstPrivateOp.setDataClause(mlir::acc::DataClause::acc_firstprivate);
1007 mlir::OpBuilder::InsertionGuard guardCase(builder);
1010 OpenACCRecipeBuilder<mlir::acc::FirstprivateRecipeOp>(cgf,
1013 cgf.getContext(), recipeInsertLocation, varExpr,
1014 varRecipe.AllocaDecl, varRecipe.InitFromTemporary,
1015 OpenACCReductionOperator::Invalid,
1018 firstPrivateOp, {});
1023 operation.addFirstPrivatization(builder.getContext(), firstPrivateOp,
1027 }
else if constexpr (isCombinedType<OpTy>) {
1030 applyToComputeOp(clause);
1032 llvm_unreachable(
"Unknown construct kind in VisitFirstPrivateClause");
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] :
1041 CIRGenFunction::OpenACCDataOperandInfo opInfo =
1042 cgf.getOpenACCDataOperandInfo(varExpr);
1044 auto reductionOp = mlir::acc::ReductionOp::create(
1047 reductionOp.setDataClause(mlir::acc::DataClause::acc_reduction);
1050 mlir::OpBuilder::InsertionGuard guardCase(builder);
1053 OpenACCRecipeBuilder<mlir::acc::ReductionRecipeOp>(cgf, builder)
1055 cgf.getContext(), recipeInsertLocation, varExpr,
1056 varRecipe.AllocaDecl,
1060 reductionOp, varRecipe.CombinerRecipes);
1062 operation.addReduction(builder.getContext(), reductionOp, recipe);
1065 }
else if constexpr (isCombinedType<OpTy>) {
1068 applyToLoopOp(clause);
1070 llvm_unreachable(
"Unknown construct kind in VisitReductionClause");
1074 void VisitDeviceResidentClause(
const OpenACCDeviceResidentClause &clause) {
1075 if constexpr (isOneOfTypes<OpTy, mlir::acc::DeclareEnterOp>) {
1077 addDataOperand<mlir::acc::DeclareDeviceResidentOp>(
1078 var, mlir::acc::DataClause::acc_declare_device_resident, {},
1082 llvm_unreachable(
"Unknown construct kind in VisitDeviceResidentClause");
1087template <
typename OpTy>
1088auto makeClauseEmitter(OpTy &op,
1089 mlir::OpBuilder::InsertPoint &recipeInsertLocation,
1093 return OpenACCClauseCIREmitter<OpTy>(op, recipeInsertLocation, cgf, builder,
1098template <
typename Op>
1099void CIRGenFunction::emitOpenACCClauses(
1102 mlir::OpBuilder::InsertionGuard guardCase(builder);
1106 builder.setInsertionPoint(op);
1107 makeClauseEmitter(op, lastRecipeLocation, *
this, builder, dirKind)
1108 .emitClauses(clauses);
1111#define EXPL_SPEC(N) \
1112 template void CIRGenFunction::emitOpenACCClauses<N>( \
1113 N &, OpenACCDirectiveKind, ArrayRef<const OpenACCClause *>);
1134template <
typename ComputeOp,
typename LoopOp>
1135void CIRGenFunction::emitOpenACCClauses(
1138 static_assert(std::is_same_v<mlir::acc::LoopOp, LoopOp>);
1140 CombinedConstructClauseInfo<ComputeOp> inf{op, loopOp};
1143 mlir::OpBuilder::InsertionGuard guardCase(builder);
1144 makeClauseEmitter(inf, lastRecipeLocation, *
this, builder, dirKind)
1145 .emitClauses(clauses);
1148#define EXPL_SPEC(N) \
1149 template void CIRGenFunction::emitOpenACCClauses<N, mlir::acc::LoopOp>( \
1150 N &, mlir::acc::LoopOp &, OpenACCDirectiveKind, \
1151 ArrayRef<const OpenACCClause *>);
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
const Expr * getIntExpr() 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()
bool hasConditionExpr() const
bool isEmptySelfClause() const
ArrayRef< Expr * > getSizeExprs()
ArrayRef< Expr * > getQueueIdExprs()
Expr * getDevNumExpr() const
bool hasDevNumExpr() const
SourceLocation getBeginLoc() const LLVM_READONLY
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.
bool isa(CodeGen::Address addr)
IdentifierLoc DeviceTypeArgument
const FunctionProtoType * T
__DEVICE__ _Tp arg(const std::complex< _Tp > &__c)
llvm::SmallVector< mlir::Value > bounds
llvm::SmallVector< QualType > boundTypes