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 = builder.create<mlir::UnrealizedConversionCastOp>(
100 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 = builder.create<mlir::arith::ConstantOp>(
111 loc, builder.getIntegerAttr(ty, value));
113 return constOp.getResult();
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);
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>) {
233 afterOp = builder.create<AfterOpTy>(
234 opInfo.
beginLoc, beforeOp.getResult(), structured, implicit,
237 afterOp = builder.create<AfterOpTy>(
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);
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<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
557 mlir::acc::KernelsOp, mlir::acc::InitOp,
558 mlir::acc::ShutdownOp, mlir::acc::SetOp,
559 mlir::acc::DataOp, mlir::acc::WaitOp,
560 mlir::acc::HostDataOp, mlir::acc::EnterDataOp,
561 mlir::acc::ExitDataOp, mlir::acc::UpdateOp>) {
562 operation.getIfCondMutable().append(
564 }
else if constexpr (isCombinedType<OpTy>) {
565 applyToComputeOp(clause);
567 llvm_unreachable(
"Unknown construct kind in VisitIfClause");
571 void VisitIfPresentClause(
const OpenACCIfPresentClause &clause) {
572 if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp,
573 mlir::acc::UpdateOp>) {
574 operation.setIfPresent(
true);
576 llvm_unreachable(
"unknown construct kind in VisitIfPresentClause");
580 void VisitDeviceNumClause(
const OpenACCDeviceNumClause &clause) {
581 if constexpr (isOneOfTypes<OpTy, mlir::acc::InitOp, mlir::acc::ShutdownOp,
583 operation.getDeviceNumMutable().append(emitIntExpr(clause.
getIntExpr()));
586 "init, shutdown, set, are only valid device_num constructs");
590 void VisitNumGangsClause(
const OpenACCNumGangsClause &clause) {
591 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
592 mlir::acc::KernelsOp>) {
593 llvm::SmallVector<mlir::Value> values;
595 values.push_back(emitIntExpr(E));
597 operation.addNumGangsOperands(builder.getContext(), values,
598 lastDeviceTypeValues);
599 }
else if constexpr (isCombinedType<OpTy>) {
600 applyToComputeOp(clause);
602 llvm_unreachable(
"Unknown construct kind in VisitNumGangsClause");
606 void VisitWaitClause(
const OpenACCWaitClause &clause) {
607 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
608 mlir::acc::KernelsOp, mlir::acc::DataOp,
609 mlir::acc::EnterDataOp, mlir::acc::ExitDataOp,
610 mlir::acc::UpdateOp>) {
612 operation.addWaitOnly(builder.getContext(), lastDeviceTypeValues);
614 llvm::SmallVector<mlir::Value> values;
618 values.push_back(emitIntExpr(E));
619 operation.addWaitOperands(builder.getContext(), clause.
hasDevNumExpr(),
620 values, lastDeviceTypeValues);
622 }
else if constexpr (isCombinedType<OpTy>) {
623 applyToComputeOp(clause);
627 return clauseNotImplemented(clause);
631 void VisitDefaultAsyncClause(
const OpenACCDefaultAsyncClause &clause) {
632 if constexpr (isOneOfTypes<OpTy, mlir::acc::SetOp>) {
633 operation.getDefaultAsyncMutable().append(
636 llvm_unreachable(
"set, is only valid device_num constructs");
640 void VisitSeqClause(
const OpenACCSeqClause &clause) {
641 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
642 operation.addSeq(builder.getContext(), lastDeviceTypeValues);
643 }
else if constexpr (isCombinedType<OpTy>) {
644 applyToLoopOp(clause);
648 return clauseNotImplemented(clause);
652 void VisitAutoClause(
const OpenACCAutoClause &clause) {
653 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
654 operation.addAuto(builder.getContext(), lastDeviceTypeValues);
655 }
else if constexpr (isCombinedType<OpTy>) {
656 applyToLoopOp(clause);
660 return clauseNotImplemented(clause);
664 void VisitIndependentClause(
const OpenACCIndependentClause &clause) {
665 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
666 operation.addIndependent(builder.getContext(), lastDeviceTypeValues);
667 }
else if constexpr (isCombinedType<OpTy>) {
668 applyToLoopOp(clause);
672 return clauseNotImplemented(clause);
676 void VisitCollapseClause(
const OpenACCCollapseClause &clause) {
677 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
681 value = value.sextOrTrunc(64);
682 operation.setCollapseForDeviceTypes(builder.getContext(),
683 lastDeviceTypeValues, value);
684 }
else if constexpr (isCombinedType<OpTy>) {
685 applyToLoopOp(clause);
687 llvm_unreachable(
"Unknown construct kind in VisitCollapseClause");
691 void VisitTileClause(
const OpenACCTileClause &clause) {
692 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
693 llvm::SmallVector<mlir::Value> values;
696 mlir::Location exprLoc = cgf.cgm.getLoc(e->getBeginLoc());
702 values.push_back(createConstantInt(exprLoc, 64, -1));
704 llvm::APInt curValue =
705 e->EvaluateKnownConstInt(cgf.cgm.getASTContext());
706 values.push_back(createConstantInt(
707 exprLoc, 64, curValue.sextOrTrunc(64).getSExtValue()));
711 operation.setTileForDeviceTypes(builder.getContext(),
712 lastDeviceTypeValues, values);
713 }
else if constexpr (isCombinedType<OpTy>) {
714 applyToLoopOp(clause);
716 llvm_unreachable(
"Unknown construct kind in VisitTileClause");
720 void VisitWorkerClause(
const OpenACCWorkerClause &clause) {
721 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
723 operation.addWorkerNumOperand(builder.getContext(),
725 lastDeviceTypeValues);
727 operation.addEmptyWorker(builder.getContext(), lastDeviceTypeValues);
729 }
else if constexpr (isCombinedType<OpTy>) {
730 applyToLoopOp(clause);
734 return clauseNotImplemented(clause);
738 void VisitVectorClause(
const OpenACCVectorClause &clause) {
739 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
741 operation.addVectorOperand(builder.getContext(),
743 lastDeviceTypeValues);
745 operation.addEmptyVector(builder.getContext(), lastDeviceTypeValues);
747 }
else if constexpr (isCombinedType<OpTy>) {
748 applyToLoopOp(clause);
752 return clauseNotImplemented(clause);
756 void VisitGangClause(
const OpenACCGangClause &clause) {
757 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
759 operation.addEmptyGang(builder.getContext(), lastDeviceTypeValues);
761 llvm::SmallVector<mlir::Value> values;
762 llvm::SmallVector<mlir::acc::GangArgType> argTypes;
763 for (
unsigned i : llvm::index_range(0u, clause.
getNumExprs())) {
765 mlir::Location exprLoc = cgf.cgm.getLoc(
expr->getBeginLoc());
766 argTypes.push_back(decodeGangType(kind));
767 if (kind == OpenACCGangKind::Dim) {
768 llvm::APInt curValue =
769 expr->EvaluateKnownConstInt(cgf.cgm.getASTContext());
772 curValue = curValue.sextOrTrunc(64);
774 createConstantInt(exprLoc, 64, curValue.getSExtValue()));
776 values.push_back(createConstantInt(exprLoc, 64, -1));
778 values.push_back(emitIntExpr(
expr));
782 operation.addGangOperands(builder.getContext(), lastDeviceTypeValues,
785 }
else if constexpr (isCombinedType<OpTy>) {
786 applyToLoopOp(clause);
788 llvm_unreachable(
"Unknown construct kind in VisitGangClause");
792 void VisitCopyClause(
const OpenACCCopyClause &clause) {
793 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
794 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
796 addDataOperand<mlir::acc::CopyinOp, mlir::acc::CopyoutOp>(
800 }
else if constexpr (isCombinedType<OpTy>) {
801 applyToComputeOp(clause);
805 return clauseNotImplemented(clause);
809 void VisitCopyInClause(
const OpenACCCopyInClause &clause) {
810 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
811 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
813 addDataOperand<mlir::acc::CopyinOp, mlir::acc::DeleteOp>(
817 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
819 addDataOperand<mlir::acc::CopyinOp>(
822 }
else if constexpr (isCombinedType<OpTy>) {
823 applyToComputeOp(clause);
827 return clauseNotImplemented(clause);
831 void VisitCopyOutClause(
const OpenACCCopyOutClause &clause) {
832 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
833 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
835 addDataOperand<mlir::acc::CreateOp, mlir::acc::CopyoutOp>(
839 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
841 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::CopyoutOp>(
845 }
else if constexpr (isCombinedType<OpTy>) {
846 applyToComputeOp(clause);
850 return clauseNotImplemented(clause);
854 void VisitCreateClause(
const OpenACCCreateClause &clause) {
855 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
856 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
858 addDataOperand<mlir::acc::CreateOp, mlir::acc::DeleteOp>(
862 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
864 addDataOperand<mlir::acc::CreateOp>(
867 }
else if constexpr (isCombinedType<OpTy>) {
868 applyToComputeOp(clause);
872 return clauseNotImplemented(clause);
876 void VisitDeleteClause(
const OpenACCDeleteClause &clause) {
877 if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
879 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::DeleteOp>(
880 var, mlir::acc::DataClause::acc_delete, {},
884 llvm_unreachable(
"Unknown construct kind in VisitDeleteClause");
888 void VisitDetachClause(
const OpenACCDetachClause &clause) {
889 if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
891 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::DetachOp>(
892 var, mlir::acc::DataClause::acc_detach, {},
896 llvm_unreachable(
"Unknown construct kind in VisitDetachClause");
900 void VisitFinalizeClause(
const OpenACCFinalizeClause &clause) {
901 if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
902 operation.setFinalize(
true);
904 llvm_unreachable(
"Unknown construct kind in VisitFinalizeClause");
908 void VisitUseDeviceClause(
const OpenACCUseDeviceClause &clause) {
909 if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp>) {
911 addDataOperand<mlir::acc::UseDeviceOp>(
912 var, mlir::acc::DataClause::acc_use_device, {},
true,
915 llvm_unreachable(
"Unknown construct kind in VisitUseDeviceClause");
919 void VisitDevicePtrClause(
const OpenACCDevicePtrClause &clause) {
920 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
921 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
923 addDataOperand<mlir::acc::DevicePtrOp>(
924 var, mlir::acc::DataClause::acc_deviceptr, {},
927 }
else if constexpr (isCombinedType<OpTy>) {
928 applyToComputeOp(clause);
932 return clauseNotImplemented(clause);
936 void VisitNoCreateClause(
const OpenACCNoCreateClause &clause) {
937 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
938 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
940 addDataOperand<mlir::acc::NoCreateOp, mlir::acc::DeleteOp>(
941 var, mlir::acc::DataClause::acc_no_create, {},
true,
943 }
else if constexpr (isCombinedType<OpTy>) {
944 applyToComputeOp(clause);
946 llvm_unreachable(
"Unknown construct kind in VisitNoCreateClause");
950 void VisitPresentClause(
const OpenACCPresentClause &clause) {
951 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
952 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
954 addDataOperand<mlir::acc::PresentOp, mlir::acc::DeleteOp>(
955 var, mlir::acc::DataClause::acc_present, {},
true,
957 }
else if constexpr (isCombinedType<OpTy>) {
958 applyToComputeOp(clause);
962 return clauseNotImplemented(clause);
966 void VisitAttachClause(
const OpenACCAttachClause &clause) {
967 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
968 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
970 addDataOperand<mlir::acc::AttachOp, mlir::acc::DetachOp>(
971 var, mlir::acc::DataClause::acc_attach, {},
true,
973 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
975 addDataOperand<mlir::acc::AttachOp>(
976 var, mlir::acc::DataClause::acc_attach, {},
978 }
else if constexpr (isCombinedType<OpTy>) {
979 applyToComputeOp(clause);
981 llvm_unreachable(
"Unknown construct kind in VisitAttachClause");
985 void VisitPrivateClause(
const OpenACCPrivateClause &clause) {
986 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
987 mlir::acc::LoopOp>) {
988 for (
const auto [varExpr, varRecipe] :
990 CIRGenFunction::OpenACCDataOperandInfo opInfo =
991 cgf.getOpenACCDataOperandInfo(varExpr);
992 auto privateOp = mlir::acc::PrivateOp::create(
995 privateOp.setDataClause(mlir::acc::DataClause::acc_private);
998 mlir::OpBuilder::InsertionGuard guardCase(builder);
1001 OpenACCRecipeBuilder<mlir::acc::PrivateRecipeOp>(cgf, builder)
1003 cgf.getContext(), recipeInsertLocation, varExpr,
1004 varRecipe.AllocaDecl, varRecipe.InitExpr,
1005 nullptr, OpenACCReductionOperator::Invalid,
1008 privateOp.getResult());
1012 operation.addPrivatization(builder.getContext(), privateOp, recipe);
1015 }
else if constexpr (isCombinedType<OpTy>) {
1018 applyToLoopOp(clause);
1020 llvm_unreachable(
"Unknown construct kind in VisitPrivateClause");
1024 void VisitFirstPrivateClause(
const OpenACCFirstPrivateClause &clause) {
1025 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
1026 mlir::acc::SerialOp>) {
1027 for (
const auto [varExpr, varRecipe] :
1029 CIRGenFunction::OpenACCDataOperandInfo opInfo =
1030 cgf.getOpenACCDataOperandInfo(varExpr);
1031 auto firstPrivateOp = mlir::acc::FirstprivateOp::create(
1035 firstPrivateOp.setDataClause(mlir::acc::DataClause::acc_firstprivate);
1038 mlir::OpBuilder::InsertionGuard guardCase(builder);
1042 VarDecl *allocaDecl = varRecipe.AllocaDecl;
1043 allocaDecl->
setInit(varRecipe.InitExpr);
1047 OpenACCRecipeBuilder<mlir::acc::FirstprivateRecipeOp>(cgf,
1050 cgf.getContext(), recipeInsertLocation, varExpr,
1051 varRecipe.AllocaDecl, varRecipe.InitExpr,
1052 varRecipe.InitFromTemporary,
1053 OpenACCReductionOperator::Invalid,
1056 firstPrivateOp.getResult());
1061 operation.addFirstPrivatization(builder.getContext(), firstPrivateOp,
1065 }
else if constexpr (isCombinedType<OpTy>) {
1068 applyToComputeOp(clause);
1070 llvm_unreachable(
"Unknown construct kind in VisitFirstPrivateClause");
1074 void VisitReductionClause(
const OpenACCReductionClause &clause) {
1075 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
1076 mlir::acc::LoopOp>) {
1077 for (
const auto [varExpr, varRecipe] :
1079 CIRGenFunction::OpenACCDataOperandInfo opInfo =
1080 cgf.getOpenACCDataOperandInfo(varExpr);
1082 auto reductionOp = mlir::acc::ReductionOp::create(
1085 reductionOp.setDataClause(mlir::acc::DataClause::acc_reduction);
1088 mlir::OpBuilder::InsertionGuard guardCase(builder);
1092 VarDecl *allocaDecl = varRecipe.AllocaDecl;
1093 allocaDecl->
setInit(varRecipe.InitExpr);
1097 OpenACCRecipeBuilder<mlir::acc::ReductionRecipeOp>(cgf, builder)
1099 cgf.getContext(), recipeInsertLocation, varExpr,
1100 varRecipe.AllocaDecl, varRecipe.InitExpr,
1104 reductionOp.getResult());
1106 operation.addReduction(builder.getContext(), reductionOp, recipe);
1109 }
else if constexpr (isCombinedType<OpTy>) {
1112 applyToLoopOp(clause);
1114 llvm_unreachable(
"Unknown construct kind in VisitReductionClause");
1119template <
typename OpTy>
1120auto makeClauseEmitter(OpTy &op,
1121 mlir::OpBuilder::InsertPoint &recipeInsertLocation,
1125 return OpenACCClauseCIREmitter<OpTy>(op, recipeInsertLocation, cgf, builder,
1130template <
typename Op>
1131void CIRGenFunction::emitOpenACCClauses(
1134 mlir::OpBuilder::InsertionGuard guardCase(builder);
1138 builder.setInsertionPoint(op);
1139 makeClauseEmitter(op, lastRecipeLocation, *
this, builder, dirKind, dirLoc)
1140 .emitClauses(clauses);
1143#define EXPL_SPEC(N) \
1144 template void CIRGenFunction::emitOpenACCClauses<N>( \
1145 N &, OpenACCDirectiveKind, SourceLocation, \
1146 ArrayRef<const OpenACCClause *>);
1162template <
typename ComputeOp,
typename LoopOp>
1163void CIRGenFunction::emitOpenACCClauses(
1166 static_assert(std::is_same_v<mlir::acc::LoopOp, LoopOp>);
1168 CombinedConstructClauseInfo<ComputeOp> inf{op, loopOp};
1171 mlir::OpBuilder::InsertionGuard guardCase(builder);
1172 makeClauseEmitter(inf, lastRecipeLocation, *
this, builder, dirKind, dirLoc)
1173 .emitClauses(clauses);
1176#define EXPL_SPEC(N) \
1177 template void CIRGenFunction::emitOpenACCClauses<N, mlir::acc::LoopOp>( \
1178 N &, mlir::acc::LoopOp &, OpenACCDirectiveKind, SourceLocation, \
1179 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
void setInitStyle(InitializationStyle Style)
@ CallInit
Call-style initialization (C++98)
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