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();
76 lastDeviceTypeValues.push_back(decodeDeviceType(
arg.getIdentifierInfo()));
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);
115 mlir::acc::DeviceType decodeDeviceType(
const IdentifierInfo *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);
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;
137 llvm_unreachable(
"unknown gang kind");
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);
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};
159 computeEmitter.lastDeviceTypeValues = lastDeviceTypeValues;
164 if (!dataOperands.empty())
165 computeEmitter.dataOperands.push_back(dataOperands.front());
166 computeEmitter.Visit(&
c);
171 dataOperands.append(computeEmitter.dataOperands);
174 template <
typename BeforeOpTy,
typename AfterOpTy>
175 void addDataOperand(
const Expr *varOperand, mlir::acc::DataClause dataClause,
178 CIRGenFunction::OpenACCDataOperandInfo opInfo =
179 cgf.getOpenACCDataOperandInfo(varOperand);
183 structured, implicit, opInfo.
name, opInfo.
bounds);
184 operation.getDataClauseOperandsMutable().append(beforeOp.getResult());
188 mlir::OpBuilder::InsertionGuard guardCase(builder);
189 builder.setInsertionPointAfter(operation);
191 if constexpr (std::is_same_v<AfterOpTy, mlir::acc::DeleteOp> ||
192 std::is_same_v<AfterOpTy, mlir::acc::DetachOp>) {
196 AfterOpTy::create(builder, opInfo.
beginLoc, beforeOp, structured,
199 afterOp = AfterOpTy::create(builder, opInfo.
beginLoc, beforeOp,
200 opInfo.
varValue, structured, implicit,
206 beforeOp.setDataClause(dataClause);
207 afterOp.setDataClause(dataClause);
212 dataOperands.push_back(beforeOp.getOperation());
213 dataOperands.push_back(afterOp.getOperation());
216 template <
typename BeforeOpTy>
217 void addDataOperand(
const Expr *varOperand, mlir::acc::DataClause dataClause,
220 CIRGenFunction::OpenACCDataOperandInfo opInfo =
221 cgf.getOpenACCDataOperandInfo(varOperand);
224 structured, implicit, opInfo.
name, opInfo.
bounds);
225 operation.getDataClauseOperandsMutable().append(beforeOp.getResult());
228 beforeOp.setDataClause(dataClause);
232 dataOperands.push_back(beforeOp.getOperation());
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{};
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();
258 llvm_unreachable(
"getting asyncOnly when clause not valid on operation?");
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{};
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();
285 "getting asyncOperandsDeviceType when clause not valid on operation?");
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();
305 "getting asyncOperandsDeviceType when clause not valid on operation?");
311 void updateDataOperandAsyncValues() {
312 if (!hasAsyncClause || dataOperands.empty())
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());
322 .
Default([&](mlir::Operation *) {
323 llvm_unreachable(
"Not a data operation?");
329 OpenACCClauseCIREmitter(OpTy &operation,
330 mlir::OpBuilder::InsertPoint &recipeInsertLocation,
331 CIRGen::CIRGenFunction &cgf,
332 CIRGen::CIRGenBuilderTy &builder,
334 : operation(operation), recipeInsertLocation(recipeInsertLocation),
335 cgf(cgf), builder(builder), dirKind(dirKind) {}
337 void VisitClause(
const OpenACCClause &clause) {
338 llvm_unreachable(
"Unknown/unhandled clause kind");
344 void emitClauses(ArrayRef<const OpenACCClause *> clauses) {
345 this->VisitClauseList(clauses);
346 updateDataOperandAsyncValues();
349 void VisitDefaultClause(
const OpenACCDefaultClause &clause) {
352 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
353 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
355 case OpenACCDefaultClauseKind::None:
356 operation.setDefaultAttr(mlir::acc::ClauseDefaultValue::None);
358 case OpenACCDefaultClauseKind::Present:
359 operation.setDefaultAttr(mlir::acc::ClauseDefaultValue::Present);
361 case OpenACCDefaultClauseKind::Invalid:
364 }
else if constexpr (isCombinedType<OpTy>) {
365 applyToComputeOp(clause);
367 llvm_unreachable(
"Unknown construct kind in VisitDefaultClause");
371 void VisitDeviceTypeClause(
const OpenACCDeviceTypeClause &clause) {
372 setLastDeviceTypeClause(clause);
374 if constexpr (isOneOfTypes<OpTy, mlir::acc::InitOp,
375 mlir::acc::ShutdownOp>) {
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?");
384 operation.setDeviceType(
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>) {
393 }
else if constexpr (isCombinedType<OpTy>) {
397 llvm_unreachable(
"Unknown construct kind in VisitDeviceTypeClause");
401 void VisitNumWorkersClause(
const OpenACCNumWorkersClause &clause) {
402 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
403 mlir::acc::KernelsOp>) {
404 operation.addNumWorkersOperand(builder.getContext(),
406 lastDeviceTypeValues);
407 }
else if constexpr (isCombinedType<OpTy>) {
408 applyToComputeOp(clause);
410 llvm_unreachable(
"Unknown construct kind in VisitNumGangsClause");
414 void VisitVectorLengthClause(
const OpenACCVectorLengthClause &clause) {
415 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
416 mlir::acc::KernelsOp>) {
417 operation.addVectorLengthOperand(builder.getContext(),
419 lastDeviceTypeValues);
420 }
else if constexpr (isCombinedType<OpTy>) {
421 applyToComputeOp(clause);
423 llvm_unreachable(
"Unknown construct kind in VisitVectorLengthClause");
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>) {
434 operation.addAsyncOnly(builder.getContext(), lastDeviceTypeValues);
442 mlir::OpBuilder::InsertionGuard guardCase(builder);
443 if (!dataOperands.empty())
444 builder.setInsertionPoint(dataOperands.front());
447 operation.addAsyncOperand(builder.getContext(), intExpr,
448 lastDeviceTypeValues);
450 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::WaitOp>) {
454 operation.setAsync(
true);
456 operation.getAsyncOperandMutable().append(
458 }
else if constexpr (isCombinedType<OpTy>) {
459 applyToComputeOp(clause);
461 llvm_unreachable(
"Unknown construct kind in VisitAsyncClause");
465 void VisitSelfClause(
const OpenACCSelfClause &clause) {
466 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
467 mlir::acc::KernelsOp>) {
469 operation.setSelfAttr(
true);
472 operation.getSelfCondMutable().append(
475 llvm_unreachable(
"var-list version of self shouldn't get here");
477 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
479 "var-list version of self required for update");
481 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::UpdateHostOp>(
482 var, mlir::acc::DataClause::acc_update_self, {},
484 }
else if constexpr (isCombinedType<OpTy>) {
485 applyToComputeOp(clause);
487 llvm_unreachable(
"Unknown construct kind in VisitSelfClause");
491 void VisitHostClause(
const OpenACCHostClause &clause) {
492 if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
494 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::UpdateHostOp>(
495 var, mlir::acc::DataClause::acc_update_host, {},
498 llvm_unreachable(
"Unknown construct kind in VisitHostClause");
502 void VisitDeviceClause(
const OpenACCDeviceClause &clause) {
503 if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
505 addDataOperand<mlir::acc::UpdateDeviceOp>(
506 var, mlir::acc::DataClause::acc_update_device, {},
509 llvm_unreachable(
"Unknown construct kind in VisitDeviceClause");
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(
525 }
else if constexpr (isCombinedType<OpTy>) {
526 applyToComputeOp(clause);
528 llvm_unreachable(
"Unknown construct kind in VisitIfClause");
532 void VisitIfPresentClause(
const OpenACCIfPresentClause &clause) {
533 if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp,
534 mlir::acc::UpdateOp>) {
535 operation.setIfPresent(
true);
537 llvm_unreachable(
"unknown construct kind in VisitIfPresentClause");
541 void VisitDeviceNumClause(
const OpenACCDeviceNumClause &clause) {
542 if constexpr (isOneOfTypes<OpTy, mlir::acc::InitOp, mlir::acc::ShutdownOp,
544 operation.getDeviceNumMutable().append(emitIntExpr(clause.
getIntExpr()));
547 "init, shutdown, set, are only valid device_num constructs");
551 void VisitNumGangsClause(
const OpenACCNumGangsClause &clause) {
552 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
553 mlir::acc::KernelsOp>) {
554 llvm::SmallVector<mlir::Value> values;
556 values.push_back(emitIntExpr(E));
558 operation.addNumGangsOperands(builder.getContext(), values,
559 lastDeviceTypeValues);
560 }
else if constexpr (isCombinedType<OpTy>) {
561 applyToComputeOp(clause);
563 llvm_unreachable(
"Unknown construct kind in VisitNumGangsClause");
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>) {
573 operation.addWaitOnly(builder.getContext(), lastDeviceTypeValues);
575 llvm::SmallVector<mlir::Value> values;
579 values.push_back(emitIntExpr(E));
580 operation.addWaitOperands(builder.getContext(), clause.
hasDevNumExpr(),
581 values, lastDeviceTypeValues);
583 }
else if constexpr (isCombinedType<OpTy>) {
584 applyToComputeOp(clause);
588 llvm_unreachable(
"Unknown construct kind in VisitWaitClause");
592 void VisitDefaultAsyncClause(
const OpenACCDefaultAsyncClause &clause) {
593 if constexpr (isOneOfTypes<OpTy, mlir::acc::SetOp>) {
594 operation.getDefaultAsyncMutable().append(
597 llvm_unreachable(
"set, is only valid device_num constructs");
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);
607 llvm_unreachable(
"Unknown construct kind in VisitSeqClause");
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);
617 llvm_unreachable(
"Unknown construct kind in VisitAutoClause");
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);
627 llvm_unreachable(
"Unknown construct kind in VisitIndependentClause");
631 void VisitCollapseClause(
const OpenACCCollapseClause &clause) {
632 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
636 value = value.sextOrTrunc(64);
637 operation.setCollapseForDeviceTypes(builder.getContext(),
638 lastDeviceTypeValues, value);
639 }
else if constexpr (isCombinedType<OpTy>) {
640 applyToLoopOp(clause);
642 llvm_unreachable(
"Unknown construct kind in VisitCollapseClause");
646 void VisitTileClause(
const OpenACCTileClause &clause) {
647 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
648 llvm::SmallVector<mlir::Value> values;
651 mlir::Location exprLoc = cgf.cgm.getLoc(e->getBeginLoc());
657 values.push_back(createConstantInt(exprLoc, 64, -1));
659 llvm::APInt curValue =
660 e->EvaluateKnownConstInt(cgf.cgm.getASTContext());
661 values.push_back(createConstantInt(
662 exprLoc, 64, curValue.sextOrTrunc(64).getSExtValue()));
666 operation.setTileForDeviceTypes(builder.getContext(),
667 lastDeviceTypeValues, values);
668 }
else if constexpr (isCombinedType<OpTy>) {
669 applyToLoopOp(clause);
671 llvm_unreachable(
"Unknown construct kind in VisitTileClause");
675 void VisitWorkerClause(
const OpenACCWorkerClause &clause) {
676 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
678 operation.addWorkerNumOperand(builder.getContext(),
680 lastDeviceTypeValues);
682 operation.addEmptyWorker(builder.getContext(), lastDeviceTypeValues);
684 }
else if constexpr (isCombinedType<OpTy>) {
685 applyToLoopOp(clause);
687 llvm_unreachable(
"Unknown construct kind in VisitWorkerClause");
691 void VisitVectorClause(
const OpenACCVectorClause &clause) {
692 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
694 operation.addVectorOperand(builder.getContext(),
696 lastDeviceTypeValues);
698 operation.addEmptyVector(builder.getContext(), lastDeviceTypeValues);
700 }
else if constexpr (isCombinedType<OpTy>) {
701 applyToLoopOp(clause);
703 llvm_unreachable(
"Unknown construct kind in VisitVectorClause");
707 void VisitGangClause(
const OpenACCGangClause &clause) {
708 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
710 operation.addEmptyGang(builder.getContext(), lastDeviceTypeValues);
712 llvm::SmallVector<mlir::Value> values;
713 llvm::SmallVector<mlir::acc::GangArgType> argTypes;
714 for (
unsigned i : llvm::index_range(0u, clause.
getNumExprs())) {
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());
723 curValue = curValue.sextOrTrunc(64);
725 createConstantInt(exprLoc, 64, curValue.getSExtValue()));
727 values.push_back(createConstantInt(exprLoc, 64, -1));
729 values.push_back(emitIntExpr(
expr));
733 operation.addGangOperands(builder.getContext(), lastDeviceTypeValues,
736 }
else if constexpr (isCombinedType<OpTy>) {
737 applyToLoopOp(clause);
739 llvm_unreachable(
"Unknown construct kind in VisitGangClause");
743 void VisitCopyClause(
const OpenACCCopyClause &clause) {
744 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
745 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
747 addDataOperand<mlir::acc::CopyinOp, mlir::acc::CopyoutOp>(
751 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::DeclareEnterOp>) {
753 addDataOperand<mlir::acc::CopyinOp>(
757 }
else if constexpr (isCombinedType<OpTy>) {
758 applyToComputeOp(clause);
760 llvm_unreachable(
"Unknown construct kind in VisitCopyClause");
764 void VisitCopyInClause(
const OpenACCCopyInClause &clause) {
765 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
766 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
768 addDataOperand<mlir::acc::CopyinOp, mlir::acc::DeleteOp>(
772 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
774 addDataOperand<mlir::acc::CopyinOp>(
777 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::DeclareEnterOp>) {
779 addDataOperand<mlir::acc::CopyinOp>(
783 }
else if constexpr (isCombinedType<OpTy>) {
784 applyToComputeOp(clause);
786 llvm_unreachable(
"Unknown construct kind in VisitCopyInClause");
790 void VisitCopyOutClause(
const OpenACCCopyOutClause &clause) {
791 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
792 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
794 addDataOperand<mlir::acc::CreateOp, mlir::acc::CopyoutOp>(
798 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
800 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::CopyoutOp>(
804 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::DeclareEnterOp>) {
806 addDataOperand<mlir::acc::CreateOp>(
810 }
else if constexpr (isCombinedType<OpTy>) {
811 applyToComputeOp(clause);
813 llvm_unreachable(
"Unknown construct kind in VisitCopyOutClause");
817 void VisitCreateClause(
const OpenACCCreateClause &clause) {
818 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
819 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
821 addDataOperand<mlir::acc::CreateOp, mlir::acc::DeleteOp>(
825 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
827 addDataOperand<mlir::acc::CreateOp>(
830 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::DeclareEnterOp>) {
832 addDataOperand<mlir::acc::CreateOp>(
836 }
else if constexpr (isCombinedType<OpTy>) {
837 applyToComputeOp(clause);
839 llvm_unreachable(
"Unknown construct kind in VisitCreateClause");
843 void VisitLinkClause(
const OpenACCLinkClause &clause) {
844 if constexpr (isOneOfTypes<OpTy, mlir::acc::DeclareEnterOp>) {
846 addDataOperand<mlir::acc::DeclareLinkOp>(
847 var, mlir::acc::DataClause::acc_declare_link, {},
851 llvm_unreachable(
"Unknown construct kind in VisitLinkClause");
855 void VisitDeleteClause(
const OpenACCDeleteClause &clause) {
856 if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
858 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::DeleteOp>(
859 var, mlir::acc::DataClause::acc_delete, {},
863 llvm_unreachable(
"Unknown construct kind in VisitDeleteClause");
867 void VisitDetachClause(
const OpenACCDetachClause &clause) {
868 if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
870 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::DetachOp>(
871 var, mlir::acc::DataClause::acc_detach, {},
875 llvm_unreachable(
"Unknown construct kind in VisitDetachClause");
879 void VisitFinalizeClause(
const OpenACCFinalizeClause &clause) {
880 if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
881 operation.setFinalize(
true);
883 llvm_unreachable(
"Unknown construct kind in VisitFinalizeClause");
887 void VisitUseDeviceClause(
const OpenACCUseDeviceClause &clause) {
888 if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp>) {
890 addDataOperand<mlir::acc::UseDeviceOp>(
891 var, mlir::acc::DataClause::acc_use_device, {},
true,
894 llvm_unreachable(
"Unknown construct kind in VisitUseDeviceClause");
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>) {
903 addDataOperand<mlir::acc::DevicePtrOp>(
904 var, mlir::acc::DataClause::acc_deviceptr, {},
907 }
else if constexpr (isCombinedType<OpTy>) {
908 applyToComputeOp(clause);
910 llvm_unreachable(
"Unknown construct kind in VisitDevicePtrClause");
914 void VisitNoCreateClause(
const OpenACCNoCreateClause &clause) {
915 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
916 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
918 addDataOperand<mlir::acc::NoCreateOp, mlir::acc::DeleteOp>(
919 var, mlir::acc::DataClause::acc_no_create, {},
true,
921 }
else if constexpr (isCombinedType<OpTy>) {
922 applyToComputeOp(clause);
924 llvm_unreachable(
"Unknown construct kind in VisitNoCreateClause");
928 void VisitPresentClause(
const OpenACCPresentClause &clause) {
929 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
930 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
932 addDataOperand<mlir::acc::PresentOp, mlir::acc::DeleteOp>(
933 var, mlir::acc::DataClause::acc_present, {},
true,
935 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::DeclareEnterOp>) {
937 addDataOperand<mlir::acc::PresentOp>(
938 var, mlir::acc::DataClause::acc_present, {},
941 }
else if constexpr (isCombinedType<OpTy>) {
942 applyToComputeOp(clause);
944 llvm_unreachable(
"Unknown construct kind in VisitPresentClause");
948 void VisitAttachClause(
const OpenACCAttachClause &clause) {
949 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
950 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
952 addDataOperand<mlir::acc::AttachOp, mlir::acc::DetachOp>(
953 var, mlir::acc::DataClause::acc_attach, {},
true,
955 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
957 addDataOperand<mlir::acc::AttachOp>(
958 var, mlir::acc::DataClause::acc_attach, {},
960 }
else if constexpr (isCombinedType<OpTy>) {
961 applyToComputeOp(clause);
963 llvm_unreachable(
"Unknown construct kind in VisitAttachClause");
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] :
972 CIRGenFunction::OpenACCDataOperandInfo opInfo =
973 cgf.getOpenACCDataOperandInfo(varExpr);
974 auto privateOp = mlir::acc::PrivateOp::create(
977 privateOp.setDataClause(mlir::acc::DataClause::acc_private);
980 mlir::OpBuilder::InsertionGuard guardCase(builder);
983 OpenACCRecipeBuilder<mlir::acc::PrivateRecipeOp>(cgf, builder)
985 cgf.getContext(), recipeInsertLocation, varExpr,
986 varRecipe.AllocaDecl,
987 nullptr, OpenACCReductionOperator::Invalid,
994 operation.addPrivatization(builder.getContext(), privateOp, recipe);
997 }
else if constexpr (isCombinedType<OpTy>) {
1000 applyToLoopOp(clause);
1002 llvm_unreachable(
"Unknown construct kind in VisitPrivateClause");
1006 void VisitFirstPrivateClause(
const OpenACCFirstPrivateClause &clause) {
1007 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
1008 mlir::acc::SerialOp>) {
1009 for (
const auto [varExpr, varRecipe] :
1011 CIRGenFunction::OpenACCDataOperandInfo opInfo =
1012 cgf.getOpenACCDataOperandInfo(varExpr);
1013 auto firstPrivateOp = mlir::acc::FirstprivateOp::create(
1017 firstPrivateOp.setDataClause(mlir::acc::DataClause::acc_firstprivate);
1020 mlir::OpBuilder::InsertionGuard guardCase(builder);
1023 OpenACCRecipeBuilder<mlir::acc::FirstprivateRecipeOp>(cgf,
1026 cgf.getContext(), recipeInsertLocation, varExpr,
1027 varRecipe.AllocaDecl, varRecipe.InitFromTemporary,
1028 OpenACCReductionOperator::Invalid,
1031 firstPrivateOp, {});
1036 operation.addFirstPrivatization(builder.getContext(), firstPrivateOp,
1040 }
else if constexpr (isCombinedType<OpTy>) {
1043 applyToComputeOp(clause);
1045 llvm_unreachable(
"Unknown construct kind in VisitFirstPrivateClause");
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] :
1054 CIRGenFunction::OpenACCDataOperandInfo opInfo =
1055 cgf.getOpenACCDataOperandInfo(varExpr);
1057 auto reductionOp = mlir::acc::ReductionOp::create(
1060 reductionOp.setDataClause(mlir::acc::DataClause::acc_reduction);
1063 mlir::OpBuilder::InsertionGuard guardCase(builder);
1066 OpenACCRecipeBuilder<mlir::acc::ReductionRecipeOp>(cgf, builder)
1068 cgf.getContext(), recipeInsertLocation, varExpr,
1069 varRecipe.AllocaDecl,
1073 reductionOp, varRecipe.CombinerRecipes);
1075 operation.addReduction(builder.getContext(), reductionOp, recipe);
1078 }
else if constexpr (isCombinedType<OpTy>) {
1081 applyToLoopOp(clause);
1083 llvm_unreachable(
"Unknown construct kind in VisitReductionClause");
1087 void VisitDeviceResidentClause(
const OpenACCDeviceResidentClause &clause) {
1088 if constexpr (isOneOfTypes<OpTy, mlir::acc::DeclareEnterOp>) {
1090 addDataOperand<mlir::acc::DeclareDeviceResidentOp>(
1091 var, mlir::acc::DataClause::acc_declare_device_resident, {},
1095 llvm_unreachable(
"Unknown construct kind in VisitDeviceResidentClause");
1100template <
typename OpTy>
1101auto makeClauseEmitter(OpTy &op,
1102 mlir::OpBuilder::InsertPoint &recipeInsertLocation,
1106 return OpenACCClauseCIREmitter<OpTy>(op, recipeInsertLocation, cgf, builder,
1111template <
typename Op>
1112void CIRGenFunction::emitOpenACCClauses(
1115 mlir::OpBuilder::InsertionGuard guardCase(builder);
1119 builder.setInsertionPoint(op);
1120 makeClauseEmitter(op, lastRecipeLocation, *
this, builder, dirKind)
1121 .emitClauses(clauses);
1124#define EXPL_SPEC(N) \
1125 template void CIRGenFunction::emitOpenACCClauses<N>( \
1126 N &, OpenACCDirectiveKind, ArrayRef<const OpenACCClause *>);
1147template <
typename ComputeOp,
typename LoopOp>
1148void CIRGenFunction::emitOpenACCClauses(
1151 static_assert(std::is_same_v<mlir::acc::LoopOp, LoopOp>);
1153 CombinedConstructClauseInfo<ComputeOp> inf{op, loopOp};
1156 mlir::OpBuilder::InsertionGuard guardCase(builder);
1157 makeClauseEmitter(inf, lastRecipeLocation, *
this, builder, dirKind)
1158 .emitClauses(clauses);
1161#define EXPL_SPEC(N) \
1162 template void CIRGenFunction::emitOpenACCClauses<N, mlir::acc::LoopOp>( \
1163 N &, mlir::acc::LoopOp &, OpenACCDirectiveKind, \
1164 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.
StringRef getName() const
Return the actual identifier string.
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::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