[OpenACC][CIR] Start work to lower 'loop' (llvm#137972)

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.

Author: erichkeane

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

@@ -1560,7 +1560,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,
@@ -1576,7 +1576,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)
 }