21#include "mlir/Dialect/Arith/IR/Arith.h"
22#include "mlir/Dialect/OpenACC/OpenACC.h"
23#include "llvm/ADT/TypeSwitch.h"
31template <
typename ToTest,
typename T,
typename... Tys>
32constexpr bool isOneOfTypes =
33 std::is_same_v<ToTest, T> || isOneOfTypes<ToTest, Tys...>;
34template <
typename ToTest,
typename T>
35constexpr bool isOneOfTypes<ToTest, T> = std::is_same_v<ToTest, T>;
40template <
typename CompOpTy>
struct CombinedConstructClauseInfo {
41 using ComputeOpTy = CompOpTy;
42 ComputeOpTy computeOp;
43 mlir::acc::LoopOp loopOp;
45template <
typename ToTest>
constexpr bool isCombinedType =
false;
47constexpr bool isCombinedType<CombinedConstructClauseInfo<T>> =
true;
49template <
typename OpTy>
50class OpenACCClauseCIREmitter final
53 template <
typename FriendOpTy>
friend class OpenACCClauseCIREmitter;
56 mlir::OpBuilder::InsertPoint &recipeInsertLocation;
57 CIRGen::CIRGenFunction &cgf;
58 CIRGen::CIRGenBuilderTy &builder;
65 SourceLocation dirLoc;
67 llvm::SmallVector<mlir::acc::DeviceType> lastDeviceTypeValues;
70 bool hasAsyncClause =
false;
72 llvm::SmallVector<mlir::Operation *> dataOperands;
74 void clauseNotImplemented(
const OpenACCClause &
c) {
75 cgf.cgm.errorNYI(
c.getSourceRange(),
"OpenACC Clause",
c.getClauseKind());
78 void setLastDeviceTypeClause(
const OpenACCDeviceTypeClause &clause) {
79 lastDeviceTypeValues.clear();
82 lastDeviceTypeValues.push_back(decodeDeviceType(
arg.getIdentifierInfo()));
85 mlir::Value emitIntExpr(
const Expr *intExpr) {
86 return cgf.emitOpenACCIntExpr(intExpr);
93 mlir::Value createCondition(
const Expr *condExpr) {
94 mlir::Value condition = cgf.evaluateExprAsBool(condExpr);
95 mlir::Location exprLoc = cgf.cgm.getLoc(condExpr->
getBeginLoc());
96 mlir::IntegerType targetType = mlir::IntegerType::get(
97 &cgf.getMLIRContext(), 1,
98 mlir::IntegerType::SignednessSemantics::Signless);
99 auto conversionOp = mlir::UnrealizedConversionCastOp::create(
100 builder, exprLoc, targetType, condition);
101 return conversionOp.getResult(0);
104 mlir::Value createConstantInt(mlir::Location loc,
unsigned width,
106 return cgf.createOpenACCConstantInt(loc, width, value);
107 mlir::IntegerType ty = mlir::IntegerType::get(
108 &cgf.getMLIRContext(), width,
109 mlir::IntegerType::SignednessSemantics::Signless);
110 auto constOp = mlir::arith::ConstantOp::create(
111 builder, loc, builder.getIntegerAttr(ty, value));
116 mlir::Value createConstantInt(SourceLocation loc,
unsigned width,
118 return createConstantInt(cgf.cgm.getLoc(loc), width, value);
121 mlir::acc::DeviceType decodeDeviceType(
const IdentifierInfo *ii) {
124 return mlir::acc::DeviceType::Star;
125 return llvm::StringSwitch<mlir::acc::DeviceType>(ii->
getName())
126 .CaseLower(
"default", mlir::acc::DeviceType::Default)
127 .CaseLower(
"host", mlir::acc::DeviceType::Host)
128 .CaseLower(
"multicore", mlir::acc::DeviceType::Multicore)
129 .CasesLower({
"nvidia",
"acc_device_nvidia"},
130 mlir::acc::DeviceType::Nvidia)
131 .CaseLower(
"radeon", mlir::acc::DeviceType::Radeon);
136 case OpenACCGangKind::Num:
137 return mlir::acc::GangArgType::Num;
138 case OpenACCGangKind::Dim:
139 return mlir::acc::GangArgType::Dim;
140 case OpenACCGangKind::Static:
141 return mlir::acc::GangArgType::Static;
143 llvm_unreachable(
"unknown gang kind");
146 template <
typename U = void,
147 typename = std::enable_if_t<isCombinedType<OpTy>, U>>
148 void applyToLoopOp(
const OpenACCClause &
c) {
149 mlir::OpBuilder::InsertionGuard guardCase(builder);
150 builder.setInsertionPoint(operation.loopOp);
151 OpenACCClauseCIREmitter<mlir::acc::LoopOp> loopEmitter{
152 operation.loopOp, recipeInsertLocation, cgf, builder, dirKind, dirLoc};
153 loopEmitter.lastDeviceTypeValues = lastDeviceTypeValues;
154 loopEmitter.Visit(&
c);
157 template <
typename U = void,
158 typename = std::enable_if_t<isCombinedType<OpTy>, U>>
159 void applyToComputeOp(
const OpenACCClause &
c) {
160 mlir::OpBuilder::InsertionGuard guardCase(builder);
161 builder.setInsertionPoint(operation.computeOp);
162 OpenACCClauseCIREmitter<typename OpTy::ComputeOpTy> computeEmitter{
164 recipeInsertLocation,
170 computeEmitter.lastDeviceTypeValues = lastDeviceTypeValues;
175 if (!dataOperands.empty())
176 computeEmitter.dataOperands.push_back(dataOperands.front());
177 computeEmitter.Visit(&
c);
182 dataOperands.append(computeEmitter.dataOperands);
185 mlir::acc::DataClauseModifier
187 using namespace mlir::acc;
188 static_assert(
static_cast<int>(OpenACCModifierKind::Zero) ==
189 static_cast<int>(DataClauseModifier::zero) &&
190 static_cast<int>(OpenACCModifierKind::Readonly) ==
191 static_cast<int>(DataClauseModifier::readonly) &&
192 static_cast<int>(OpenACCModifierKind::AlwaysIn) ==
193 static_cast<int>(DataClauseModifier::alwaysin) &&
194 static_cast<int>(OpenACCModifierKind::AlwaysOut) ==
195 static_cast<int>(DataClauseModifier::alwaysout) &&
196 static_cast<int>(OpenACCModifierKind::Capture) ==
197 static_cast<int>(DataClauseModifier::capture));
199 DataClauseModifier mlirModifiers{};
204 mlirModifiers = mlirModifiers | DataClauseModifier::always;
205 modifiers &= ~OpenACCModifierKind
::Always;
208 mlirModifiers = mlirModifiers |
static_cast<DataClauseModifier
>(modifiers);
209 return mlirModifiers;
212 template <
typename BeforeOpTy,
typename AfterOpTy>
213 void addDataOperand(
const Expr *varOperand, mlir::acc::DataClause dataClause,
216 CIRGenFunction::OpenACCDataOperandInfo opInfo =
217 cgf.getOpenACCDataOperandInfo(varOperand);
221 structured, implicit, opInfo.
name, opInfo.
bounds);
222 operation.getDataClauseOperandsMutable().append(beforeOp.getResult());
226 mlir::OpBuilder::InsertionGuard guardCase(builder);
227 builder.setInsertionPointAfter(operation);
229 if constexpr (std::is_same_v<AfterOpTy, mlir::acc::DeleteOp> ||
230 std::is_same_v<AfterOpTy, mlir::acc::DetachOp>) {
234 AfterOpTy::create(builder, opInfo.
beginLoc, beforeOp, structured,
237 afterOp = AfterOpTy::create(builder, opInfo.
beginLoc, beforeOp,
238 opInfo.
varValue, structured, implicit,
244 beforeOp.setDataClause(dataClause);
245 afterOp.setDataClause(dataClause);
246 beforeOp.setModifiers(convertModifiers(modifiers));
247 afterOp.setModifiers(convertModifiers(modifiers));
250 dataOperands.push_back(beforeOp.getOperation());
251 dataOperands.push_back(afterOp.getOperation());
254 template <
typename BeforeOpTy>
255 void addDataOperand(
const Expr *varOperand, mlir::acc::DataClause dataClause,
258 CIRGenFunction::OpenACCDataOperandInfo opInfo =
259 cgf.getOpenACCDataOperandInfo(varOperand);
262 structured, implicit, opInfo.
name, opInfo.
bounds);
263 operation.getDataClauseOperandsMutable().append(beforeOp.getResult());
266 beforeOp.setDataClause(dataClause);
267 beforeOp.setModifiers(convertModifiers(modifiers));
270 dataOperands.push_back(beforeOp.getOperation());
275 mlir::ArrayAttr getAsyncOnlyAttr() {
276 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
277 mlir::acc::KernelsOp, mlir::acc::DataOp,
278 mlir::acc::UpdateOp>) {
279 return operation.getAsyncOnlyAttr();
280 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp,
281 mlir::acc::ExitDataOp>) {
282 if (!operation.getAsyncAttr())
283 return mlir::ArrayAttr{};
285 llvm::SmallVector<mlir::Attribute> devTysTemp;
286 devTysTemp.push_back(mlir::acc::DeviceTypeAttr::get(
287 builder.getContext(), mlir::acc::DeviceType::None));
288 return mlir::ArrayAttr::get(builder.getContext(), devTysTemp);
289 }
else if constexpr (isCombinedType<OpTy>) {
290 return operation.computeOp.getAsyncOnlyAttr();
296 llvm_unreachable(
"getting asyncOnly when clause not valid on operation?");
301 mlir::ArrayAttr getAsyncOperandsDeviceTypeAttr() {
302 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
303 mlir::acc::KernelsOp, mlir::acc::DataOp,
304 mlir::acc::UpdateOp>) {
305 return operation.getAsyncOperandsDeviceTypeAttr();
306 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp,
307 mlir::acc::ExitDataOp>) {
308 if (!operation.getAsyncOperand())
309 return mlir::ArrayAttr{};
311 llvm::SmallVector<mlir::Attribute> devTysTemp;
312 devTysTemp.push_back(mlir::acc::DeviceTypeAttr::get(
313 builder.getContext(), mlir::acc::DeviceType::None));
314 return mlir::ArrayAttr::get(builder.getContext(), devTysTemp);
315 }
else if constexpr (isCombinedType<OpTy>) {
316 return operation.computeOp.getAsyncOperandsDeviceTypeAttr();
323 "getting asyncOperandsDeviceType when clause not valid on operation?");
328 mlir::OperandRange getAsyncOperands() {
329 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
330 mlir::acc::KernelsOp, mlir::acc::DataOp,
331 mlir::acc::UpdateOp>)
332 return operation.getAsyncOperands();
333 else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp,
334 mlir::acc::ExitDataOp>)
335 return operation.getAsyncOperandMutable();
336 else if constexpr (isCombinedType<OpTy>)
337 return operation.computeOp.getAsyncOperands();
343 "getting asyncOperandsDeviceType when clause not valid on operation?");
349 void updateDataOperandAsyncValues() {
350 if (!hasAsyncClause || dataOperands.empty())
353 for (mlir::Operation *dataOp : dataOperands) {
354 llvm::TypeSwitch<mlir::Operation *, void>(dataOp)
355 .Case<ACC_DATA_ENTRY_OPS, ACC_DATA_EXIT_OPS>([&](
auto op) {
356 op.setAsyncOnlyAttr(getAsyncOnlyAttr());
357 op.setAsyncOperandsDeviceTypeAttr(getAsyncOperandsDeviceTypeAttr());
358 op.getAsyncOperandsMutable().assign(getAsyncOperands());
360 .
Default([&](mlir::Operation *) {
361 llvm_unreachable(
"Not a data operation?");
367 OpenACCClauseCIREmitter(OpTy &operation,
368 mlir::OpBuilder::InsertPoint &recipeInsertLocation,
369 CIRGen::CIRGenFunction &cgf,
370 CIRGen::CIRGenBuilderTy &builder,
372 : operation(operation), recipeInsertLocation(recipeInsertLocation),
373 cgf(cgf), builder(builder), dirKind(dirKind), dirLoc(dirLoc) {}
375 void VisitClause(
const OpenACCClause &clause) {
376 clauseNotImplemented(clause);
382 void emitClauses(ArrayRef<const OpenACCClause *> clauses) {
383 this->VisitClauseList(clauses);
384 updateDataOperandAsyncValues();
387 void VisitDefaultClause(
const OpenACCDefaultClause &clause) {
390 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
391 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
393 case OpenACCDefaultClauseKind::None:
394 operation.setDefaultAttr(mlir::acc::ClauseDefaultValue::None);
396 case OpenACCDefaultClauseKind::Present:
397 operation.setDefaultAttr(mlir::acc::ClauseDefaultValue::Present);
399 case OpenACCDefaultClauseKind::Invalid:
402 }
else if constexpr (isCombinedType<OpTy>) {
403 applyToComputeOp(clause);
405 llvm_unreachable(
"Unknown construct kind in VisitDefaultClause");
409 void VisitDeviceTypeClause(
const OpenACCDeviceTypeClause &clause) {
410 setLastDeviceTypeClause(clause);
412 if constexpr (isOneOfTypes<OpTy, mlir::acc::InitOp,
413 mlir::acc::ShutdownOp>) {
415 operation.addDeviceType(builder.getContext(),
416 decodeDeviceType(
arg.getIdentifierInfo()));
417 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::SetOp>) {
418 assert(!operation.getDeviceTypeAttr() &&
"already have device-type?");
422 operation.setDeviceType(
424 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
425 mlir::acc::SerialOp, mlir::acc::KernelsOp,
426 mlir::acc::DataOp, mlir::acc::LoopOp,
427 mlir::acc::UpdateOp>) {
431 }
else if constexpr (isCombinedType<OpTy>) {
437 return clauseNotImplemented(clause);
441 void VisitNumWorkersClause(
const OpenACCNumWorkersClause &clause) {
442 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
443 mlir::acc::KernelsOp>) {
444 operation.addNumWorkersOperand(builder.getContext(),
446 lastDeviceTypeValues);
447 }
else if constexpr (isCombinedType<OpTy>) {
448 applyToComputeOp(clause);
450 llvm_unreachable(
"Unknown construct kind in VisitNumGangsClause");
454 void VisitVectorLengthClause(
const OpenACCVectorLengthClause &clause) {
455 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
456 mlir::acc::KernelsOp>) {
457 operation.addVectorLengthOperand(builder.getContext(),
459 lastDeviceTypeValues);
460 }
else if constexpr (isCombinedType<OpTy>) {
461 applyToComputeOp(clause);
463 llvm_unreachable(
"Unknown construct kind in VisitVectorLengthClause");
467 void VisitAsyncClause(
const OpenACCAsyncClause &clause) {
468 hasAsyncClause =
true;
469 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
470 mlir::acc::KernelsOp, mlir::acc::DataOp,
471 mlir::acc::EnterDataOp, mlir::acc::ExitDataOp,
472 mlir::acc::UpdateOp>) {
474 operation.addAsyncOnly(builder.getContext(), lastDeviceTypeValues);
482 mlir::OpBuilder::InsertionGuard guardCase(builder);
483 if (!dataOperands.empty())
484 builder.setInsertionPoint(dataOperands.front());
487 operation.addAsyncOperand(builder.getContext(), intExpr,
488 lastDeviceTypeValues);
490 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::WaitOp>) {
494 operation.setAsync(
true);
496 operation.getAsyncOperandMutable().append(
498 }
else if constexpr (isCombinedType<OpTy>) {
499 applyToComputeOp(clause);
503 return clauseNotImplemented(clause);
507 void VisitSelfClause(
const OpenACCSelfClause &clause) {
508 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
509 mlir::acc::KernelsOp>) {
511 operation.setSelfAttr(
true);
514 operation.getSelfCondMutable().append(
517 llvm_unreachable(
"var-list version of self shouldn't get here");
519 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
521 "var-list version of self required for update");
523 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::UpdateHostOp>(
524 var, mlir::acc::DataClause::acc_update_self, {},
526 }
else if constexpr (isCombinedType<OpTy>) {
527 applyToComputeOp(clause);
529 llvm_unreachable(
"Unknown construct kind in VisitSelfClause");
533 void VisitHostClause(
const OpenACCHostClause &clause) {
534 if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
536 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::UpdateHostOp>(
537 var, mlir::acc::DataClause::acc_update_host, {},
540 llvm_unreachable(
"Unknown construct kind in VisitHostClause");
544 void VisitDeviceClause(
const OpenACCDeviceClause &clause) {
545 if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
547 addDataOperand<mlir::acc::UpdateDeviceOp>(
548 var, mlir::acc::DataClause::acc_update_device, {},
551 llvm_unreachable(
"Unknown construct kind in VisitDeviceClause");
555 void VisitIfClause(
const OpenACCIfClause &clause) {
556 if constexpr (isOneOfTypes<
557 OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
558 mlir::acc::KernelsOp, mlir::acc::InitOp,
559 mlir::acc::ShutdownOp, mlir::acc::SetOp,
560 mlir::acc::DataOp, mlir::acc::WaitOp,
561 mlir::acc::HostDataOp, mlir::acc::EnterDataOp,
562 mlir::acc::ExitDataOp, mlir::acc::UpdateOp,
563 mlir::acc::AtomicReadOp, mlir::acc::AtomicWriteOp,
564 mlir::acc::AtomicUpdateOp, mlir::acc::AtomicCaptureOp>) {
565 operation.getIfCondMutable().append(
567 }
else if constexpr (isCombinedType<OpTy>) {
568 applyToComputeOp(clause);
570 llvm_unreachable(
"Unknown construct kind in VisitIfClause");
574 void VisitIfPresentClause(
const OpenACCIfPresentClause &clause) {
575 if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp,
576 mlir::acc::UpdateOp>) {
577 operation.setIfPresent(
true);
579 llvm_unreachable(
"unknown construct kind in VisitIfPresentClause");
583 void VisitDeviceNumClause(
const OpenACCDeviceNumClause &clause) {
584 if constexpr (isOneOfTypes<OpTy, mlir::acc::InitOp, mlir::acc::ShutdownOp,
586 operation.getDeviceNumMutable().append(emitIntExpr(clause.
getIntExpr()));
589 "init, shutdown, set, are only valid device_num constructs");
593 void VisitNumGangsClause(
const OpenACCNumGangsClause &clause) {
594 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
595 mlir::acc::KernelsOp>) {
596 llvm::SmallVector<mlir::Value> values;
598 values.push_back(emitIntExpr(E));
600 operation.addNumGangsOperands(builder.getContext(), values,
601 lastDeviceTypeValues);
602 }
else if constexpr (isCombinedType<OpTy>) {
603 applyToComputeOp(clause);
605 llvm_unreachable(
"Unknown construct kind in VisitNumGangsClause");
609 void VisitWaitClause(
const OpenACCWaitClause &clause) {
610 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
611 mlir::acc::KernelsOp, mlir::acc::DataOp,
612 mlir::acc::EnterDataOp, mlir::acc::ExitDataOp,
613 mlir::acc::UpdateOp>) {
615 operation.addWaitOnly(builder.getContext(), lastDeviceTypeValues);
617 llvm::SmallVector<mlir::Value> values;
621 values.push_back(emitIntExpr(E));
622 operation.addWaitOperands(builder.getContext(), clause.
hasDevNumExpr(),
623 values, lastDeviceTypeValues);
625 }
else if constexpr (isCombinedType<OpTy>) {
626 applyToComputeOp(clause);
630 return clauseNotImplemented(clause);
634 void VisitDefaultAsyncClause(
const OpenACCDefaultAsyncClause &clause) {
635 if constexpr (isOneOfTypes<OpTy, mlir::acc::SetOp>) {
636 operation.getDefaultAsyncMutable().append(
639 llvm_unreachable(
"set, is only valid device_num constructs");
643 void VisitSeqClause(
const OpenACCSeqClause &clause) {
644 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
645 operation.addSeq(builder.getContext(), lastDeviceTypeValues);
646 }
else if constexpr (isCombinedType<OpTy>) {
647 applyToLoopOp(clause);
651 return clauseNotImplemented(clause);
655 void VisitAutoClause(
const OpenACCAutoClause &clause) {
656 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
657 operation.addAuto(builder.getContext(), lastDeviceTypeValues);
658 }
else if constexpr (isCombinedType<OpTy>) {
659 applyToLoopOp(clause);
663 return clauseNotImplemented(clause);
667 void VisitIndependentClause(
const OpenACCIndependentClause &clause) {
668 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
669 operation.addIndependent(builder.getContext(), lastDeviceTypeValues);
670 }
else if constexpr (isCombinedType<OpTy>) {
671 applyToLoopOp(clause);
675 return clauseNotImplemented(clause);
679 void VisitCollapseClause(
const OpenACCCollapseClause &clause) {
680 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
684 value = value.sextOrTrunc(64);
685 operation.setCollapseForDeviceTypes(builder.getContext(),
686 lastDeviceTypeValues, value);
687 }
else if constexpr (isCombinedType<OpTy>) {
688 applyToLoopOp(clause);
690 llvm_unreachable(
"Unknown construct kind in VisitCollapseClause");
694 void VisitTileClause(
const OpenACCTileClause &clause) {
695 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
696 llvm::SmallVector<mlir::Value> values;
699 mlir::Location exprLoc = cgf.cgm.getLoc(e->getBeginLoc());
705 values.push_back(createConstantInt(exprLoc, 64, -1));
707 llvm::APInt curValue =
708 e->EvaluateKnownConstInt(cgf.cgm.getASTContext());
709 values.push_back(createConstantInt(
710 exprLoc, 64, curValue.sextOrTrunc(64).getSExtValue()));
714 operation.setTileForDeviceTypes(builder.getContext(),
715 lastDeviceTypeValues, values);
716 }
else if constexpr (isCombinedType<OpTy>) {
717 applyToLoopOp(clause);
719 llvm_unreachable(
"Unknown construct kind in VisitTileClause");
723 void VisitWorkerClause(
const OpenACCWorkerClause &clause) {
724 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
726 operation.addWorkerNumOperand(builder.getContext(),
728 lastDeviceTypeValues);
730 operation.addEmptyWorker(builder.getContext(), lastDeviceTypeValues);
732 }
else if constexpr (isCombinedType<OpTy>) {
733 applyToLoopOp(clause);
737 return clauseNotImplemented(clause);
741 void VisitVectorClause(
const OpenACCVectorClause &clause) {
742 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
744 operation.addVectorOperand(builder.getContext(),
746 lastDeviceTypeValues);
748 operation.addEmptyVector(builder.getContext(), lastDeviceTypeValues);
750 }
else if constexpr (isCombinedType<OpTy>) {
751 applyToLoopOp(clause);
755 return clauseNotImplemented(clause);
759 void VisitGangClause(
const OpenACCGangClause &clause) {
760 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
762 operation.addEmptyGang(builder.getContext(), lastDeviceTypeValues);
764 llvm::SmallVector<mlir::Value> values;
765 llvm::SmallVector<mlir::acc::GangArgType> argTypes;
766 for (
unsigned i : llvm::index_range(0u, clause.
getNumExprs())) {
768 mlir::Location exprLoc = cgf.cgm.getLoc(
expr->getBeginLoc());
769 argTypes.push_back(decodeGangType(kind));
770 if (kind == OpenACCGangKind::Dim) {
771 llvm::APInt curValue =
772 expr->EvaluateKnownConstInt(cgf.cgm.getASTContext());
775 curValue = curValue.sextOrTrunc(64);
777 createConstantInt(exprLoc, 64, curValue.getSExtValue()));
779 values.push_back(createConstantInt(exprLoc, 64, -1));
781 values.push_back(emitIntExpr(
expr));
785 operation.addGangOperands(builder.getContext(), lastDeviceTypeValues,
788 }
else if constexpr (isCombinedType<OpTy>) {
789 applyToLoopOp(clause);
791 llvm_unreachable(
"Unknown construct kind in VisitGangClause");
795 void VisitCopyClause(
const OpenACCCopyClause &clause) {
796 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
797 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
799 addDataOperand<mlir::acc::CopyinOp, mlir::acc::CopyoutOp>(
803 }
else if constexpr (isCombinedType<OpTy>) {
804 applyToComputeOp(clause);
808 return clauseNotImplemented(clause);
812 void VisitCopyInClause(
const OpenACCCopyInClause &clause) {
813 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
814 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
816 addDataOperand<mlir::acc::CopyinOp, mlir::acc::DeleteOp>(
820 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
822 addDataOperand<mlir::acc::CopyinOp>(
825 }
else if constexpr (isCombinedType<OpTy>) {
826 applyToComputeOp(clause);
830 return clauseNotImplemented(clause);
834 void VisitCopyOutClause(
const OpenACCCopyOutClause &clause) {
835 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
836 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
838 addDataOperand<mlir::acc::CreateOp, mlir::acc::CopyoutOp>(
842 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
844 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::CopyoutOp>(
848 }
else if constexpr (isCombinedType<OpTy>) {
849 applyToComputeOp(clause);
853 return clauseNotImplemented(clause);
857 void VisitCreateClause(
const OpenACCCreateClause &clause) {
858 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
859 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
861 addDataOperand<mlir::acc::CreateOp, mlir::acc::DeleteOp>(
865 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
867 addDataOperand<mlir::acc::CreateOp>(
870 }
else if constexpr (isCombinedType<OpTy>) {
871 applyToComputeOp(clause);
875 return clauseNotImplemented(clause);
879 void VisitDeleteClause(
const OpenACCDeleteClause &clause) {
880 if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
882 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::DeleteOp>(
883 var, mlir::acc::DataClause::acc_delete, {},
887 llvm_unreachable(
"Unknown construct kind in VisitDeleteClause");
891 void VisitDetachClause(
const OpenACCDetachClause &clause) {
892 if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
894 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::DetachOp>(
895 var, mlir::acc::DataClause::acc_detach, {},
899 llvm_unreachable(
"Unknown construct kind in VisitDetachClause");
903 void VisitFinalizeClause(
const OpenACCFinalizeClause &clause) {
904 if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
905 operation.setFinalize(
true);
907 llvm_unreachable(
"Unknown construct kind in VisitFinalizeClause");
911 void VisitUseDeviceClause(
const OpenACCUseDeviceClause &clause) {
912 if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp>) {
914 addDataOperand<mlir::acc::UseDeviceOp>(
915 var, mlir::acc::DataClause::acc_use_device, {},
true,
918 llvm_unreachable(
"Unknown construct kind in VisitUseDeviceClause");
922 void VisitDevicePtrClause(
const OpenACCDevicePtrClause &clause) {
923 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
924 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
926 addDataOperand<mlir::acc::DevicePtrOp>(
927 var, mlir::acc::DataClause::acc_deviceptr, {},
930 }
else if constexpr (isCombinedType<OpTy>) {
931 applyToComputeOp(clause);
935 return clauseNotImplemented(clause);
939 void VisitNoCreateClause(
const OpenACCNoCreateClause &clause) {
940 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
941 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
943 addDataOperand<mlir::acc::NoCreateOp, mlir::acc::DeleteOp>(
944 var, mlir::acc::DataClause::acc_no_create, {},
true,
946 }
else if constexpr (isCombinedType<OpTy>) {
947 applyToComputeOp(clause);
949 llvm_unreachable(
"Unknown construct kind in VisitNoCreateClause");
953 void VisitPresentClause(
const OpenACCPresentClause &clause) {
954 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
955 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
957 addDataOperand<mlir::acc::PresentOp, mlir::acc::DeleteOp>(
958 var, mlir::acc::DataClause::acc_present, {},
true,
960 }
else if constexpr (isCombinedType<OpTy>) {
961 applyToComputeOp(clause);
965 return clauseNotImplemented(clause);
969 void VisitAttachClause(
const OpenACCAttachClause &clause) {
970 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
971 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
973 addDataOperand<mlir::acc::AttachOp, mlir::acc::DetachOp>(
974 var, mlir::acc::DataClause::acc_attach, {},
true,
976 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
978 addDataOperand<mlir::acc::AttachOp>(
979 var, mlir::acc::DataClause::acc_attach, {},
981 }
else if constexpr (isCombinedType<OpTy>) {
982 applyToComputeOp(clause);
984 llvm_unreachable(
"Unknown construct kind in VisitAttachClause");
988 void VisitPrivateClause(
const OpenACCPrivateClause &clause) {
989 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
990 mlir::acc::LoopOp>) {
991 for (
const auto [varExpr, varRecipe] :
993 CIRGenFunction::OpenACCDataOperandInfo opInfo =
994 cgf.getOpenACCDataOperandInfo(varExpr);
995 auto privateOp = mlir::acc::PrivateOp::create(
998 privateOp.setDataClause(mlir::acc::DataClause::acc_private);
1001 mlir::OpBuilder::InsertionGuard guardCase(builder);
1004 OpenACCRecipeBuilder<mlir::acc::PrivateRecipeOp>(cgf, builder)
1006 cgf.getContext(), recipeInsertLocation, varExpr,
1007 varRecipe.AllocaDecl,
1008 nullptr, OpenACCReductionOperator::Invalid,
1015 operation.addPrivatization(builder.getContext(), privateOp, recipe);
1018 }
else if constexpr (isCombinedType<OpTy>) {
1021 applyToLoopOp(clause);
1023 llvm_unreachable(
"Unknown construct kind in VisitPrivateClause");
1027 void VisitFirstPrivateClause(
const OpenACCFirstPrivateClause &clause) {
1028 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
1029 mlir::acc::SerialOp>) {
1030 for (
const auto [varExpr, varRecipe] :
1032 CIRGenFunction::OpenACCDataOperandInfo opInfo =
1033 cgf.getOpenACCDataOperandInfo(varExpr);
1034 auto firstPrivateOp = mlir::acc::FirstprivateOp::create(
1038 firstPrivateOp.setDataClause(mlir::acc::DataClause::acc_firstprivate);
1041 mlir::OpBuilder::InsertionGuard guardCase(builder);
1044 OpenACCRecipeBuilder<mlir::acc::FirstprivateRecipeOp>(cgf,
1047 cgf.getContext(), recipeInsertLocation, varExpr,
1048 varRecipe.AllocaDecl, varRecipe.InitFromTemporary,
1049 OpenACCReductionOperator::Invalid,
1052 firstPrivateOp, {});
1057 operation.addFirstPrivatization(builder.getContext(), firstPrivateOp,
1061 }
else if constexpr (isCombinedType<OpTy>) {
1064 applyToComputeOp(clause);
1066 llvm_unreachable(
"Unknown construct kind in VisitFirstPrivateClause");
1070 void VisitReductionClause(
const OpenACCReductionClause &clause) {
1071 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
1072 mlir::acc::LoopOp>) {
1073 for (
const auto [varExpr, varRecipe] :
1075 CIRGenFunction::OpenACCDataOperandInfo opInfo =
1076 cgf.getOpenACCDataOperandInfo(varExpr);
1078 auto reductionOp = mlir::acc::ReductionOp::create(
1081 reductionOp.setDataClause(mlir::acc::DataClause::acc_reduction);
1084 mlir::OpBuilder::InsertionGuard guardCase(builder);
1087 OpenACCRecipeBuilder<mlir::acc::ReductionRecipeOp>(cgf, builder)
1089 cgf.getContext(), recipeInsertLocation, varExpr,
1090 varRecipe.AllocaDecl,
1094 reductionOp, varRecipe.CombinerRecipes);
1096 operation.addReduction(builder.getContext(), reductionOp, recipe);
1099 }
else if constexpr (isCombinedType<OpTy>) {
1102 applyToLoopOp(clause);
1104 llvm_unreachable(
"Unknown construct kind in VisitReductionClause");
1109template <
typename OpTy>
1110auto makeClauseEmitter(OpTy &op,
1111 mlir::OpBuilder::InsertPoint &recipeInsertLocation,
1115 return OpenACCClauseCIREmitter<OpTy>(op, recipeInsertLocation, cgf, builder,
1120template <
typename Op>
1121void CIRGenFunction::emitOpenACCClauses(
1124 mlir::OpBuilder::InsertionGuard guardCase(builder);
1128 builder.setInsertionPoint(op);
1129 makeClauseEmitter(op, lastRecipeLocation, *
this, builder, dirKind, dirLoc)
1130 .emitClauses(clauses);
1133#define EXPL_SPEC(N) \
1134 template void CIRGenFunction::emitOpenACCClauses<N>( \
1135 N &, OpenACCDirectiveKind, SourceLocation, \
1136 ArrayRef<const OpenACCClause *>);
1156template <
typename ComputeOp,
typename LoopOp>
1157void CIRGenFunction::emitOpenACCClauses(
1160 static_assert(std::is_same_v<mlir::acc::LoopOp, LoopOp>);
1162 CombinedConstructClauseInfo<ComputeOp> inf{op, loopOp};
1165 mlir::OpBuilder::InsertionGuard guardCase(builder);
1166 makeClauseEmitter(inf, lastRecipeLocation, *
this, builder, dirKind, dirLoc)
1167 .emitClauses(clauses);
1170#define EXPL_SPEC(N) \
1171 template void CIRGenFunction::emitOpenACCClauses<N, mlir::acc::LoopOp>( \
1172 N &, mlir::acc::LoopOp &, OpenACCDirectiveKind, SourceLocation, \
1173 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
Encodes a location in the source.
SourceLocation getBeginLoc() const LLVM_READONLY
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
bool isOpenACCModifierBitSet(OpenACCModifierKind List, OpenACCModifierKind Bit)
__DEVICE__ _Tp arg(const std::complex< _Tp > &__c)
llvm::SmallVector< mlir::Value > bounds
llvm::SmallVector< QualType > boundTypes