diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACCLoop.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACCLoop.cpp index b01ff85607939..2b78bc1a6d4a5 100644 --- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACCLoop.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACCLoop.cpp @@ -13,10 +13,12 @@ #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 +26,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 retTy; + llvm::SmallVector operands; + auto op = builder.create(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' equivalent 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 equivalent for the operator-call, as + // well as member-call. + // + // For builtin operators, we perhaps should refine the assignment here. It + // doesn't really 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(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 loc{{.*}}, %[[ARG_B:.*]]: !cir.ptr loc{{.*}}, %[[ARG_C:.*]]: !cir.ptr loc{{.*}}, %[[ARG_N:.*]]: !s32i loc{{.*}}) { + // CHECK-NEXT: %[[ALLOCA_A:.*]] = cir.alloca !cir.ptr, !cir.ptr>, ["A", init] + // CHECK-NEXT: %[[ALLOCA_B:.*]] = cir.alloca !cir.ptr, !cir.ptr>, ["B", init] + // CHECK-NEXT: %[[ALLOCA_C:.*]] = cir.alloca !cir.ptr, !cir.ptr>, ["C", init] + // CHECK-NEXT: %[[ALLOCA_N:.*]] = cir.alloca !s32i, !cir.ptr, ["N", init] + // CHECK-NEXT: cir.store %[[ARG_A]], %[[ALLOCA_A]] : !cir.ptr, !cir.ptr> + // CHECK-NEXT: cir.store %[[ARG_B]], %[[ALLOCA_B]] : !cir.ptr, !cir.ptr> + // CHECK-NEXT: cir.store %[[ARG_C]], %[[ALLOCA_C]] : !cir.ptr, !cir.ptr> + // CHECK-NEXT: cir.store %[[ARG_N]], %[[ALLOCA_N]] : !s32i, !cir.ptr + + +#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) }