-
Notifications
You must be signed in to change notification settings - Fork 15.2k
[OpenACC][CIR] Implement 'reduction' combiner lowering for 5 ops #162906
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
base: main
Are you sure you want to change the base?
[OpenACC][CIR] Implement 'reduction' combiner lowering for 5 ops #162906
Conversation
Following on the Sema changes, this does the lowering for all of the operators that can be done as a compound operator. Lowering is very simply looping through the objects based on array/compound/etc, and doing a call to the operation.
@llvm/pr-subscribers-clangir @llvm/pr-subscribers-clang Author: Erich Keane (erichkeane) ChangesFollowing on the Sema changes, this does the lowering for all of the operators that can be done as a compound operator. Lowering is very simply looping through the objects based on array/compound/etc, and doing a call to the operation. Patch is 719.05 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/162906.diff 24 Files Affected:
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index 3d86f71b077e2..59f09942019d4 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -1005,7 +1005,7 @@ class OpenACCClauseCIREmitter final
/*temporary=*/nullptr, OpenACCReductionOperator::Invalid,
Decl::castToDeclContext(cgf.curFuncDecl), opInfo.origType,
opInfo.bounds.size(), opInfo.boundTypes, opInfo.baseType,
- privateOp);
+ privateOp, /*reductionCombinerRecipe=*/{});
// TODO: OpenACC: The dialect is going to change in the near future to
// have these be on a different operation, so when that changes, we
// probably need to change these here.
@@ -1046,7 +1046,7 @@ class OpenACCClauseCIREmitter final
OpenACCReductionOperator::Invalid,
Decl::castToDeclContext(cgf.curFuncDecl), opInfo.origType,
opInfo.bounds.size(), opInfo.boundTypes, opInfo.baseType,
- firstPrivateOp);
+ firstPrivateOp, /*reductionCombinerRecipe=*/{});
// TODO: OpenACC: The dialect is going to change in the near future to
// have these be on a different operation, so when that changes, we
@@ -1088,7 +1088,7 @@ class OpenACCClauseCIREmitter final
/*temporary=*/nullptr, clause.getReductionOp(),
Decl::castToDeclContext(cgf.curFuncDecl), opInfo.origType,
opInfo.bounds.size(), opInfo.boundTypes, opInfo.baseType,
- reductionOp);
+ reductionOp, varRecipe.CombinerRecipes);
operation.addReduction(builder.getContext(), reductionOp, recipe);
}
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCRecipe.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCRecipe.cpp
index 24a5fc27c4775..92a8338b70ee7 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCRecipe.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCRecipe.cpp
@@ -527,16 +527,140 @@ void OpenACCRecipeBuilderBase::createFirstprivateRecipeCopy(
// doesn't restore it aftewards.
void OpenACCRecipeBuilderBase::createReductionRecipeCombiner(
mlir::Location loc, mlir::Location locEnd, mlir::Value mainOp,
- mlir::acc::ReductionRecipeOp recipe, size_t numBounds) {
+ mlir::acc::ReductionRecipeOp recipe, size_t numBounds, QualType origType,
+ llvm::ArrayRef<OpenACCReductionRecipe::CombinerRecipe> combinerRecipes) {
mlir::Block *block =
createRecipeBlock(recipe.getCombinerRegion(), mainOp.getType(), loc,
numBounds, /*isInit=*/false);
builder.setInsertionPointToEnd(&recipe.getCombinerRegion().back());
CIRGenFunction::LexicalScope ls(cgf, loc, block);
- mlir::BlockArgument lhsArg = block->getArgument(0);
+ mlir::Value lhsArg = block->getArgument(0);
+ mlir::Value rhsArg = block->getArgument(1);
+ llvm::MutableArrayRef<mlir::BlockArgument> boundsRange =
+ block->getArguments().drop_front(2);
+
+ if (llvm::any_of(combinerRecipes, [](auto &r) { return r.Op == nullptr; })) {
+ cgf.cgm.errorNYI(loc, "OpenACC Reduction combiner not generated");
+ mlir::acc::YieldOp::create(builder, locEnd, block->getArgument(0));
+ return;
+ }
+
+ // apply the bounds so that we can get our bounds emitted correctly.
+ for (mlir::BlockArgument boundArg : llvm::reverse(boundsRange))
+ std::tie(lhsArg, rhsArg) =
+ createBoundsLoop(lhsArg, rhsArg, boundArg, loc, /*inverse=*/false);
+
+ // Emitter for when we know this isn't a struct or array we have to loop
+ // through. This should work for the 'field' once the get-element call has
+ // been made.
+ auto emitSingleCombiner =
+ [&](mlir::Value lhsArg, mlir::Value rhsArg,
+ const OpenACCReductionRecipe::CombinerRecipe &combiner) {
+ mlir::Type elementTy =
+ mlir::cast<cir::PointerType>(lhsArg.getType()).getPointee();
+ CIRGenFunction::DeclMapRevertingRAII declMapRAIILhs{cgf, combiner.LHS};
+ cgf.setAddrOfLocalVar(
+ combiner.LHS, Address{lhsArg, elementTy,
+ cgf.getContext().getDeclAlign(combiner.LHS)});
+ CIRGenFunction::DeclMapRevertingRAII declMapRAIIRhs{cgf, combiner.RHS};
+ cgf.setAddrOfLocalVar(
+ combiner.RHS, Address{rhsArg, elementTy,
+ cgf.getContext().getDeclAlign(combiner.RHS)});
+
+ [[maybe_unused]] mlir::LogicalResult stmtRes =
+ cgf.emitStmt(combiner.Op, /*useCurrentScope=*/true);
+ };
+
+ // Emitter for when we know this is either a non-array or element of an array
+ // (which also shouldn't be an array type?). This function should generate the
+ // loop to do this on each individual array or struct element (if necessary).
+ auto emitCombiner = [&](mlir::Value lhsArg, mlir::Value rhsArg, QualType Ty) {
+ if (const auto *RD = Ty->getAsRecordDecl()) {
+ if (combinerRecipes.size() == 1 &&
+ cgf.getContext().hasSameType(Ty, combinerRecipes[0].LHS->getType())) {
+ // If this is a 'top level' operator on the type we can just emit this
+ // as a simple one.
+ emitSingleCombiner(lhsArg, rhsArg, combinerRecipes[0]);
+ } else {
+ // else we have to handle each individual field after after a
+ // get-element.
+ for (const auto &[field, combiner] :
+ llvm::zip_equal(RD->fields(), combinerRecipes)) {
+ mlir::Type fieldType = cgf.convertType(field->getType());
+ auto fieldPtr = cir::PointerType::get(fieldType);
+
+ mlir::Value lhsField = builder.createGetMember(
+ loc, fieldPtr, lhsArg, field->getName(), field->getFieldIndex());
+ mlir::Value rhsField = builder.createGetMember(
+ loc, fieldPtr, rhsArg, field->getName(), field->getFieldIndex());
+
+ emitSingleCombiner(lhsField, rhsField, combiner);
+ }
+ }
+
+ } else {
+ // if this is a single-thing (because we should know this isn't an array,
+ // as Sema wouldn't let us get here), we can just do a normal emit call.
+ emitSingleCombiner(lhsArg, rhsArg, combinerRecipes[0]);
+ }
+ };
+
+ if (const auto *cat = cgf.getContext().getAsConstantArrayType(origType)) {
+ // If we're in an array, we have to emit the combiner for each element of
+ // the array.
+ auto itrTy = mlir::cast<cir::IntType>(cgf.PtrDiffTy);
+ auto itrPtrTy = cir::PointerType::get(itrTy);
+
+ mlir::Value zero =
+ builder.getConstInt(loc, mlir::cast<cir::IntType>(cgf.PtrDiffTy), 0);
+ mlir::Value itr =
+ cir::AllocaOp::create(builder, loc, itrPtrTy, itrTy, "itr",
+ cgf.cgm.getSize(cgf.getPointerAlign()));
+ builder.CIRBaseBuilderTy::createStore(loc, zero, itr);
+
+ builder.setInsertionPointAfter(builder.createFor(
+ loc,
+ /*condBuilder=*/
+ [&](mlir::OpBuilder &b, mlir::Location loc) {
+ auto loadItr = cir::LoadOp::create(builder, loc, {itr});
+ mlir::Value arraySize = builder.getConstInt(
+ loc, mlir::cast<cir::IntType>(cgf.PtrDiffTy), cat->getZExtSize());
+ auto cmp = builder.createCompare(loc, cir::CmpOpKind::lt, loadItr,
+ arraySize);
+ builder.createCondition(cmp);
+ },
+ /*bodyBuilder=*/
+ [&](mlir::OpBuilder &b, mlir::Location loc) {
+ auto loadItr = cir::LoadOp::create(builder, loc, {itr});
+ auto lhsElt = builder.getArrayElement(
+ loc, loc, lhsArg, cgf.convertType(cat->getElementType()), loadItr,
+ /*shouldDecay=*/true);
+ auto rhsElt = builder.getArrayElement(
+ loc, loc, rhsArg, cgf.convertType(cat->getElementType()), loadItr,
+ /*shouldDecay=*/true);
+
+ emitCombiner(lhsElt, rhsElt, cat->getElementType());
+ builder.createYield(loc);
+ },
+ /*stepBuilder=*/
+ [&](mlir::OpBuilder &b, mlir::Location loc) {
+ auto loadItr = cir::LoadOp::create(builder, loc, {itr});
+ auto inc = cir::UnaryOp::create(builder, loc, loadItr.getType(),
+ cir::UnaryOpKind::Inc, loadItr);
+ builder.CIRBaseBuilderTy::createStore(loc, inc, itr);
+ builder.createYield(loc);
+ }));
- mlir::acc::YieldOp::create(builder, locEnd, lhsArg);
+ } else if (origType->isArrayType()) {
+ cgf.cgm.errorNYI(loc,
+ "OpenACC Reduction combiner non-constant array recipe");
+ } else {
+ emitCombiner(lhsArg, rhsArg, origType);
+ }
+
+ builder.setInsertionPointToEnd(&recipe.getCombinerRegion().back());
+ mlir::acc::YieldOp::create(builder, locEnd, block->getArgument(0));
}
} // namespace clang::CIRGen
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCRecipe.h b/clang/lib/CIR/CodeGen/CIRGenOpenACCRecipe.h
index a5da7444ebf96..745d42446896e 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCRecipe.h
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCRecipe.h
@@ -64,10 +64,10 @@ class OpenACCRecipeBuilderBase {
// that this function is not 'insertion point' clean, in that it alters the
// insertion point to be inside of the 'combiner' section of the recipe, but
// doesn't restore it aftewards.
- void createReductionRecipeCombiner(mlir::Location loc, mlir::Location locEnd,
- mlir::Value mainOp,
- mlir::acc::ReductionRecipeOp recipe,
- size_t numBounds);
+ void createReductionRecipeCombiner(
+ mlir::Location loc, mlir::Location locEnd, mlir::Value mainOp,
+ mlir::acc::ReductionRecipeOp recipe, size_t numBounds, QualType origType,
+ llvm::ArrayRef<OpenACCReductionRecipe::CombinerRecipe> combinerRecipes);
void createInitRecipe(mlir::Location loc, mlir::Location locEnd,
SourceRange exprRange, mlir::Value mainOp,
@@ -169,7 +169,9 @@ class OpenACCRecipeBuilder : OpenACCRecipeBuilderBase {
const Expr *varRef, const VarDecl *varRecipe, const VarDecl *temporary,
OpenACCReductionOperator reductionOp, DeclContext *dc, QualType origType,
size_t numBounds, llvm::ArrayRef<QualType> boundTypes, QualType baseType,
- mlir::Value mainOp) {
+ mlir::Value mainOp,
+ llvm::ArrayRef<OpenACCReductionRecipe::CombinerRecipe>
+ reductionCombinerRecipes) {
assert(!varRecipe->getType()->isSpecificBuiltinType(
BuiltinType::ArraySection) &&
"array section shouldn't make it to recipe creation");
@@ -208,7 +210,8 @@ class OpenACCRecipeBuilder : OpenACCRecipeBuilderBase {
createInitRecipe(loc, locEnd, varRef->getSourceRange(), mainOp,
recipe.getInitRegion(), numBounds, boundTypes, varRecipe,
origType, /*emitInitExpr=*/true);
- createReductionRecipeCombiner(loc, locEnd, mainOp, recipe, numBounds);
+ createReductionRecipeCombiner(loc, locEnd, mainOp, recipe, numBounds,
+ origType, reductionCombinerRecipes);
} else {
static_assert(std::is_same_v<RecipeTy, mlir::acc::FirstprivateRecipeOp>);
createInitRecipe(loc, locEnd, varRef->getSourceRange(), mainOp,
diff --git a/clang/lib/Sema/SemaOpenACCClause.cpp b/clang/lib/Sema/SemaOpenACCClause.cpp
index 881e960e5a24f..ca7b78a55e7f4 100644
--- a/clang/lib/Sema/SemaOpenACCClause.cpp
+++ b/clang/lib/Sema/SemaOpenACCClause.cpp
@@ -1924,7 +1924,7 @@ bool SemaOpenACC::CheckReductionVarType(Expr *VarExpr) {
// off here. This will result in CurType being the actual 'type' of the
// expression, which is what we are looking to check.
QualType CurType = isa<ArraySectionExpr>(VarExpr)
- ? ArraySectionExpr::getBaseOriginalType(VarExpr)
+ ? cast<ArraySectionExpr>(VarExpr)->getElementType()
: VarExpr->getType();
// This can happen when we have a dependent type in an array element that the
diff --git a/clang/test/CIR/CodeGenOpenACC/combined-reduction-clause-default-ops.cpp b/clang/test/CIR/CodeGenOpenACC/combined-reduction-clause-default-ops.cpp
index 040ddd3ca458c..ee4fffef971e0 100644
--- a/clang/test/CIR/CodeGenOpenACC/combined-reduction-clause-default-ops.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/combined-reduction-clause-default-ops.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fopenacc -triple x86_64-linux-gnu -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir -triple x86_64-linux-pc %s -o - | FileCheck %s
+// RUN: not %clang_cc1 -fopenacc -triple x86_64-linux-gnu -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir -triple x86_64-linux-pc %s -o - | FileCheck %s
struct DefaultOperators {
int i;
@@ -43,12 +43,43 @@ void acc_combined() {
//
// CHECK-NEXT: } combiner {
// CHECK-NEXT: ^bb0(%[[LHSARG:.*]]: !cir.ptr<!rec_DefaultOperators> {{.*}}, %[[RHSARG:.*]]: !cir.ptr<!rec_DefaultOperators> {{.*}})
-// TODO OpenACC: Expecting combination operation here
+// CHECK-NEXT: %[[GET_MEM_LHS:.*]] = cir.get_member %[[LHSARG]][0] {name = "i"} : !cir.ptr<!rec_DefaultOperators> -> !cir.ptr<!s32i>
+// CHECK-NEXT: %[[GET_MEM_RHS:.*]] = cir.get_member %[[RHSARG]][0] {name = "i"} : !cir.ptr<!rec_DefaultOperators> -> !cir.ptr<!s32i>
+// CHECK-NEXT: %[[RHS_LOAD:.*]] = cir.load{{.*}} %[[GET_MEM_RHS]] : !cir.ptr<!s32i>, !s32i
+// CHECK-NEXT: %[[LHS_LOAD:.*]] = cir.load{{.*}} %[[GET_MEM_LHS]] : !cir.ptr<!s32i>, !s32i
+// CHECK-NEXT: %[[ADD:.*]] = cir.binop(add, %[[LHS_LOAD]], %[[RHS_LOAD]]) nsw : !s32i
+// CHECK-NEXT: cir.store {{.*}} %[[ADD]], %[[GET_MEM_LHS]] : !s32i, !cir.ptr<!s32i>
+// CHECK-NEXT: %[[GET_MEM_LHS:.*]] = cir.get_member %[[LHSARG]][1] {name = "u"} : !cir.ptr<!rec_DefaultOperators> -> !cir.ptr<!u32i>
+// CHECK-NEXT: %[[GET_MEM_RHS:.*]] = cir.get_member %[[RHSARG]][1] {name = "u"} : !cir.ptr<!rec_DefaultOperators> -> !cir.ptr<!u32i>
+// CHECK-NEXT: %[[RHS_LOAD:.*]] = cir.load{{.*}} %[[GET_MEM_RHS]] : !cir.ptr<!u32i>, !u32i
+// CHECK-NEXT: %[[LHS_LOAD:.*]] = cir.load{{.*}} %[[GET_MEM_LHS]] : !cir.ptr<!u32i>, !u32i
+// CHECK-NEXT: %[[ADD:.*]] = cir.binop(add, %[[LHS_LOAD]], %[[RHS_LOAD]]) : !u32i
+// CHECK-NEXT: cir.store {{.*}} %[[ADD]], %[[GET_MEM_LHS]] : !u32i, !cir.ptr<!u32i>
+// CHECK-NEXT: %[[GET_MEM_LHS:.*]] = cir.get_member %[[LHSARG]][2] {name = "f"} : !cir.ptr<!rec_DefaultOperators> -> !cir.ptr<!cir.float>
+// CHECK-NEXT: %[[GET_MEM_RHS:.*]] = cir.get_member %[[RHSARG]][2] {name = "f"} : !cir.ptr<!rec_DefaultOperators> -> !cir.ptr<!cir.float>
+// CHECK-NEXT: %[[RHS_LOAD:.*]] = cir.load{{.*}} %[[GET_MEM_RHS]] : !cir.ptr<!cir.float>, !cir.float
+// CHECK-NEXT: %[[LHS_LOAD:.*]] = cir.load{{.*}} %[[GET_MEM_LHS]] : !cir.ptr<!cir.float>, !cir.float
+// CHECK-NEXT: %[[ADD:.*]] = cir.binop(add, %[[LHS_LOAD]], %[[RHS_LOAD]]) : !cir.float
+// CHECK-NEXT: cir.store {{.*}} %[[ADD]], %[[GET_MEM_LHS]] : !cir.float, !cir.ptr<!cir.float>
+// CHECK-NEXT: %[[GET_MEM_LHS:.*]] = cir.get_member %[[LHSARG]][3] {name = "d"} : !cir.ptr<!rec_DefaultOperators> -> !cir.ptr<!cir.double>
+// CHECK-NEXT: %[[GET_MEM_RHS:.*]] = cir.get_member %[[RHSARG]][3] {name = "d"} : !cir.ptr<!rec_DefaultOperators> -> !cir.ptr<!cir.double>
+// CHECK-NEXT: %[[RHS_LOAD:.*]] = cir.load{{.*}} %[[GET_MEM_RHS]] : !cir.ptr<!cir.double>, !cir.double
+// CHECK-NEXT: %[[LHS_LOAD:.*]] = cir.load{{.*}} %[[GET_MEM_LHS]] : !cir.ptr<!cir.double>, !cir.double
+// CHECK-NEXT: %[[ADD:.*]] = cir.binop(add, %[[LHS_LOAD]], %[[RHS_LOAD]]) : !cir.double
+// CHECK-NEXT: cir.store {{.*}} %[[ADD]], %[[GET_MEM_LHS]] : !cir.double, !cir.ptr<!cir.double>
+// CHECK-NEXT: %[[GET_MEM_LHS:.*]] = cir.get_member %[[LHSARG]][4] {name = "b"} : !cir.ptr<!rec_DefaultOperators> -> !cir.ptr<!cir.bool>
+// CHECK-NEXT: %[[GET_MEM_RHS:.*]] = cir.get_member %[[RHSARG]][4] {name = "b"} : !cir.ptr<!rec_DefaultOperators> -> !cir.ptr<!cir.bool>
+// CHECK-NEXT: %[[RHS_LOAD:.*]] = cir.load{{.*}} %[[GET_MEM_RHS]] : !cir.ptr<!cir.bool>, !cir.bool
+// CHECK-NEXT: %[[RHS_INT_CAST:.*]] = cir.cast bool_to_int %[[RHS_LOAD]] : !cir.bool -> !s32i
+// CHECK-NEXT: %[[LHS_LOAD:.*]] = cir.load{{.*}} %[[GET_MEM_LHS]] : !cir.ptr<!cir.bool>, !cir.bool
+// CHECK-NEXT: %[[LHS_INT_CAST:.*]] = cir.cast bool_to_int %[[LHS_LOAD]] : !cir.bool -> !s32i
+// CHECK-NEXT: %[[ADD:.*]] = cir.binop(add, %[[LHS_INT_CAST]], %[[RHS_INT_CAST]]) nsw : !s32i
+// CHECK-NEXT: %[[RES_TO_BOOL_CAST:.*]] = cir.cast int_to_bool %[[ADD]] : !s32i -> !cir.bool
+// CHECK-NEXT: cir.store {{.*}} %[[RES_TO_BOOL_CAST]], %[[GET_MEM_LHS]] : !cir.bool, !cir.ptr<!cir.bool>
// CHECK-NEXT: acc.yield %[[LHSARG]] : !cir.ptr<!rec_DefaultOperators>
// CHECK-NEXT: }
- for(int i=0;i < 5; ++i);
+ for(int i = 0; i < 5; ++i);
#pragma acc parallel loop reduction(*:someVar)
-
// CHECK-NEXT: acc.reduction.recipe @reduction_mul__ZTS16DefaultOperators : !cir.ptr<!rec_DefaultOperators> reduction_operator <mul> init {
// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_DefaultOperators>{{.*}})
// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !rec_DefaultOperators, !cir.ptr<!rec_DefaultOperators>, ["openacc.reduction.init", init]
@@ -71,10 +102,42 @@ void acc_combined() {
//
// CHECK-NEXT: } combiner {
// CHECK-NEXT: ^bb0(%[[LHSARG:.*]]: !cir.ptr<!rec_DefaultOperators> {{.*}}, %[[RHSARG:.*]]: !cir.ptr<!rec_DefaultOperators> {{.*}})
-// TODO OpenACC: Expecting combination operation here
+// CHECK-NEXT: %[[GET_MEM_LHS:.*]] = cir.get_member %[[LHSARG]][0] {name = "i"} : !cir.ptr<!rec_DefaultOperators> -> !cir.ptr<!s32i>
+// CHECK-NEXT: %[[GET_MEM_RHS:.*]] = cir.get_member %[[RHSARG]][0] {name = "i"} : !cir.ptr<!rec_DefaultOperators> -> !cir.ptr<!s32i>
+// CHECK-NEXT: %[[RHS_LOAD:.*]] = cir.load{{.*}} %[[GET_MEM_RHS]] : !cir.ptr<!s32i>, !s32i
+// CHECK-NEXT: %[[LHS_LOAD:.*]] = cir.load{{.*}} %[[GET_MEM_LHS]] : !cir.ptr<!s32i>, !s32i
+// CHECK-NEXT: %[[MUL:.*]] = cir.binop(mul, %[[LHS_LOAD]], %[[RHS_LOAD]]) nsw : !s32i
+// CHECK-NEXT: cir.store {{.*}} %[[MUL]], %[[GET_MEM_LHS]] : !s32i, !cir.ptr<!s32i>
+// CHECK-NEXT: %[[GET_MEM_LHS:.*]] = cir.get_member %[[LHSARG]][1] {name = "u"} : !cir.ptr<!rec_DefaultOperators> -> !cir.ptr<!u32i>
+// CHECK-NEXT: %[[GET_MEM_RHS:.*]] = cir.get_member %[[RHSARG]][1] {name = "u"} : !cir.ptr<!rec_DefaultOperators> -> !cir.ptr<!u32i>
+// CHECK-NEXT: %[[RHS_LOAD:.*]] = cir.load{{.*}} %[[GET_MEM_RHS]] : !cir.ptr<!u32i>, !u32i
+// CHECK-NEXT: %[[LHS_LOAD:.*]] = cir.load{{.*}} %[[GET_MEM_LHS]] : !cir.ptr<!u32i>, !u32i
+// CHECK-NEXT: %[[MUL:.*]] = cir.binop(mul, %[[LHS_LOAD]], %[[RHS_LOAD]]) : !u32i
+// CHECK-NEXT: cir.store {{.*}} %[[MUL]], %[[GET_MEM_LHS]] : !u32i, !cir.ptr<!u32i>
+// CHECK-NEXT: %[[GET_MEM_LHS:.*]] = cir.get_member %[[LHSARG]][2] {name = "f"} : !cir.ptr<!rec_DefaultOperators> -> !cir.ptr<!cir.float>
+// CHECK-NEXT: %[[GET_MEM_RHS:.*]] = cir.get_member %[[RHSARG]][2] {name = "f"} : !cir.ptr<!rec_DefaultOperators> -> !cir.ptr<!cir.float>
+// CHECK-NEXT: %[[RHS_LOAD:.*]] = cir.load{{.*}} %[[GET_MEM_RHS]] : !cir.ptr<!cir.float>, !cir.float
+// CHECK-NEXT: %[[LHS_LOAD:.*]] = cir.load{{.*}} %[[GET_MEM_LHS]] : !cir.ptr<!cir.float>, !cir.float
+// CHECK-NEXT: %[[MUL:.*]] = cir.binop(mul, %[[LHS_LOAD]], %[[RHS_LOAD]]) : !cir.float
+// CHECK-NEXT: cir.store {{.*}} %[[MUL]], %[[GET_MEM_LHS]] : !cir.float, !cir.ptr<!cir.float>
+// CHECK-NEXT: %[[GET_MEM_LHS:.*]] = cir.get_member %[[LHSARG]][3] {name = "d"} : !cir.ptr<!rec_DefaultOperators> -> !cir.ptr<!cir.double>
+// CHECK-NEXT: %[[GET_MEM_RHS:.*]] = cir.get_member %[[RHSARG]][3] {name = "d"} : !cir.ptr<!rec_DefaultOperators> -> !cir.ptr<!cir.double>
+// CHECK-NEXT: %[[RHS_LOAD:.*]] = cir.load{{.*}} %[[GET_MEM_RHS]] : !cir.ptr<!cir.double>, !cir.double
+// CHECK-NEXT: %[[LHS_LOAD:.*]] = cir.load{{.*}} %[[GET_MEM_LHS]] : !cir.ptr<!cir.double>, !cir.double
+// CHECK-NEXT: %[[MUL:.*]] = cir.binop(mul, %[[LHS_LOAD]], %[[RHS_LOAD]]) : !cir.double
+// CHECK-NEXT: cir.store {{.*}} %[[MUL]], %[[GET_MEM_LHS]] : !cir.double, !cir.ptr<!cir.double>
+// CHECK-NEXT: %[[GET_MEM_LHS:.*]] = cir.get_member %[[LHSARG]][4] {name = "b"} : !cir.ptr<!rec_DefaultOperators> -> !cir.ptr<!cir.bool>
+// CHECK-NEXT: %[[GET_MEM_RHS:.*]] = cir.get_member %[[RHSARG]][4] {name = "b"} : !cir.ptr<!rec_DefaultOperators> -> !cir.ptr<!cir.bool>
+// CH...
[truncated]
|
Sema patch this depends on reverted here #162920 So we'll have to wait to merge this until that one gets re-added. |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Looks good, just a bunch of nits.
Decl::castToDeclContext(cgf.curFuncDecl), opInfo.origType, | ||
opInfo.bounds.size(), opInfo.boundTypes, opInfo.baseType, | ||
privateOp); | ||
privateOp, /*reductionCombinerRecipe=*/{}); |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
privateOp, /*reductionCombinerRecipe=*/{}); | |
privateOp, /*reductionCombinerRecipes=*/{}); |
// (which also shouldn't be an array type?). This function should generate the | ||
// loop to do this on each individual array or struct element (if necessary). | ||
auto emitCombiner = [&](mlir::Value lhsArg, mlir::Value rhsArg, QualType Ty) { | ||
if (const auto *RD = Ty->getAsRecordDecl()) { |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
if (const auto *RD = Ty->getAsRecordDecl()) { | |
if (const auto *rd = Ty->getAsRecordDecl()) { |
// Emitter for when we know this is either a non-array or element of an array | ||
// (which also shouldn't be an array type?). This function should generate the | ||
// loop to do this on each individual array or struct element (if necessary). | ||
auto emitCombiner = [&](mlir::Value lhsArg, mlir::Value rhsArg, QualType Ty) { |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
auto emitCombiner = [&](mlir::Value lhsArg, mlir::Value rhsArg, QualType Ty) { | |
auto emitCombiner = [&](mlir::Value lhsArg, mlir::Value rhsArg, QualType ty) { |
|
||
// Emitter for when we know this is either a non-array or element of an array | ||
// (which also shouldn't be an array type?). This function should generate the | ||
// loop to do this on each individual array or struct element (if necessary). |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I don't understand the comment. If we know this is a non-array, what does the part about a loop to do this on each individual array element mean?
|
||
} else { | ||
// if this is a single-thing (because we should know this isn't an array, | ||
// as Sema wouldn't let us get here), we can just do a normal emit call. |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Maybe assert that it's not an array?
Following on the Sema changes, this does the lowering for all of the operators that can be done as a compound operator. Lowering is very simply looping through the objects based on array/compound/etc, and doing a call to the operation.