20#include "mlir/Dialect/Arith/IR/Arith.h"
21#include "mlir/Dialect/OpenACC/OpenACC.h"
22#include "llvm/ADT/TypeSwitch.h"
30template <
typename ToTest,
typename T,
typename... Tys>
31constexpr bool isOneOfTypes =
32 std::is_same_v<ToTest, T> || isOneOfTypes<ToTest, Tys...>;
33template <
typename ToTest,
typename T>
34constexpr bool isOneOfTypes<ToTest, T> = std::is_same_v<ToTest, T>;
39template <
typename CompOpTy>
struct CombinedConstructClauseInfo {
40 using ComputeOpTy = CompOpTy;
41 ComputeOpTy computeOp;
42 mlir::acc::LoopOp loopOp;
44template <
typename ToTest>
constexpr bool isCombinedType =
false;
46constexpr bool isCombinedType<CombinedConstructClauseInfo<T>> =
true;
48template <
typename OpTy>
49class OpenACCClauseCIREmitter final
52 template <
typename FriendOpTy>
friend class OpenACCClauseCIREmitter;
68 bool hasAsyncClause =
false;
73 cgf.
cgm.
errorNYI(
c.getSourceRange(),
"OpenACC Clause",
c.getClauseKind());
77 lastDeviceTypeValues.clear();
80 lastDeviceTypeValues.push_back(decodeDeviceType(
arg.getIdentifierInfo()));
83 mlir::Value emitIntExpr(
const Expr *intExpr) {
91 mlir::Value createCondition(
const Expr *condExpr) {
94 mlir::IntegerType targetType = mlir::IntegerType::get(
96 mlir::IntegerType::SignednessSemantics::Signless);
97 auto conversionOp = builder.create<mlir::UnrealizedConversionCastOp>(
98 exprLoc, targetType, condition);
99 return conversionOp.getResult(0);
102 mlir::Value createConstantInt(mlir::Location loc,
unsigned width,
105 mlir::IntegerType ty = mlir::IntegerType::get(
107 mlir::IntegerType::SignednessSemantics::Signless);
108 auto constOp = builder.create<mlir::arith::ConstantOp>(
109 loc, builder.getIntegerAttr(ty, value));
111 return constOp.getResult();
116 return createConstantInt(cgf.
cgm.
getLoc(loc), width, value);
119 mlir::acc::DeviceType decodeDeviceType(
const IdentifierInfo *ii) {
122 return mlir::acc::DeviceType::Star;
123 return llvm::StringSwitch<mlir::acc::DeviceType>(ii->
getName())
124 .CaseLower(
"default", mlir::acc::DeviceType::Default)
125 .CaseLower(
"host", mlir::acc::DeviceType::Host)
126 .CaseLower(
"multicore", mlir::acc::DeviceType::Multicore)
127 .CasesLower(
"nvidia",
"acc_device_nvidia",
128 mlir::acc::DeviceType::Nvidia)
129 .CaseLower(
"radeon", mlir::acc::DeviceType::Radeon);
135 return mlir::acc::GangArgType::Num;
137 return mlir::acc::GangArgType::Dim;
139 return mlir::acc::GangArgType::Static;
141 llvm_unreachable(
"unknown gang kind");
144 template <
typename U = void,
145 typename = std::enable_if_t<isCombinedType<OpTy>,
U>>
147 mlir::OpBuilder::InsertionGuard guardCase(builder);
148 builder.setInsertionPoint(operation.loopOp);
149 OpenACCClauseCIREmitter<mlir::acc::LoopOp> loopEmitter{
150 operation.loopOp, cgf, builder, dirKind, dirLoc};
151 loopEmitter.lastDeviceTypeValues = lastDeviceTypeValues;
152 loopEmitter.Visit(&
c);
155 template <
typename U = void,
156 typename = std::enable_if_t<isCombinedType<OpTy>,
U>>
158 mlir::OpBuilder::InsertionGuard guardCase(builder);
159 builder.setInsertionPoint(operation.computeOp);
160 OpenACCClauseCIREmitter<typename OpTy::ComputeOpTy> computeEmitter{
161 operation.computeOp, cgf, builder, dirKind, dirLoc};
163 computeEmitter.lastDeviceTypeValues = lastDeviceTypeValues;
168 if (!dataOperands.empty())
169 computeEmitter.dataOperands.push_back(dataOperands.front());
170 computeEmitter.Visit(&
c);
175 dataOperands.append(computeEmitter.dataOperands);
178 mlir::acc::DataClauseModifier
182 static_cast<int>(DataClauseModifier::zero) &&
184 static_cast<int>(DataClauseModifier::readonly) &&
186 static_cast<int>(DataClauseModifier::alwaysin) &&
188 static_cast<int>(DataClauseModifier::alwaysout) &&
190 static_cast<int>(DataClauseModifier::capture));
192 DataClauseModifier mlirModifiers{};
197 mlirModifiers = mlirModifiers | DataClauseModifier::always;
201 mlirModifiers = mlirModifiers |
static_cast<DataClauseModifier
>(modifiers);
202 return mlirModifiers;
205 template <
typename BeforeOpTy,
typename AfterOpTy>
206 void addDataOperand(
const Expr *varOperand, mlir::acc::DataClause dataClause,
215 operation.getDataClauseOperandsMutable().append(beforeOp.getResult());
219 mlir::OpBuilder::InsertionGuard guardCase(builder);
220 builder.setInsertionPointAfter(operation);
222 if constexpr (std::is_same_v<AfterOpTy, mlir::acc::DeleteOp> ||
223 std::is_same_v<AfterOpTy, mlir::acc::DetachOp>) {
226 afterOp = builder.create<AfterOpTy>(
227 opInfo.
beginLoc, beforeOp.getResult(), structured, implicit,
230 afterOp = builder.create<AfterOpTy>(
237 beforeOp.setDataClause(dataClause);
238 afterOp.setDataClause(dataClause);
239 beforeOp.setModifiers(convertModifiers(modifiers));
240 afterOp.setModifiers(convertModifiers(modifiers));
243 dataOperands.push_back(beforeOp.getOperation());
244 dataOperands.push_back(afterOp.getOperation());
247 template <
typename BeforeOpTy>
248 void addDataOperand(
const Expr *varOperand, mlir::acc::DataClause dataClause,
256 operation.getDataClauseOperandsMutable().append(beforeOp.getResult());
259 beforeOp.setDataClause(dataClause);
260 beforeOp.setModifiers(convertModifiers(modifiers));
263 dataOperands.push_back(beforeOp.getOperation());
268 mlir::ArrayAttr getAsyncOnlyAttr() {
269 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
270 mlir::acc::KernelsOp, mlir::acc::DataOp,
271 mlir::acc::UpdateOp>) {
272 return operation.getAsyncOnlyAttr();
273 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp,
274 mlir::acc::ExitDataOp>) {
275 if (!operation.getAsyncAttr())
276 return mlir::ArrayAttr{};
279 devTysTemp.push_back(mlir::acc::DeviceTypeAttr::get(
280 builder.getContext(), mlir::acc::DeviceType::None));
281 return mlir::ArrayAttr::get(builder.getContext(), devTysTemp);
282 }
else if constexpr (isCombinedType<OpTy>) {
283 return operation.computeOp.getAsyncOnlyAttr();
289 llvm_unreachable(
"getting asyncOnly when clause not valid on operation?");
294 mlir::ArrayAttr getAsyncOperandsDeviceTypeAttr() {
295 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
296 mlir::acc::KernelsOp, mlir::acc::DataOp,
297 mlir::acc::UpdateOp>) {
298 return operation.getAsyncOperandsDeviceTypeAttr();
299 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp,
300 mlir::acc::ExitDataOp>) {
301 if (!operation.getAsyncOperand())
302 return mlir::ArrayAttr{};
305 devTysTemp.push_back(mlir::acc::DeviceTypeAttr::get(
306 builder.getContext(), mlir::acc::DeviceType::None));
307 return mlir::ArrayAttr::get(builder.getContext(), devTysTemp);
308 }
else if constexpr (isCombinedType<OpTy>) {
309 return operation.computeOp.getAsyncOperandsDeviceTypeAttr();
316 "getting asyncOperandsDeviceType when clause not valid on operation?");
321 mlir::OperandRange getAsyncOperands() {
322 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
323 mlir::acc::KernelsOp, mlir::acc::DataOp,
324 mlir::acc::UpdateOp>)
325 return operation.getAsyncOperands();
326 else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp,
327 mlir::acc::ExitDataOp>)
328 return operation.getAsyncOperandMutable();
329 else if constexpr (isCombinedType<OpTy>)
330 return operation.computeOp.getAsyncOperands();
336 "getting asyncOperandsDeviceType when clause not valid on operation?");
342 void updateDataOperandAsyncValues() {
343 if (!hasAsyncClause || dataOperands.empty())
346 for (mlir::Operation *dataOp : dataOperands) {
347 llvm::TypeSwitch<mlir::Operation *, void>(dataOp)
348 .Case<ACC_DATA_ENTRY_OPS, ACC_DATA_EXIT_OPS>([&](
auto op) {
349 op.setAsyncOnlyAttr(getAsyncOnlyAttr());
350 op.setAsyncOperandsDeviceTypeAttr(getAsyncOperandsDeviceTypeAttr());
351 op.getAsyncOperandsMutable().assign(getAsyncOperands());
353 .
Default([&](mlir::Operation *) {
354 llvm_unreachable(
"Not a data operation?");
359 template <
typename RecipeTy>
362 std::string recipeName;
364 llvm::raw_string_ostream stream(recipeName);
366 if constexpr (std::is_same_v<RecipeTy, mlir::acc::PrivateRecipeOp>) {
367 stream <<
"privatization_";
368 }
else if constexpr (std::is_same_v<RecipeTy,
369 mlir::acc::FirstprivateRecipeOp>) {
370 stream <<
"firstprivatization_";
372 }
else if constexpr (std::is_same_v<RecipeTy,
373 mlir::acc::ReductionRecipeOp>) {
374 stream <<
"reduction_";
378 switch (reductionOp) {
407 llvm_unreachable(
"invalid reduction operator");
410 static_assert(!
sizeof(RecipeTy),
"Unknown Recipe op kind");
419 void createFirstprivateRecipeCopy(
420 mlir::Location loc, mlir::Location locEnd, mlir::Value mainOp,
422 mlir::acc::FirstprivateRecipeOp recipe,
const VarDecl *varRecipe,
424 mlir::Block *block = builder.createBlock(
425 &recipe.getCopyRegion(), recipe.getCopyRegion().end(),
426 {mainOp.getType(), mainOp.getType()}, {loc, loc});
427 builder.setInsertionPointToEnd(&recipe.getCopyRegion().back());
430 mlir::BlockArgument fromArg = block->getArgument(0);
431 mlir::BlockArgument toArg = block->getArgument(1);
433 mlir::Type elementTy =
434 mlir::cast<cir::PointerType>(mainOp.getType()).getPointee();
448 mlir::acc::YieldOp::create(builder, locEnd);
455 template <
typename RecipeTy>
456 void createRecipeInitCopy(mlir::Location loc, mlir::Location locEnd,
458 RecipeTy recipe,
const VarDecl *varRecipe,
460 assert(varRecipe &&
"Required recipe variable not set?");
468 mlir::Block *block = builder.createBlock(&recipe.getInitRegion(),
469 recipe.getInitRegion().end(),
470 {mainOp.getType()}, {loc});
471 builder.setInsertionPointToEnd(&recipe.getInitRegion().back());
482 if constexpr (std::is_same_v<RecipeTy, mlir::acc::PrivateRecipeOp>) {
495 cgf.
cgm.
errorNYI(exprRange,
"private default-init recipe");
498 }
else if constexpr (std::is_same_v<RecipeTy,
499 mlir::acc::ReductionRecipeOp>) {
503 cgf.
cgm.
errorNYI(exprRange,
"reduction init recipe");
507 mlir::acc::YieldOp::create(builder, locEnd);
509 if constexpr (std::is_same_v<RecipeTy, mlir::acc::FirstprivateRecipeOp>) {
517 exprRange,
"firstprivate copy-init recipe not properly generated");
520 createFirstprivateRecipeCopy(loc, locEnd, mainOp, tempDeclEmission,
521 recipe, varRecipe, temporary);
529 void createReductionRecipeCombiner(mlir::Location loc, mlir::Location locEnd,
531 mlir::acc::ReductionRecipeOp recipe) {
532 mlir::Block *block = builder.createBlock(
533 &recipe.getCombinerRegion(), recipe.getCombinerRegion().end(),
534 {mainOp.getType(), mainOp.getType()}, {loc, loc});
535 builder.setInsertionPointToEnd(&recipe.getCombinerRegion().back());
538 mlir::BlockArgument lhsArg = block->getArgument(0);
540 mlir::acc::YieldOp::create(builder, locEnd, lhsArg);
547 void createRecipeDestroySection(mlir::Location loc, mlir::Location locEnd,
550 mlir::Region &destroyRegion) {
551 mlir::Block *block = builder.createBlock(
552 &destroyRegion, destroyRegion.end(), {mainOp.getType()}, {loc});
553 builder.setInsertionPointToEnd(&destroyRegion.back());
556 mlir::Type elementTy =
557 mlir::cast<cir::PointerType>(mainOp.getType()).getPointee();
558 Address addr{block->getArgument(0), elementTy, alignment};
562 mlir::acc::YieldOp::create(builder, locEnd);
568 return mlir::acc::ReductionOperator::AccAdd;
570 return mlir::acc::ReductionOperator::AccMul;
572 return mlir::acc::ReductionOperator::AccMax;
574 return mlir::acc::ReductionOperator::AccMin;
576 return mlir::acc::ReductionOperator::AccIand;
578 return mlir::acc::ReductionOperator::AccIor;
580 return mlir::acc::ReductionOperator::AccXor;
582 return mlir::acc::ReductionOperator::AccLand;
584 return mlir::acc::ReductionOperator::AccLor;
586 llvm_unreachable(
"invalid reduction operator");
589 llvm_unreachable(
"invalid reduction operator");
592 template <
typename RecipeTy>
593 RecipeTy getOrCreateRecipe(
ASTContext &astCtx,
const Expr *varRef,
597 mlir::Value mainOp) {
598 mlir::ModuleOp mod = builder.getBlock()
600 ->template getParentOfType<mlir::ModuleOp>();
602 std::string recipeName = getRecipeName<RecipeTy>(varRef->
getSourceRange(),
603 baseType, reductionOp);
604 if (
auto recipe = mod.lookupSymbol<RecipeTy>(recipeName))
610 mlir::OpBuilder modBuilder(mod.getBodyRegion());
613 if constexpr (std::is_same_v<RecipeTy, mlir::acc::ReductionRecipeOp>) {
614 recipe = RecipeTy::create(modBuilder, loc, recipeName, mainOp.getType(),
615 convertReductionOp(reductionOp));
617 recipe = RecipeTy::create(modBuilder, loc, recipeName, mainOp.getType());
620 createRecipeInitCopy(loc, locEnd, varRef->
getSourceRange(), mainOp, recipe,
621 varRecipe, temporary);
623 if constexpr (std::is_same_v<RecipeTy, mlir::acc::ReductionRecipeOp>) {
624 createReductionRecipeCombiner(loc, locEnd, mainOp, recipe);
628 createRecipeDestroySection(loc, locEnd, mainOp,
630 baseType, recipe.getDestroyRegion());
638 : operation(operation), cgf(cgf), builder(builder), dirKind(dirKind),
642 clauseNotImplemented(clause);
650 updateDataOperandAsyncValues();
656 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
657 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
660 operation.setDefaultAttr(mlir::acc::ClauseDefaultValue::None);
663 operation.setDefaultAttr(mlir::acc::ClauseDefaultValue::Present);
668 }
else if constexpr (isCombinedType<OpTy>) {
669 applyToComputeOp(clause);
671 llvm_unreachable(
"Unknown construct kind in VisitDefaultClause");
676 setLastDeviceTypeClause(clause);
678 if constexpr (isOneOfTypes<OpTy, mlir::acc::InitOp,
679 mlir::acc::ShutdownOp>) {
681 operation.addDeviceType(builder.getContext(),
682 decodeDeviceType(
arg.getIdentifierInfo()));
683 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::SetOp>) {
684 assert(!operation.getDeviceTypeAttr() &&
"already have device-type?");
688 operation.setDeviceType(
690 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
691 mlir::acc::SerialOp, mlir::acc::KernelsOp,
692 mlir::acc::DataOp, mlir::acc::LoopOp,
693 mlir::acc::UpdateOp>) {
697 }
else if constexpr (isCombinedType<OpTy>) {
703 return clauseNotImplemented(clause);
708 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
709 mlir::acc::KernelsOp>) {
710 operation.addNumWorkersOperand(builder.getContext(),
712 lastDeviceTypeValues);
713 }
else if constexpr (isCombinedType<OpTy>) {
714 applyToComputeOp(clause);
716 llvm_unreachable(
"Unknown construct kind in VisitNumGangsClause");
721 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
722 mlir::acc::KernelsOp>) {
723 operation.addVectorLengthOperand(builder.getContext(),
725 lastDeviceTypeValues);
726 }
else if constexpr (isCombinedType<OpTy>) {
727 applyToComputeOp(clause);
729 llvm_unreachable(
"Unknown construct kind in VisitVectorLengthClause");
734 hasAsyncClause =
true;
735 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
736 mlir::acc::KernelsOp, mlir::acc::DataOp,
737 mlir::acc::EnterDataOp, mlir::acc::ExitDataOp,
738 mlir::acc::UpdateOp>) {
740 operation.addAsyncOnly(builder.getContext(), lastDeviceTypeValues);
748 mlir::OpBuilder::InsertionGuard guardCase(builder);
749 if (!dataOperands.empty())
750 builder.setInsertionPoint(dataOperands.front());
753 operation.addAsyncOperand(builder.getContext(), intExpr,
754 lastDeviceTypeValues);
756 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::WaitOp>) {
760 operation.setAsync(
true);
762 operation.getAsyncOperandMutable().append(
764 }
else if constexpr (isCombinedType<OpTy>) {
765 applyToComputeOp(clause);
769 return clauseNotImplemented(clause);
774 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
775 mlir::acc::KernelsOp>) {
777 operation.setSelfAttr(
true);
780 operation.getSelfCondMutable().append(
783 llvm_unreachable(
"var-list version of self shouldn't get here");
785 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
787 "var-list version of self required for update");
789 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::UpdateHostOp>(
790 var, mlir::acc::DataClause::acc_update_self, {},
792 }
else if constexpr (isCombinedType<OpTy>) {
793 applyToComputeOp(clause);
795 llvm_unreachable(
"Unknown construct kind in VisitSelfClause");
800 if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
802 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::UpdateHostOp>(
803 var, mlir::acc::DataClause::acc_update_host, {},
806 llvm_unreachable(
"Unknown construct kind in VisitHostClause");
811 if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
813 addDataOperand<mlir::acc::UpdateDeviceOp>(
814 var, mlir::acc::DataClause::acc_update_device, {},
817 llvm_unreachable(
"Unknown construct kind in VisitDeviceClause");
822 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
823 mlir::acc::KernelsOp, mlir::acc::InitOp,
824 mlir::acc::ShutdownOp, mlir::acc::SetOp,
825 mlir::acc::DataOp, mlir::acc::WaitOp,
826 mlir::acc::HostDataOp, mlir::acc::EnterDataOp,
827 mlir::acc::ExitDataOp, mlir::acc::UpdateOp>) {
828 operation.getIfCondMutable().append(
830 }
else if constexpr (isCombinedType<OpTy>) {
831 applyToComputeOp(clause);
833 llvm_unreachable(
"Unknown construct kind in VisitIfClause");
838 if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp,
839 mlir::acc::UpdateOp>) {
840 operation.setIfPresent(
true);
842 llvm_unreachable(
"unknown construct kind in VisitIfPresentClause");
847 if constexpr (isOneOfTypes<OpTy, mlir::acc::InitOp, mlir::acc::ShutdownOp,
849 operation.getDeviceNumMutable().append(emitIntExpr(clause.
getIntExpr()));
852 "init, shutdown, set, are only valid device_num constructs");
857 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
858 mlir::acc::KernelsOp>) {
861 values.push_back(emitIntExpr(
E));
863 operation.addNumGangsOperands(builder.getContext(), values,
864 lastDeviceTypeValues);
865 }
else if constexpr (isCombinedType<OpTy>) {
866 applyToComputeOp(clause);
868 llvm_unreachable(
"Unknown construct kind in VisitNumGangsClause");
873 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
874 mlir::acc::KernelsOp, mlir::acc::DataOp,
875 mlir::acc::EnterDataOp, mlir::acc::ExitDataOp,
876 mlir::acc::UpdateOp>) {
878 operation.addWaitOnly(builder.getContext(), lastDeviceTypeValues);
884 values.push_back(emitIntExpr(
E));
885 operation.addWaitOperands(builder.getContext(), clause.
hasDevNumExpr(),
886 values, lastDeviceTypeValues);
888 }
else if constexpr (isCombinedType<OpTy>) {
889 applyToComputeOp(clause);
893 return clauseNotImplemented(clause);
898 if constexpr (isOneOfTypes<OpTy, mlir::acc::SetOp>) {
899 operation.getDefaultAsyncMutable().append(
902 llvm_unreachable(
"set, is only valid device_num constructs");
907 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
908 operation.addSeq(builder.getContext(), lastDeviceTypeValues);
909 }
else if constexpr (isCombinedType<OpTy>) {
910 applyToLoopOp(clause);
914 return clauseNotImplemented(clause);
919 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
920 operation.addAuto(builder.getContext(), lastDeviceTypeValues);
921 }
else if constexpr (isCombinedType<OpTy>) {
922 applyToLoopOp(clause);
926 return clauseNotImplemented(clause);
931 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
932 operation.addIndependent(builder.getContext(), lastDeviceTypeValues);
933 }
else if constexpr (isCombinedType<OpTy>) {
934 applyToLoopOp(clause);
938 return clauseNotImplemented(clause);
943 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
947 value = value.sextOrTrunc(64);
948 operation.setCollapseForDeviceTypes(builder.getContext(),
949 lastDeviceTypeValues, value);
950 }
else if constexpr (isCombinedType<OpTy>) {
951 applyToLoopOp(clause);
953 llvm_unreachable(
"Unknown construct kind in VisitCollapseClause");
958 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
962 mlir::Location exprLoc = cgf.
cgm.
getLoc(e->getBeginLoc());
967 if (isa<OpenACCAsteriskSizeExpr>(e)) {
968 values.push_back(createConstantInt(exprLoc, 64, -1));
970 llvm::APInt curValue =
972 values.push_back(createConstantInt(
973 exprLoc, 64, curValue.sextOrTrunc(64).getSExtValue()));
977 operation.setTileForDeviceTypes(builder.getContext(),
978 lastDeviceTypeValues, values);
979 }
else if constexpr (isCombinedType<OpTy>) {
980 applyToLoopOp(clause);
982 llvm_unreachable(
"Unknown construct kind in VisitTileClause");
987 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
989 operation.addWorkerNumOperand(builder.getContext(),
991 lastDeviceTypeValues);
993 operation.addEmptyWorker(builder.getContext(), lastDeviceTypeValues);
995 }
else if constexpr (isCombinedType<OpTy>) {
996 applyToLoopOp(clause);
1000 return clauseNotImplemented(clause);
1005 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
1007 operation.addVectorOperand(builder.getContext(),
1009 lastDeviceTypeValues);
1011 operation.addEmptyVector(builder.getContext(), lastDeviceTypeValues);
1013 }
else if constexpr (isCombinedType<OpTy>) {
1014 applyToLoopOp(clause);
1018 return clauseNotImplemented(clause);
1023 if constexpr (isOneOfTypes<OpTy, mlir::acc::LoopOp>) {
1025 operation.addEmptyGang(builder.getContext(), lastDeviceTypeValues);
1029 for (
unsigned i : llvm::index_range(0u, clause.
getNumExprs())) {
1031 mlir::Location exprLoc = cgf.
cgm.
getLoc(
expr->getBeginLoc());
1032 argTypes.push_back(decodeGangType(kind));
1033 if (kind == OpenACCGangKind::Dim) {
1034 llvm::APInt curValue =
1038 curValue = curValue.sextOrTrunc(64);
1040 createConstantInt(exprLoc, 64, curValue.getSExtValue()));
1041 }
else if (isa<OpenACCAsteriskSizeExpr>(
expr)) {
1042 values.push_back(createConstantInt(exprLoc, 64, -1));
1044 values.push_back(emitIntExpr(
expr));
1048 operation.addGangOperands(builder.getContext(), lastDeviceTypeValues,
1051 }
else if constexpr (isCombinedType<OpTy>) {
1052 applyToLoopOp(clause);
1054 llvm_unreachable(
"Unknown construct kind in VisitGangClause");
1059 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
1060 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
1062 addDataOperand<mlir::acc::CopyinOp, mlir::acc::CopyoutOp>(
1066 }
else if constexpr (isCombinedType<OpTy>) {
1067 applyToComputeOp(clause);
1071 return clauseNotImplemented(clause);
1076 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
1077 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
1079 addDataOperand<mlir::acc::CopyinOp, mlir::acc::DeleteOp>(
1083 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
1085 addDataOperand<mlir::acc::CopyinOp>(
1088 }
else if constexpr (isCombinedType<OpTy>) {
1089 applyToComputeOp(clause);
1093 return clauseNotImplemented(clause);
1098 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
1099 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
1101 addDataOperand<mlir::acc::CreateOp, mlir::acc::CopyoutOp>(
1105 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
1107 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::CopyoutOp>(
1111 }
else if constexpr (isCombinedType<OpTy>) {
1112 applyToComputeOp(clause);
1116 return clauseNotImplemented(clause);
1121 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
1122 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
1124 addDataOperand<mlir::acc::CreateOp, mlir::acc::DeleteOp>(
1128 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
1130 addDataOperand<mlir::acc::CreateOp>(
1133 }
else if constexpr (isCombinedType<OpTy>) {
1134 applyToComputeOp(clause);
1138 return clauseNotImplemented(clause);
1143 if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
1145 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::DeleteOp>(
1146 var, mlir::acc::DataClause::acc_delete, {},
1150 llvm_unreachable(
"Unknown construct kind in VisitDeleteClause");
1155 if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
1157 addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::DetachOp>(
1158 var, mlir::acc::DataClause::acc_detach, {},
1162 llvm_unreachable(
"Unknown construct kind in VisitDetachClause");
1167 if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
1168 operation.setFinalize(
true);
1170 llvm_unreachable(
"Unknown construct kind in VisitFinalizeClause");
1175 if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp>) {
1177 addDataOperand<mlir::acc::UseDeviceOp>(
1178 var, mlir::acc::DataClause::acc_use_device, {},
true,
1181 llvm_unreachable(
"Unknown construct kind in VisitUseDeviceClause");
1186 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
1187 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
1189 addDataOperand<mlir::acc::DevicePtrOp>(
1190 var, mlir::acc::DataClause::acc_deviceptr, {},
1193 }
else if constexpr (isCombinedType<OpTy>) {
1194 applyToComputeOp(clause);
1198 return clauseNotImplemented(clause);
1203 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
1204 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
1206 addDataOperand<mlir::acc::NoCreateOp, mlir::acc::DeleteOp>(
1207 var, mlir::acc::DataClause::acc_no_create, {},
true,
1209 }
else if constexpr (isCombinedType<OpTy>) {
1210 applyToComputeOp(clause);
1212 llvm_unreachable(
"Unknown construct kind in VisitNoCreateClause");
1217 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
1218 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
1220 addDataOperand<mlir::acc::PresentOp, mlir::acc::DeleteOp>(
1221 var, mlir::acc::DataClause::acc_present, {},
true,
1223 }
else if constexpr (isCombinedType<OpTy>) {
1224 applyToComputeOp(clause);
1228 return clauseNotImplemented(clause);
1233 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
1234 mlir::acc::KernelsOp, mlir::acc::DataOp>) {
1236 addDataOperand<mlir::acc::AttachOp, mlir::acc::DetachOp>(
1237 var, mlir::acc::DataClause::acc_attach, {},
true,
1239 }
else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
1241 addDataOperand<mlir::acc::AttachOp>(
1242 var, mlir::acc::DataClause::acc_attach, {},
1244 }
else if constexpr (isCombinedType<OpTy>) {
1245 applyToComputeOp(clause);
1247 llvm_unreachable(
"Unknown construct kind in VisitAttachClause");
1252 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
1253 mlir::acc::LoopOp>) {
1254 for (
const auto [varExpr, varRecipe] :
1258 auto privateOp = mlir::acc::PrivateOp::create(
1261 privateOp.setDataClause(mlir::acc::DataClause::acc_private);
1264 mlir::OpBuilder::InsertionGuard guardCase(builder);
1265 auto recipe = getOrCreateRecipe<mlir::acc::PrivateRecipeOp>(
1266 cgf.
getContext(), varExpr, varRecipe,
nullptr,
1267 OpenACCReductionOperator::Invalid,
1270 privateOp.getResult());
1274 operation.addPrivatization(builder.getContext(), privateOp, recipe);
1277 }
else if constexpr (isCombinedType<OpTy>) {
1280 applyToLoopOp(clause);
1282 llvm_unreachable(
"Unknown construct kind in VisitPrivateClause");
1287 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
1288 mlir::acc::SerialOp>) {
1289 for (
const auto [varExpr, varRecipe] :
1293 auto firstPrivateOp = mlir::acc::FirstprivateOp::create(
1297 firstPrivateOp.setDataClause(mlir::acc::DataClause::acc_firstprivate);
1300 mlir::OpBuilder::InsertionGuard guardCase(builder);
1301 auto recipe = getOrCreateRecipe<mlir::acc::FirstprivateRecipeOp>(
1302 cgf.
getContext(), varExpr, varRecipe.RecipeDecl,
1303 varRecipe.InitFromTemporary, OpenACCReductionOperator::Invalid,
1305 firstPrivateOp.getResult());
1310 operation.addFirstPrivatization(builder.getContext(), firstPrivateOp,
1314 }
else if constexpr (isCombinedType<OpTy>) {
1317 applyToComputeOp(clause);
1319 llvm_unreachable(
"Unknown construct kind in VisitFirstPrivateClause");
1324 if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
1325 mlir::acc::LoopOp>) {
1326 for (
const auto [varExpr, varRecipe] :
1331 auto reductionOp = mlir::acc::ReductionOp::create(
1334 reductionOp.setDataClause(mlir::acc::DataClause::acc_reduction);
1337 mlir::OpBuilder::InsertionGuard guardCase(builder);
1339 auto recipe = getOrCreateRecipe<mlir::acc::ReductionRecipeOp>(
1340 cgf.
getContext(), varExpr, varRecipe.RecipeDecl,
1343 reductionOp.getResult());
1345 operation.addReduction(builder.getContext(), reductionOp, recipe);
1348 }
else if constexpr (isCombinedType<OpTy>) {
1351 applyToLoopOp(clause);
1353 llvm_unreachable(
"Unknown construct kind in VisitReductionClause");
1358template <
typename OpTy>
1362 return OpenACCClauseCIREmitter<OpTy>(op, cgf, builder, dirKind, dirLoc);
1366template <
typename Op>
1367void CIRGenFunction::emitOpenACCClauses(
1370 mlir::OpBuilder::InsertionGuard guardCase(builder);
1374 builder.setInsertionPoint(op);
1375 makeClauseEmitter(op, *
this, builder, dirKind, dirLoc).emitClauses(clauses);
1378#define EXPL_SPEC(N) \
1379 template void CIRGenFunction::emitOpenACCClauses<N>( \
1380 N &, OpenACCDirectiveKind, SourceLocation, \
1381 ArrayRef<const OpenACCClause *>);
1397template <
typename ComputeOp,
typename LoopOp>
1398void CIRGenFunction::emitOpenACCClauses(
1401 static_assert(std::is_same_v<mlir::acc::LoopOp, LoopOp>);
1403 CombinedConstructClauseInfo<ComputeOp> inf{op, loopOp};
1406 mlir::OpBuilder::InsertionGuard guardCase(builder);
1407 makeClauseEmitter(inf, *
this, builder, dirKind, dirLoc).emitClauses(clauses);
1410#define EXPL_SPEC(N) \
1411 template void CIRGenFunction::emitOpenACCClauses<N, mlir::acc::LoopOp>( \
1412 N &, mlir::acc::LoopOp &, OpenACCDirectiveKind, SourceLocation, \
1413 ArrayRef<const OpenACCClause *>);
Defines the clang::Expr interface and subclasses for C++ expressions.
__device__ __2f16 float c
Holds long-lived AST nodes (such as types and decls) that can be referred to throughout the semantic ...
const LangOptions & getLangOpts() const
CharUnits getDeclAlign(const Decl *D, bool ForAlignof=false) const
Return a conservative estimate of the alignment of the specified decl D.
clang::MangleContext & getMangleContext()
Gets the mangle context.
AutoVarEmission emitAutoVarAlloca(const clang::VarDecl &d, mlir::OpBuilder::InsertPoint ip={})
const clang::Decl * curFuncDecl
mlir::Value evaluateExprAsBool(const clang::Expr *e)
Perform the usual unary conversions on the specified expression and compare the result against zero,...
OpenACCDataOperandInfo getOpenACCDataOperandInfo(const Expr *e)
void emitAutoVarInit(const AutoVarEmission &emission)
Emit the initializer for an allocated variable.
mlir::Value emitOpenACCIntExpr(const Expr *intExpr)
void emitDestroy(Address addr, QualType type, Destroyer *destroyer)
Immediately perform the destruction of the given object.
mlir::MLIRContext & getMLIRContext()
Destroyer * getDestroyer(clang::QualType::DestructionKind kind)
mlir::Value createOpenACCConstantInt(mlir::Location loc, unsigned width, int64_t value)
clang::ASTContext & getContext() const
void setAddrOfLocalVar(const clang::VarDecl *vd, Address addr)
Set the address of a local variable.
DiagnosticBuilder errorNYI(SourceLocation, llvm::StringRef)
Helpers to emit "not yet implemented" error diagnostics.
clang::ASTContext & getASTContext() const
mlir::Location getLoc(clang::SourceLocation cLoc)
Helpers to convert the presumed location of Clang's SourceLocation to an MLIR Location.
CIRGenCXXABI & getCXXABI() const
CharUnits - This is an opaque type for sizes expressed in character units.
DeclContext - This is used only as base class of specific decl types that can act as declaration cont...
static DeclContext * castToDeclContext(const Decl *)
This represents one expression.
llvm::APSInt EvaluateKnownConstInt(const ASTContext &Ctx, SmallVectorImpl< PartialDiagnosticAt > *Diag=nullptr) const
EvaluateKnownConstInt - Call EvaluateAsRValue and return the folded integer.
One of these records is kept for each identifier that is lexed.
StringRef getName() const
Return the actual identifier string.
A simple pair of identifier info and location.
MangleContext - Context for tracking state which persists across multiple calls to the C++ name mangl...
virtual void mangleCanonicalTypeName(QualType T, raw_ostream &, bool NormalizeIntegers=false)=0
Generates a unique string for an externally visible type for use with TBAA or type uniquing.
void VisitClauseList(ArrayRef< const OpenACCClause * > List)
const Expr * getConditionExpr() const
const Expr * getIntExpr() const
ArrayRef< Expr * > getVarList()
This is the base type for all OpenACC Clauses.
Represents a 'collapse' clause on a 'loop' construct.
OpenACCModifierKind getModifierList() const
OpenACCModifierKind getModifierList() const
OpenACCModifierKind getModifierList() const
OpenACCModifierKind getModifierList() const
A 'default' clause, has the optional 'none' or 'present' argument.
OpenACCDefaultClauseKind getDefaultClauseKind() const
A 'device_type' or 'dtype' clause, takes a list of either an 'asterisk' or an identifier.
ArrayRef< DeviceTypeArgument > getArchitectures() const
ArrayRef< OpenACCFirstPrivateRecipe > getInitRecipes()
unsigned getNumExprs() const
std::pair< OpenACCGangKind, const Expr * > getExpr(unsigned I) const
An 'if' clause, which has a required condition expression.
ArrayRef< Expr * > getIntExprs()
ArrayRef< VarDecl * > getInitRecipes()
ArrayRef< OpenACCReductionRecipe > getRecipes()
OpenACCReductionOperator getReductionOp() const
A 'self' clause, which has an optional condition expression, or, in the event of an 'update' directiv...
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
A (possibly-)qualified type.
Encodes a location in the source.
A trivial tuple used to represent a source range.
SourceLocation getEndLoc() const LLVM_READONLY
SourceRange getSourceRange() const LLVM_READONLY
SourceLocation tokens are not useful in isolation - they are low level value objects created/interpre...
SourceLocation getBeginLoc() const LLVM_READONLY
const Type * getPointeeOrArrayElementType() const
If this is a pointer type, return the pointee type.
bool isPointerType() const
bool isBuiltinType() const
Helper methods to distinguish type categories.
Represents a variable declaration or definition.
QualType::DestructionKind needsDestruction(const ASTContext &Ctx) const
Would the destruction of this variable have any effect, and if so, what kind?
const Expr * getInit() const
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.
@ Invalid
Invalid Reduction Clause Kind.
@ Invalid
Not a valid option.
@ Present
'present' option.
bool isOpenACCModifierBitSet(OpenACCModifierKind List, OpenACCModifierKind Bit)
const FunctionProtoType * T
__DEVICE__ _Tp arg(const std::complex< _Tp > &__c)
void setAllocatedAddress(Address A)
bool EmittedAsOffload
True if the variable was emitted as an offload recipe, and thus doesn't have the same sort of alloca ...
static AutoVarEmission invalid()
Represents a scope, including function bodies, compound statements, and the substatements of if/while...
llvm::SmallVector< mlir::Value > bounds