[OpenACC][CIR] Start work to lower 'loop'

clang/lib/CIR/CodeGen/CIRGenStmtOpenACCLoop.cpp

@@ -13,17 +13,103 @@
 #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;
 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' 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<mlir::acc::YieldOp>(end);
+  }
+
+  return stmtRes;
 }

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,

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
+}

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)
 }