llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: Erich Keane (erichkeane) <details> <summary>Changes</summary> As can be seen by the comment, this ends up being a construct that is going to be quite a lot of work in the future to make sure we properly identify the upperbound, lowerbound, and step. For now, we just treat the 'loop' as container so that we can put the 'for' loop into it. In the future, we'll have to teach the OpenACC dialect how to derive the upperbound, lowerbound, and step from the cir.for loop. Additionally, we'll probably have to add a few more options to it so that we can give it the recipes it needs to determine these for random access iterators. For Integer and Pointer values, these should already be known. --- Full diff: https://github.com/llvm/llvm-project/pull/137972.diff 4 Files Affected: - (modified) clang/lib/CIR/CodeGen/CIRGenStmtOpenACCLoop.cpp (+90-3) - (modified) clang/lib/Sema/SemaOpenACC.cpp (+2-2) - (added) clang/test/CIR/CodeGenOpenACC/loop.cpp (+33) - (modified) clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp (-6) ``````````diff diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACCLoop.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACCLoop.cpp index b01ff85607939..f26096f708f99 100644 --- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACCLoop.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACCLoop.cpp @@ -13,10 +13,13 @@ #include "CIRGenBuilder.h" #include "CIRGenFunction.h" #include "CIRGenOpenACCClause.h" -#include "mlir/Dialect/OpenACC/OpenACC.h" + #include "clang/AST/OpenACCClause.h" #include "clang/AST/StmtOpenACC.h" +#include "mlir/Dialect/OpenACC/OpenACC.h" + + using namespace clang; using namespace clang::CIRGen; using namespace cir; @@ -24,6 +27,90 @@ using namespace mlir::acc; mlir::LogicalResult CIRGenFunction::emitOpenACCLoopConstruct(const OpenACCLoopConstruct &s) { - cgm.errorNYI(s.getSourceRange(), "OpenACC Loop Construct"); - return mlir::failure(); + mlir::Location start = getLoc(s.getSourceRange().getBegin()); + mlir::Location end = getLoc(s.getSourceRange().getEnd()); + llvm::SmallVector<mlir::Type> retTy; + llvm::SmallVector<mlir::Value> operands; + auto op = builder.create<LoopOp>(start, retTy, operands); + + // TODO(OpenACC): In the future we are going to need to come up with a + // transformation here that can teach the acc.loop how to figure out the + // 'lowerbound', 'upperbound', and 'step'. + // + // -'upperbound' should fortunately be pretty easy as it should be + // in the initialization section of the cir.for loop. In Sema, we limit to + // just the forms 'Var = init', `Type Var = init`, or `Var = init` (where it + // is an operator= call)`. However, as those are all necessary to emit for + // the init section of the for loop, they should be inside the initial + // cir.scope. + // + // -'upperbound' should be somewhat easy to determine. Sema is limiting this + // to: ==, <, >, !=, <=, >= builtin operators, the overloaded 'comparison' + // operations, and member-call expressions. + // + // For the builtin comparison operators, we can pretty well deduce based on + // the comparison what the 'end' object is going to be, and the inclusive + // nature of it. + // + // For the overloaded operators, Sema will ensure that at least one side of + // the operator is the init variable, so we can deduce the comparison there + // too. The standard places no real bounds on WHAT the comparison operators do + // for a `RandomAccessIterator` however, so we'll have to just 'assume' they + // do the right thing? Note that this might be incrementing by a different + // 'object', not an integral, so it isn't really clear to me what we can do to + // determine the other side. + // + // Member-call expressions are the difficult ones. I don't think there is + // anything we can deduce from this to determine the 'end', so we might end up + // having to go back to Sema and make this ill-formed. + // + // HOWEVER: What ACC dialect REALLY cares about is the tripcount, which you + // cannot get (in the case of `RandomAccessIterator`) from JUST 'upperbound' + // and 'lowerbound'. We will likely have to provide a 'recipe' equivilent to + // `std::distance` instead. In the case of integer/pointers, it is fairly + // simple to find: it is just the mathematical subtraction. Howver, in the + // case of `RandomAccessIterator`, we have to enable the use of `operator-`. + // FORTUNATELY the standard requires this to work correctly for + // `RandomAccessIterator`, so we don't have to implement a `std::distance` + // that loops through, like we would for a forward/etc iterator. + // + // 'step': Sema is currently allowing builtin ++,--, +=, -=, *=, /=, and = + // operators. Additionally, it allows the equivilent for the operator-call, as + // well as member-call. + // + // For builtin operators, we perhaps should refine the assignment here. It + // doesn't reallly help us know the 'step' count at all, but we could perhaps + // do one more step of analysis in Sema to allow something like Var = Var + 1. + // For the others, this should get us the step reasonably well. + // + // For the overloaded operators, we have the same problems as for + // 'upperbound', plus not really knowing what they do. Member-call expressions + // are again difficult, and we might want to reconsider allowing these in + // Sema. + // + + // Emit all clauses. + { + mlir::OpBuilder::InsertionGuard guardCase(builder); + // Sets insertion point before the 'op', since every new expression needs to + // be before the operation. + builder.setInsertionPoint(op); + makeClauseEmitter(op, *this, builder, s.getDirectiveKind(), + s.getDirectiveLoc()) + .VisitClauseList(s.clauses()); + } + + mlir::LogicalResult stmtRes = mlir::success(); + // Emit body. + { + mlir::Block &block = op.getRegion().emplaceBlock(); + mlir::OpBuilder::InsertionGuard guardCase(builder); + builder.setInsertionPointToEnd(&block); + LexicalScope ls{*this, start, builder.getInsertionBlock()}; + + stmtRes = emitStmt(s.getLoop(), /*useCurrentScope=*/true); + builder.create<mlir::acc::YieldOp>(end); + } + + return stmtRes; } diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp index fc191bd95a35e..1fb9fa4d63c16 100644 --- a/clang/lib/Sema/SemaOpenACC.cpp +++ b/clang/lib/Sema/SemaOpenACC.cpp @@ -1517,7 +1517,7 @@ void SemaOpenACC::ActOnForStmtBegin(SourceLocation ForLoc, const Stmt *First, void SemaOpenACC::ActOnRangeForStmtBegin(SourceLocation ForLoc, const Stmt *OldRangeFor, const Stmt *RangeFor) { - if (!getLangOpts().OpenACC) + if (!getLangOpts().OpenACC || OldRangeFor == nullptr || RangeFor == nullptr) return; ForStmtBeginChecker FSBC{*this, ForLoc, @@ -1533,7 +1533,7 @@ void SemaOpenACC::ActOnRangeForStmtBegin(SourceLocation ForLoc, void SemaOpenACC::ActOnRangeForStmtBegin(SourceLocation ForLoc, const Stmt *RangeFor) { - if (!getLangOpts().OpenACC) + if (!getLangOpts().OpenACC || RangeFor == nullptr) return; ForStmtBeginChecker FSBC = {*this, ForLoc, diff --git a/clang/test/CIR/CodeGenOpenACC/loop.cpp b/clang/test/CIR/CodeGenOpenACC/loop.cpp new file mode 100644 index 0000000000000..792edfedaacc6 --- /dev/null +++ b/clang/test/CIR/CodeGenOpenACC/loop.cpp @@ -0,0 +1,33 @@ +// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s + +extern "C" void acc_loop(int *A, int *B, int *C, int N) { + // CHECK: cir.func @acc_loop(%[[ARG_A:.*]]: !cir.ptr<!s32i> loc{{.*}}, %[[ARG_B:.*]]: !cir.ptr<!s32i> loc{{.*}}, %[[ARG_C:.*]]: !cir.ptr<!s32i> loc{{.*}}, %[[ARG_N:.*]]: !s32i loc{{.*}}) { + // CHECK-NEXT: %[[ALLOCA_A:.*]] = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["A", init] + // CHECK-NEXT: %[[ALLOCA_B:.*]] = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["B", init] + // CHECK-NEXT: %[[ALLOCA_C:.*]] = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["C", init] + // CHECK-NEXT: %[[ALLOCA_N:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["N", init] + // CHECK-NEXT: cir.store %[[ARG_A]], %[[ALLOCA_A]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>> + // CHECK-NEXT: cir.store %[[ARG_B]], %[[ALLOCA_B]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>> + // CHECK-NEXT: cir.store %[[ARG_C]], %[[ALLOCA_C]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>> + // CHECK-NEXT: cir.store %[[ARG_N]], %[[ALLOCA_N]] : !s32i, !cir.ptr<!s32i> + + +#pragma acc loop + for (unsigned I = 0u; I < N; ++I) { + A[I] = B[I] + C[I]; + } + // CHECK-NEXT: acc.loop { + // CHECK-NEXT: cir.scope { + // CHECK: cir.for : cond { + // CHECK: cir.condition + // CHECK-NEXT: } body { + // CHECK-NEXT: cir.scope { + // CHECK: } + // CHECK-NEXT: cir.yield + // CHECK-NEXT: } step { + // CHECK: cir.yield + // CHECK-NEXT: } loc + // CHECK-NEXT: } loc + // CHECK-NEXT: acc.yield + // CHECK-NEXT: } loc +} diff --git a/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp b/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp index a7a179c0b2e3c..e95d4b8bfacbd 100644 --- a/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp +++ b/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp @@ -9,12 +9,6 @@ void HelloWorld(int *A, int *B, int *C, int N) { for (unsigned I = 0; I < N; ++I) A[I] = B[I] + C[I]; -// expected-error@+2{{ClangIR code gen Not Yet Implemented: OpenACC Loop Construct}} -// expected-error@+1{{ClangIR code gen Not Yet Implemented: statement}} -#pragma acc loop - for (unsigned I = 0; I < N; ++I) - A[I] = B[I] + C[I]; - // expected-error@+1{{ClangIR code gen Not Yet Implemented: OpenACC Declare Construct}} #pragma acc declare create(A) } `````````` </details> https://github.com/llvm/llvm-project/pull/137972 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
