[mlir] NFC: Rename LoopOps dialect to SCF (Structured Control Flow)

[ftynse]

Summary

This dialect contains various structured control flow operaitons, not only
loops, reflect this in the name. Drop the Ops suffix for consistency with other
dialects.

Note that this only moves the files and changes the C++ namespace from 'loop'
to 'scf'. The visible IR prefix remains the same and will be updated
separately. The conversions will also be updated separately.

[PeteSteinfeld]

@ftynse and @antiagainst , note that one of the clients of mlir is the flang project. Since this revision changed the interface of mlir, all clients of mlir should be checked to see if they still build and test correctly before delivering such changes. Note that the person who knows most about flang's dependency on mlir is @schweitz, but any member of the flang team should be able to verify if a change works for flang.

[ftynse]

In D79578#2029705, @PeteSteinfeld wrote:

@ftynse and @antiagainst , note that one of the clients of mlir is the flang project. Since this revision changed the interface of mlir, all clients of mlir should be checked to see if they still build and test correctly before delivering such changes. Note that the person who knows most about flang's dependency on mlir is @schweitz, but any member of the flang team should be able to verify if a change works for flang.

I'm sorry this broke your build, I had approved the fix that was mailed before this message.

I don't know who are the members of the flang team (is there a list?) and would feel concerned if we need to block any changes to MLIR APIs for review by people other than regular MLIR reviewers. If something is broken by an identifiable commit, and the fix isn't obvious, the usual policy in LLVM is to revert and report the problem to the author of the commit. Sometimes, it may be faster to just write the author a quick direct email. Generally, this is something that can and should be solved with a pre-merge check, but the one we have is broken for random reasons so often that it becomes useless.

[mehdi_amini]

In D79578#2029762, @ftynse wrote:

I don't know who are the members of the flang team (is there a list?) and would feel concerned if we need to block any changes to MLIR APIs for review by people other than regular MLIR reviewers.

I don't think this is the request. The request is just "don't break the build".
The pre-merge testing will test flang as well, but you would have had to rebase the patch after you landed the dependent ones first so that the pre-merge testing could work (I think marking the dependencies in Phabricator would have worked also).

In general trying to get the pre-merge testing green before pushing changes is safer.

[PeteSteinfeld]

I don't think this is the request. The request is just "don't break the build".
The pre-merge testing will test flang as well, but you would have had to rebase the patch after you landed the dependent ones first so that the pre-merge testing could work (I think marking the dependencies in Phabricator would have worked also).

In general trying to get the pre-merge testing green before pushing changes is safer.

Yes, the request is to not break the build. And in the case where the pre-merge testing is not working, I encourage you to contact @schweitz or @PeteSteinfeld or any other member of the flang team to verify that our builds succeed.

[ftynse]

In D79578#2029832, @PeteSteinfeld wrote:

I don't think this is the request. The request is just "don't break the build".
The pre-merge testing will test flang as well, but you would have had to rebase the patch after you landed the dependent ones first so that the pre-merge testing could work (I think marking the dependencies in Phabricator would have worked also).

In general trying to get the pre-merge testing green before pushing changes is safer.

Yes, the request is to not break the build.

This is standard policy. So is reverting the breaking changes. Sometimes we do break stuff and there are technical solutions to decrease the probability of this happening.

I am however opposed to having to contact someone other than people who would normally review the patch. If something is not covered by check-flang (which I forgot to run, I admit, that's why we need automation!) or any other simple command, I think it is unreasonable to block us from working and wait for somebody to think about how the changes might affect dependent projects.

And in the case where the pre-merge testing is not working, I encourage you to contact @schweitz or @PeteSteinfeld or any other member of the flang team to verify that our builds succeed.

Thanks for offering help! Note that I still don't know who "any other member"s are, is there a list? a chat channel?

[PeteSteinfeld]

Thanks for offering help! Note that I still don't know who "any other member"s are, is there a list? a chat channel?

Thanks, @ftynse. It would have been sufficient to run check-flang.

There's a Slack channel for Fortran development at https://flang-compiler.slack.com/. There's also a mailing list at https://lists.llvm.org/pipermail/flang-dev/.

Pete

[ftynse]

In D79578#2031354, @PeteSteinfeld wrote:

Thanks for offering help! Note that I still don't know who "any other member"s are, is there a list? a chat channel?

Thanks, @ftynse. It would have been sufficient to run check-flang.

Great! I'll just make an alias that runs both mlir and flang tests.

There's a Slack channel for Fortran development at https://flang-compiler.slack.com/. There's also a mailing list at https://lists.llvm.org/pipermail/flang-dev/.

Thanks for the information!

[mehdi_amini]

In D79578#2031204, @ftynse wrote:

In D79578#2029832, @PeteSteinfeld wrote:

I don't think this is the request. The request is just "don't break the build".
The pre-merge testing will test flang as well, but you would have had to rebase the patch after you landed the dependent ones first so that the pre-merge testing could work (I think marking the dependencies in Phabricator would have worked also).

In general trying to get the pre-merge testing green before pushing changes is safer.

Yes, the request is to not break the build.

This is standard policy. So is reverting the breaking changes. Sometimes we do break stuff and there are technical solutions to decrease the probability of this happening.

I am however opposed to having to contact someone other than people who would normally review the patch. If something is not covered by check-flang (which I forgot to run, I admit, that's why we need automation!) or any other simple command, I think it is unreasonable to block us from working and wait for somebody to think about how the changes might affect dependent projects.

This is not what was asked: it was mentioned that some people can *help*, I think if you can't figure out yourself how to not break flang or if you don't understand a flang failure in the pre-merge testing for example.

[kiranchandramohan]

There is now a buildbot for flang.
http://lab.llvm.org:8011/builders/flang-aarch64-ubuntu

The flang projects page lists some of the members.
https://reviews.llvm.org/project/members/73/

mlir/examples/toy/Ch6/mlir/LowerToLLVM.cpp

 #include "toy/Passes.h"
 
 #include "mlir/Conversion/AffineToStandard/AffineToStandard.h"
 #include "mlir/Conversion/LoopToStandard/ConvertLoopToStandard.h"
 #include "mlir/Conversion/StandardToLLVM/ConvertStandardToLLVM.h"
 #include "mlir/Conversion/StandardToLLVM/ConvertStandardToLLVMPass.h"
 #include "mlir/Dialect/Affine/IR/AffineOps.h"
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
-#include "mlir/Dialect/LoopOps/LoopOps.h"
+#include "mlir/Dialect/SCF/SCF.h"
 #include "mlir/Dialect/StandardOps/IR/Ops.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Transforms/DialectConversion.h"
 #include "llvm/ADT/Sequence.h"
 
 using namespace mlir;
 
 //===----------------------------------------------------------------------===//
matchAndRewrite(Operation *op, ArrayRef<Value> operands,
 
     // Create a loop for each of the dimensions within the shape.
     SmallVector<Value, 4> loopIvs;
     for (unsigned i = 0, e = memRefShape.size(); i != e; ++i) {
       auto lowerBound = rewriter.create<ConstantIndexOp>(loc, 0);
       auto upperBound = rewriter.create<ConstantIndexOp>(loc, memRefShape[i]);
       auto step = rewriter.create<ConstantIndexOp>(loc, 1);
       auto loop =
-          rewriter.create<loop::ForOp>(loc, lowerBound, upperBound, step);
+          rewriter.create<scf::ForOp>(loc, lowerBound, upperBound, step);
       loop.getBody()->clear();
       loopIvs.push_back(loop.getInductionVar());
 
       // Terminate the loop body.
       rewriter.setInsertionPointToStart(loop.getBody());
 
       // Insert a newline after each of the inner dimensions of the shape.
       if (i != e - 1)
         rewriter.create<CallOp>(loc, printfRef, rewriter.getIntegerType(32),
                                 newLineCst);
-      rewriter.create<loop::YieldOp>(loc);
+      rewriter.create<scf::YieldOp>(loc);
       rewriter.setInsertionPointToStart(loop.getBody());
     }
 
     // Generate a call to printf for the current element of the loop.
     auto printOp = cast<toy::PrintOp>(op);
     auto elementLoad = rewriter.create<LoadOp>(loc, printOp.input(), loopIvs);
     rewriter.create<CallOp>(loc, printfRef, rewriter.getIntegerType(32),
                             ArrayRef<Value>({formatSpecifierCst, elementLoad}));

mlir/examples/toy/Ch7/mlir/LowerToLLVM.cpp

 #include "toy/Passes.h"
 
 #include "mlir/Conversion/AffineToStandard/AffineToStandard.h"
 #include "mlir/Conversion/LoopToStandard/ConvertLoopToStandard.h"
 #include "mlir/Conversion/StandardToLLVM/ConvertStandardToLLVM.h"
 #include "mlir/Conversion/StandardToLLVM/ConvertStandardToLLVMPass.h"
 #include "mlir/Dialect/Affine/IR/AffineOps.h"
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
-#include "mlir/Dialect/LoopOps/LoopOps.h"
+#include "mlir/Dialect/SCF/SCF.h"
 #include "mlir/Dialect/StandardOps/IR/Ops.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Transforms/DialectConversion.h"
 #include "llvm/ADT/Sequence.h"
 
 using namespace mlir;
 
 //===----------------------------------------------------------------------===//
matchAndRewrite(Operation *op, ArrayRef<Value> operands,
 
     // Create a loop for each of the dimensions within the shape.
     SmallVector<Value, 4> loopIvs;
     for (unsigned i = 0, e = memRefShape.size(); i != e; ++i) {
       auto lowerBound = rewriter.create<ConstantIndexOp>(loc, 0);
       auto upperBound = rewriter.create<ConstantIndexOp>(loc, memRefShape[i]);
       auto step = rewriter.create<ConstantIndexOp>(loc, 1);
       auto loop =
-          rewriter.create<loop::ForOp>(loc, lowerBound, upperBound, step);
+          rewriter.create<scf::ForOp>(loc, lowerBound, upperBound, step);
       loop.getBody()->clear();
       loopIvs.push_back(loop.getInductionVar());
 
       // Terminate the loop body.
       rewriter.setInsertionPointToStart(loop.getBody());
 
       // Insert a newline after each of the inner dimensions of the shape.
       if (i != e - 1)
         rewriter.create<CallOp>(loc, printfRef, rewriter.getIntegerType(32),
                                 newLineCst);
-      rewriter.create<loop::YieldOp>(loc);
+      rewriter.create<scf::YieldOp>(loc);
       rewriter.setInsertionPointToStart(loop.getBody());
     }
 
     // Generate a call to printf for the current element of the loop.
     auto printOp = cast<toy::PrintOp>(op);
     auto elementLoad = rewriter.create<LoadOp>(loc, printOp.input(), loopIvs);
     rewriter.create<CallOp>(loc, printfRef, rewriter.getIntegerType(32),
                             ArrayRef<Value>({formatSpecifierCst, elementLoad}));

mlir/include/mlir/Conversion/LoopsToGPU/LoopsToGPU.h

 
 namespace mlir {
 class AffineForOp;
 class MLIRContext;
 class OwningRewritePatternList;
 struct LogicalResult;
 class Value;
 
-namespace loop {
+namespace scf {
 class ForOp;
-} // end namespace loop
+} // end namespace scf
 
 /// Convert a perfect affine loop nest with the outermost loop identified by
 /// `forOp` into a gpu::Launch operation.  Map `numBlockDims` outer loops to
 /// GPU blocks and `numThreadDims` to GPU threads.  The bounds of the loops that
 /// are mapped should be independent of the induction variables of the other
 /// mapped loops.
 ///
 /// No check on the size of the block or grid, or on the validity of

 /// GPU blocks and `numThreadDims` to GPU threads.  The bounds of the loops that
 /// are mapped should be independent of the induction variables of the other
 /// mapped loops.
 ///
 /// No check on the size of the block or grid, or on the validity of
 /// parallelization is performed, it is under the responsibility of the caller
 /// to strip-mine the loops and to perform the dependence analysis before
 /// calling the conversion.
-LogicalResult convertLoopNestToGPULaunch(loop::ForOp forOp,
+LogicalResult convertLoopNestToGPULaunch(scf::ForOp forOp,
                                          unsigned numBlockDims,
                                          unsigned numThreadDims);
 
 /// Convert a loop operation into a GPU launch with the values provided in
 /// `numWorkGroups` as the grid size and the values provided in `workGroupSizes`
 /// as the block size. Size of `numWorkGroups` and workGroupSizes` must be less
 /// than or equal to 3. The loop operation can be an imperfectly nested
 /// computation with the following restrictions:
 /// 1) The loop nest must contain as many perfectly nested loops as the number
 /// of values passed in through `numWorkGroups`. This corresponds to the number
 /// of grid dimensions of the launch. All loops within the loop nest must be
 /// parallel.
 /// 2) The body of the innermost loop of the above perfectly nested loops, must
 /// contain statements that satisfy one of the two conditions below:
 ///   a) A perfect loop nest of depth greater than or equal to the number of
 ///   values passed in through `workGroupSizes`, i.e. the number of thread
 ///   dimensions of the launch. Loops at depth less than or equal to size of
 ///   `workGroupSizes` must be parallel. Loops nested deeper can be sequential
 ///   and are retained as such in the generated GPU launch code.
 ///   b) Statements that are safe to be executed by all threads within the
 ///   workgroup. No checks are performed that this is indeed the case.
 ///   TODO(ravishankarm) : Add checks that verify 2(b) above.
 /// The above conditions are assumed to be satisfied by the computation rooted
 /// at `forOp`.
-LogicalResult convertLoopToGPULaunch(loop::ForOp forOp,
+LogicalResult convertLoopToGPULaunch(scf::ForOp forOp,
                                      ArrayRef<Value> numWorkGroups,
                                      ArrayRef<Value> workGroupSizes);
 
 /// Adds the conversion pattern from `loop.parallel` to `gpu.launch` to the
 /// provided pattern list.
 void populateParallelLoopToGPUPatterns(OwningRewritePatternList &patterns,
                                        MLIRContext *ctx);
 
 } // namespace mlir
 
 #endif // MLIR_CONVERSION_LOOPSTOGPU_LOOPSTOGPU_H_

mlir/include/mlir/Conversion/Passes.td

let options = [
     Option<"numBlockDims", "gpu-block-dims", "unsigned", /*default=*/"1u",
            "Number of GPU block dimensions for mapping">,
     Option<"numThreadDims", "gpu-thread-dims", "unsigned", /*default=*/"1u",
            "Number of GPU thread dimensions for mapping">
   ];
 }
 
 def ConvertLoopsToGPU : FunctionPass<"convert-loop-op-to-gpu"> {
-  let summary = "Convert top-level loop::ForOp to GPU kernels";
+  let summary = "Convert top-level scf::ForOp to GPU kernels";
   let constructor = "mlir::createLoopToGPUPass()";
   let options = [
     ListOption<"numWorkGroups", "gpu-num-workgroups", "int64_t",
                "Num workgroups in the GPU launch",
                "llvm::cl::ZeroOrMore, llvm::cl::MiscFlags::CommaSeparated">,
     ListOption<"workGroupSize", "gpu-workgroup-size", "int64_t",
                "Workgroup Size in the GPU launch",
                "llvm::cl::ZeroOrMore, llvm::cl::MiscFlags::CommaSeparated">

mlir/include/mlir/Dialect/CMakeLists.txt

 add_subdirectory(Affine)
 add_subdirectory(AVX512)
 add_subdirectory(GPU)
 add_subdirectory(Linalg)
 add_subdirectory(LLVMIR)
-add_subdirectory(LoopOps)
 add_subdirectory(OpenMP)
 add_subdirectory(Quant)
+add_subdirectory(SCF)
 add_subdirectory(Shape)
 add_subdirectory(SPIRV)
 add_subdirectory(StandardOps)
 add_subdirectory(Vector)

mlir/include/mlir/Dialect/GPU/ParallelLoopMapper.h

 
 class AffineMap;
 struct LogicalResult;
 class Operation;
 class Region;
 
 #include "mlir/Dialect/GPU/ParallelLoopMapperAttr.h.inc"
 
-namespace loop {
+namespace scf {
 class ParallelOp;
 }
 
 namespace gpu {
 
 /// Name of the mapping attribute produced by loop mappers.
 StringRef getMappingAttrName();
 
 /// Get the value of the processor in the ParallelLoopDimMapping attribute.
 inline Processor getProcessor(ParallelLoopDimMapping attr) {
   return static_cast<Processor>(attr.processor().getInt());
 }
 
 /// Helper function to create a ParallelDimMapperAttr.
 /// TODO(ravishankarm/antiagainst): Replace its uses with an auto-gened method.
 ParallelLoopDimMapping getParallelLoopDimMappingAttr(Processor processor,
                                                      AffineMap map,
                                                      AffineMap bound);
 
 /// Sets the mapping attribute of a loop.parallel operation. Verifies that the
 /// mapping passed is valid.
 /// - the number of DimMapperAttr provided is same as the number of loops of
 ///   the `ploopOp`.
 /// - the mapping does not map multiple loops to the same processor.
-LogicalResult setMappingAttr(loop::ParallelOp ploopOp,
+LogicalResult setMappingAttr(scf::ParallelOp ploopOp,
                              ArrayRef<ParallelLoopDimMapping> mapping);
 } // end namespace gpu
 
 /// Maps the parallel loops found in the given function to workgroups. The first
 /// loop encountered will be mapped to the global workgroup and the second loop
 /// encountered to the local workgroup. Within each mapping, the first three
 /// dimensions are mapped to x/y/z hardware ids and all following dimensions are
 /// mapped to sequential loops.
 void greedilyMapParallelLoopsToGPU(Region &region);
 
 } // end namespace mlir
 #endif // MLIR_DIALECT_GPU_PARALLELLOOPMAPPER_H

mlir/include/mlir/Dialect/Linalg/EDSC/Builders.h

 #include "mlir/IR/AffineExpr.h"
 #include "mlir/IR/Builders.h"
 
 namespace mlir {
 class AffineForOp;
 class BlockArgument;
 class SubViewOp;
 
-namespace loop {
+namespace scf {
 class ParallelOp;
-} // namespace loop
+} // namespace scf
 
 namespace edsc {
 class AffineLoopNestBuilder;
 class ParallelLoopNestBuilder;
 
 /// A LoopRangeBuilder is a generic NestedBuilder for loop.for operations.
 /// More specifically it is meant to be used as a temporary object for
 /// representing any nested MLIR construct that is "related to" an mlir::Value
GenericLoopNestRangeBuilder(MutableArrayRef<Value> ivs,
                               ArrayRef<Value> ranges);
   void operator()(std::function<void(void)> fun = nullptr) { (*builder)(fun); }
 
 private:
   using LoopOrAffineLoopBuilder =
       typename std::conditional_t<std::is_same<LoopTy, AffineForOp>::value,
                                   AffineLoopNestBuilder, LoopNestRangeBuilder>;
   using BuilderType =
-      typename std::conditional_t<std::is_same<LoopTy, loop::ParallelOp>::value,
+      typename std::conditional_t<std::is_same<LoopTy, scf::ParallelOp>::value,
                                   ParallelLoopNestBuilder,
                                   LoopOrAffineLoopBuilder>;
 
   std::unique_ptr<BuilderType> builder;
 };
 
 inline void defaultRegionBuilder(ArrayRef<BlockArgument> args) {}
 

mlir/include/mlir/Dialect/Linalg/Utils/Utils.h

 //===- Utils.h - Utilities to support the Linalg dialect --------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef MLIR_DIALECT_LINALG_UTILS_H_
 #define MLIR_DIALECT_LINALG_UTILS_H_
 
 #include "mlir/Dialect/Linalg/IR/LinalgOps.h"
-#include "mlir/Dialect/LoopOps/LoopOps.h"
+#include "mlir/Dialect/SCF/SCF.h"
 #include "mlir/Dialect/StandardOps/IR/Ops.h"
 
 #include "llvm/ADT/SetVector.h"
 
 namespace mlir {
 class AffineExpr;
 class AffineMap;
 class OperationFolder;

mlir/include/mlir/Dialect/LoopOps/CMakeLists.txt

-add_mlir_dialect(LoopOps loop)
-add_mlir_doc(LoopOps -gen-dialect-doc LoopDialect Dialects/)
-
-set(LLVM_TARGET_DEFINITIONS Passes.td)
-mlir_tablegen(Passes.h.inc -gen-pass-decls)
-add_public_tablegen_target(MLIRLoopPassIncGen)
-add_dependencies(mlir-headers MLIRLoopPassIncGen)
-
-add_mlir_doc(Passes -gen-pass-doc LoopPasses ./)

mlir/include/mlir/Dialect/SCF/CMakeLists.txt

+add_mlir_dialect(SCFOps loop Ops)
+add_mlir_doc(SCFOps -gen-dialect-doc SCFDialect Dialects/)
+
+set(LLVM_TARGET_DEFINITIONS Passes.td)
+mlir_tablegen(Passes.h.inc -gen-pass-decls)
+add_public_tablegen_target(MLIRSCFPassIncGen)
+add_dependencies(mlir-headers MLIRSCFPassIncGen)
+
+add_mlir_doc(Passes -gen-pass-doc SCFPasses ./)

mlir/include/mlir/Dialect/SCF/EDSC/Builders.h

 //===- Builders.h - MLIR Declarative Builder Classes ------------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // Provides intuitive composable interfaces for building structured MLIR
 // snippets in a declarative fashion.
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef MLIR_DIALECT_LOOPOPS_EDSC_BUILDERS_H_
-#define MLIR_DIALECT_LOOPOPS_EDSC_BUILDERS_H_
+#ifndef MLIR_DIALECT_SCF_EDSC_BUILDERS_H_
+#define MLIR_DIALECT_SCF_EDSC_BUILDERS_H_
 
-#include "mlir/Dialect/LoopOps/LoopOps.h"
+#include "mlir/Dialect/SCF/SCF.h"
 #include "mlir/EDSC/Builders.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/Types.h"
 
 namespace mlir {
 namespace edsc {
 
-/// Constructs a new loop::ParallelOp and captures the associated induction
+/// Constructs a new scf::ParallelOp and captures the associated induction
 /// variables. An array of Value pointers is passed as the first
 /// argument and is the *only* way to capture loop induction variables.
 LoopBuilder makeParallelLoopBuilder(MutableArrayRef<Value> ivs,
                                     ArrayRef<Value> lbs, ArrayRef<Value> ubs,
                                     ArrayRef<Value> steps);
-/// Constructs a new loop::ForOp and captures the associated induction
+/// Constructs a new scf::ForOp and captures the associated induction
 /// variable. A Value pointer is passed as the first argument and is the
 /// *only* way to capture the loop induction variable.
 LoopBuilder makeLoopBuilder(Value *iv, Value lb, Value ub, Value step,
                             MutableArrayRef<Value> iterArgsHandles,
                             ValueRange iterArgsInitValues);
 LoopBuilder makeLoopBuilder(Value *iv, Value lb, Value ub, Value step,
                             MutableArrayRef<Value> iterArgsHandles,
                             ValueRange iterArgsInitValues);

 
 private:
   SmallVector<LoopBuilder, 4> loops;
 };
 
 } // namespace edsc
 } // namespace mlir
 
-#endif // MLIR_DIALECT_LOOPOPS_EDSC_BUILDERS_H_
+#endif // MLIR_DIALECT_SCF_EDSC_BUILDERS_H_

mlir/include/mlir/Dialect/SCF/EDSC/Intrinsics.h

-//===- Intrinsics.h - MLIR EDSC Intrinsics for Linalg -----------*- C++ -*-===//
+//===- Intrinsics.h - MLIR EDSC Intrinsics for SCF --------------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM
 // Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
-#ifndef MLIR_DIALECT_LOOPOPS_EDSC_INTRINSICS_H_
-#define MLIR_DIALECT_LOOPOPS_EDSC_INTRINSICS_H_
+#ifndef MLIR_DIALECT_SCF_EDSC_INTRINSICS_H_
+#define MLIR_DIALECT_SCF_EDSC_INTRINSICS_H_
 
-#include "mlir/Dialect/LoopOps/EDSC/Builders.h"
+#include "mlir/Dialect/SCF/EDSC/Builders.h"
 
 namespace mlir {
 namespace edsc {
 namespace intrinsics {
 
-using loop_yield = OperationBuilder<loop::YieldOp>;
+using loop_yield = OperationBuilder<scf::YieldOp>;
 
 } // namespace intrinsics
 } // namespace edsc
 } // namespace mlir
 
-#endif // MLIR_DIALECT_LOOPOPS_EDSC_INTRINSICS_H_
+#endif // MLIR_DIALECT_SCF_EDSC_INTRINSICS_H_

mlir/include/mlir/Dialect/SCF/Passes.h

 //===- Passes.h - Pass Entrypoints ------------------------------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This header file defines prototypes that expose pass constructors.
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef MLIR_DIALECT_LOOPOPS_PASSES_H_
-#define MLIR_DIALECT_LOOPOPS_PASSES_H_
+#ifndef MLIR_DIALECT_SCF_PASSES_H_
+#define MLIR_DIALECT_SCF_PASSES_H_
 
 #include "llvm/ADT/ArrayRef.h"
 #include <memory>
 
 namespace mlir {
 
 class Pass;
 
 /// Creates a loop fusion pass which fuses parallel loops.
 std::unique_ptr<Pass> createParallelLoopFusionPass();
 
 /// Creates a pass that specializes parallel loop for unrolling and
 /// vectorization.
 std::unique_ptr<Pass> createParallelLoopSpecializationPass();
 
 /// Creates a pass which tiles innermost parallel loops.
 std::unique_ptr<Pass>
 createParallelLoopTilingPass(llvm::ArrayRef<int64_t> tileSize = {});
 
 } // namespace mlir
 
-#endif // MLIR_DIALECT_LOOPOPS_PASSES_H_
+#endif // MLIR_DIALECT_SCF_PASSES_H_

mlir/include/mlir/Dialect/SCF/Passes.td

 //===-- Passes.td - Loop pass definition file --------------*- tablegen -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef MLIR_DIALECT_LOOP_PASSES
-#define MLIR_DIALECT_LOOP_PASSES
+#ifndef MLIR_DIALECT_SCF_PASSES
+#define MLIR_DIALECT_SCF_PASSES
 
 include "mlir/Pass/PassBase.td"
 
 def LoopParallelLoopFusion : Pass<"parallel-loop-fusion"> {
   let summary = "Fuse adjacent parallel loops";
   let constructor = "mlir::createParallelLoopFusionPass()";
 }
 
 def LoopParallelLoopSpecialization
     : FunctionPass<"parallel-loop-specialization"> {
   let summary = "Specialize parallel loops for vectorization";
   let constructor = "mlir::createParallelLoopSpecializationPass()";
 }
 
 def LoopParallelLoopTiling : FunctionPass<"parallel-loop-tiling"> {
   let summary = "Tile parallel loops";
   let constructor = "mlir::createParallelLoopTilingPass()";
   let options = [
     ListOption<"tileSizes", "parallel-loop-tile-sizes", "int64_t",
                "Factors to tile parallel loops by",
                "llvm::cl::ZeroOrMore, llvm::cl::MiscFlags::CommaSeparated">
   ];
 }
 
-#endif // MLIR_DIALECT_LOOP_PASSES
+#endif // MLIR_DIALECT_SCF_PASSES

mlir/include/mlir/Dialect/SCF/SCF.h

-//===- Ops.h - Loop MLIR Operations -----------------------------*- C++ -*-===//
+//===- SCFOps.h - Structured Control Flow -----------------------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
-// This file defines convenience types for working with loop operations.
+// This file defines structured control flow operations.
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef MLIR_LOOPOPS_OPS_H_
-#define MLIR_LOOPOPS_OPS_H_
+#ifndef MLIR_DIALECT_SCF_H_
+#define MLIR_DIALECT_SCF_H_
 
 #include "mlir/IR/Attributes.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/Dialect.h"
 #include "mlir/IR/OpDefinition.h"
 #include "mlir/Interfaces/ControlFlowInterfaces.h"
 #include "mlir/Interfaces/LoopLikeInterface.h"
 #include "mlir/Interfaces/SideEffects.h"
 
 namespace mlir {
-namespace loop {
+namespace scf {
 
 class TerminatorOp;
 
-#include "mlir/Dialect/LoopOps/LoopOpsDialect.h.inc"
+#include "mlir/Dialect/SCF/SCFOpsDialect.h.inc"
 
 #define GET_OP_CLASSES
-#include "mlir/Dialect/LoopOps/LoopOps.h.inc"
+#include "mlir/Dialect/SCF/SCFOps.h.inc"
 
 // Insert `loop.terminator` at the end of the only region's only block if it
 // does not have a terminator already.  If a new `loop.terminator` is inserted,
 // the location is specified by `loc`. If the region is empty, insert a new
 // block first.
 void ensureLoopTerminator(Region &region, Builder &builder, Location loc);
 
 /// Returns the loop parent of an induction variable. If the provided value is
 /// not an induction variable, then return nullptr.
 ForOp getForInductionVarOwner(Value val);
 
 /// Returns the parallel loop parent of an induction variable. If the provided
 // value is not an induction variable, then return nullptr.
 ParallelOp getParallelForInductionVarOwner(Value val);
 
-} // end namespace loop
+} // end namespace scf
 } // end namespace mlir
-#endif // MLIR_LOOPOPS_OPS_H_
+#endif // MLIR_DIALECT_SCF_H_

mlir/include/mlir/Dialect/SCF/SCFOps.td

-//===- Ops.td - Loop operation definitions ---------------*- tablegen -*-===//
+//===- SCFOps.td - Structured Control Flow operations ------*- tablegen -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
-// Defines MLIR loop operations.
+// Defines MLIR structured control flow operations.
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef LOOP_OPS
-#define LOOP_OPS
+#ifndef MLIR_DIALECT_SCF_SCFOPS
+#define MLIR_DIALECT_SCF_SCFOPS
 
 include "mlir/Interfaces/ControlFlowInterfaces.td"
 include "mlir/Interfaces/LoopLikeInterface.td"
 include "mlir/Interfaces/SideEffects.td"
 
-def LoopOps_Dialect : Dialect {
+def SCF_Dialect : Dialect {
   let name = "loop";
   let cppNamespace = "";
 }
 
-// Base class for Loop dialect ops.
-class Loop_Op<string mnemonic, list<OpTrait> traits = []> :
-    Op<LoopOps_Dialect, mnemonic, traits> {
+// Base class for SCF dialect ops.
+class SCF_Op<string mnemonic, list<OpTrait> traits = []> :
+    Op<SCF_Dialect, mnemonic, traits> {
   // For every standard op, there needs to be a:
   //   * void print(OpAsmPrinter &p, ${C++ class of Op} op)
   //   * LogicalResult verify(${C++ class of Op} op)
   //   * ParseResult parse${C++ class of Op}(OpAsmParser &parser,
   //                                         OperationState &result)
   // functions.
   let printer = [{ return ::print(p, *this); }];
   let verifier = [{ return ::verify(*this); }];
   let parser = [{ return ::parse$cppClass(parser, result); }];
 }
 
-def ForOp : Loop_Op<"for",
+def ForOp : SCF_Op<"for",
       [DeclareOpInterfaceMethods<LoopLikeOpInterface>,
        DeclareOpInterfaceMethods<RegionBranchOpInterface>,
        SingleBlockImplicitTerminator<"YieldOp">,
        RecursiveSideEffects]> {
   let summary = "for operation";
   let description = [{
     The "loop.for" operation represents a loop taking 3 SSA value as operands
     that represent the lower bound, upper bound and step respectively.  The
let extraClassDeclaration = [{
     /// Return operands used when entering the region at 'index'. These operands
     /// correspond to the loop iterator operands, i.e., those exclusing the
     /// induction variable. LoopOp only has one region, so 0 is the only valid
     /// value for `index`.
     OperandRange getSuccessorEntryOperands(unsigned index);
   }];
 }
 
-def IfOp : Loop_Op<"if",
+def IfOp : SCF_Op<"if",
       [DeclareOpInterfaceMethods<RegionBranchOpInterface>,
        SingleBlockImplicitTerminator<"YieldOp">, RecursiveSideEffects]> {
   let summary = "if-then-else operation";
   let description = [{
     The `loop.if` operation represents an if-then-else construct for
     conditionally executing two regions of code. The operand to an if operation
     is a boolean value. For example:
 
let extraClassDeclaration = [{
     OpBuilder getElseBodyBuilder() {
       Block* body = getBody(1);
       return results().empty() ? OpBuilder::atBlockTerminator(body)
                                : OpBuilder::atBlockEnd(body);
     }
   }];
 }
 
-def ParallelOp : Loop_Op<"parallel",
+def ParallelOp : SCF_Op<"parallel",
     [AttrSizedOperandSegments,
      DeclareOpInterfaceMethods<LoopLikeOpInterface>,
      RecursiveSideEffects,
      SingleBlockImplicitTerminator<"YieldOp">]> {
   let summary = "parallel for operation";
   let description = [{
     The "loop.parallel" operation represents a loop nest taking 4 groups of SSA
     values as operands that represent the lower bounds, upper bounds, steps and
let extraClassDeclaration = [{
     ValueRange getInductionVars() {
       return getBody()->getArguments();
     }
     unsigned getNumLoops() { return step().size(); }
     unsigned getNumReductions() { return initVals().size(); }
   }];
 }
 
-def ReduceOp : Loop_Op<"reduce", [HasParent<"ParallelOp">]> {
+def ReduceOp : SCF_Op<"reduce", [HasParent<"ParallelOp">]> {
   let summary = "reduce operation for parallel for";
   let description = [{
     "loop.reduce" is an operation occurring inside "loop.parallel" operations.
     It consists of one block with two arguments which have the same type as the
     operand of "loop.reduce".
 
     "loop.reduce" is used to model the value for reduction computations of a
     "loop.parallel" operation. It has to appear as an immediate child of a
OpBuilder<"OpBuilder &builder, OperationState &result, "
               "Value operand">
   ];
 
   let arguments = (ins AnyType:$operand);
   let regions = (region SizedRegion<1>:$reductionOperator);
 }
 
 def ReduceReturnOp :
-    Loop_Op<"reduce.return", [HasParent<"ReduceOp">, NoSideEffect,
+    SCF_Op<"reduce.return", [HasParent<"ReduceOp">, NoSideEffect,
                               Terminator]> {
   let summary = "terminator for reduce operation";
   let description = [{
     "loop.reduce.return" is a special terminator operation for the block inside
     "loop.reduce". It terminates the region. It should have the same type as
     the operand of "loop.reduce". Example for the custom format:
 
     ```mlir
     loop.reduce.return %res : f32
     ```
   }];
 
   let arguments = (ins AnyType:$result);
   let assemblyFormat = "$result attr-dict `:` type($result)";
 }
 
-def YieldOp : Loop_Op<"yield", [NoSideEffect, ReturnLike, Terminator]> {
+def YieldOp : SCF_Op<"yield", [NoSideEffect, ReturnLike, Terminator]> {
   let summary = "loop yield and termination operation";
   let description = [{
     "loop.yield" yields an SSA value from a loop dialect op region and
     terminates the regions. The semantics of how the values are yielded is
     defined by the parent operation.
     If "loop.yield" has any operands, the operands must match the parent
     operation's results.
     If the parent operation defines no values, then the "loop.yield" may be
     left out in the custom syntax and the builders will insert one implicitly.
     Otherwise, it has to be present in the syntax to indicate which values are
     yielded.
   }];
 
   let arguments = (ins Variadic<AnyType>:$results);
   let builders = [
     OpBuilder<"OpBuilder &builder, OperationState &result",
               [{ /* nothing to do */ }]>
   ];
 }
-#endif // LOOP_OPS
+#endif // MLIR_DIALECT_SCF_SCFOPS

mlir/include/mlir/Dialect/SCF/Transforms.h

-//===- Transforms.h - Pass Entrypoints --------------------------*- C++ -*-===//
+//===- Transforms.h - SCF dialect transformation utilities ------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
-// This header file defines transformations on loop operations.
+// This header file defines transformations on SCF operations.
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef MLIR_DIALECT_LOOPOPS_TRANSFORMS_H_
-#define MLIR_DIALECT_LOOPOPS_TRANSFORMS_H_
+#ifndef MLIR_DIALECT_SCF_TRANSFORMS_H_
+#define MLIR_DIALECT_SCF_TRANSFORMS_H_
 
 #include "llvm/ADT/ArrayRef.h"
 
 namespace mlir {
 
 class Region;
 
-namespace loop {
+namespace scf {
 
 class ParallelOp;
 
 /// Fuses all adjacent loop.parallel operations with identical bounds and step
 /// into one loop.parallel operations. Uses a naive aliasing and dependency
 /// analysis.
 void naivelyFuseParallelOps(Region &region);
 
 /// Tile a parallel loop of the form
 ///   loop.parallel (%i0, %i1) = (%arg0, %arg1) to (%arg2, %arg3)
 ///                                             step (%arg4, %arg5)
 ///
 /// into
 ///   loop.parallel (%i0, %i1) = (%arg0, %arg1) to (%arg2, %arg3)
 ///                                             step (%arg4*tileSize[0],
 ///                                                   %arg5*tileSize[1])
 ///     loop.parallel (%j0, %j1) = (0, 0) to (min(tileSize[0], %arg2-%j0)
 ///                                           min(tileSize[1], %arg3-%j1))
 ///                                        step (%arg4, %arg5)
 /// The old loop is replaced with the new one.
 void tileParallelLoop(ParallelOp op, llvm::ArrayRef<int64_t> tileSizes);
 
-} // namespace loop
+} // namespace scf
 } // namespace mlir
 
-#endif // MLIR_DIALECT_LOOPOPS_TRANSFORMS_H_
+#endif // MLIR_DIALECT_SCF_TRANSFORMS_H_

mlir/include/mlir/InitAllDialects.h

 #include "mlir/Dialect/AVX512/AVX512Dialect.h"
 #include "mlir/Dialect/Affine/IR/AffineOps.h"
 #include "mlir/Dialect/GPU/GPUDialect.h"
 #include "mlir/Dialect/LLVMIR/LLVMAVX512Dialect.h"
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
 #include "mlir/Dialect/LLVMIR/NVVMDialect.h"
 #include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
 #include "mlir/Dialect/Linalg/IR/LinalgOps.h"
-#include "mlir/Dialect/LoopOps/LoopOps.h"
 #include "mlir/Dialect/OpenMP/OpenMPDialect.h"
 #include "mlir/Dialect/Quant/QuantOps.h"
+#include "mlir/Dialect/SCF/SCF.h"
 #include "mlir/Dialect/SDBM/SDBMDialect.h"
 #include "mlir/Dialect/SPIRV/SPIRVDialect.h"
 #include "mlir/Dialect/Shape/IR/Shape.h"
 #include "mlir/Dialect/StandardOps/IR/Ops.h"
 #include "mlir/Dialect/Vector/VectorOps.h"
 #include "mlir/IR/Dialect.h"
 
 namespace mlir {
 
 // This function should be called before creating any MLIRContext if one expect
 // all the possible dialects to be made available to the context automatically.
 inline void registerAllDialects() {
   static bool init_once = []() {
     registerDialect<AffineDialect>();
     registerDialect<avx512::AVX512Dialect>();
     registerDialect<gpu::GPUDialect>();
     registerDialect<LLVM::LLVMAVX512Dialect>();
     registerDialect<LLVM::LLVMDialect>();
     registerDialect<linalg::LinalgDialect>();
-    registerDialect<loop::LoopOpsDialect>();
+    registerDialect<scf::SCFDialect>();
     registerDialect<omp::OpenMPDialect>();
     registerDialect<quant::QuantizationDialect>();
     registerDialect<spirv::SPIRVDialect>();
     registerDialect<StandardOpsDialect>();
     registerDialect<vector::VectorDialect>();
     registerDialect<NVVM::NVVMDialect>();
     registerDialect<ROCDL::ROCDLDialect>();
     registerDialect<SDBMDialect>();
     registerDialect<shape::ShapeDialect>();
     return true;
   }();
   (void)init_once;
 }
 } // namespace mlir
 
 #endif // MLIR_INITALLDIALECTS_H_

mlir/include/mlir/InitAllPasses.h

 #include "mlir/Conversion/LoopsToGPU/LoopsToGPUPass.h"
 #include "mlir/Conversion/StandardToLLVM/ConvertStandardToLLVMPass.h"
 #include "mlir/Conversion/StandardToSPIRV/ConvertStandardToSPIRVPass.h"
 #include "mlir/Conversion/VectorToLLVM/ConvertVectorToLLVM.h"
 #include "mlir/Dialect/Affine/Passes.h"
 #include "mlir/Dialect/GPU/Passes.h"
 #include "mlir/Dialect/LLVMIR/Transforms/LegalizeForExport.h"
 #include "mlir/Dialect/Linalg/Passes.h"
-#include "mlir/Dialect/LoopOps/Passes.h"
 #include "mlir/Dialect/Quant/Passes.h"
+#include "mlir/Dialect/SCF/Passes.h"
 #include "mlir/Dialect/SPIRV/Passes.h"
 #include "mlir/Transforms/LocationSnapshot.h"
 #include "mlir/Transforms/Passes.h"
 #include "mlir/Transforms/ViewOpGraph.h"
 #include "mlir/Transforms/ViewRegionGraph.h"
 
 #include <cstdlib>
 

 #include "mlir/Dialect/Linalg/Passes.h.inc"
 
   // LLVM
 #define GEN_PASS_REGISTRATION
 #include "mlir/Dialect/LLVMIR/Transforms/Passes.h.inc"
 
   // Loop
 #define GEN_PASS_REGISTRATION
-#include "mlir/Dialect/LoopOps/Passes.h.inc"
+#include "mlir/Dialect/SCF/Passes.h.inc"
 
   // Quant
 #define GEN_PASS_REGISTRATION
 #include "mlir/Dialect/Quant/Passes.h.inc"
 
   // SPIR-V
 #define GEN_PASS_REGISTRATION
 #include "mlir/Dialect/SPIRV/Passes.h.inc"
 }
 
 } // namespace mlir
 
 #endif // MLIR_INITALLPASSES_H_

mlir/include/mlir/Transforms/LoopUtils.h

 namespace mlir {
 class AffineForOp;
 class FuncOp;
 class OpBuilder;
 class Value;
 class ValueRange;
 struct MemRefRegion;
 
-namespace loop {
+namespace scf {
 class ForOp;
 class ParallelOp;
-} // end namespace loop
+} // end namespace scf
 
 /// Unrolls this for operation completely if the trip count is known to be
 /// constant. Returns failure otherwise.
 LogicalResult loopUnrollFull(AffineForOp forOp);
 
 /// Unrolls this for operation by the specified unroll factor. Returns failure
 /// if the loop cannot be unrolled either due to restrictions or due to invalid
 /// unroll factors.
 LogicalResult loopUnrollByFactor(AffineForOp forOp, uint64_t unrollFactor);
 
 /// Unrolls this loop by the specified unroll factor or its trip count,
 /// whichever is lower.
 LogicalResult loopUnrollUpToFactor(AffineForOp forOp, uint64_t unrollFactor);
 
 /// Returns true if `loops` is a perfectly nested loop nest, where loops appear
 /// in it from outermost to innermost.
 bool LLVM_ATTRIBUTE_UNUSED isPerfectlyNested(ArrayRef<AffineForOp> loops);
 
 /// Get perfectly nested sequence of loops starting at root of loop nest
 /// (the first op being another AffineFor, and the second op - a terminator).
 /// A loop is perfectly nested iff: the first op in the loop's body is another
 /// AffineForOp, and the second op is a terminator).
 void getPerfectlyNestedLoops(SmallVectorImpl<AffineForOp> &nestedLoops,
                              AffineForOp root);
-void getPerfectlyNestedLoops(SmallVectorImpl<loop::ForOp> &nestedLoops,
-                             loop::ForOp root);
+void getPerfectlyNestedLoops(SmallVectorImpl<scf::ForOp> &nestedLoops,
+                             scf::ForOp root);
 
 /// Unrolls and jams this loop by the specified factor. Returns success if the
 /// loop is successfully unroll-jammed.
 LogicalResult loopUnrollJamByFactor(AffineForOp forOp,
                                     uint64_t unrollJamFactor);
 
 /// Unrolls and jams this loop by the specified factor or by the trip count (if
 /// constant), whichever is lower.

 // Returns AffineForOp of the root of the new loop nest after loop interchanges.
 AffineForOp sinkSequentialLoops(AffineForOp forOp);
 
 /// Performs tiling fo imperfectly nested loops (with interchange) by
 /// strip-mining the `forOps` by `sizes` and sinking them, in their order of
 /// occurrence in `forOps`, under each of the `targets`.
 /// Returns the new AffineForOps, one per each of (`forOps`, `targets`) pair,
 /// nested immediately under each of `targets`.
-using Loops = SmallVector<loop::ForOp, 8>;
+using Loops = SmallVector<scf::ForOp, 8>;
 using TileLoops = std::pair<Loops, Loops>;
 SmallVector<SmallVector<AffineForOp, 8>, 8> tile(ArrayRef<AffineForOp> forOps,
                                                  ArrayRef<uint64_t> sizes,
                                                  ArrayRef<AffineForOp> targets);
-SmallVector<Loops, 8> tile(ArrayRef<loop::ForOp> forOps, ArrayRef<Value> sizes,
-                           ArrayRef<loop::ForOp> targets);
+SmallVector<Loops, 8> tile(ArrayRef<scf::ForOp> forOps, ArrayRef<Value> sizes,
+                           ArrayRef<scf::ForOp> targets);
 
 /// Performs tiling (with interchange) by strip-mining the `forOps` by `sizes`
 /// and sinking them, in their order of occurrence in `forOps`, under `target`.
 /// Returns the new AffineForOps, one per `forOps`, nested immediately under
 /// `target`.
 SmallVector<AffineForOp, 8> tile(ArrayRef<AffineForOp> forOps,
                                  ArrayRef<uint64_t> sizes, AffineForOp target);
-Loops tile(ArrayRef<loop::ForOp> forOps, ArrayRef<Value> sizes,
-           loop::ForOp target);
+Loops tile(ArrayRef<scf::ForOp> forOps, ArrayRef<Value> sizes,
+           scf::ForOp target);
 
-/// Tile a nest of loop::ForOp loops rooted at `rootForOp` with the given
+/// Tile a nest of scf::ForOp loops rooted at `rootForOp` with the given
 /// (parametric) sizes. Sizes are expected to be strictly positive values at
 /// runtime.  If more sizes than loops are provided, discard the trailing values
 /// in sizes.  Assumes the loop nest is permutable.
 /// Returns the newly created intra-tile loops.
-Loops tilePerfectlyNested(loop::ForOp rootForOp, ArrayRef<Value> sizes);
+Loops tilePerfectlyNested(scf::ForOp rootForOp, ArrayRef<Value> sizes);
 
 /// Explicit copy / DMA generation options for mlir::affineDataCopyGenerate.
 struct AffineCopyOptions {
   // True if DMAs should be generated instead of point-wise copies.
   bool generateDma;
   // The slower memory space from which data is to be moved.
   unsigned slowMemorySpace;
   // Memory space of the faster one (typically a scratchpad).

 LogicalResult generateCopyForMemRegion(const MemRefRegion &memrefRegion,
                                        Operation *analyzedOp,
                                        const AffineCopyOptions &copyOptions,
                                        CopyGenerateResult &result);
 
 /// Tile a nest of standard for loops rooted at `rootForOp` by finding such
 /// parametric tile sizes that the outer loops have a fixed number of iterations
 /// as defined in `sizes`.
-TileLoops extractFixedOuterLoops(loop::ForOp rootFOrOp,
-                                 ArrayRef<int64_t> sizes);
+TileLoops extractFixedOuterLoops(scf::ForOp rootFOrOp, ArrayRef<int64_t> sizes);
 
 /// Replace a perfect nest of "for" loops with a single linearized loop. Assumes
 /// `loops` contains a list of perfectly nested loops with bounds and steps
 /// independent of any loop induction variable involved in the nest.
-void coalesceLoops(MutableArrayRef<loop::ForOp> loops);
+void coalesceLoops(MutableArrayRef<scf::ForOp> loops);
 
 /// Take the ParallelLoop and for each set of dimension indices, combine them
 /// into a single dimension. combinedDimensions must contain each index into
 /// loops exactly once.
-void collapseParallelLoops(loop::ParallelOp loops,
+void collapseParallelLoops(scf::ParallelOp loops,
                            ArrayRef<std::vector<unsigned>> combinedDimensions);
 
 /// Maps `forOp` for execution on a parallel grid of virtual `processorIds` of
 /// size given by `numProcessors`. This is achieved by embedding the SSA values
 /// corresponding to `processorIds` and `numProcessors` into the bounds and step
 /// of the `forOp`. No check is performed on the legality of the rewrite, it is
 /// the caller's responsibility to ensure legality.
 ///

 /// is rewritten into a version resembling the following pseudo-IR:
 ///
 /// ```
 ///    loop.for %i = %lb + %step * (threadIdx.x + blockIdx.x * blockDim.x)
 ///       to %ub step %gridDim.x * blockDim.x * %step {
 ///      ...
 ///    }
 /// ```
-void mapLoopToProcessorIds(loop::ForOp forOp, ArrayRef<Value> processorId,
+void mapLoopToProcessorIds(scf::ForOp forOp, ArrayRef<Value> processorId,
                            ArrayRef<Value> numProcessors);
 
 /// Gathers all AffineForOps in 'func' grouped by loop depth.
 void gatherLoops(FuncOp func,
                  std::vector<SmallVector<AffineForOp, 2>> &depthToLoops);
 
 /// Creates an AffineForOp while ensuring that the lower and upper bounds are
 /// canonicalized, i.e., unused and duplicate operands are removed, any constant

mlir/lib/Analysis/CMakeLists.txt

add_mlir_library(MLIRAnalysis
   )
 
 target_link_libraries(MLIRAnalysis
   PUBLIC
   MLIRAffineOps
   MLIRCallInterfaces
   MLIRControlFlowInterfaces
   MLIRInferTypeOpInterface
-  MLIRLoopOps
+  MLIRSCF
   )
 
 add_mlir_library(MLIRLoopAnalysis
   AffineAnalysis.cpp
   AffineStructures.cpp
   LoopAnalysis.cpp
   NestedMatcher.cpp
   Utils.cpp
 
   ADDITIONAL_HEADER_DIRS
   ${MLIR_MAIN_INCLUDE_DIR}/mlir/Analysis
 
   DEPENDS
   mlir-generic-headers
   )
 
 target_link_libraries(MLIRLoopAnalysis
   PUBLIC
   MLIRAffineOps
   MLIRCallInterfaces
   MLIRControlFlowInterfaces
   MLIRInferTypeOpInterface
-  MLIRLoopOps)
+  MLIRSCF)

mlir/lib/Analysis/SliceAnalysis.cpp

 //===- UseDefAnalysis.cpp - Analysis for Transitive UseDef chains ---------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements Analysis functions specific to slicing in Function.
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Analysis/SliceAnalysis.h"
 #include "mlir/Dialect/Affine/IR/AffineOps.h"
-#include "mlir/Dialect/LoopOps/LoopOps.h"
+#include "mlir/Dialect/SCF/SCF.h"
 #include "mlir/IR/Function.h"
 #include "mlir/IR/Operation.h"
 #include "mlir/Support/LLVM.h"
 #include "llvm/ADT/SetVector.h"
 
 ///
 /// Implements Analysis functions specific to slicing in Function.
 ///
static void getForwardSliceImpl(Operation *op,
   if (!filter(op)) {
     return;
   }
 
   if (auto forOp = dyn_cast<AffineForOp>(op)) {
     for (auto *ownerInst : forOp.getInductionVar().getUsers())
       if (forwardSlice->count(ownerInst) == 0)
         getForwardSliceImpl(ownerInst, forwardSlice, filter);
-  } else if (auto forOp = dyn_cast<loop::ForOp>(op)) {
+  } else if (auto forOp = dyn_cast<scf::ForOp>(op)) {
     for (auto *ownerInst : forOp.getInductionVar().getUsers())
       if (forwardSlice->count(ownerInst) == 0)
         getForwardSliceImpl(ownerInst, forwardSlice, filter);
   } else {
     assert(op->getNumRegions() == 0 && "unexpected generic op with regions");
     assert(op->getNumResults() <= 1 && "unexpected multiple results");
     if (op->getNumResults() > 0) {
       for (auto *ownerInst : op->getResult(0).getUsers())

 
 static void getBackwardSliceImpl(Operation *op,
                                  SetVector<Operation *> *backwardSlice,
                                  TransitiveFilter filter) {
   if (!op)
     return;
 
   assert((op->getNumRegions() == 0 || isa<AffineForOp>(op) ||
-          isa<loop::ForOp>(op)) &&
+          isa<scf::ForOp>(op)) &&
          "unexpected generic op with regions");
 
   // Evaluate whether we should keep this def.
   // This is useful in particular to implement scoping; i.e. return the
   // transitive forwardSlice in the current scope.
   if (!filter(op)) {
     return;
   }
 
   for (auto en : llvm::enumerate(op->getOperands())) {
     auto operand = en.value();
     if (auto blockArg = operand.dyn_cast<BlockArgument>()) {
       if (auto affIv = getForInductionVarOwner(operand)) {
         auto *affOp = affIv.getOperation();
         if (backwardSlice->count(affOp) == 0)
           getBackwardSliceImpl(affOp, backwardSlice, filter);
-      } else if (auto loopIv = loop::getForInductionVarOwner(operand)) {
+      } else if (auto loopIv = scf::getForInductionVarOwner(operand)) {
         auto *loopOp = loopIv.getOperation();
         if (backwardSlice->count(loopOp) == 0)
           getBackwardSliceImpl(loopOp, backwardSlice, filter);
       } else if (blockArg.getOwner() !=
                  &op->getParentOfType<FuncOp>().getBody().front()) {
         op->emitError("unsupported CF for operand ") << en.index();
         llvm_unreachable("Unsupported control flow");
       }

mlir/lib/Conversion/AffineToStandard/AffineToStandard.cpp

 // operations) within a function into their standard If and For equivalent ops.
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Conversion/AffineToStandard/AffineToStandard.h"
 
 #include "../PassDetail.h"
 #include "mlir/Dialect/Affine/IR/AffineOps.h"
-#include "mlir/Dialect/LoopOps/LoopOps.h"
+#include "mlir/Dialect/SCF/SCF.h"
 #include "mlir/Dialect/StandardOps/IR/Ops.h"
 #include "mlir/IR/AffineExprVisitor.h"
 #include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/IntegerSet.h"
 #include "mlir/IR/MLIRContext.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Transforms/DialectConversion.h"

 
 /// Affine terminators are removed.
 class AffineTerminatorLowering : public OpRewritePattern<AffineTerminatorOp> {
 public:
   using OpRewritePattern<AffineTerminatorOp>::OpRewritePattern;
 
   LogicalResult matchAndRewrite(AffineTerminatorOp op,
                                 PatternRewriter &rewriter) const override {
-    rewriter.replaceOpWithNewOp<loop::YieldOp>(op);
+    rewriter.replaceOpWithNewOp<scf::YieldOp>(op);
     return success();
   }
 };
 
 class AffineForLowering : public OpRewritePattern<AffineForOp> {
 public:
   using OpRewritePattern<AffineForOp>::OpRewritePattern;
 
   LogicalResult matchAndRewrite(AffineForOp op,
                                 PatternRewriter &rewriter) const override {
     Location loc = op.getLoc();
     Value lowerBound = lowerAffineLowerBound(op, rewriter);
     Value upperBound = lowerAffineUpperBound(op, rewriter);
     Value step = rewriter.create<ConstantIndexOp>(loc, op.getStep());
-    auto f = rewriter.create<loop::ForOp>(loc, lowerBound, upperBound, step);
+    auto f = rewriter.create<scf::ForOp>(loc, lowerBound, upperBound, step);
     f.region().getBlocks().clear();
     rewriter.inlineRegionBefore(op.region(), f.region(), f.region().end());
     rewriter.eraseOp(op);
     return success();
   }
 };
 
 class AffineIfLowering : public OpRewritePattern<AffineIfOp> {
for (unsigned i = 0, e = integerSet.getNumConstraints(); i < e; ++i) {
           rewriter.create<CmpIOp>(loc, pred, affResult, zeroConstant);
       cond =
           cond ? rewriter.create<AndOp>(loc, cond, cmpVal).getResult() : cmpVal;
     }
     cond = cond ? cond
                 : rewriter.create<ConstantIntOp>(loc, /*value=*/1, /*width=*/1);
 
     bool hasElseRegion = !op.elseRegion().empty();
-    auto ifOp = rewriter.create<loop::IfOp>(loc, cond, hasElseRegion);
+    auto ifOp = rewriter.create<scf::IfOp>(loc, cond, hasElseRegion);
     rewriter.inlineRegionBefore(op.thenRegion(), &ifOp.thenRegion().back());
     ifOp.thenRegion().back().erase();
     if (hasElseRegion) {
       rewriter.inlineRegionBefore(op.elseRegion(), &ifOp.elseRegion().back());
       ifOp.elseRegion().back().erase();
     }
 
     // Ok, we're done!

 }
 
 namespace {
 class LowerAffinePass : public ConvertAffineToStandardBase<LowerAffinePass> {
   void runOnFunction() override {
     OwningRewritePatternList patterns;
     populateAffineToStdConversionPatterns(patterns, &getContext());
     ConversionTarget target(getContext());
-    target.addLegalDialect<loop::LoopOpsDialect, StandardOpsDialect>();
+    target.addLegalDialect<scf::SCFDialect, StandardOpsDialect>();
     if (failed(applyPartialConversion(getFunction(), target, patterns)))
       signalPassFailure();
   }
 };
 } // namespace
 
 /// Lowers If and For operations within a function into their lower level CFG
 /// equivalent blocks.
 std::unique_ptr<OperationPass<FuncOp>> mlir::createLowerAffinePass() {
   return std::make_unique<LowerAffinePass>();
 }

mlir/lib/Conversion/AffineToStandard/CMakeLists.txt

 add_mlir_conversion_library(MLIRAffineToStandard
   AffineToStandard.cpp
 
   ADDITIONAL_HEADER_DIRS
   ${MLIR_MAIN_INCLUDE_DIR}/mlir/Conversion/AffineToStandard
 
   DEPENDS
   MLIRConversionPassIncGen
 )
 target_link_libraries(
   MLIRAffineToStandard
   PUBLIC
   MLIRAffineOps
-  MLIRLoopOps
+  MLIRSCF
   MLIRPass
   MLIRStandardOps
   MLIRTransforms
   MLIRIR
   LLVMCore
   LLVMSupport
   )

mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.cpp

 //===- ConvertGPUToSPIRV.cpp - Convert GPU ops to SPIR-V dialect ----------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements the conversion patterns from GPU ops to SPIR-V dialect.
 //
 //===----------------------------------------------------------------------===//
 #include "mlir/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.h"
 #include "mlir/Dialect/GPU/GPUDialect.h"
-#include "mlir/Dialect/LoopOps/LoopOps.h"
+#include "mlir/Dialect/SCF/SCF.h"
 #include "mlir/Dialect/SPIRV/SPIRVDialect.h"
 #include "mlir/Dialect/SPIRV/SPIRVLowering.h"
 #include "mlir/Dialect/SPIRV/SPIRVOps.h"
 #include "mlir/IR/Module.h"
 
 using namespace mlir;
 
 namespace {
 
-/// Pattern to convert a loop::ForOp within kernel functions into spirv::LoopOp.
-class ForOpConversion final : public SPIRVOpLowering<loop::ForOp> {
+/// Pattern to convert a scf::ForOp within kernel functions into spirv::LoopOp.
+class ForOpConversion final : public SPIRVOpLowering<scf::ForOp> {
 public:
-  using SPIRVOpLowering<loop::ForOp>::SPIRVOpLowering;
+  using SPIRVOpLowering<scf::ForOp>::SPIRVOpLowering;
 
   LogicalResult
-  matchAndRewrite(loop::ForOp forOp, ArrayRef<Value> operands,
+  matchAndRewrite(scf::ForOp forOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
-/// Pattern to convert a loop::IfOp within kernel functions into
+/// Pattern to convert a scf::IfOp within kernel functions into
 /// spirv::SelectionOp.
-class IfOpConversion final : public SPIRVOpLowering<loop::IfOp> {
+class IfOpConversion final : public SPIRVOpLowering<scf::IfOp> {
 public:
-  using SPIRVOpLowering<loop::IfOp>::SPIRVOpLowering;
+  using SPIRVOpLowering<scf::IfOp>::SPIRVOpLowering;
 
   LogicalResult
-  matchAndRewrite(loop::IfOp IfOp, ArrayRef<Value> operands,
+  matchAndRewrite(scf::IfOp IfOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
-/// Pattern to erase a loop::YieldOp.
-class TerminatorOpConversion final : public SPIRVOpLowering<loop::YieldOp> {
+/// Pattern to erase a scf::YieldOp.
+class TerminatorOpConversion final : public SPIRVOpLowering<scf::YieldOp> {
 public:
-  using SPIRVOpLowering<loop::YieldOp>::SPIRVOpLowering;
+  using SPIRVOpLowering<scf::YieldOp>::SPIRVOpLowering;
 
   LogicalResult
-  matchAndRewrite(loop::YieldOp terminatorOp, ArrayRef<Value> operands,
+  matchAndRewrite(scf::YieldOp terminatorOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     rewriter.eraseOp(terminatorOp);
     return success();
   }
 };
 
 /// Pattern lowering GPU block/thread size/id to loading SPIR-V invocation
 /// builtin variables.
public:
   LogicalResult
   matchAndRewrite(gpu::ReturnOp returnOp, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override;
 };
 
 } // namespace
 
 //===----------------------------------------------------------------------===//
-// loop::ForOp.
+// scf::ForOp.
 //===----------------------------------------------------------------------===//
 
 LogicalResult
-ForOpConversion::matchAndRewrite(loop::ForOp forOp, ArrayRef<Value> operands,
+ForOpConversion::matchAndRewrite(scf::ForOp forOp, ArrayRef<Value> operands,
                                  ConversionPatternRewriter &rewriter) const {
-  // loop::ForOp can be lowered to the structured control flow represented by
+  // scf::ForOp can be lowered to the structured control flow represented by
   // spirv::LoopOp by making the continue block of the spirv::LoopOp the loop
   // latch and the merge block the exit block. The resulting spirv::LoopOp has a
   // single back edge from the continue to header block, and a single exit from
   // header to merge.
-  loop::ForOpOperandAdaptor forOperands(operands);
+  scf::ForOpOperandAdaptor forOperands(operands);
   auto loc = forOp.getLoc();
   auto loopControl = rewriter.getI32IntegerAttr(
       static_cast<uint32_t>(spirv::LoopControl::None));
   auto loopOp = rewriter.create<spirv::LoopOp>(loc, loopControl);
   loopOp.addEntryAndMergeBlock();
 
   OpBuilder::InsertionGuard guard(rewriter);
   // Create the block for the header.
Value updatedIndVar = rewriter.create<spirv::IAddOp>(
       loc, newIndVar.getType(), newIndVar, forOperands.step());
   rewriter.create<spirv::BranchOp>(loc, header, updatedIndVar);
 
   rewriter.eraseOp(forOp);
   return success();
 }
 
 //===----------------------------------------------------------------------===//
-// loop::IfOp.
+// scf::IfOp.
 //===----------------------------------------------------------------------===//
 
 LogicalResult
-IfOpConversion::matchAndRewrite(loop::IfOp ifOp, ArrayRef<Value> operands,
+IfOpConversion::matchAndRewrite(scf::IfOp ifOp, ArrayRef<Value> operands,
                                 ConversionPatternRewriter &rewriter) const {
-  // When lowering `loop::IfOp` we explicitly create a selection header block
+  // When lowering `scf::IfOp` we explicitly create a selection header block
   // before the control flow diverges and a merge block where control flow
   // subsequently converges.
-  loop::IfOpOperandAdaptor ifOperands(operands);
+  scf::IfOpOperandAdaptor ifOperands(operands);
   auto loc = ifOp.getLoc();
 
   // Create `spv.selection` operation, selection header block and merge block.
   auto selectionControl = rewriter.getI32IntegerAttr(
       static_cast<uint32_t>(spirv::SelectionControl::None));
   auto selectionOp = rewriter.create<spirv::SelectionOp>(loc, selectionControl);
   selectionOp.addMergeBlock();
   auto *mergeBlock = selectionOp.getMergeBlock();

mlir/lib/Conversion/GPUToSPIRV/ConvertGPUToSPIRVPass.cpp

 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Conversion/GPUToSPIRV/ConvertGPUToSPIRVPass.h"
 #include "../PassDetail.h"
 #include "mlir/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.h"
 #include "mlir/Conversion/StandardToSPIRV/ConvertStandardToSPIRV.h"
 #include "mlir/Dialect/GPU/GPUDialect.h"
-#include "mlir/Dialect/LoopOps/LoopOps.h"
+#include "mlir/Dialect/SCF/SCF.h"
 #include "mlir/Dialect/SPIRV/SPIRVDialect.h"
 #include "mlir/Dialect/SPIRV/SPIRVLowering.h"
 #include "mlir/Dialect/SPIRV/SPIRVOps.h"
 
 using namespace mlir;
 
 namespace {
 /// Pass to lower GPU Dialect to SPIR-V. The pass only converts the gpu.func ops

mlir/lib/Conversion/LoopToStandard/CMakeLists.txt

 add_mlir_conversion_library(MLIRLoopToStandard
   LoopToStandard.cpp
 
   ADDITIONAL_HEADER_DIRS
   ${MLIR_MAIN_INCLUDE_DIR}/mlir/Conversion/LoopToStandard
 
   DEPENDS
   MLIRConversionPassIncGen
 )
 target_link_libraries(
   MLIRLoopToStandard
   PUBLIC
-  MLIRLoopOps
+  MLIRSCF
   MLIRTransforms
   LLVMCore
   LLVMSupport
   )

mlir/lib/Conversion/LoopToStandard/LoopToStandard.cpp

 //===- LoopToStandard.cpp - ControlFlow to CFG conversion -----------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements a pass to convert loop.for, loop.if and loop.terminator
 // ops into standard CFG ops.
 //
 //===----------------------------------------------------------------------===//
 
 #include "../PassDetail.h"
 #include "mlir/Conversion/LoopToStandard/ConvertLoopToStandard.h"
-#include "mlir/Dialect/LoopOps/LoopOps.h"
+#include "mlir/Dialect/SCF/SCF.h"
 #include "mlir/Dialect/StandardOps/IR/Ops.h"
 #include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/MLIRContext.h"
 #include "mlir/IR/Module.h"
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/Transforms/DialectConversion.h"
 #include "mlir/Transforms/Passes.h"
 #include "mlir/Transforms/Utils.h"
 
 using namespace mlir;
-using namespace mlir::loop;
+using namespace mlir::scf;
 
 namespace {
 
 struct LoopToStandardPass
     : public ConvertLoopToStandardBase<LoopToStandardPass> {
   void runOnOperation() override;
 };
 

 //
 struct IfLowering : public OpRewritePattern<IfOp> {
   using OpRewritePattern<IfOp>::OpRewritePattern;
 
   LogicalResult matchAndRewrite(IfOp ifOp,
                                 PatternRewriter &rewriter) const override;
 };
 
-struct ParallelLowering : public OpRewritePattern<mlir::loop::ParallelOp> {
-  using OpRewritePattern<mlir::loop::ParallelOp>::OpRewritePattern;
+struct ParallelLowering : public OpRewritePattern<mlir::scf::ParallelOp> {
+  using OpRewritePattern<mlir::scf::ParallelOp>::OpRewritePattern;
 
-  LogicalResult matchAndRewrite(mlir::loop::ParallelOp parallelOp,
+  LogicalResult matchAndRewrite(mlir::scf::ParallelOp parallelOp,
                                 PatternRewriter &rewriter) const override;
 };
 } // namespace
 
 LogicalResult ForLowering::matchAndRewrite(ForOp forOp,
                                            PatternRewriter &rewriter) const {
   Location loc = forOp.getLoc();
 

mlir/lib/Conversion/LoopsToGPU/LoopsToGPU.cpp

 //===----------------------------------------------------------------------===//
 
 #include "mlir/Conversion/LoopsToGPU/LoopsToGPU.h"
 
 #include "mlir/Conversion/AffineToStandard/AffineToStandard.h"
 #include "mlir/Dialect/Affine/IR/AffineOps.h"
 #include "mlir/Dialect/GPU/GPUDialect.h"
 #include "mlir/Dialect/GPU/ParallelLoopMapper.h"
-#include "mlir/Dialect/LoopOps/LoopOps.h"
+#include "mlir/Dialect/SCF/SCF.h"
 #include "mlir/Dialect/StandardOps/IR/Ops.h"
 #include "mlir/IR/AffineExpr.h"
 #include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Transforms/DialectConversion.h"
 #include "mlir/Transforms/LoopUtils.h"
 #include "mlir/Transforms/Passes.h"
 #include "mlir/Transforms/RegionUtils.h"
 #include "llvm/ADT/Sequence.h"
 #include "llvm/Support/Debug.h"
 
 #define DEBUG_TYPE "loops-to-gpu"
 
 using namespace mlir;
-using namespace mlir::loop;
+using namespace mlir::scf;
 
 using llvm::seq;
 
 // Extract an indexed value from KernelDim3.
 static Value getDim3Value(const gpu::KernelDim3 &dim3, unsigned pos) {
   switch (pos) {
   case 0:
     return dim3.x;

 }
 
 LogicalResult mlir::convertLoopNestToGPULaunch(ForOp forOp,
                                                unsigned numBlockDims,
                                                unsigned numThreadDims) {
   return ::convertLoopNestToGPULaunch(forOp, numBlockDims, numThreadDims);
 }
 
-LogicalResult mlir::convertLoopToGPULaunch(loop::ForOp forOp,
+LogicalResult mlir::convertLoopToGPULaunch(scf::ForOp forOp,
                                            ArrayRef<Value> numWorkGroups,
                                            ArrayRef<Value> workGroupSizes) {
   return ::convertLoopToGPULaunch(forOp, numWorkGroups, workGroupSizes);
 }
 
 namespace {
 struct ParallelToGpuLaunchLowering : public OpRewritePattern<ParallelOp> {
   using OpRewritePattern<ParallelOp>::OpRewritePattern;
if (isMappedToProcessor(processor)) {
           bounds[processor] = launchBound;
         }
         if (!boundIsPrecise) {
           // We are using an approximation, create a surrounding conditional.
           Value originalBound = std::get<3>(config);
           CmpIOp pred = rewriter.create<CmpIOp>(
               loc, CmpIPredicate::slt, newIndex,
               cloningMap.lookupOrDefault(originalBound));
-          loop::IfOp ifOp = rewriter.create<loop::IfOp>(loc, pred, false);
+          scf::IfOp ifOp = rewriter.create<scf::IfOp>(loc, pred, false);
           rewriter.setInsertionPointToStart(&ifOp.thenRegion().front());
           // Put a sentinel into the worklist so we know when to pop out of the
           // if body again. We use the launchOp here, as that cannot be part of
           // the bodies instruction.
           worklist.push_back(launchOp.getOperation());
         }
       }
     } else {
       // Create a sequential for loop.
-      auto loopOp = rewriter.create<loop::ForOp>(
+      auto loopOp = rewriter.create<scf::ForOp>(
           loc, cloningMap.lookupOrDefault(lowerBound),
           cloningMap.lookupOrDefault(upperBound),
           cloningMap.lookupOrDefault(step));
       newIndex = loopOp.getInductionVar();
       rewriter.setInsertionPointToStart(loopOp.getBody());
       // Put a sentinel into the worklist so we know when to pop out of the loop
       // body again. We use the launchOp here, as that cannot be part of the
       // bodies instruction.

mlir/lib/Conversion/LoopsToGPU/LoopsToGPUPass.cpp

 //===- LoopsToGPUPass.cpp - Convert a loop nest to a GPU kernel -----------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Conversion/LoopsToGPU/LoopsToGPUPass.h"
 #include "../PassDetail.h"
 #include "mlir/Conversion/LoopsToGPU/LoopsToGPU.h"
 #include "mlir/Dialect/Affine/IR/AffineOps.h"
 #include "mlir/Dialect/GPU/GPUDialect.h"
-#include "mlir/Dialect/LoopOps/LoopOps.h"
+#include "mlir/Dialect/SCF/SCF.h"
 #include "mlir/Dialect/StandardOps/IR/Ops.h"
 #include "mlir/Transforms/DialectConversion.h"
 
 #include "llvm/ADT/ArrayRef.h"
 #include "llvm/Support/CommandLine.h"
 
 #define PASS_NAME "convert-loops-to-gpu"
 #define LOOPOP_TO_GPU_PASS_NAME "convert-loop-op-to-gpu"
 
 using namespace mlir;
-using namespace mlir::loop;
+using namespace mlir::scf;
 
 namespace {
 // A pass that traverses top-level loops in the function and converts them to
 // GPU launch operations.  Nested launches are not allowed, so this does not
 // walk the function recursively to avoid considering nested loops.
 struct ForLoopMapper : public ConvertSimpleLoopsToGPUBase<ForLoopMapper> {
   ForLoopMapper() = default;
   ForLoopMapper(unsigned numBlockDims, unsigned numThreadDims) {
struct ParallelLoopToGpuPass
     : public ConvertParallelLoopToGpuBase<ParallelLoopToGpuPass> {
   void runOnOperation() override {
     OwningRewritePatternList patterns;
     populateParallelLoopToGPUPatterns(patterns, &getContext());
     ConversionTarget target(getContext());
     target.addLegalDialect<StandardOpsDialect>();
     target.addLegalDialect<AffineDialect>();
     target.addLegalDialect<gpu::GPUDialect>();
-    target.addLegalDialect<loop::LoopOpsDialect>();
-    target.addIllegalOp<loop::ParallelOp>();
+    target.addLegalDialect<scf::SCFDialect>();
+    target.addIllegalOp<scf::ParallelOp>();
     if (failed(applyPartialConversion(getOperation(), target, patterns)))
       signalPassFailure();
   }
 };
 
 } // namespace
 
 std::unique_ptr<OperationPass<FuncOp>>

mlir/lib/Conversion/VectorToLoops/ConvertVectorToLoops.cpp

 //===- VectorToLoops.cpp - Conversion from Vector to mix of Loops and Std -===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements target-dependent lowering of vector transfer operations.
 //
 //===----------------------------------------------------------------------===//
 
 #include <type_traits>
 
 #include "mlir/Conversion/VectorToLoops/ConvertVectorToLoops.h"
 #include "mlir/Dialect/Affine/EDSC/Intrinsics.h"
-#include "mlir/Dialect/LoopOps/EDSC/Builders.h"
-#include "mlir/Dialect/LoopOps/EDSC/Intrinsics.h"
+#include "mlir/Dialect/SCF/EDSC/Builders.h"
+#include "mlir/Dialect/SCF/EDSC/Intrinsics.h"
 #include "mlir/Dialect/StandardOps/EDSC/Intrinsics.h"
 #include "mlir/Dialect/Vector/EDSC/Intrinsics.h"
 #include "mlir/Dialect/Vector/VectorOps.h"
 #include "mlir/IR/AffineExpr.h"
 #include "mlir/IR/AffineMap.h"
 #include "mlir/IR/Attributes.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/Location.h"
void NDTransferOpHelper<ConcreteOp>::emitInBounds(
   for (auto it : llvm::zip(majorIvs, majorOffsets, memrefBounds.getUbs())) {
     Value iv = std::get<0>(it), off = std::get<1>(it), ub = std::get<2>(it);
     using namespace mlir::edsc::op;
     majorIvsPlusOffsets.push_back(iv + off);
     Value inBounds2 = majorIvsPlusOffsets.back() < ub;
     inBounds = inBounds && inBounds2;
   }
 
-  auto ifOp = ScopedContext::getBuilderRef().create<loop::IfOp>(
+  auto ifOp = ScopedContext::getBuilderRef().create<scf::IfOp>(
       ScopedContext::getLocation(), TypeRange{}, inBounds,
       /*withElseRegion=*/std::is_same<ConcreteOp, TransferReadOp>());
   BlockBuilder(&ifOp.thenRegion().front(),
                Append())([&] { thenBlockBuilder(majorIvsPlusOffsets); });
   if (std::is_same<ConcreteOp, TransferReadOp>())
     BlockBuilder(&ifOp.elseRegion().front(),
                  Append())([&] { elseBlockBuilder(majorIvsPlusOffsets); });
 }

mlir/lib/Dialect/CMakeLists.txt

 add_subdirectory(Affine)
 add_subdirectory(AVX512)
 add_subdirectory(GPU)
 add_subdirectory(Linalg)
 add_subdirectory(LLVMIR)
-add_subdirectory(LoopOps)
 add_subdirectory(OpenMP)
 add_subdirectory(Quant)
+add_subdirectory(SCF)
 add_subdirectory(SDBM)
 add_subdirectory(Shape)
 add_subdirectory(SPIRV)
 add_subdirectory(StandardOps)
 add_subdirectory(Vector)
 
 set(LLVM_OPTIONAL_SOURCES
   Traits.cpp

mlir/lib/Dialect/GPU/CMakeLists.txt

add_mlir_dialect_library(MLIRGPU
   MLIRParallelLoopMapperAttrGen
   MLIRParallelLoopMapperEnumsGen
   )
 target_link_libraries(MLIRGPU
   PUBLIC
   MLIREDSC
   MLIRIR
   MLIRLLVMIR
-  MLIRLoopOps
+  MLIRSCF
   MLIRPass
   MLIRSideEffects
   MLIRStandardOps
   MLIRSupport
   MLIRTransformUtils
   LLVMSupport
   )

mlir/lib/Dialect/GPU/Transforms/MemoryPromotion.cpp

 //===- MemoryPromotion.cpp - Utilities for moving data across GPU memories ===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements utilities that allow one to create IR moving the data
 // across different levels of the GPU memory hierarchy.
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/GPU/MemoryPromotion.h"
 #include "mlir/Dialect/GPU/GPUDialect.h"
-#include "mlir/Dialect/LoopOps/EDSC/Builders.h"
+#include "mlir/Dialect/SCF/EDSC/Builders.h"
 #include "mlir/Dialect/StandardOps/EDSC/Intrinsics.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Transforms/LoopUtils.h"
 
 using namespace mlir;
 using namespace mlir::edsc;
 using namespace mlir::edsc::intrinsics;
 using namespace mlir::gpu;
LoopNestBuilder(ivs, lbs, ubs, steps)([&]() {
     toHandle(activeIvs) = fromHandle(activeIvs);
   });
 
   // Map the innermost loops to threads in reverse order.
   for (auto en :
        llvm::enumerate(llvm::reverse(llvm::makeArrayRef(ivs).take_back(
            GPUDialect::getNumWorkgroupDimensions())))) {
     Value v = en.value();
-    auto loop = cast<loop::ForOp>(v.getParentRegion()->getParentOp());
+    auto loop = cast<scf::ForOp>(v.getParentRegion()->getParentOp());
     mapLoopToProcessorIds(loop, {threadIds[en.index()]},
                           {blockDims[en.index()]});
   }
 }
 
 /// Emits the loop nests performing the copy to the designated location in the
 /// beginning of the region, and from the designated location immediately before
 /// the terminator of the first block of the region. The region is expected to

mlir/lib/Dialect/GPU/Transforms/ParallelLoopMapper.cpp

 // GPU devices.
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/GPU/ParallelLoopMapper.h"
 
 #include "mlir/Dialect/GPU/GPUDialect.h"
 #include "mlir/Dialect/GPU/Passes.h"
-#include "mlir/Dialect/LoopOps/LoopOps.h"
+#include "mlir/Dialect/SCF/SCF.h"
 #include "mlir/IR/AffineMap.h"
 #include "mlir/Pass/Pass.h"
 
 using namespace mlir;
 using namespace mlir::gpu;
-using namespace mlir::loop;
+using namespace mlir::scf;
 
 #include "mlir/Dialect/GPU/ParallelLoopMapperEnums.cpp.inc"
 namespace mlir {
 
 #include "mlir/Dialect/GPU/ParallelLoopMapperAttr.cpp.inc"
 namespace gpu {
 
 StringRef getMappingAttrName() { return "mapping"; }
 
 ParallelLoopDimMapping getParallelLoopDimMappingAttr(Processor processor,
                                                      AffineMap map,
                                                      AffineMap bound) {
   MLIRContext *context = map.getContext();
   OpBuilder builder(context);
   return ParallelLoopDimMapping::get(
       builder.getI64IntegerAttr(static_cast<int32_t>(processor)),
       AffineMapAttr::get(map), AffineMapAttr::get(bound), context);
 }
 
-LogicalResult setMappingAttr(loop::ParallelOp ploopOp,
+LogicalResult setMappingAttr(scf::ParallelOp ploopOp,
                              ArrayRef<ParallelLoopDimMapping> mapping) {
   // Verify that each processor is mapped to only once.
   llvm::DenseSet<gpu::Processor> specifiedMappings;
   for (auto dimAttr : mapping) {
     gpu::Processor processor = getProcessor(dimAttr);
     if (processor != gpu::Processor::Sequential &&
         specifiedMappings.count(processor))
       return ploopOp.emitError(

mlir/lib/Dialect/Linalg/EDSC/Builders.cpp

 //===- Builders.cpp - MLIR Declarative Linalg Builders --------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/IR/Builders.h"
 #include "mlir/Dialect/Affine/EDSC/Intrinsics.h"
 #include "mlir/Dialect/Linalg/EDSC/Builders.h"
 #include "mlir/Dialect/Linalg/EDSC/Intrinsics.h"
-#include "mlir/Dialect/LoopOps/EDSC/Builders.h"
+#include "mlir/Dialect/SCF/EDSC/Builders.h"
 #include "mlir/Dialect/StandardOps/EDSC/Intrinsics.h"
 #include "mlir/Dialect/Utils/StructuredOpsUtils.h"
 #include "mlir/IR/AffineExpr.h"
 
 using namespace mlir;
 using namespace mlir::edsc;
 using namespace mlir::edsc::intrinsics;
 using namespace mlir::linalg;
-using namespace mlir::loop;
+using namespace mlir::scf;
 
 mlir::edsc::LoopRangeBuilder::LoopRangeBuilder(Value *iv, Value range) {
   assert(range.getType() && "expected !linalg.range type");
   assert(range.getDefiningOp() && "need operations to extract range parts");
   auto rangeOp = cast<RangeOp>(range.getDefiningOp());
   auto lb = rangeOp.min();
   auto ub = rangeOp.max();
   auto step = rangeOp.step();
Value LoopNestRangeBuilder::LoopNestRangeBuilder::operator()(
   }
   return Value();
 }
 
 namespace mlir {
 namespace edsc {
 
 template <>
-GenericLoopNestRangeBuilder<loop::ForOp>::GenericLoopNestRangeBuilder(
+GenericLoopNestRangeBuilder<scf::ForOp>::GenericLoopNestRangeBuilder(
     MutableArrayRef<Value> ivs, ArrayRef<Value> ranges) {
   builder = std::make_unique<LoopNestRangeBuilder>(ivs, ranges);
 }
 
 template <>
 GenericLoopNestRangeBuilder<AffineForOp>::GenericLoopNestRangeBuilder(
     MutableArrayRef<Value> ivs, ArrayRef<Value> ranges) {
   SmallVector<Value, 4> lbs;
   SmallVector<Value, 4> ubs;
   SmallVector<int64_t, 4> steps;
   for (Value range : ranges) {
     assert(range.getType() && "expected linalg.range type");
     assert(range.getDefiningOp() && "need operations to extract range parts");
     RangeOp rangeOp = cast<RangeOp>(range.getDefiningOp());
     lbs.emplace_back(rangeOp.min());
     ubs.emplace_back(rangeOp.max());
     steps.emplace_back(rangeOp.step());
   }
   builder = std::make_unique<AffineLoopNestBuilder>(ivs, lbs, ubs, steps);
 }
 
 template <>
-GenericLoopNestRangeBuilder<loop::ParallelOp>::GenericLoopNestRangeBuilder(
+GenericLoopNestRangeBuilder<scf::ParallelOp>::GenericLoopNestRangeBuilder(
     MutableArrayRef<Value> ivs, ArrayRef<Value> ranges) {
   SmallVector<Value, 4> lbs, ubs, steps;
   for (Value range : ranges) {
     assert(range.getType() && "expected linalg.range type");
     assert(range.getDefiningOp() && "need operations to extract range parts");
     RangeOp rangeOp = cast<RangeOp>(range.getDefiningOp());
     lbs.emplace_back(rangeOp.min());
     ubs.emplace_back(rangeOp.max());

mlir/lib/Dialect/Linalg/EDSC/CMakeLists.txt

 
 target_link_libraries(MLIRLinalgEDSC
   PUBLIC
   MLIREDSC
   MLIRIR
   MLIRAffineOps
   MLIRAffineEDSC
   MLIRLinalgOps
-  MLIRLoopOps
+  MLIRSCF
   MLIRStandardOps
   LLVMSupport
   )

mlir/lib/Dialect/Linalg/Transforms/CMakeLists.txt

target_link_libraries(MLIRLinalgTransforms
   MLIRAffineOps
   MLIRAnalysis
   MLIREDSC
   MLIRIR
   MLIRLinalgAnalysis
   MLIRLinalgEDSC
   MLIRLinalgOps
   MLIRLinalgUtils
-  MLIRLoopOps
+  MLIRSCF
   MLIRPass
   MLIRStandardOps
   MLIRStandardToLLVM
   MLIRTransformUtils
   MLIRVector
   )

mlir/lib/Dialect/Linalg/Transforms/Loops.cpp

 //===- Loops.cpp - conversion from Linalg named and generic ops to loops --===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #include "PassDetail.h"
 #include "mlir/Dialect/Affine/EDSC/Intrinsics.h"
 #include "mlir/Dialect/Linalg/EDSC/FoldedIntrinsics.h"
 #include "mlir/Dialect/Linalg/IR/LinalgOps.h"
 #include "mlir/Dialect/Linalg/IR/LinalgTypes.h"
 #include "mlir/Dialect/Linalg/Passes.h"
 #include "mlir/Dialect/Linalg/Transforms/Transforms.h"
 #include "mlir/Dialect/Linalg/Utils/Utils.h"
-#include "mlir/Dialect/LoopOps/EDSC/Builders.h"
+#include "mlir/Dialect/SCF/EDSC/Builders.h"
 #include "mlir/Dialect/StandardOps/EDSC/Intrinsics.h"
 #include "mlir/IR/AffineExpr.h"
 #include "mlir/IR/AffineMap.h"
 #include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/Support/LLVM.h"
 #include "mlir/Transforms/DialectConversion.h"
 #include "mlir/Transforms/FoldUtils.h"
 
static void doit(ConcreteOpTy linalgOp, ArrayRef<Value> loopRanges,
     });
   }
 };
 
 /// Generates loop nest using loop.parallel. loop.parallel is only used for the
 /// outer parallel loops. All other loops are generated using loop.for
 /// operation.
 template <typename ConcreteOpTy>
-class GenerateLoopNest<loop::ParallelOp, ConcreteOpTy> {
+class GenerateLoopNest<scf::ParallelOp, ConcreteOpTy> {
 public:
   using IndexedValueTy = StdIndexedValue;
 
   static void doit(ConcreteOpTy linalgOp, ArrayRef<Value> loopRanges,
                    MutableArrayRef<Value> allIvs) {
     // Only generate loop.parallel for outer consecutive "parallel"
     // iterator_types.
     // TODO(ravishankarm): Generate loop.parallel for all "parallel" iterator
     // types, not just the outer most ones. Also handle "reduction" iterator
     // types.
     auto nOuterPar = linalgOp.iterator_types()
                          .getValue()
                          .take_while([](Attribute attr) {
                            return attr.cast<StringAttr>().getValue() ==
                                   getParallelIteratorTypeName();
                          })
                          .size();
     // If there are no outer parallel loops, then number of loop ops is same as
     // the number of loops, and they are all loop.for ops.
     if (nOuterPar) {
-      GenericLoopNestRangeBuilder<loop::ParallelOp>(
+      GenericLoopNestRangeBuilder<scf::ParallelOp>(
           allIvs.take_front(nOuterPar), loopRanges.take_front(nOuterPar))([&] {
-        GenericLoopNestRangeBuilder<loop::ForOp>(
+        GenericLoopNestRangeBuilder<scf::ForOp>(
             allIvs.drop_front(nOuterPar),
             loopRanges.drop_front(nOuterPar))([&] {
           SmallVector<Value, 4> allIvValues(allIvs.begin(), allIvs.end());
           LinalgScopedEmitter<StdIndexedValue, ConcreteOpTy>::
               emitScalarImplementation(allIvValues, linalgOp);
         });
       });
     } else {
       // If there are no parallel loops then fallback to generating all loop.for
       // operations.
-      GenericLoopNestRangeBuilder<loop::ForOp>(allIvs, loopRanges)([&] {
+      GenericLoopNestRangeBuilder<scf::ForOp>(allIvs, loopRanges)([&] {
         SmallVector<Value, 4> allIvValues(allIvs.begin(), allIvs.end());
         LinalgScopedEmitter<StdIndexedValue,
                             ConcreteOpTy>::emitScalarImplementation(allIvValues,
                                                                     linalgOp);
       });
     }
   }
 };

 struct LowerToAffineLoops
     : public LinalgLowerToAffineLoopsBase<LowerToAffineLoops> {
   void runOnFunction() override {
     lowerLinalgToLoopsImpl<AffineForOp>(getFunction(), &getContext());
   }
 };
 struct LowerToLoops : public LinalgLowerToLoopsBase<LowerToLoops> {
   void runOnFunction() override {
-    lowerLinalgToLoopsImpl<loop::ForOp>(getFunction(), &getContext());
+    lowerLinalgToLoopsImpl<scf::ForOp>(getFunction(), &getContext());
   }
 };
 struct LowerToParallelLoops
     : public LinalgLowerToParallelLoopsBase<LowerToParallelLoops> {
   void runOnFunction() override {
-    lowerLinalgToLoopsImpl<loop::ParallelOp>(getFunction(), &getContext());
+    lowerLinalgToLoopsImpl<scf::ParallelOp>(getFunction(), &getContext());
   }
 };
 } // namespace
 
 std::unique_ptr<OperationPass<FuncOp>> mlir::createConvertLinalgToLoopsPass() {
   return std::make_unique<LowerToLoops>();
 }
 
Optional<LinalgLoops> mlir::linalg::linalgLowerOpToLoops(OpBuilder &builder,
                                                          Operation *op) {
   return linalgOpToLoopsImpl<LoopTy, ConcreteOp>(op, builder);
 }
 
 /// Emits a loop nest of `loop.for` with the proper body for `op`.
 template <typename ConcreteOp>
 LogicalResult mlir::linalg::linalgOpToLoops(OpBuilder &builder, Operation *op) {
   Optional<LinalgLoops> loops =
-      linalgLowerOpToLoops<loop::ForOp, ConcreteOp>(builder, op);
+      linalgLowerOpToLoops<scf::ForOp, ConcreteOp>(builder, op);
   return loops ? success() : failure();
 }
 
 /// Emits a loop nest of `affine.for` with the proper body for `op`.
 template <typename ConcreteOp>
 LogicalResult mlir::linalg::linalgOpToAffineLoops(OpBuilder &builder,
                                                   Operation *op) {
   Optional<LinalgLoops> loops =
       linalgLowerOpToLoops<AffineForOp, ConcreteOp>(builder, op);
   return loops ? success() : failure();
 }
 
 /// Emits a loop nest of `loop.parallel` with the proper body for `op`.
 template <typename ConcreteOp>
 LogicalResult mlir::linalg::linalgOpToParallelLoops(OpBuilder &builder,
                                                     Operation *op) {
   Optional<LinalgLoops> loops =
-      linalgLowerOpToLoops<loop::ParallelOp, ConcreteOp>(builder, op);
+      linalgLowerOpToLoops<scf::ParallelOp, ConcreteOp>(builder, op);
   return loops ? success() : failure();
 }
 
 // TODO Need to make these instantiations more future-proof to avoid the need to
 // update as soon as we add new ops.
 #define INSTANTIATE_LINALG_OP_TO_LOOPS(OP_TYPE)                                \
   template LogicalResult mlir::linalg::linalgOpToLoops<OP_TYPE>(               \
       OpBuilder & builder, Operation * op);                                    \
   template LogicalResult mlir::linalg::linalgOpToAffineLoops<OP_TYPE>(         \
       OpBuilder & builder, Operation * op);                                    \
   template LogicalResult mlir::linalg::linalgOpToParallelLoops<OP_TYPE>(       \
       OpBuilder & builder, Operation * op);                                    \
   template Optional<LinalgLoops>                                               \
-      mlir::linalg::linalgLowerOpToLoops<loop::ParallelOp, OP_TYPE>(           \
+      mlir::linalg::linalgLowerOpToLoops<scf::ParallelOp, OP_TYPE>(            \
           OpBuilder & builder, Operation * op);
 
 INSTANTIATE_LINALG_OP_TO_LOOPS(CopyOp)
 INSTANTIATE_LINALG_OP_TO_LOOPS(FillOp)
 INSTANTIATE_LINALG_OP_TO_LOOPS(DotOp)
 INSTANTIATE_LINALG_OP_TO_LOOPS(MatvecOp)
 INSTANTIATE_LINALG_OP_TO_LOOPS(MatmulOp)
 INSTANTIATE_LINALG_OP_TO_LOOPS(ConvOp)
 INSTANTIATE_LINALG_OP_TO_LOOPS(PoolingMaxOp)
 INSTANTIATE_LINALG_OP_TO_LOOPS(PoolingMinOp)
 INSTANTIATE_LINALG_OP_TO_LOOPS(PoolingSumOp)
 INSTANTIATE_LINALG_OP_TO_LOOPS(GenericOp)
 INSTANTIATE_LINALG_OP_TO_LOOPS(IndexedGenericOp)

mlir/lib/Dialect/Linalg/Transforms/Promotion.cpp

 #include "PassDetail.h"
 #include "mlir/Dialect/Affine/EDSC/Intrinsics.h"
 #include "mlir/Dialect/Linalg/EDSC/FoldedIntrinsics.h"
 #include "mlir/Dialect/Linalg/IR/LinalgOps.h"
 #include "mlir/Dialect/Linalg/IR/LinalgTypes.h"
 #include "mlir/Dialect/Linalg/Passes.h"
 #include "mlir/Dialect/Linalg/Transforms/Transforms.h"
 #include "mlir/Dialect/Linalg/Utils/Utils.h"
-#include "mlir/Dialect/LoopOps/LoopOps.h"
+#include "mlir/Dialect/SCF/SCF.h"
 #include "mlir/Dialect/StandardOps/EDSC/Intrinsics.h"
 #include "mlir/IR/AffineExpr.h"
 #include "mlir/IR/AffineExprVisitor.h"
 #include "mlir/IR/AffineMap.h"
 #include "mlir/Support/LLVM.h"
 #include "mlir/Transforms/FoldUtils.h"
 
 #include "llvm/ADT/SetVector.h"
 #include "llvm/Support/CommandLine.h"
 
 using namespace mlir;
 using namespace mlir::edsc;
 using namespace mlir::edsc::intrinsics;
 using namespace mlir::linalg;
-using namespace mlir::loop;
+using namespace mlir::scf;
 
 using llvm::SetVector;
 
 using folded_affine_min = FoldedValueBuilder<AffineMinOp>;
 using folded_linalg_range = FoldedValueBuilder<linalg::RangeOp>;
 using folded_std_dim = FoldedValueBuilder<DimOp>;
 using folded_std_subview = FoldedValueBuilder<SubViewOp>;
 using folded_std_view = FoldedValueBuilder<ViewOp>;

mlir/lib/Dialect/Linalg/Transforms/Tiling.cpp

 
 #include "PassDetail.h"
 #include "mlir/Dialect/Affine/EDSC/Intrinsics.h"
 #include "mlir/Dialect/Linalg/EDSC/FoldedIntrinsics.h"
 #include "mlir/Dialect/Linalg/IR/LinalgTypes.h"
 #include "mlir/Dialect/Linalg/Passes.h"
 #include "mlir/Dialect/Linalg/Transforms/Transforms.h"
 #include "mlir/Dialect/Linalg/Utils/Utils.h"
-#include "mlir/Dialect/LoopOps/EDSC/Builders.h"
+#include "mlir/Dialect/SCF/EDSC/Builders.h"
 #include "mlir/Dialect/StandardOps/EDSC/Intrinsics.h"
 #include "mlir/IR/AffineExpr.h"
 #include "mlir/IR/AffineExprVisitor.h"
 #include "mlir/IR/AffineMap.h"
 #include "mlir/Support/LLVM.h"
 #include "mlir/Transforms/FoldUtils.h"
 
 #include "llvm/Support/CommandLine.h"
 
 using namespace mlir;
 using namespace mlir::edsc;
 using namespace mlir::edsc::intrinsics;
 using namespace mlir::linalg;
-using namespace mlir::loop;
+using namespace mlir::scf;
 
 using folded_affine_min = FoldedValueBuilder<AffineMinOp>;
 
 #define DEBUG_TYPE "linalg-tiling"
 
 static bool isZero(Value v) {
   return isa_and_nonnull<ConstantIndexOp>(v.getDefiningOp()) &&
          cast<ConstantIndexOp>(v.getDefiningOp()).getValue() == 0;
tileLinalgOpImpl(OpBuilder &b, LinalgOp op, ArrayRef<int64_t> tileSizes,
   return tileLinalgOpImpl<LoopTy>(b, op, tileSizeValues, interchangeVector,
                                   folder);
 }
 
 Optional<TiledLinalgOp>
 mlir::linalg::tileLinalgOp(OpBuilder &b, LinalgOp op, ArrayRef<Value> tileSizes,
                            ArrayRef<unsigned> interchangeVector,
                            OperationFolder *folder) {
-  return tileLinalgOpImpl<loop::ForOp>(b, op, tileSizes, interchangeVector,
-                                       folder);
+  return tileLinalgOpImpl<scf::ForOp>(b, op, tileSizes, interchangeVector,
 }
 
 Optional<TiledLinalgOp> mlir::linalg::tileLinalgOpToParallelLoops(
     OpBuilder &b, LinalgOp op, ArrayRef<Value> tileSizes,
     ArrayRef<unsigned> interchangeVector, OperationFolder *folder) {
-  return tileLinalgOpImpl<loop::ParallelOp>(b, op, tileSizes, interchangeVector,
-                                            folder);
+  return tileLinalgOpImpl<scf::ParallelOp>(b, op, tileSizes, interchangeVector,
 }
 
 Optional<TiledLinalgOp> mlir::linalg::tileLinalgOp(
     OpBuilder &b, LinalgOp op, ArrayRef<int64_t> tileSizes,
     ArrayRef<unsigned> interchangeVector, OperationFolder *folder) {
-  return tileLinalgOpImpl<loop::ForOp>(b, op, tileSizes, interchangeVector,
-                                       folder);
+  return tileLinalgOpImpl<scf::ForOp>(b, op, tileSizes, interchangeVector,
 }
 
 Optional<TiledLinalgOp> mlir::linalg::tileLinalgOpToParallelLoops(
     OpBuilder &b, LinalgOp op, ArrayRef<int64_t> tileSizes,
     ArrayRef<unsigned> interchangeVector, OperationFolder *folder) {
-  return tileLinalgOpImpl<loop::ParallelOp>(b, op, tileSizes, interchangeVector,
-                                            folder);
+  return tileLinalgOpImpl<scf::ParallelOp>(b, op, tileSizes, interchangeVector,
 }
 
 template <typename LoopTy>
 static void tileLinalgOps(FuncOp f, ArrayRef<int64_t> tileSizes) {
   OpBuilder b(f);
   OperationFolder folder(f.getContext());
   f.walk([tileSizes, &b, &folder](LinalgOp op) {
     if (!op.hasBufferSemantics())

 }
 
 namespace {
 struct LinalgTilingPass : public LinalgTilingBase<LinalgTilingPass> {
   LinalgTilingPass() = default;
   LinalgTilingPass(ArrayRef<int64_t> sizes) { tileSizes = sizes; }
 
   void runOnFunction() override {
-    tileLinalgOps<loop::ForOp>(getFunction(), tileSizes);
+    tileLinalgOps<scf::ForOp>(getFunction(), tileSizes);
   }
 };
 
 struct LinalgTilingToParallelLoopsPass
     : public LinalgTilingToParallelLoopsBase<LinalgTilingToParallelLoopsPass> {
   LinalgTilingToParallelLoopsPass() = default;
   LinalgTilingToParallelLoopsPass(ArrayRef<int64_t> sizes) {
     tileSizes = sizes;
   }
 
   void runOnFunction() override {
-    tileLinalgOps<loop::ParallelOp>(getFunction(), tileSizes);
+    tileLinalgOps<scf::ParallelOp>(getFunction(), tileSizes);
   }
 };
 
 } // namespace
 
 std::unique_ptr<OperationPass<FuncOp>>
 mlir::createLinalgTilingPass(ArrayRef<int64_t> tileSizes) {
   return std::make_unique<LinalgTilingPass>(tileSizes);
 }
 
 std::unique_ptr<OperationPass<FuncOp>>
 mlir::createLinalgTilingToParallelLoopsPass(ArrayRef<int64_t> tileSizes) {
   return std::make_unique<LinalgTilingToParallelLoopsPass>(tileSizes);
 }

mlir/lib/Dialect/Linalg/Utils/CMakeLists.txt

 add_mlir_dialect_library(MLIRLinalgUtils
   Utils.cpp
 
   ADDITIONAL_HEADER_DIRS
   ${MLIR_MAIN_INCLUDE_DIR}/mlir/Dialect/Linalg
   DEPENDS
   intrinsics_gen
   )
 
 target_link_libraries(MLIRLinalgUtils
   PUBLIC
   MLIRAffineOps
   MLIREDSC
   MLIRIR
   MLIRLinalgOps
-  MLIRLoopOps
+  MLIRSCF
   MLIRPass
   MLIRStandardOps
   MLIRTransformUtils
   )

mlir/lib/Dialect/Linalg/Utils/Utils.cpp

 //===- Utils.cpp - Utilities to support the Linalg dialect ----------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements utilities for the Linalg dialect.
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/Linalg/Utils/Utils.h"
 #include "mlir/Dialect/Affine/IR/AffineOps.h"
 #include "mlir/Dialect/Linalg/IR/LinalgOps.h"
 #include "mlir/Dialect/Linalg/IR/LinalgTypes.h"
-#include "mlir/Dialect/LoopOps/LoopOps.h"
+#include "mlir/Dialect/SCF/SCF.h"
 #include "mlir/Dialect/StandardOps/IR/Ops.h"
 #include "mlir/IR/AffineExpr.h"
 #include "mlir/IR/AffineMap.h"
 #include "mlir/IR/Matchers.h"
 #include "mlir/IR/OpImplementation.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Transforms/FoldUtils.h"
 
 using namespace mlir;
 using namespace mlir::linalg;
-using namespace mlir::loop;
+using namespace mlir::scf;
 
 Optional<RegionMatcher::BinaryOpKind>
 RegionMatcher::matchAsScalarBinaryOp(GenericOp op) {
   auto &region = op.region();
   if (!llvm::hasSingleElement(region))
     return llvm::None;
 
   Block &block = region.front();

mlir/lib/Dialect/SCF/CMakeLists.txt

 file(GLOB globbed *.c *.cpp)
-add_mlir_dialect_library(MLIRLoopOps
+add_mlir_dialect_library(MLIRSCF
   ${globbed}
   EDSC/Builders.cpp
 
   ADDITIONAL_HEADER_DIRS
   ${MLIR_MAIN_INCLUDE_DIR}/mlir/LoopOps
 
   DEPENDS
-  MLIRLoopOpsIncGen
+  MLIRSCFOpsIncGen
   )
-target_link_libraries(MLIRLoopOps
+target_link_libraries(MLIRSCF
   PUBLIC
   MLIREDSC
   MLIRIR
   MLIRLoopLikeInterface
   MLIRSideEffects
   MLIRStandardOps
   LLVMSupport
   )
 
 add_subdirectory(Transforms)

mlir/lib/Dialect/SCF/EDSC/Builders.cpp

 //===- Builders.cpp - MLIR Declarative Builder Classes --------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#include "mlir/Dialect/LoopOps/EDSC/Builders.h"
+#include "mlir/Dialect/SCF/EDSC/Builders.h"
 #include "mlir/IR/AffineExpr.h"
 #include "mlir/IR/AffineMap.h"
 
 using namespace mlir;
 using namespace mlir::edsc;
 
 mlir::edsc::ParallelLoopNestBuilder::ParallelLoopNestBuilder(
     MutableArrayRef<Value> ivs, ArrayRef<Value> lbs, ArrayRef<Value> ubs,
mlir::edsc::LoopNestBuilder::LoopNestBuilder::operator()(
 
   return loops[0].getOp()->getResults();
 }
 
 LoopBuilder mlir::edsc::makeParallelLoopBuilder(MutableArrayRef<Value> ivs,
                                                 ArrayRef<Value> lbs,
                                                 ArrayRef<Value> ubs,
                                                 ArrayRef<Value> steps) {
-  loop::ParallelOp parallelOp = OperationBuilder<loop::ParallelOp>(
+  scf::ParallelOp parallelOp = OperationBuilder<scf::ParallelOp>(
       SmallVector<Value, 4>(lbs.begin(), lbs.end()),
       SmallVector<Value, 4>(ubs.begin(), ubs.end()),
       SmallVector<Value, 4>(steps.begin(), steps.end()));
   for (size_t i = 0, e = ivs.size(); i < e; ++i)
     ivs[i] = parallelOp.getBody()->getArgument(i);
   LoopBuilder result;
   result.enter(parallelOp.getBody());
   return result;
 }
 
 mlir::edsc::LoopBuilder
 mlir::edsc::makeLoopBuilder(Value *iv, Value lb, Value ub, Value step,
                             MutableArrayRef<Value> iterArgsHandles,
                             ValueRange iterArgsInitValues) {
   mlir::edsc::LoopBuilder result;
-  loop::ForOp forOp =
-      OperationBuilder<loop::ForOp>(lb, ub, step, iterArgsInitValues);
+  scf::ForOp forOp =
+      OperationBuilder<scf::ForOp>(lb, ub, step, iterArgsInitValues);
   *iv = Value(forOp.getInductionVar());
-  auto *body = loop::getForInductionVarOwner(*iv).getBody();
+  auto *body = scf::getForInductionVarOwner(*iv).getBody();
   for (size_t i = 0, e = iterArgsHandles.size(); i < e; ++i) {
     // Skipping the induction variable.
     iterArgsHandles[i] = body->getArgument(i + 1);
   }
   result.setOp(forOp);
   result.enter(body);
   return result;
 }

mlir/lib/Dialect/SCF/SCF.cpp

-//===- Ops.cpp - Loop MLIR Operations -------------------------------------===//
+//===- SCF.cpp - Structured Control Flow Operations -----------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#include "mlir/Dialect/LoopOps/LoopOps.h"
+#include "mlir/Dialect/SCF/SCF.h"
 #include "mlir/Dialect/StandardOps/IR/Ops.h"
 #include "mlir/IR/AffineExpr.h"
 #include "mlir/IR/AffineMap.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/Function.h"
 #include "mlir/IR/Matchers.h"
 #include "mlir/IR/Module.h"
 #include "mlir/IR/OpImplementation.h"
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/IR/StandardTypes.h"
 #include "mlir/IR/Value.h"
 #include "mlir/Support/MathExtras.h"
 
 using namespace mlir;
-using namespace mlir::loop;
+using namespace mlir::scf;
 
 //===----------------------------------------------------------------------===//
-// LoopOpsDialect
+// SCFDialect
 //===----------------------------------------------------------------------===//
 
-LoopOpsDialect::LoopOpsDialect(MLIRContext *context)
+SCFDialect::SCFDialect(MLIRContext *context)
     : Dialect(getDialectNamespace(), context) {
   addOperations<
 #define GET_OP_LIST
-#include "mlir/Dialect/LoopOps/LoopOps.cpp.inc"
+#include "mlir/Dialect/SCF/SCFOps.cpp.inc"
       >();
 }
 
 //===----------------------------------------------------------------------===//
 // ForOp
 //===----------------------------------------------------------------------===//
 
 void ForOp::build(OpBuilder &builder, OperationState &result, Value lb,

 }
 
 LogicalResult ForOp::moveOutOfLoop(ArrayRef<Operation *> ops) {
   for (auto op : ops)
     op->moveBefore(*this);
   return success();
 }
 
-ForOp mlir::loop::getForInductionVarOwner(Value val) {
+ForOp mlir::scf::getForInductionVarOwner(Value val) {
   auto ivArg = val.dyn_cast<BlockArgument>();
   if (!ivArg)
     return ForOp();
   assert(ivArg.getOwner() && "unlinked block argument");
   auto *containingOp = ivArg.getOwner()->getParentOp();
   return dyn_cast_or_null<ForOp>(containingOp);
 }
 

 }
 
 LogicalResult ParallelOp::moveOutOfLoop(ArrayRef<Operation *> ops) {
   for (auto op : ops)
     op->moveBefore(*this);
   return success();
 }
 
-ParallelOp mlir::loop::getParallelForInductionVarOwner(Value val) {
+ParallelOp mlir::scf::getParallelForInductionVarOwner(Value val) {
   auto ivArg = val.dyn_cast<BlockArgument>();
   if (!ivArg)
     return ParallelOp();
   assert(ivArg.getOwner() && "unlinked block argument");
   auto *containingOp = ivArg.getOwner()->getParentOp();
   return dyn_cast<ParallelOp>(containingOp);
 }
 
if (op.getNumOperands() != 0)
     p << ' ' << op.getOperands() << " : " << op.getOperandTypes();
 }
 
 //===----------------------------------------------------------------------===//
 // TableGen'd op method definitions
 //===----------------------------------------------------------------------===//
 
 #define GET_OP_CLASSES
-#include "mlir/Dialect/LoopOps/LoopOps.cpp.inc"
+#include "mlir/Dialect/SCF/SCFOps.cpp.inc"

mlir/lib/Dialect/SCF/Transforms/CMakeLists.txt

-add_mlir_dialect_library(MLIRLoopOpsTransforms
+add_mlir_dialect_library(MLIRSCFTransforms
   ParallelLoopFusion.cpp
   ParallelLoopSpecialization.cpp
   ParallelLoopTiling.cpp
 
   ADDITIONAL_HEADER_DIRS
-  ${MLIR_MAIN_INCLUDE_DIR}/mlir/Dialect/LoopOps
+  ${MLIR_MAIN_INCLUDE_DIR}/mlir/Dialect/SCF
 
   DEPENDS
-  MLIRLoopPassIncGen
+  MLIRSCFPassIncGen
   )
-target_link_libraries(MLIRLoopOpsTransforms
+target_link_libraries(MLIRSCFTransforms
   PUBLIC
   MLIRAffineOps
   MLIRIR
   MLIRPass
-  MLIRLoopOps
+  MLIRSCF
   MLIRStandardOps
   MLIRSupport
   LLVMSupport
   )

mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp

 //===- ParallelLoopFusion.cpp - Code to perform loop fusion ---------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements loop fusion on parallel loops.
 //
 //===----------------------------------------------------------------------===//
 
 #include "PassDetail.h"
-#include "mlir/Dialect/LoopOps/LoopOps.h"
-#include "mlir/Dialect/LoopOps/Passes.h"
-#include "mlir/Dialect/LoopOps/Transforms.h"
+#include "mlir/Dialect/SCF/Passes.h"
+#include "mlir/Dialect/SCF/SCF.h"
+#include "mlir/Dialect/SCF/Transforms.h"
 #include "mlir/Dialect/StandardOps/IR/Ops.h"
 #include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/OpDefinition.h"
 
 using namespace mlir;
-using namespace mlir::loop;
+using namespace mlir::scf;
 
 /// Verify there are no nested ParallelOps.
 static bool hasNestedParallelOp(ParallelOp ploop) {
   auto walkResult =
       ploop.getBody()->walk([](ParallelOp) { return WalkResult::interrupt(); });
   return walkResult.wasInterrupted();
 }
 
if (!isFusionLegal(firstPloop, secondPloop, firstToSecondPloopIndices))
     return;
 
   b.setInsertionPointToStart(secondPloop.getBody());
   for (auto &op : firstPloop.getBody()->without_terminator())
     b.clone(op, firstToSecondPloopIndices);
   firstPloop.erase();
 }
 
-void mlir::loop::naivelyFuseParallelOps(Region &region) {
+void mlir::scf::naivelyFuseParallelOps(Region &region) {
   OpBuilder b(region);
   // Consider every single block and attempt to fuse adjacent loops.
   for (auto &block : region) {
     SmallVector<SmallVector<ParallelOp, 8>, 1> ploopChains{{}};
     // Not using `walk()` to traverse only top-level parallel loops and also
     // make sure that there are no side-effecting ops between the parallel
     // loops.
     bool noSideEffects = true;

mlir/lib/Dialect/SCF/Transforms/ParallelLoopSpecialization.cpp

 //===- ParallelLoopSpecialization.cpp - loop.parallel specialization ------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // Specializes parallel loops for easier unrolling and vectorization.
 //
 //===----------------------------------------------------------------------===//
 
 #include "PassDetail.h"
 #include "mlir/Dialect/Affine/IR/AffineOps.h"
-#include "mlir/Dialect/LoopOps/LoopOps.h"
-#include "mlir/Dialect/LoopOps/Passes.h"
+#include "mlir/Dialect/SCF/Passes.h"
+#include "mlir/Dialect/SCF/SCF.h"
 #include "mlir/Dialect/StandardOps/IR/Ops.h"
 #include "mlir/IR/AffineExpr.h"
 #include "mlir/IR/BlockAndValueMapping.h"
 
 using namespace mlir;
-using loop::ParallelOp;
+using scf::ParallelOp;
 
 /// Rewrite a loop with bounds defined by an affine.min with a constant into 2
 /// loops after checking if the bounds are equal to that constant. This is
 /// beneficial if the loop will almost always have the constant bound and that
 /// version can be fully unrolled and vectorized.
 static void specializeLoopForUnrolling(ParallelOp op) {
   SmallVector<int64_t, 2> constantIndices;
   constantIndices.reserve(op.upperBound().size());
static void specializeLoopForUnrolling(ParallelOp op) {
   Value cond;
   for (auto bound : llvm::zip(op.upperBound(), constantIndices)) {
     Value constant = b.create<ConstantIndexOp>(op.getLoc(), std::get<1>(bound));
     Value cmp = b.create<CmpIOp>(op.getLoc(), CmpIPredicate::eq,
                                  std::get<0>(bound), constant);
     cond = cond ? b.create<AndOp>(op.getLoc(), cond, cmp) : cmp;
     map.map(std::get<0>(bound), constant);
   }
-  auto ifOp = b.create<loop::IfOp>(op.getLoc(), cond, /*withElseRegion=*/true);
+  auto ifOp = b.create<scf::IfOp>(op.getLoc(), cond, /*withElseRegion=*/true);
   ifOp.getThenBodyBuilder().clone(*op.getOperation(), map);
   ifOp.getElseBodyBuilder().clone(*op.getOperation());
   op.erase();
 }
 
 namespace {
 struct ParallelLoopSpecialization
     : public LoopParallelLoopSpecializationBase<ParallelLoopSpecialization> {

mlir/lib/Dialect/SCF/Transforms/ParallelLoopTiling.cpp

 //===- ParallelLoopTiling.cpp - Tiles loop.parallel ---------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements loop tiling on parallel loops.
 //
 //===----------------------------------------------------------------------===//
 
 #include "PassDetail.h"
 #include "mlir/Dialect/Affine/IR/AffineOps.h"
-#include "mlir/Dialect/LoopOps/LoopOps.h"
-#include "mlir/Dialect/LoopOps/Passes.h"
-#include "mlir/Dialect/LoopOps/Transforms.h"
+#include "mlir/Dialect/SCF/Passes.h"
+#include "mlir/Dialect/SCF/SCF.h"
+#include "mlir/Dialect/SCF/Transforms.h"
 #include "mlir/Dialect/StandardOps/IR/Ops.h"
 #include "mlir/Transforms/RegionUtils.h"
 #include "llvm/Support/CommandLine.h"
 
 using namespace mlir;
-using namespace mlir::loop;
+using namespace mlir::scf;
 
 /// Tile a parallel loop of the form
 ///   loop.parallel (%i0, %i1) = (%arg0, %arg1) to (%arg2, %arg3)
 ///                                             step (%arg4, %arg5)
 ///
 /// into
 ///   loop.parallel (%i0, %i1) = (%arg0, %arg1) to (%arg2, %arg3)
 ///                                             step (%arg4*tileSize[0],
 ///                                                   %arg5*tileSize[1])
 ///     loop.parallel (%j0, %j1) = (0, 0) to (min(tileSize[0], %arg2-%j0)
 ///                                           min(tileSize[1], %arg3-%j1))
 ///                                        step (%arg4, %arg5)
 /// The old loop is replaced with the new one.
-void mlir::loop::tileParallelLoop(ParallelOp op, ArrayRef<int64_t> tileSizes) {
+void mlir::scf::tileParallelLoop(ParallelOp op, ArrayRef<int64_t> tileSizes) {
   OpBuilder b(op);
   auto zero = b.create<ConstantIndexOp>(op.getLoc(), 0);
   SmallVector<Value, 2> tileSizeConstants;
   tileSizeConstants.reserve(op.upperBound().size());
   for (size_t i = 0, end = op.upperBound().size(); i != end; ++i) {
     if (i < tileSizes.size())
       tileSizeConstants.push_back(
           b.create<ConstantIndexOp>(op.getLoc(), tileSizes[i]));

mlir/lib/Dialect/SCF/Transforms/PassDetail.h

 //===- PassDetail.h - Loop Pass class details -------------------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef DIALECT_LOOPOPS_TRANSFORMS_PASSDETAIL_H_
 #define DIALECT_LOOPOPS_TRANSFORMS_PASSDETAIL_H_
 
 #include "mlir/Pass/Pass.h"
 
 namespace mlir {
 
 #define GEN_PASS_CLASSES
-#include "mlir/Dialect/LoopOps/Passes.h.inc"
+#include "mlir/Dialect/SCF/Passes.h.inc"
 
 } // end namespace mlir
 
 #endif // DIALECT_LOOPOPS_TRANSFORMS_PASSDETAIL_H_

mlir/lib/Dialect/Vector/CMakeLists.txt

add_mlir_dialect_library(MLIRVector
   MLIRVectorTransformPatternsIncGen
   )
 target_link_libraries(MLIRVector
   PUBLIC
   MLIREDSC
   MLIRIR
   MLIRStandardOps
   MLIRAffineOps
-  MLIRLoopOps
+  MLIRSCF
   MLIRLoopAnalysis
   MLIRSideEffects
   )

mlir/lib/Transforms/CMakeLists.txt

add_mlir_library(MLIRTransforms
   MLIRTransformsPassIncGen
   )
 
 target_link_libraries(MLIRTransforms
   PUBLIC
   MLIRAffineOps
   MLIRAnalysis
   MLIRLoopLikeInterface
-  MLIRLoopOps
+  MLIRSCF
   MLIRPass
   MLIRTransformUtils
   MLIRVector
   )

mlir/lib/Transforms/LoopCoalescing.cpp

 //===- LoopCoalescing.cpp - Pass transforming loop nests into single loops-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #include "PassDetail.h"
-#include "mlir/Dialect/LoopOps/LoopOps.h"
+#include "mlir/Dialect/SCF/SCF.h"
 #include "mlir/Transforms/LoopUtils.h"
 #include "mlir/Transforms/Passes.h"
 #include "mlir/Transforms/RegionUtils.h"
 #include "llvm/Support/Debug.h"
 
 #define PASS_NAME "loop-coalescing"
 #define DEBUG_TYPE PASS_NAME
 
 using namespace mlir;
 
 namespace {
 struct LoopCoalescingPass : public LoopCoalescingBase<LoopCoalescingPass> {
   void runOnFunction() override {
     FuncOp func = getFunction();
 
-    func.walk([](loop::ForOp op) {
+    func.walk([](scf::ForOp op) {
       // Ignore nested loops.
-      if (op.getParentOfType<loop::ForOp>())
+      if (op.getParentOfType<scf::ForOp>())
         return;
 
-      SmallVector<loop::ForOp, 4> loops;
+      SmallVector<scf::ForOp, 4> loops;
       getPerfectlyNestedLoops(loops, op);
       LLVM_DEBUG(llvm::dbgs()
                  << "found a perfect nest of depth " << loops.size() << '\n');
 
       // Look for a band of loops that can be coalesced, i.e. perfectly nested
       // loops with bounds defined above some loop.
       // 1. For each loop, find above which parent loop its operands are
       // defined.

mlir/lib/Transforms/ParallelLoopCollapsing.cpp

 //===- ParallelLoopCollapsing.cpp - Pass collapsing parallel loop indices -===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #include "PassDetail.h"
-#include "mlir/Dialect/LoopOps/LoopOps.h"
+#include "mlir/Dialect/SCF/SCF.h"
 #include "mlir/Transforms/LoopUtils.h"
 #include "mlir/Transforms/Passes.h"
 #include "mlir/Transforms/RegionUtils.h"
 #include "llvm/Support/CommandLine.h"
 #include "llvm/Support/Debug.h"
 
 #define DEBUG_TYPE "parallel-loop-collapsing"
 
 using namespace mlir;
 
 namespace {
 struct ParallelLoopCollapsing
     : public ParallelLoopCollapsingBase<ParallelLoopCollapsing> {
   void runOnOperation() override {
     Operation *module = getOperation();
 
-    module->walk([&](loop::ParallelOp op) {
+    module->walk([&](scf::ParallelOp op) {
       // The common case for GPU dialect will be simplifying the ParallelOp to 3
       // arguments, so we do that here to simplify things.
       llvm::SmallVector<std::vector<unsigned>, 3> combinedLoops;
       if (clCollapsedIndices0.size())
         combinedLoops.push_back(clCollapsedIndices0);
       if (clCollapsedIndices1.size())
         combinedLoops.push_back(clCollapsedIndices1);
       if (clCollapsedIndices2.size())

mlir/lib/Transforms/Utils/CMakeLists.txt

add_mlir_library(MLIRTransformUtils
   MLIRStandardOpsIncGen
   )
 
 target_link_libraries(MLIRTransformUtils
   PUBLIC
   MLIRAffineOps
   MLIRAnalysis
   MLIRLoopAnalysis
-  MLIRLoopOps
+  MLIRSCF
   MLIRPass
   MLIRStandardOps
   )

mlir/lib/Transforms/Utils/LoopUtils.cpp

 #include "mlir/Transforms/LoopUtils.h"
 
 #include "mlir/Analysis/AffineAnalysis.h"
 #include "mlir/Analysis/LoopAnalysis.h"
 #include "mlir/Analysis/SliceAnalysis.h"
 #include "mlir/Analysis/Utils.h"
 #include "mlir/Dialect/Affine/IR/AffineOps.h"
 #include "mlir/Dialect/Affine/IR/AffineValueMap.h"
-#include "mlir/Dialect/LoopOps/LoopOps.h"
+#include "mlir/Dialect/SCF/SCF.h"
 #include "mlir/IR/AffineMap.h"
 #include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Function.h"
 #include "mlir/IR/IntegerSet.h"
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/Transforms/RegionUtils.h"
 #include "mlir/Transforms/Utils.h"
 #include "llvm/ADT/DenseMap.h"

 /// (the first op being another AffineFor, and the second op - a terminator).
 /// A loop is perfectly nested iff: the first op in the loop's body is another
 /// AffineForOp, and the second op is a terminator).
 void mlir::getPerfectlyNestedLoops(SmallVectorImpl<AffineForOp> &nestedLoops,
                                    AffineForOp root) {
   getPerfectlyNestedLoopsImpl(nestedLoops, root);
 }
 
-void mlir::getPerfectlyNestedLoops(SmallVectorImpl<loop::ForOp> &nestedLoops,
-                                   loop::ForOp root) {
+void mlir::getPerfectlyNestedLoops(SmallVectorImpl<scf::ForOp> &nestedLoops,
+                                   scf::ForOp root) {
   getPerfectlyNestedLoopsImpl(nestedLoops, root);
 }
 
 /// Unrolls this loop completely.
 LogicalResult mlir::loopUnrollFull(AffineForOp forOp) {
   Optional<uint64_t> mayBeConstantTripCount = getConstantTripCount(forOp);
   if (mayBeConstantTripCount.hasValue()) {
     uint64_t tripCount = mayBeConstantTripCount.getValue();
for (auto t : targets) {
     replaceAllUsesInRegionWith(iv, newForOp.getInductionVar(),
                                newForOp.region());
     innerLoops.push_back(newForOp);
   }
 
   return innerLoops;
 }
 
-static Loops stripmineSink(loop::ForOp forOp, Value factor,
-                           ArrayRef<loop::ForOp> targets) {
+static Loops stripmineSink(scf::ForOp forOp, Value factor,
+                           ArrayRef<scf::ForOp> targets) {
   auto originalStep = forOp.step();
   auto iv = forOp.getInductionVar();
 
   OpBuilder b(forOp);
   forOp.setStep(b.create<MulIOp>(forOp.getLoc(), originalStep, factor));
 
   Loops innerLoops;
   for (auto t : targets) {
     // Save information for splicing ops out of t when done
     auto begin = t.getBody()->begin();
     auto nOps = t.getBody()->getOperations().size();
 
     // Insert newForOp before the terminator of `t`.
     auto b = OpBuilder::atBlockTerminator((t.getBody()));
     Value stepped = b.create<AddIOp>(t.getLoc(), iv, forOp.step());
     Value less = b.create<CmpIOp>(t.getLoc(), CmpIPredicate::slt,
                                   forOp.upperBound(), stepped);
     Value ub =
         b.create<SelectOp>(t.getLoc(), less, forOp.upperBound(), stepped);
 
     // Splice [begin, begin + nOps - 1) into `newForOp` and replace uses.
-    auto newForOp = b.create<loop::ForOp>(t.getLoc(), iv, ub, originalStep);
+    auto newForOp = b.create<scf::ForOp>(t.getLoc(), iv, ub, originalStep);
     newForOp.getBody()->getOperations().splice(
         newForOp.getBody()->getOperations().begin(),
         t.getBody()->getOperations(), begin, std::next(begin, nOps - 1));
     replaceAllUsesInRegionWith(iv, newForOp.getInductionVar(),
                                newForOp.region());
 
     innerLoops.push_back(newForOp);
   }

 }
 
 SmallVector<SmallVector<AffineForOp, 8>, 8>
 mlir::tile(ArrayRef<AffineForOp> forOps, ArrayRef<uint64_t> sizes,
            ArrayRef<AffineForOp> targets) {
   return tileImpl(forOps, sizes, targets);
 }
 
-SmallVector<Loops, 8> mlir::tile(ArrayRef<loop::ForOp> forOps,
+SmallVector<Loops, 8> mlir::tile(ArrayRef<scf::ForOp> forOps,
                                  ArrayRef<Value> sizes,
-                                 ArrayRef<loop::ForOp> targets) {
+                                 ArrayRef<scf::ForOp> targets) {
   return tileImpl(forOps, sizes, targets);
 }
 
 template <typename ForType, typename SizeType>
 static SmallVector<ForType, 8>
 tileImpl(ArrayRef<ForType> forOps, ArrayRef<SizeType> sizes, ForType target) {
   SmallVector<ForType, 8> res;
   for (auto loops : tile(forOps, sizes, ArrayRef<ForType>{target})) {
     assert(loops.size() == 1);
     res.push_back(loops[0]);
   }
   return res;
 }
 
 SmallVector<AffineForOp, 8> mlir::tile(ArrayRef<AffineForOp> forOps,
                                        ArrayRef<uint64_t> sizes,
                                        AffineForOp target) {
   return tileImpl(forOps, sizes, target);
 }
 
-Loops mlir::tile(ArrayRef<loop::ForOp> forOps, ArrayRef<Value> sizes,
-                 loop::ForOp target) {
+Loops mlir::tile(ArrayRef<scf::ForOp> forOps, ArrayRef<Value> sizes,
+                 scf::ForOp target) {
   return tileImpl(forOps, sizes, target);
 }
 
-Loops mlir::tilePerfectlyNested(loop::ForOp rootForOp, ArrayRef<Value> sizes) {
+Loops mlir::tilePerfectlyNested(scf::ForOp rootForOp, ArrayRef<Value> sizes) {
   // Collect perfectly nested loops.  If more size values provided than nested
   // loops available, truncate `sizes`.
-  SmallVector<loop::ForOp, 4> forOps;
+  SmallVector<scf::ForOp, 4> forOps;
   forOps.reserve(sizes.size());
   getPerfectlyNestedLoopsImpl(forOps, rootForOp, sizes.size());
   if (forOps.size() < sizes.size())
     sizes = sizes.take_front(forOps.size());
 
   return ::tile(forOps, sizes, forOps.back());
 }
 
static Value ceilDivPositive(OpBuilder &builder, Location loc, Value dividend,
   return builder.create<SignedDivIOp>(loc, sum, divisor);
 }
 
 // Hoist the ops within `outer` that appear before `inner`.
 // Such ops include the ops that have been introduced by parametric tiling.
 // Ops that come from triangular loops (i.e. that belong to the program slice
 // rooted at `outer`) and ops that have side effects cannot be hoisted.
 // Return failure when any op fails to hoist.
-static LogicalResult hoistOpsBetween(loop::ForOp outer, loop::ForOp inner) {
+static LogicalResult hoistOpsBetween(scf::ForOp outer, scf::ForOp inner) {
   SetVector<Operation *> forwardSlice;
   getForwardSlice(outer.getOperation(), &forwardSlice, [&inner](Operation *op) {
     return op != inner.getOperation();
   });
   LogicalResult status = success();
   SmallVector<Operation *, 8> toHoist;
   for (auto &op : outer.getBody()->without_terminator()) {
     // Stop when encountering the inner loop.
     if (&op == inner.getOperation())
       break;
     // Skip over non-hoistable ops.
     if (forwardSlice.count(&op) > 0) {
       status = failure();
       continue;
     }
-    // Skip loop::ForOp, these are not considered a failure.
+    // Skip scf::ForOp, these are not considered a failure.
     if (op.getNumRegions() > 0)
       continue;
     // Skip other ops with regions.
     if (op.getNumRegions() > 0) {
       status = failure();
       continue;
     }
     // Skip if op has side effects.
for (unsigned s = 1; s < size; ++s)
     status = succeeded(status) ? hoistOpsBetween(intraTile[0], intraTile[s])
                                : failure();
   for (unsigned s = 1; s < size; ++s)
     status = succeeded(status) ? hoistOpsBetween(interTile[0], interTile[s])
                                : failure();
   return status;
 }
 
-TileLoops mlir::extractFixedOuterLoops(loop::ForOp rootForOp,
+TileLoops mlir::extractFixedOuterLoops(scf::ForOp rootForOp,
                                        ArrayRef<int64_t> sizes) {
   // Collect perfectly nested loops.  If more size values provided than nested
   // loops available, truncate `sizes`.
-  SmallVector<loop::ForOp, 4> forOps;
+  SmallVector<scf::ForOp, 4> forOps;
   forOps.reserve(sizes.size());
   getPerfectlyNestedLoopsImpl(forOps, rootForOp, sizes.size());
   if (forOps.size() < sizes.size())
     sizes = sizes.take_front(forOps.size());
 
   // Compute the tile sizes such that i-th outer loop executes size[i]
   // iterations.  Given that the loop current executes
   //   numIterations = ceildiv((upperBound - lowerBound), step)

 ///   for %i = %lb to %ub step %s
 /// into a 0-based loop with step 1
 ///   for %ii = 0 to ceildiv(%ub - %lb, %s) step 1 {
 ///     %i = %ii * %s + %lb
 /// Insert the induction variable remapping in the body of `inner`, which is
 /// expected to be either `loop` or another loop perfectly nested under `loop`.
 /// Insert the definition of new bounds immediate before `outer`, which is
 /// expected to be either `loop` or its parent in the loop nest.
-static void normalizeLoop(loop::ForOp loop, loop::ForOp outer,
-                          loop::ForOp inner) {
+static void normalizeLoop(scf::ForOp loop, scf::ForOp outer, scf::ForOp inner) {
   OpBuilder builder(outer);
   OpBuilder innerBuilder = OpBuilder::atBlockBegin(inner.getBody());
   auto loopPieces =
       normalizeLoop(builder, innerBuilder, loop.getLoc(), loop.lowerBound(),
                     loop.upperBound(), loop.step(), loop.getInductionVar());
 
   loop.setLowerBound(loopPieces.lowerBound);
   loop.setUpperBound(loopPieces.upperBound);
   loop.setStep(loopPieces.step);
 }
 
-void mlir::coalesceLoops(MutableArrayRef<loop::ForOp> loops) {
+void mlir::coalesceLoops(MutableArrayRef<scf::ForOp> loops) {
   if (loops.size() < 2)
     return;
 
-  loop::ForOp innermost = loops.back();
-  loop::ForOp outermost = loops.front();
+  scf::ForOp innermost = loops.back();
+  scf::ForOp outermost = loops.front();
 
   // 1. Make sure all loops iterate from 0 to upperBound with step 1.  This
   // allows the following code to assume upperBound is the number of iterations.
   for (auto loop : loops)
     normalizeLoop(loop, outermost, innermost);
 
   // 2. Emit code computing the upper bound of the coalesced loop as product
   // of the number of iterations of all loops.
Value iv = (i == e - 1) ? previous
                             : builder.create<SignedRemIOp>(
                                   loc, previous, loops[idx].upperBound());
     replaceAllUsesInRegionWith(loops[idx].getInductionVar(), iv,
                                loops.back().region());
   }
 
   // 4. Move the operations from the innermost just above the second-outermost
   // loop, delete the extra terminator and the second-outermost loop.
-  loop::ForOp second = loops[1];
+  scf::ForOp second = loops[1];
   innermost.getBody()->back().erase();
   outermost.getBody()->getOperations().splice(
       Block::iterator(second.getOperation()),
       innermost.getBody()->getOperations());
   second.erase();
 }
 
 void mlir::collapseParallelLoops(
-    loop::ParallelOp loops,
-    ArrayRef<std::vector<unsigned>> combinedDimensions) {
+    scf::ParallelOp loops, ArrayRef<std::vector<unsigned>> combinedDimensions) {
   OpBuilder outsideBuilder(loops);
   Location loc = loops.getLoc();
 
   // Normalize ParallelOp's iteration pattern.
   SmallVector<Value, 3> normalizedLowerBounds;
   SmallVector<Value, 3> normalizedSteps;
   SmallVector<Value, 3> normalizedUpperBounds;
   for (unsigned i = 0, e = loops.getNumLoops(); i < e; ++i) {
void mlir::collapseParallelLoops(
   }
 
   // Create new ParallelLoop with conversions to the original induction values.
   // The loop below uses divisions to get the relevant range of values in the
   // new induction value that represent each range of the original induction
   // value. The remainders then determine based on that range, which iteration
   // of the original induction value this represents. This is a normalized value
   // that is un-normalized already by the previous logic.
-  auto newPloop = outsideBuilder.create<loop::ParallelOp>(loc, lowerBounds,
-                                                          upperBounds, steps);
+  auto newPloop = outsideBuilder.create<scf::ParallelOp>(loc, lowerBounds,
   OpBuilder insideBuilder(newPloop.region());
   for (unsigned i = 0, e = combinedDimensions.size(); i < e; ++i) {
     Value previous = newPloop.getBody()->getArgument(i);
     unsigned numberCombinedDimensions = combinedDimensions[i].size();
     // Iterate over all except the last induction value.
     for (unsigned j = 0, e = numberCombinedDimensions - 1; j < e; ++j) {
       unsigned idx = combinedDimensions[i][j];
 
void mlir::collapseParallelLoops(
   // Replace the old loop with the new loop.
   loops.getBody()->back().erase();
   newPloop.getBody()->getOperations().splice(
       Block::iterator(newPloop.getBody()->back()),
       loops.getBody()->getOperations());
   loops.erase();
 }
 
-void mlir::mapLoopToProcessorIds(loop::ForOp forOp, ArrayRef<Value> processorId,
+void mlir::mapLoopToProcessorIds(scf::ForOp forOp, ArrayRef<Value> processorId,
                                  ArrayRef<Value> numProcessors) {
   assert(processorId.size() == numProcessors.size());
   if (processorId.empty())
     return;
 
   OpBuilder b(forOp);
   Location loc(forOp.getLoc());
   Value mul = processorId.front();

mlir/test/Dialect/SCF/invalid.mlir

 // RUN: mlir-opt -allow-unregistered-dialect %s -split-input-file -verify-diagnostics
 
 func @loop_for_lb(%arg0: f32, %arg1: index) {
   // expected-error@+1 {{operand #0 must be index}}
   "loop.for"(%arg0, %arg1, %arg1) ({}) : (f32, index, index) -> ()
   return
 }
 
 // -----
 
 func @loop_for_ub(%arg0: f32, %arg1: index) {
   // expected-error@+1 {{operand #1 must be index}}
   "loop.for"(%arg1, %arg0, %arg1) ({}) : (index, f32, index) -> ()
   return
 }
 
 // -----
 
 func @loop_for_step(%arg0: f32, %arg1: index) {
   // expected-error@+1 {{operand #2 must be index}}
   "loop.for"(%arg1, %arg1, %arg0) ({}) : (index, index, f32) -> ()
   return
 }
 
 // -----
 
 func @loop_for_step_positive(%arg0: index) {
   // expected-error@+2 {{constant step operand must be positive}}
   %c0 = constant 0 : index
   "loop.for"(%arg0, %arg0, %c0) ({
     ^bb0(%arg1: index):
       loop.yield
   }) : (index, index, index) -> ()
   return
 }
 
 // -----
 
 func @loop_for_one_region(%arg0: index) {
   // expected-error@+1 {{requires one region}}
   "loop.for"(%arg0, %arg0, %arg0) (
     {loop.yield},
     {loop.yield}
   ) : (index, index, index) -> ()
   return
 }
 
 // -----
 
 func @loop_for_single_block(%arg0: index) {
   // expected-error@+1 {{expects region #0 to have 0 or 1 blocks}}
   "loop.for"(%arg0, %arg0, %arg0) (
     {
     ^bb1:
       loop.yield
     ^bb2:
       loop.yield
     }
   ) : (index, index, index) -> ()
   return
 }
 
 // -----
 
 func @loop_for_single_index_argument(%arg0: index) {
   // expected-error@+1 {{op expected body first argument to be an index argument for the induction variable}}
   "loop.for"(%arg0, %arg0, %arg0) (
     {
     ^bb0(%i0 : f32):
       loop.yield
     }
   ) : (index, index, index) -> ()
   return
 }
 
 // -----
 
 func @loop_if_not_i1(%arg0: index) {
   // expected-error@+1 {{operand #0 must be 1-bit signless integer}}
   "loop.if"(%arg0) ({}, {}) : (index) -> ()
   return
 }
 
 // -----
 
 func @loop_if_more_than_2_regions(%arg0: i1) {
   // expected-error@+1 {{expected 2 regions}}
   "loop.if"(%arg0) ({}, {}, {}): (i1) -> ()
   return
 }
 
 // -----
 
 func @loop_if_not_one_block_per_region(%arg0: i1) {
   // expected-error@+1 {{expects region #0 to have 0 or 1 blocks}}
   "loop.if"(%arg0) ({
     ^bb0:
       loop.yield
     ^bb1:
       loop.yield
   }, {}): (i1) -> ()
   return
 }
 
 // -----
 
 func @loop_if_illegal_block_argument(%arg0: i1) {
   // expected-error@+1 {{requires that child entry blocks have no arguments}}
   "loop.if"(%arg0) ({
     ^bb0(%0 : index):
       loop.yield
   }, {}): (i1) -> ()
   return
 }
 
 // -----
 
 func @parallel_arguments_different_tuple_size(
     %arg0: index, %arg1: index, %arg2: index) {
   // expected-error@+1 {{custom op 'loop.parallel' expected 1 operands}}
   loop.parallel (%i0) = (%arg0) to (%arg1, %arg2) step () {
   }
   return
 }
 
 // -----
 
 func @parallel_body_arguments_wrong_type(
     %arg0: index, %arg1: index, %arg2: index) {
   // expected-error@+1 {{'loop.parallel' op expects arguments for the induction variable to be of index type}}
   "loop.parallel"(%arg0, %arg1, %arg2) ({
     ^bb0(%i0: f32):
       loop.yield
   }) {operand_segment_sizes = dense<[1, 1, 1, 0]>: vector<4xi32>}: (index, index, index) -> ()
   return
 }
 
 // -----
 
 func @parallel_body_wrong_number_of_arguments(
     %arg0: index, %arg1: index, %arg2: index) {
   // expected-error@+1 {{'loop.parallel' op expects the same number of induction variables: 2 as bound and step values: 1}}
   "loop.parallel"(%arg0, %arg1, %arg2) ({
     ^bb0(%i0: index, %i1: index):
       loop.yield
   }) {operand_segment_sizes = dense<[1, 1, 1, 0]>: vector<4xi32>}: (index, index, index) -> ()
   return
 }
 
 // -----
 
 func @parallel_no_tuple_elements() {
   // expected-error@+1 {{'loop.parallel' op needs at least one tuple element for lowerBound, upperBound and step}}
   loop.parallel () = () to () step () {
   }
   return
 }
 
 // -----
 
 func @parallel_step_not_positive(
     %arg0: index, %arg1: index, %arg2: index, %arg3: index) {
   // expected-error@+3 {{constant step operand must be positive}}
   %c0 = constant 1 : index
   %c1 = constant 0 : index
   loop.parallel (%i0, %i1) = (%arg0, %arg1) to (%arg2, %arg3) step (%c0, %c1) {
   }
   return
 }
 
 // -----
 
 func @parallel_fewer_results_than_reduces(
     %arg0 : index, %arg1: index, %arg2: index) {
   // expected-error@+1 {{expects number of results: 0 to be the same as number of reductions: 1}}
   loop.parallel (%i0) = (%arg0) to (%arg1) step (%arg2) {
     %c0 = constant 1.0 : f32
     loop.reduce(%c0) : f32 {
       ^bb0(%lhs: f32, %rhs: f32):
         loop.reduce.return %lhs : f32
     }
   }
   return
 }
 
 // -----
 
 func @parallel_more_results_than_reduces(
     %arg0 : index, %arg1 : index, %arg2 : index) {
   // expected-error@+2 {{expects number of results: 1 to be the same as number of reductions: 0}}
   %zero = constant 1.0 : f32
   %res = loop.parallel (%i0) = (%arg0) to (%arg1) step (%arg2) init (%zero) -> f32 {
   }
 
   return
 }
 
 // -----
 
 func @parallel_more_results_than_initial_values(
     %arg0 : index, %arg1: index, %arg2: index) {
   // expected-error@+1 {{expects number of results: 1 to be the same as number of initial values: 0}}
   %res = loop.parallel (%i0) = (%arg0) to (%arg1) step (%arg2) -> f32 {
     loop.reduce(%arg0) : index {
       ^bb0(%lhs: index, %rhs: index):
         loop.reduce.return %lhs : index
     }
   }
 }
 
 // -----
 
 func @parallel_different_types_of_results_and_reduces(
     %arg0 : index, %arg1: index, %arg2: index) {
   %zero = constant 0.0 : f32
   %res = loop.parallel (%i0) = (%arg0) to (%arg1)
                                        step (%arg2) init (%zero) -> f32 {
     // expected-error@+1 {{expects type of reduce: 'index' to be the same as result type: 'f32'}}
     loop.reduce(%arg0) : index {
       ^bb0(%lhs: index, %rhs: index):
         loop.reduce.return %lhs : index
     }
   }
   return
 }
 
 // -----
 
 func @top_level_reduce(%arg0 : f32) {
   // expected-error@+1 {{expects parent op 'loop.parallel'}}
   loop.reduce(%arg0) : f32 {
     ^bb0(%lhs : f32, %rhs : f32):
       loop.reduce.return %lhs : f32
   }
   return
 }
 
 // -----
 
 func @reduce_empty_block(%arg0 : index, %arg1 : f32) {
   %zero = constant 0.0 : f32
   %res = loop.parallel (%i0) = (%arg0) to (%arg0)
                                        step (%arg0) init (%zero) -> f32 {
     // expected-error@+1 {{the block inside reduce should not be empty}}
     loop.reduce(%arg1) : f32 {
       ^bb0(%lhs : f32, %rhs : f32):
     }
   }
   return
 }
 
 // -----
 
 func @reduce_too_many_args(%arg0 : index, %arg1 : f32) {
   %zero = constant 0.0 : f32
   %res = loop.parallel (%i0) = (%arg0) to (%arg0)
                                        step (%arg0) init (%zero) -> f32 {
     // expected-error@+1 {{expects two arguments to reduce block of type 'f32'}}
     loop.reduce(%arg1) : f32 {
       ^bb0(%lhs : f32, %rhs : f32, %other : f32):
         loop.reduce.return %lhs : f32
     }
   }
   return
 }
 
 // -----
 
 func @reduce_wrong_args(%arg0 : index, %arg1 : f32) {
   %zero = constant 0.0 : f32
   %res = loop.parallel (%i0) = (%arg0) to (%arg0)
                                        step (%arg0) init (%zero) -> f32 {
     // expected-error@+1 {{expects two arguments to reduce block of type 'f32'}}
     loop.reduce(%arg1) : f32 {
       ^bb0(%lhs : f32, %rhs : i32):
         loop.reduce.return %lhs : f32
     }
   }
   return
 }
 
 
 // -----
 
 func @reduce_wrong_terminator(%arg0 : index, %arg1 : f32) {
   %zero = constant 0.0 : f32
   %res = loop.parallel (%i0) = (%arg0) to (%arg0)
                                        step (%arg0) init (%zero) -> f32 {
     // expected-error@+1 {{the block inside reduce should be terminated with a 'loop.reduce.return' op}}
     loop.reduce(%arg1) : f32 {
       ^bb0(%lhs : f32, %rhs : f32):
         loop.yield
     }
   }
   return
 }
 
 // -----
 
 func @reduceReturn_wrong_type(%arg0 : index, %arg1: f32) {
   %zero = constant 0.0 : f32
   %res = loop.parallel (%i0) = (%arg0) to (%arg0)
                                        step (%arg0) init (%zero) -> f32 {
     loop.reduce(%arg1) : f32 {
       ^bb0(%lhs : f32, %rhs : f32):
         %c0 = constant 1 : index
         // expected-error@+1 {{needs to have type 'f32' (the type of the enclosing ReduceOp)}}
         loop.reduce.return %c0 : index
     }
   }
   return
 }
 
 // -----
 
 func @reduceReturn_not_inside_reduce(%arg0 : f32) {
   "foo.region"() ({
     // expected-error@+1 {{expects parent op 'loop.reduce'}}
     loop.reduce.return %arg0 : f32
   }): () -> ()
   return
 }
 
 // -----
 
 func @std_if_incorrect_yield(%arg0: i1, %arg1: f32)
 {
   %x, %y = loop.if %arg0 -> (f32, f32) {
     %0 = addf %arg1, %arg1 : f32
     // expected-error@+1 {{parent of yield must have same number of results as the yield operands}}
     loop.yield %0 : f32
   } else {
     %0 = subf %arg1, %arg1 : f32
     loop.yield %0 : f32
   }
   return
 }
 
 // -----
 
 func @std_if_missing_else(%arg0: i1, %arg1: f32)
 {
   // expected-error@+1 {{must have an else block if defining values}}
   %x = loop.if %arg0 -> (f32) {
     %0 = addf %arg1, %arg1 : f32
     loop.yield %0 : f32
   }
   return
 }
 
 // -----
 
 func @std_for_operands_mismatch(%arg0 : index, %arg1 : index, %arg2 : index) {
   %s0 = constant 0.0 : f32
   %t0 = constant 1 : i32
   // expected-error@+1 {{mismatch in number of loop-carried values and defined values}}
   %result1:3 = loop.for %i0 = %arg0 to %arg1 step %arg2
                     iter_args(%si = %s0, %ti = %t0) -> (f32, i32, f32) {
     %sn = addf %si, %si : f32
     %tn = addi %ti, %ti : i32
     loop.yield %sn, %tn, %sn : f32, i32, f32
   }
   return
 }
 
 // -----
 
 func @std_for_operands_mismatch_2(%arg0 : index, %arg1 : index, %arg2 : index) {
   %s0 = constant 0.0 : f32
   %t0 = constant 1 : i32
   %u0 = constant 1.0 : f32
   // expected-error@+1 {{mismatch in number of loop-carried values and defined values}}
   %result1:2 = loop.for %i0 = %arg0 to %arg1 step %arg2
                     iter_args(%si = %s0, %ti = %t0, %ui = %u0) -> (f32, i32) {
     %sn = addf %si, %si : f32
     %tn = addi %ti, %ti : i32
     %un = subf %ui, %ui : f32
     loop.yield %sn, %tn, %un : f32, i32, f32
   }
   return
 }
 
 // -----
 
 func @std_for_operands_mismatch_3(%arg0 : index, %arg1 : index, %arg2 : index) {
   // expected-note@+1 {{prior use here}}
   %s0 = constant 0.0 : f32
   %t0 = constant 1.0 : f32
   // expected-error@+2 {{expects different type than prior uses: 'i32' vs 'f32'}}
   %result1:2 = loop.for %i0 = %arg0 to %arg1 step %arg2
                     iter_args(%si = %s0, %ti = %t0) -> (i32, i32) {
     %sn = addf %si, %si : i32
     %tn = addf %ti, %ti : i32
     loop.yield %sn, %tn : i32, i32
   }
   return
 }
 
 // -----
 
 func @parallel_invalid_yield(
     %arg0: index, %arg1: index, %arg2: index) {
   loop.parallel (%i0) = (%arg0) to (%arg1) step (%arg2) {
     %c0 = constant 1.0 : f32
     // expected-error@+1 {{yield inside loop.parallel is not allowed to have operands}}
     loop.yield %c0 : f32
   }
   return
 }

mlir/test/Dialect/SCF/ops.mlir

 // RUN: mlir-opt %s | FileCheck %s --dump-input-on-failure
 // Verify the printed output can be parsed.
 // RUN: mlir-opt %s | mlir-opt | FileCheck %s --dump-input-on-failure
 // Verify the generic form can be parsed.
 // RUN: mlir-opt -mlir-print-op-generic %s | mlir-opt | FileCheck %s --dump-input-on-failure
 
 func @std_for(%arg0 : index, %arg1 : index, %arg2 : index) {
   loop.for %i0 = %arg0 to %arg1 step %arg2 {
     loop.for %i1 = %arg0 to %arg1 step %arg2 {
       %min_cmp = cmpi "slt", %i0, %i1 : index
       %min = select %min_cmp, %i0, %i1 : index
       %max_cmp = cmpi "sge", %i0, %i1 : index
       %max = select %max_cmp, %i0, %i1 : index
       loop.for %i2 = %min to %max step %i1 {
       }
     }
   }
   return
 }
 // CHECK-LABEL: func @std_for(
 //  CHECK-NEXT:   loop.for %{{.*}} = %{{.*}} to %{{.*}} step %{{.*}} {
 //  CHECK-NEXT:     loop.for %{{.*}} = %{{.*}} to %{{.*}} step %{{.*}} {
 //  CHECK-NEXT:       %{{.*}} = cmpi "slt", %{{.*}}, %{{.*}} : index
 //  CHECK-NEXT:       %{{.*}} = select %{{.*}}, %{{.*}}, %{{.*}} : index
 //  CHECK-NEXT:       %{{.*}} = cmpi "sge", %{{.*}}, %{{.*}} : index
 //  CHECK-NEXT:       %{{.*}} = select %{{.*}}, %{{.*}}, %{{.*}} : index
 //  CHECK-NEXT:       loop.for %{{.*}} = %{{.*}} to %{{.*}} step %{{.*}} {
 
 func @std_if(%arg0: i1, %arg1: f32) {
   loop.if %arg0 {
     %0 = addf %arg1, %arg1 : f32
   }
   return
 }
 // CHECK-LABEL: func @std_if(
 //  CHECK-NEXT:   loop.if %{{.*}} {
 //  CHECK-NEXT:     %{{.*}} = addf %{{.*}}, %{{.*}} : f32
 
 func @std_if_else(%arg0: i1, %arg1: f32) {
   loop.if %arg0 {
     %0 = addf %arg1, %arg1 : f32
   } else {
     %1 = addf %arg1, %arg1 : f32
   }
   return
 }
 // CHECK-LABEL: func @std_if_else(
 //  CHECK-NEXT:   loop.if %{{.*}} {
 //  CHECK-NEXT:     %{{.*}} = addf %{{.*}}, %{{.*}} : f32
 //  CHECK-NEXT:   } else {
 //  CHECK-NEXT:     %{{.*}} = addf %{{.*}}, %{{.*}} : f32
 
 func @std_parallel_loop(%arg0 : index, %arg1 : index, %arg2 : index,
                         %arg3 : index, %arg4 : index) {
   %step = constant 1 : index
   loop.parallel (%i0, %i1) = (%arg0, %arg1) to (%arg2, %arg3)
                                           step (%arg4, %step) {
     %min_cmp = cmpi "slt", %i0, %i1 : index
     %min = select %min_cmp, %i0, %i1 : index
     %max_cmp = cmpi "sge", %i0, %i1 : index
     %max = select %max_cmp, %i0, %i1 : index
     %zero = constant 0.0 : f32
     %int_zero = constant 0 : i32
     %red:2 = loop.parallel (%i2) = (%min) to (%max) step (%i1)
                                       init (%zero, %int_zero) -> (f32, i32) {
       %one = constant 1.0 : f32
       loop.reduce(%one) : f32 {
         ^bb0(%lhs : f32, %rhs: f32):
           %res = addf %lhs, %rhs : f32
           loop.reduce.return %res : f32
       }
       %int_one = constant 1 : i32
       loop.reduce(%int_one) : i32 {
         ^bb0(%lhs : i32, %rhs: i32):
           %res = muli %lhs, %rhs : i32
           loop.reduce.return %res : i32
       }
     }
   }
   return
 }
 // CHECK-LABEL: func @std_parallel_loop(
 //  CHECK-SAME: %[[ARG0:[A-Za-z0-9]+]]:
 //  CHECK-SAME: %[[ARG1:[A-Za-z0-9]+]]:
 //  CHECK-SAME: %[[ARG2:[A-Za-z0-9]+]]:
 //  CHECK-SAME: %[[ARG3:[A-Za-z0-9]+]]:
 //  CHECK-SAME: %[[ARG4:[A-Za-z0-9]+]]:
 //       CHECK:   %[[STEP:.*]] = constant 1 : index
 //  CHECK-NEXT:   loop.parallel (%[[I0:.*]], %[[I1:.*]]) = (%[[ARG0]], %[[ARG1]]) to
 //       CHECK:   (%[[ARG2]], %[[ARG3]]) step (%[[ARG4]], %[[STEP]]) {
 //  CHECK-NEXT:     %[[MIN_CMP:.*]] = cmpi "slt", %[[I0]], %[[I1]] : index
 //  CHECK-NEXT:     %[[MIN:.*]] = select %[[MIN_CMP]], %[[I0]], %[[I1]] : index
 //  CHECK-NEXT:     %[[MAX_CMP:.*]] = cmpi "sge", %[[I0]], %[[I1]] : index
 //  CHECK-NEXT:     %[[MAX:.*]] = select %[[MAX_CMP]], %[[I0]], %[[I1]] : index
 //  CHECK-NEXT:     %[[ZERO:.*]] = constant 0.000000e+00 : f32
 //  CHECK-NEXT:     %[[INT_ZERO:.*]] = constant 0 : i32
 //  CHECK-NEXT:     loop.parallel (%{{.*}}) = (%[[MIN]]) to (%[[MAX]])
 //  CHECK-SAME:          step (%[[I1]])
 //  CHECK-SAME:          init (%[[ZERO]], %[[INT_ZERO]]) -> (f32, i32) {
 //  CHECK-NEXT:       %[[ONE:.*]] = constant 1.000000e+00 : f32
 //  CHECK-NEXT:       loop.reduce(%[[ONE]]) : f32 {
 //  CHECK-NEXT:       ^bb0(%[[LHS:.*]]: f32, %[[RHS:.*]]: f32):
 //  CHECK-NEXT:         %[[RES:.*]] = addf %[[LHS]], %[[RHS]] : f32
 //  CHECK-NEXT:         loop.reduce.return %[[RES]] : f32
 //  CHECK-NEXT:       }
 //  CHECK-NEXT:       %[[INT_ONE:.*]] = constant 1 : i32
 //  CHECK-NEXT:       loop.reduce(%[[INT_ONE]]) : i32 {
 //  CHECK-NEXT:       ^bb0(%[[LHS:.*]]: i32, %[[RHS:.*]]: i32):
 //  CHECK-NEXT:         %[[RES:.*]] = muli %[[LHS]], %[[RHS]] : i32
 //  CHECK-NEXT:         loop.reduce.return %[[RES]] : i32
 //  CHECK-NEXT:       }
 //  CHECK-NEXT:       loop.yield
 //  CHECK-NEXT:     }
 //  CHECK-NEXT:     loop.yield
 
 func @parallel_explicit_yield(
     %arg0: index, %arg1: index, %arg2: index) {
   loop.parallel (%i0) = (%arg0) to (%arg1) step (%arg2) {
     loop.yield
   }
   return
 }
 
 // CHECK-LABEL: func @parallel_explicit_yield(
 //  CHECK-SAME: %[[ARG0:[A-Za-z0-9]+]]:
 //  CHECK-SAME: %[[ARG1:[A-Za-z0-9]+]]:
 //  CHECK-SAME: %[[ARG2:[A-Za-z0-9]+]]:
 //  CHECK-NEXT: loop.parallel (%{{.*}}) = (%[[ARG0]]) to (%[[ARG1]]) step (%[[ARG2]])
 //  CHECK-NEXT: loop.yield
 //  CHECK-NEXT: }
 //  CHECK-NEXT: return
 //  CHECK-NEXT: }
 
 func @std_if_yield(%arg0: i1, %arg1: f32)
 {
   %x, %y = loop.if %arg0 -> (f32, f32) {
     %0 = addf %arg1, %arg1 : f32
     %1 = subf %arg1, %arg1 : f32
     loop.yield %0, %1 : f32, f32
   } else {
     %0 = subf %arg1, %arg1 : f32
     %1 = addf %arg1, %arg1 : f32
     loop.yield %0, %1 : f32, f32
   }
   return
 }
 // CHECK-LABEL: func @std_if_yield(
 //  CHECK-SAME: %[[ARG0:[A-Za-z0-9]+]]:
 //  CHECK-SAME: %[[ARG1:[A-Za-z0-9]+]]:
 //  CHECK-NEXT: %{{.*}}:2 = loop.if %[[ARG0]] -> (f32, f32) {
 //  CHECK-NEXT: %[[T1:.*]] = addf %[[ARG1]], %[[ARG1]]
 //  CHECK-NEXT: %[[T2:.*]] = subf %[[ARG1]], %[[ARG1]]
 //  CHECK-NEXT: loop.yield %[[T1]], %[[T2]] : f32, f32
 //  CHECK-NEXT: } else {
 //  CHECK-NEXT: %[[T3:.*]] = subf %[[ARG1]], %[[ARG1]]
 //  CHECK-NEXT: %[[T4:.*]] = addf %[[ARG1]], %[[ARG1]]
 //  CHECK-NEXT: loop.yield %[[T3]], %[[T4]] : f32, f32
 //  CHECK-NEXT: }
 
 func @std_for_yield(%arg0 : index, %arg1 : index, %arg2 : index) {
   %s0 = constant 0.0 : f32
   %result = loop.for %i0 = %arg0 to %arg1 step %arg2 iter_args(%si = %s0) -> (f32) {
     %sn = addf %si, %si : f32
     loop.yield %sn : f32
   }
   return
 }
 // CHECK-LABEL: func @std_for_yield(
 // CHECK-SAME: %[[ARG0:[A-Za-z0-9]+]]:
 // CHECK-SAME: %[[ARG1:[A-Za-z0-9]+]]:
 // CHECK-SAME: %[[ARG2:[A-Za-z0-9]+]]:
 // CHECK-NEXT: %[[INIT:.*]] = constant
 // CHECK-NEXT: %{{.*}} = loop.for %{{.*}} = %[[ARG0]] to %[[ARG1]] step %[[ARG2]]
 // CHECK-SAME: iter_args(%[[ITER:.*]] = %[[INIT]]) -> (f32) {
 // CHECK-NEXT: %[[NEXT:.*]] = addf %[[ITER]], %[[ITER]] : f32
 // CHECK-NEXT: loop.yield %[[NEXT]] : f32
 // CHECK-NEXT: }
 
 
 func @std_for_yield_multi(%arg0 : index, %arg1 : index, %arg2 : index) {
   %s0 = constant 0.0 : f32
   %t0 = constant 1 : i32
   %u0 = constant 1.0 : f32
   %result1:3 = loop.for %i0 = %arg0 to %arg1 step %arg2 iter_args(%si = %s0, %ti = %t0, %ui = %u0) -> (f32, i32, f32) {
     %sn = addf %si, %si : f32
     %tn = addi %ti, %ti : i32
     %un = subf %ui, %ui : f32
     loop.yield %sn, %tn, %un : f32, i32, f32
   }
   return
 }
 // CHECK-LABEL: func @std_for_yield_multi(
 // CHECK-SAME: %[[ARG0:[A-Za-z0-9]+]]:
 // CHECK-SAME: %[[ARG1:[A-Za-z0-9]+]]:
 // CHECK-SAME: %[[ARG2:[A-Za-z0-9]+]]:
 // CHECK-NEXT: %[[INIT1:.*]] = constant
 // CHECK-NEXT: %[[INIT2:.*]] = constant
 // CHECK-NEXT: %[[INIT3:.*]] = constant
 // CHECK-NEXT: %{{.*}}:3 = loop.for %{{.*}} = %[[ARG0]] to %[[ARG1]] step %[[ARG2]]
 // CHECK-SAME: iter_args(%[[ITER1:.*]] = %[[INIT1]], %[[ITER2:.*]] = %[[INIT2]], %[[ITER3:.*]] = %[[INIT3]]) -> (f32, i32, f32) {
 // CHECK-NEXT: %[[NEXT1:.*]] = addf %[[ITER1]], %[[ITER1]] : f32
 // CHECK-NEXT: %[[NEXT2:.*]] = addi %[[ITER2]], %[[ITER2]] : i32
 // CHECK-NEXT: %[[NEXT3:.*]] = subf %[[ITER3]], %[[ITER3]] : f32
 // CHECK-NEXT: loop.yield %[[NEXT1]], %[[NEXT2]], %[[NEXT3]] : f32, i32, f32
 
 
 func @conditional_reduce(%buffer: memref<1024xf32>, %lb: index, %ub: index, %step: index) -> (f32) {
   %sum_0 = constant 0.0 : f32
   %c0 = constant 0.0 : f32
   %sum = loop.for %iv = %lb to %ub step %step iter_args(%sum_iter = %sum_0) -> (f32) {
     %t = load %buffer[%iv] : memref<1024xf32>
     %cond = cmpf "ugt", %t, %c0 : f32
     %sum_next = loop.if %cond -> (f32) {
       %new_sum = addf %sum_iter, %t : f32
       loop.yield %new_sum : f32
     } else {
       loop.yield %sum_iter : f32
     }
     loop.yield %sum_next : f32
   }
   return %sum : f32
 }
 // CHECK-LABEL: func @conditional_reduce(
 //  CHECK-SAME: %[[ARG0:[A-Za-z0-9]+]]
 //  CHECK-SAME: %[[ARG1:[A-Za-z0-9]+]]
 //  CHECK-SAME: %[[ARG2:[A-Za-z0-9]+]]
 //  CHECK-SAME: %[[ARG3:[A-Za-z0-9]+]]
 //  CHECK-NEXT: %[[INIT:.*]] = constant
 //  CHECK-NEXT: %[[ZERO:.*]] = constant
 //  CHECK-NEXT: %[[RESULT:.*]] = loop.for %[[IV:.*]] = %[[ARG1]] to %[[ARG2]] step %[[ARG3]]
 //  CHECK-SAME: iter_args(%[[ITER:.*]] = %[[INIT]]) -> (f32) {
 //  CHECK-NEXT: %[[T:.*]] = load %[[ARG0]][%[[IV]]]
 //  CHECK-NEXT: %[[COND:.*]] = cmpf "ugt", %[[T]], %[[ZERO]]
 //  CHECK-NEXT: %[[IFRES:.*]] = loop.if %[[COND]] -> (f32) {
 //  CHECK-NEXT: %[[THENRES:.*]] = addf %[[ITER]], %[[T]]
 //  CHECK-NEXT: loop.yield %[[THENRES]] : f32
 //  CHECK-NEXT: } else {
 //  CHECK-NEXT: loop.yield %[[ITER]] : f32
 //  CHECK-NEXT: }
 //  CHECK-NEXT: loop.yield %[[IFRES]] : f32
 //  CHECK-NEXT: }
 //  CHECK-NEXT: return %[[RESULT]]

mlir/test/Dialect/SCF/parallel-loop-fusion.mlir

 // RUN: mlir-opt -allow-unregistered-dialect %s -pass-pipeline='func(parallel-loop-fusion)' -split-input-file | FileCheck %s --dump-input-on-failure
 
 func @fuse_empty_loops() {
   %c2 = constant 2 : index
   %c0 = constant 0 : index
   %c1 = constant 1 : index
   loop.parallel (%i, %j) = (%c0, %c0) to (%c2, %c2) step (%c1, %c1) {
     loop.yield
   }
   loop.parallel (%i, %j) = (%c0, %c0) to (%c2, %c2) step (%c1, %c1) {
     loop.yield
   }
   "xla_lhlo.terminator"() : () -> ()
 }
 // CHECK-LABEL: func @fuse_empty_loops
 // CHECK:        [[C2:%.*]] = constant 2 : index
 // CHECK:        [[C0:%.*]] = constant 0 : index
 // CHECK:        [[C1:%.*]] = constant 1 : index
 // CHECK:        loop.parallel ([[I:%.*]], [[J:%.*]]) = ([[C0]], [[C0]])
 // CHECK-SAME:       to ([[C2]], [[C2]]) step ([[C1]], [[C1]]) {
 // CHECK:          loop.yield
 // CHECK:        }
 // CHECK-NOT:    loop.parallel
 
 // -----
 
 func @fuse_two(%A: memref<2x2xf32>, %B: memref<2x2xf32>,
                     %C: memref<2x2xf32>, %result: memref<2x2xf32>) {
   %c2 = constant 2 : index
   %c0 = constant 0 : index
   %c1 = constant 1 : index
   %sum = alloc()  : memref<2x2xf32>
   loop.parallel (%i, %j) = (%c0, %c0) to (%c2, %c2) step (%c1, %c1) {
     %B_elem = load %B[%i, %j] : memref<2x2xf32>
     %C_elem = load %C[%i, %j] : memref<2x2xf32>
     %sum_elem = addf %B_elem, %C_elem : f32
     store %sum_elem, %sum[%i, %j] : memref<2x2xf32>
     loop.yield
   }
   loop.parallel (%i, %j) = (%c0, %c0) to (%c2, %c2) step (%c1, %c1) {
     %sum_elem = load %sum[%i, %j] : memref<2x2xf32>
     %A_elem = load %A[%i, %j] : memref<2x2xf32>
     %product_elem = mulf %sum_elem, %A_elem : f32
     store %product_elem, %result[%i, %j] : memref<2x2xf32>
     loop.yield
   }
   dealloc %sum : memref<2x2xf32>
   return
 }
 // CHECK-LABEL: func @fuse_two
 // CHECK-SAME:   ([[A:%.*]]: {{.*}}, [[B:%.*]]: {{.*}}, [[C:%.*]]: {{.*}},
 // CHECK-SAME:    [[RESULT:%.*]]: {{.*}}) {
 // CHECK:      [[C2:%.*]] = constant 2 : index
 // CHECK:      [[C0:%.*]] = constant 0 : index
 // CHECK:      [[C1:%.*]] = constant 1 : index
 // CHECK:      [[SUM:%.*]] = alloc()
 // CHECK:      loop.parallel ([[I:%.*]], [[J:%.*]]) = ([[C0]], [[C0]])
 // CHECK-SAME:     to ([[C2]], [[C2]]) step ([[C1]], [[C1]]) {
 // CHECK:        [[B_ELEM:%.*]] = load [[B]]{{\[}}[[I]], [[J]]]
 // CHECK:        [[C_ELEM:%.*]] = load [[C]]{{\[}}[[I]], [[J]]]
 // CHECK:        [[SUM_ELEM:%.*]] = addf [[B_ELEM]], [[C_ELEM]]
 // CHECK:        store [[SUM_ELEM]], [[SUM]]{{\[}}[[I]], [[J]]]
 // CHECK:        [[SUM_ELEM_:%.*]] = load [[SUM]]{{\[}}[[I]], [[J]]]
 // CHECK:        [[A_ELEM:%.*]] = load [[A]]{{\[}}[[I]], [[J]]]
 // CHECK:        [[PRODUCT_ELEM:%.*]] = mulf [[SUM_ELEM_]], [[A_ELEM]]
 // CHECK:        store [[PRODUCT_ELEM]], [[RESULT]]{{\[}}[[I]], [[J]]]
 // CHECK:        loop.yield
 // CHECK:      }
 // CHECK:      dealloc [[SUM]]
 
 // -----
 
 func @fuse_three(%lhs: memref<100x10xf32>, %rhs: memref<100xf32>,
                       %result: memref<100x10xf32>) {
   %c100 = constant 100 : index
   %c10 = constant 10 : index
   %c0 = constant 0 : index
   %c1 = constant 1 : index
   %broadcast_rhs = alloc() : memref<100x10xf32>
   %diff = alloc() : memref<100x10xf32>
   loop.parallel (%i, %j) = (%c0, %c0) to (%c100, %c10) step (%c1, %c1) {
     %rhs_elem = load %rhs[%i] : memref<100xf32>
     store %rhs_elem, %broadcast_rhs[%i, %j] : memref<100x10xf32>
     loop.yield
   }
   loop.parallel (%i, %j) = (%c0, %c0) to (%c100, %c10) step (%c1, %c1) {
     %lhs_elem = load %lhs[%i, %j] : memref<100x10xf32>
     %broadcast_rhs_elem = load %broadcast_rhs[%i, %j] : memref<100x10xf32>
     %diff_elem = subf %lhs_elem, %broadcast_rhs_elem : f32
     store %diff_elem, %diff[%i, %j] : memref<100x10xf32>
     loop.yield
   }
   loop.parallel (%i, %j) = (%c0, %c0) to (%c100, %c10) step (%c1, %c1) {
     %diff_elem = load %diff[%i, %j] : memref<100x10xf32>
     %exp_elem = exp %diff_elem : f32
     store %exp_elem, %result[%i, %j] : memref<100x10xf32>
     loop.yield
   }
   dealloc %broadcast_rhs : memref<100x10xf32>
   dealloc %diff : memref<100x10xf32>
   return
 }
 // CHECK-LABEL: func @fuse_three
 // CHECK-SAME: ([[LHS:%.*]]: memref<100x10xf32>, [[RHS:%.*]]: memref<100xf32>,
 // CHECK-SAME:  [[RESULT:%.*]]: memref<100x10xf32>) {
 // CHECK:      [[C100:%.*]] = constant 100 : index
 // CHECK:      [[C10:%.*]] = constant 10 : index
 // CHECK:      [[C0:%.*]] = constant 0 : index
 // CHECK:      [[C1:%.*]] = constant 1 : index
 // CHECK:      [[BROADCAST_RHS:%.*]] = alloc()
 // CHECK:      [[DIFF:%.*]] = alloc()
 // CHECK:      loop.parallel ([[I:%.*]], [[J:%.*]]) = ([[C0]], [[C0]])
 // CHECK-SAME:     to ([[C100]], [[C10]]) step ([[C1]], [[C1]]) {
 // CHECK:        [[RHS_ELEM:%.*]] = load [[RHS]]{{\[}}[[I]]]
 // CHECK:        store [[RHS_ELEM]], [[BROADCAST_RHS]]{{\[}}[[I]], [[J]]]
 // CHECK:        [[LHS_ELEM:%.*]] = load [[LHS]]{{\[}}[[I]], [[J]]]
 // CHECK:        [[BROADCAST_RHS_ELEM:%.*]] = load [[BROADCAST_RHS]]
 // CHECK:        [[DIFF_ELEM:%.*]] = subf [[LHS_ELEM]], [[BROADCAST_RHS_ELEM]]
 // CHECK:        store [[DIFF_ELEM]], [[DIFF]]{{\[}}[[I]], [[J]]]
 // CHECK:        [[DIFF_ELEM_:%.*]] = load [[DIFF]]{{\[}}[[I]], [[J]]]
 // CHECK:        [[EXP_ELEM:%.*]] = exp [[DIFF_ELEM_]]
 // CHECK:        store [[EXP_ELEM]], [[RESULT]]{{\[}}[[I]], [[J]]]
 // CHECK:        loop.yield
 // CHECK:      }
 // CHECK:      dealloc [[BROADCAST_RHS]]
 // CHECK:      dealloc [[DIFF]]
 
 // -----
 
 func @do_not_fuse_nested_ploop1() {
   %c2 = constant 2 : index
   %c0 = constant 0 : index
   %c1 = constant 1 : index
   loop.parallel (%i, %j) = (%c0, %c0) to (%c2, %c2) step (%c1, %c1) {
     loop.parallel (%k, %l) = (%c0, %c0) to (%c2, %c2) step (%c1, %c1) {
       loop.yield
     }
     loop.yield
   }
   loop.parallel (%i, %j) = (%c0, %c0) to (%c2, %c2) step (%c1, %c1) {
     loop.yield
   }
   "xla_lhlo.terminator"() : () -> ()
 }
 // CHECK-LABEL: func @do_not_fuse_nested_ploop1
 // CHECK:        loop.parallel
 // CHECK:          loop.parallel
 // CHECK:        loop.parallel
 
 // -----
 
 func @do_not_fuse_nested_ploop2() {
   %c2 = constant 2 : index
   %c0 = constant 0 : index
   %c1 = constant 1 : index
   loop.parallel (%i, %j) = (%c0, %c0) to (%c2, %c2) step (%c1, %c1) {
     loop.yield
   }
   loop.parallel (%i, %j) = (%c0, %c0) to (%c2, %c2) step (%c1, %c1) {
     loop.parallel (%k, %l) = (%c0, %c0) to (%c2, %c2) step (%c1, %c1) {
       loop.yield
     }
     loop.yield
   }
   "xla_lhlo.terminator"() : () -> ()
 }
 // CHECK-LABEL: func @do_not_fuse_nested_ploop2
 // CHECK:        loop.parallel
 // CHECK:        loop.parallel
 // CHECK:          loop.parallel
 
 // -----
 
 func @do_not_fuse_loops_unmatching_num_loops() {
   %c2 = constant 2 : index
   %c0 = constant 0 : index
   %c1 = constant 1 : index
   loop.parallel (%i, %j) = (%c0, %c0) to (%c2, %c2) step (%c1, %c1) {
     loop.yield
   }
   loop.parallel (%i) = (%c0) to (%c2) step (%c1) {
     loop.yield
   }
   "xla_lhlo.terminator"() : () -> ()
 }
 // CHECK-LABEL: func @do_not_fuse_loops_unmatching_num_loops
 // CHECK:        loop.parallel
 // CHECK:        loop.parallel
 
 // -----
 
 func @do_not_fuse_loops_with_side_effecting_ops_in_between() {
   %c2 = constant 2 : index
   %c0 = constant 0 : index
   %c1 = constant 1 : index
   loop.parallel (%i, %j) = (%c0, %c0) to (%c2, %c2) step (%c1, %c1) {
     loop.yield
   }
   %buffer  = alloc() : memref<2x2xf32>
   loop.parallel (%i, %j) = (%c0, %c0) to (%c2, %c2) step (%c1, %c1) {
     loop.yield
   }
   "xla_lhlo.terminator"() : () -> ()
 }
 // CHECK-LABEL: func @do_not_fuse_loops_with_side_effecting_ops_in_between
 // CHECK:        loop.parallel
 // CHECK:        loop.parallel
 
 // -----
 
 func @do_not_fuse_loops_unmatching_iteration_space() {
   %c0 = constant 0 : index
   %c1 = constant 1 : index
   %c2 = constant 2 : index
   %c4 = constant 4 : index
   loop.parallel (%i, %j) = (%c0, %c0) to (%c4, %c4) step (%c2, %c2) {
     loop.yield
   }
   loop.parallel (%i, %j) = (%c0, %c0) to (%c2, %c2) step (%c1, %c1) {
     loop.yield
   }
   "xla_lhlo.terminator"() : () -> ()
 }
 // CHECK-LABEL: func @do_not_fuse_loops_unmatching_iteration_space
 // CHECK:        loop.parallel
 // CHECK:        loop.parallel
 
 // -----
 
 func @do_not_fuse_unmatching_write_read_patterns(
     %A: memref<2x2xf32>, %B: memref<2x2xf32>,
     %C: memref<2x2xf32>, %result: memref<2x2xf32>) {
   %c2 = constant 2 : index
   %c0 = constant 0 : index
   %c1 = constant 1 : index
   %common_buf = alloc() : memref<2x2xf32>
   loop.parallel (%i, %j) = (%c0, %c0) to (%c2, %c2) step (%c1, %c1) {
     %B_elem = load %B[%i, %j] : memref<2x2xf32>
     %C_elem = load %C[%i, %j] : memref<2x2xf32>
     %sum_elem = addf %B_elem, %C_elem : f32
     store %sum_elem, %common_buf[%i, %j] : memref<2x2xf32>
     loop.yield
   }
   loop.parallel (%i, %j) = (%c0, %c0) to (%c2, %c2) step (%c1, %c1) {
     %k = addi %i, %c1 : index
     %sum_elem = load %common_buf[%k, %j] : memref<2x2xf32>
     %A_elem = load %A[%i, %j] : memref<2x2xf32>
     %product_elem = mulf %sum_elem, %A_elem : f32
     store %product_elem, %result[%i, %j] : memref<2x2xf32>
     loop.yield
   }
   dealloc %common_buf : memref<2x2xf32>
   return
 }
 // CHECK-LABEL: func @do_not_fuse_unmatching_write_read_patterns
 // CHECK:        loop.parallel
 // CHECK:        loop.parallel
 
 // -----
 
 func @do_not_fuse_unmatching_read_write_patterns(
     %A: memref<2x2xf32>, %B: memref<2x2xf32>, %common_buf: memref<2x2xf32>) {
   %c2 = constant 2 : index
   %c0 = constant 0 : index
   %c1 = constant 1 : index
   %sum = alloc() : memref<2x2xf32>
   loop.parallel (%i, %j) = (%c0, %c0) to (%c2, %c2) step (%c1, %c1) {
     %B_elem = load %B[%i, %j] : memref<2x2xf32>
     %C_elem = load %common_buf[%i, %j] : memref<2x2xf32>
     %sum_elem = addf %B_elem, %C_elem : f32
     store %sum_elem, %sum[%i, %j] : memref<2x2xf32>
     loop.yield
   }
   loop.parallel (%i, %j) = (%c0, %c0) to (%c2, %c2) step (%c1, %c1) {
     %k = addi %i, %c1 : index
     %sum_elem = load %sum[%k, %j] : memref<2x2xf32>
     %A_elem = load %A[%i, %j] : memref<2x2xf32>
     %product_elem = mulf %sum_elem, %A_elem : f32
     store %product_elem, %common_buf[%j, %i] : memref<2x2xf32>
     loop.yield
   }
   dealloc %sum : memref<2x2xf32>
   return
 }
 // CHECK-LABEL: func @do_not_fuse_unmatching_read_write_patterns
 // CHECK:        loop.parallel
 // CHECK:        loop.parallel
 
 // -----
 
 func @do_not_fuse_loops_with_memref_defined_in_loop_bodies() {
   %c2 = constant 2 : index
   %c0 = constant 0 : index
   %c1 = constant 1 : index
   %buffer  = alloc() : memref<2x2xf32>
   loop.parallel (%i, %j) = (%c0, %c0) to (%c2, %c2) step (%c1, %c1) {
     loop.yield
   }
   loop.parallel (%i, %j) = (%c0, %c0) to (%c2, %c2) step (%c1, %c1) {
     %A = subview %buffer[%c0, %c0][%c2, %c2][%c1, %c1]
       : memref<2x2xf32> to memref<?x?xf32, offset: ?, strides:[?, ?]>
     %A_elem = load %A[%i, %j] : memref<?x?xf32, offset: ?, strides:[?, ?]>
     loop.yield
   }
   "xla_lhlo.terminator"() : () -> ()
 }
 // CHECK-LABEL: func @do_not_fuse_loops_with_memref_defined_in_loop_bodies
 // CHECK:        loop.parallel
 // CHECK:        loop.parallel

mlir/test/Dialect/SCF/parallel-loop-specialization.mlir

 // RUN: mlir-opt %s -parallel-loop-specialization -split-input-file | FileCheck %s --dump-input-on-failure
 
 #map0 = affine_map<()[s0, s1] -> (1024, s0 - s1)>
 #map1 = affine_map<()[s0, s1] -> (64, s0 - s1)>
 
 func @parallel_loop(%outer_i0: index, %outer_i1: index, %A: memref<?x?xf32>, %B: memref<?x?xf32>,
                     %C: memref<?x?xf32>, %result: memref<?x?xf32>) {
   %c0 = constant 0 : index
   %c1 = constant 1 : index
   %d0 = dim %A, 0 : memref<?x?xf32>
   %d1 = dim %A, 1 : memref<?x?xf32>
   %b0 = affine.min #map0()[%d0, %outer_i0]
   %b1 = affine.min #map1()[%d1, %outer_i1]
   loop.parallel (%i0, %i1) = (%c0, %c0) to (%b0, %b1) step (%c1, %c1) {
     %B_elem = load %B[%i0, %i1] : memref<?x?xf32>
     %C_elem = load %C[%i0, %i1] : memref<?x?xf32>
     %sum_elem = addf %B_elem, %C_elem : f32
     store %sum_elem, %result[%i0, %i1] : memref<?x?xf32>
   }
   return
 }
 
 // CHECK-LABEL:   func @parallel_loop(
 // CHECK-SAME:                        [[VAL_0:%.*]]: index, [[VAL_1:%.*]]: index, [[VAL_2:%.*]]: memref<?x?xf32>, [[VAL_3:%.*]]: memref<?x?xf32>, [[VAL_4:%.*]]: memref<?x?xf32>, [[VAL_5:%.*]]: memref<?x?xf32>) {
 // CHECK:           [[VAL_6:%.*]] = constant 0 : index
 // CHECK:           [[VAL_7:%.*]] = constant 1 : index
 // CHECK:           [[VAL_8:%.*]] = dim [[VAL_2]], 0 : memref<?x?xf32>
 // CHECK:           [[VAL_9:%.*]] = dim [[VAL_2]], 1 : memref<?x?xf32>
 // CHECK:           [[VAL_10:%.*]] = affine.min #map0(){{\[}}[[VAL_8]], [[VAL_0]]]
 // CHECK:           [[VAL_11:%.*]] = affine.min #map1(){{\[}}[[VAL_9]], [[VAL_1]]]
 // CHECK:           [[VAL_12:%.*]] = constant 1024 : index
 // CHECK:           [[VAL_13:%.*]] = cmpi "eq", [[VAL_10]], [[VAL_12]] : index
 // CHECK:           [[VAL_14:%.*]] = constant 64 : index
 // CHECK:           [[VAL_15:%.*]] = cmpi "eq", [[VAL_11]], [[VAL_14]] : index
 // CHECK:           [[VAL_16:%.*]] = and [[VAL_13]], [[VAL_15]] : i1
 // CHECK:           loop.if [[VAL_16]] {
 // CHECK:             loop.parallel ([[VAL_17:%.*]], [[VAL_18:%.*]]) = ([[VAL_6]], [[VAL_6]]) to ([[VAL_12]], [[VAL_14]]) step ([[VAL_7]], [[VAL_7]]) {
 // CHECK:               store
 // CHECK:             }
 // CHECK:           } else {
 // CHECK:             loop.parallel ([[VAL_22:%.*]], [[VAL_23:%.*]]) = ([[VAL_6]], [[VAL_6]]) to ([[VAL_10]], [[VAL_11]]) step ([[VAL_7]], [[VAL_7]]) {
 // CHECK:               store
 // CHECK:             }
 // CHECK:           }
 // CHECK:           return
 // CHECK:         }

mlir/test/Dialect/SCF/parallel-loop-tiling.mlir

 // RUN: mlir-opt %s -pass-pipeline='func(parallel-loop-tiling{parallel-loop-tile-sizes=1,4})' -split-input-file | FileCheck %s --dump-input-on-failure
 
 func @parallel_loop(%arg0 : index, %arg1 : index, %arg2 : index,
                     %arg3 : index, %arg4 : index, %arg5 : index,
         %A: memref<?x?xf32>, %B: memref<?x?xf32>,
                     %C: memref<?x?xf32>, %result: memref<?x?xf32>) {
   loop.parallel (%i0, %i1) = (%arg0, %arg1) to (%arg2, %arg3) step (%arg4, %arg5) {
     %B_elem = load %B[%i0, %i1] : memref<?x?xf32>
     %C_elem = load %C[%i0, %i1] : memref<?x?xf32>
     %sum_elem = addf %B_elem, %C_elem : f32
     store %sum_elem, %result[%i0, %i1] : memref<?x?xf32>
   }
   return
 }
 
 // CHECK:       #map0 = affine_map<(d0, d1, d2) -> (d0, d1 - d2)>
 // CHECK-LABEL:   func @parallel_loop(
 // CHECK-SAME:                        [[VAL_0:%.*]]: index, [[VAL_1:%.*]]: index, [[VAL_2:%.*]]: index, [[VAL_3:%.*]]: index, [[VAL_4:%.*]]: index, [[VAL_5:%.*]]: index, [[VAL_6:%.*]]: memref<?x?xf32>, [[VAL_7:%.*]]: memref<?x?xf32>, [[VAL_8:%.*]]: memref<?x?xf32>, [[VAL_9:%.*]]: memref<?x?xf32>) {
 // CHECK:           [[VAL_10:%.*]] = constant 0 : index
 // CHECK:           [[VAL_11:%.*]] = constant 1 : index
 // CHECK:           [[VAL_12:%.*]] = constant 4 : index
 // CHECK:           [[VAL_13:%.*]] = muli [[VAL_4]], [[VAL_11]] : index
 // CHECK:           [[VAL_14:%.*]] = muli [[VAL_5]], [[VAL_12]] : index
 // CHECK:           loop.parallel ([[VAL_15:%.*]], [[VAL_16:%.*]]) = ([[VAL_0]], [[VAL_1]]) to ([[VAL_2]], [[VAL_3]]) step ([[VAL_13]], [[VAL_14]]) {
 // CHECK:             [[VAL_17:%.*]] = affine.min #map0([[VAL_11]], [[VAL_2]], [[VAL_15]])
 // CHECK:             [[VAL_18:%.*]] = affine.min #map0([[VAL_12]], [[VAL_3]], [[VAL_16]])
 // CHECK:             loop.parallel ([[VAL_19:%.*]], [[VAL_20:%.*]]) = ([[VAL_10]], [[VAL_10]]) to ([[VAL_17]], [[VAL_18]]) step ([[VAL_4]], [[VAL_5]]) {
 // CHECK:               [[VAL_21:%.*]] = load [[VAL_7]]{{\[}}[[VAL_19]], [[VAL_20]]] : memref<?x?xf32>
 // CHECK:               [[VAL_22:%.*]] = load [[VAL_8]]{{\[}}[[VAL_19]], [[VAL_20]]] : memref<?x?xf32>
 // CHECK:               [[VAL_23:%.*]] = addf [[VAL_21]], [[VAL_22]] : f32
 // CHECK:               store [[VAL_23]], [[VAL_9]]{{\[}}[[VAL_19]], [[VAL_20]]] : memref<?x?xf32>
 // CHECK:             }
 // CHECK:           }
 // CHECK:           return
 
 // -----
 
 func @tile_nested_innermost() {
   %c2 = constant 2 : index
   %c0 = constant 0 : index
   %c1 = constant 1 : index
   loop.parallel (%i, %j) = (%c0, %c0) to (%c2, %c2) step (%c1, %c1) {
     loop.parallel (%k, %l) = (%c0, %c0) to (%c2, %c2) step (%c1, %c1) {
     }
   }
   loop.parallel (%i, %j) = (%c0, %c0) to (%c2, %c2) step (%c1, %c1) {
   }
   return
 }
 
 // CHECK-LABEL:   func @tile_nested_innermost() {
 // CHECK:           [[VAL_24:%.*]] = constant 2 : index
 // CHECK:           [[VAL_25:%.*]] = constant 0 : index
 // CHECK:           [[VAL_26:%.*]] = constant 1 : index
 // CHECK:           loop.parallel ([[VAL_27:%.*]], [[VAL_28:%.*]]) = ([[VAL_25]], [[VAL_25]]) to ([[VAL_24]], [[VAL_24]]) step ([[VAL_26]], [[VAL_26]]) {
 // CHECK:             [[VAL_29:%.*]] = constant 0 : index
 // CHECK:             [[VAL_30:%.*]] = constant 1 : index
 // CHECK:             [[VAL_31:%.*]] = constant 4 : index
 // CHECK:             [[VAL_32:%.*]] = muli [[VAL_26]], [[VAL_30]] : index
 // CHECK:             [[VAL_33:%.*]] = muli [[VAL_26]], [[VAL_31]] : index
 // CHECK:             loop.parallel ([[VAL_34:%.*]], [[VAL_35:%.*]]) = ([[VAL_25]], [[VAL_25]]) to ([[VAL_24]], [[VAL_24]]) step ([[VAL_32]], [[VAL_33]]) {
 // CHECK:               [[VAL_36:%.*]] = affine.min #map0([[VAL_30]], [[VAL_24]], [[VAL_34]])
 // CHECK:               [[VAL_37:%.*]] = affine.min #map0([[VAL_31]], [[VAL_24]], [[VAL_35]])
 // CHECK:               loop.parallel ([[VAL_38:%.*]], [[VAL_39:%.*]]) = ([[VAL_29]], [[VAL_29]]) to ([[VAL_36]], [[VAL_37]]) step ([[VAL_26]], [[VAL_26]]) {
 // CHECK:               }
 // CHECK:             }
 // CHECK:           }
 // CHECK:           [[VAL_40:%.*]] = constant 0 : index
 // CHECK:           [[VAL_41:%.*]] = constant 1 : index
 // CHECK:           [[VAL_42:%.*]] = constant 4 : index
 // CHECK:           [[VAL_43:%.*]] = muli [[VAL_26]], [[VAL_41]] : index
 // CHECK:           [[VAL_44:%.*]] = muli [[VAL_26]], [[VAL_42]] : index
 // CHECK:           loop.parallel ([[VAL_45:%.*]], [[VAL_46:%.*]]) = ([[VAL_25]], [[VAL_25]]) to ([[VAL_24]], [[VAL_24]]) step ([[VAL_43]], [[VAL_44]]) {
 // CHECK:             [[VAL_47:%.*]] = affine.min #map0([[VAL_41]], [[VAL_24]], [[VAL_45]])
 // CHECK:             [[VAL_48:%.*]] = affine.min #map0([[VAL_42]], [[VAL_24]], [[VAL_46]])
 // CHECK:             loop.parallel ([[VAL_49:%.*]], [[VAL_50:%.*]]) = ([[VAL_40]], [[VAL_40]]) to ([[VAL_47]], [[VAL_48]]) step ([[VAL_26]], [[VAL_26]]) {
 // CHECK:             }
 // CHECK:           }
 // CHECK:           return
 // CHECK:         }

mlir/test/EDSC/CMakeLists.txt

 add_llvm_executable(mlir-edsc-builder-api-test
   builder-api-test.cpp
 )
 
 llvm_update_compile_flags(mlir-edsc-builder-api-test)
 
 target_link_libraries(mlir-edsc-builder-api-test
   PRIVATE
   MLIRAffineOps
   MLIRAffineEDSC
   MLIREDSC
   MLIRIR
   MLIRLinalgEDSC
   MLIRLinalgOps
-  MLIRLoopOps
+  MLIRSCF
   MLIRStandardOps
   MLIRTransforms
   MLIRVector
   LLVMCore
   LLVMSupport
 )
 
 target_include_directories(mlir-edsc-builder-api-test PRIVATE ..)

mlir/test/EDSC/builder-api-test.cpp

 //===- builder-api-test.cpp - Tests for Declarative Builder APIs ----------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 // RUN: mlir-edsc-builder-api-test | FileCheck %s -dump-input-on-failure
 
 #include "mlir/Dialect/Affine/EDSC/Intrinsics.h"
 #include "mlir/Dialect/Linalg/EDSC/Builders.h"
 #include "mlir/Dialect/Linalg/EDSC/Intrinsics.h"
-#include "mlir/Dialect/LoopOps/EDSC/Intrinsics.h"
+#include "mlir/Dialect/SCF/EDSC/Intrinsics.h"
 #include "mlir/Dialect/StandardOps/EDSC/Intrinsics.h"
 #include "mlir/Dialect/Vector/EDSC/Intrinsics.h"
 #include "mlir/EDSC/Builders.h"
 #include "mlir/IR/AffineExpr.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/IntegerSet.h"
 #include "mlir/IR/MLIRContext.h"
 #include "mlir/IR/Module.h"

 using namespace mlir;
 using namespace mlir::edsc;
 using namespace mlir::edsc::intrinsics;
 
 static MLIRContext &globalContext() {
   static bool init_once = []() {
     registerDialect<AffineDialect>();
     registerDialect<linalg::LinalgDialect>();
-    registerDialect<loop::LoopOpsDialect>();
+    registerDialect<scf::SCFDialect>();
     registerDialect<StandardOpsDialect>();
     registerDialect<vector::VectorDialect>();
     return true;
   }();
   (void)init_once;
   static thread_local MLIRContext context;
   context.allowUnregisteredDialects();
   return context;

mlir/test/lib/Transforms/CMakeLists.txt

target_link_libraries(MLIRTestTransforms
   PUBLIC
   MLIRAffineOps
   MLIRAnalysis
   MLIREDSC
   MLIRGPU
   MLIRGPUtoCUDATransforms
   MLIRLinalgOps
   MLIRLinalgTransforms
-  MLIRLoopOps
+  MLIRSCF
   MLIRGPU
   MLIRPass
   MLIRStandardToStandard
   MLIRTestDialect
   MLIRTransformUtils
   MLIRVectorToLoops
   MLIRVector
   )

mlir/test/lib/Transforms/TestLoopMapping.cpp

 //===- TestLoopMapping.cpp --- Parametric loop mapping pass ---------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements a pass to parametrically map loop.for loops to virtual
 // processing element dimensions.
 //
 //===----------------------------------------------------------------------===//
 
-#include "mlir/Dialect/LoopOps/LoopOps.h"
+#include "mlir/Dialect/SCF/SCF.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Transforms/LoopUtils.h"
 #include "mlir/Transforms/Passes.h"
 
 #include "llvm/ADT/SetVector.h"
 
 using namespace mlir;
void runOnFunction() override {
     SmallVector<Value, 8> processorIds, numProcessors;
     func.walk([&processorIds, &numProcessors](Operation *op) {
       if (op->getName().getStringRef() != "new_processor_id_and_range")
         return;
       processorIds.push_back(op->getResult(0));
       numProcessors.push_back(op->getResult(1));
     });
 
-    func.walk([&processorIds, &numProcessors](loop::ForOp op) {
+    func.walk([&processorIds, &numProcessors](scf::ForOp op) {
       // Ignore nested loops.
-      if (op.getParentRegion()->getParentOfType<loop::ForOp>())
+      if (op.getParentRegion()->getParentOfType<scf::ForOp>())
         return;
       mapLoopToProcessorIds(op, processorIds, numProcessors);
     });
   }
 };
 } // end namespace
 
 namespace mlir {
 void registerTestLoopMappingPass() {
   PassRegistration<TestLoopMappingPass>(
       "test-mapping-to-processing-elements",
       "test mapping a single loop on a virtual processor grid");
 }
 } // namespace mlir

mlir/test/lib/Transforms/TestLoopParametricTiling.cpp

 //===- TestLoopParametricTiling.cpp --- Parametric loop tiling pass -------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements a pass to parametrically tile nests of standard loops.
 //
 //===----------------------------------------------------------------------===//
 
-#include "mlir/Dialect/LoopOps/LoopOps.h"
+#include "mlir/Dialect/SCF/SCF.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Transforms/LoopUtils.h"
 #include "mlir/Transforms/Passes.h"
 
 using namespace mlir;
 
 namespace {
 
 // Extracts fixed-range loops for top-level loop nests with ranges defined in
 // the pass constructor.  Assumes loops are permutable.
 class SimpleParametricLoopTilingPass
     : public PassWrapper<SimpleParametricLoopTilingPass, FunctionPass> {
 public:
   SimpleParametricLoopTilingPass() = default;
   SimpleParametricLoopTilingPass(const SimpleParametricLoopTilingPass &) {}
   explicit SimpleParametricLoopTilingPass(ArrayRef<int64_t> outerLoopSizes) {
     sizes = outerLoopSizes;
   }
 
   void runOnFunction() override {
     FuncOp func = getFunction();
-    func.walk([this](loop::ForOp op) {
+    func.walk([this](scf::ForOp op) {
       // Ignore nested loops.
-      if (op.getParentRegion()->getParentOfType<loop::ForOp>())
+      if (op.getParentRegion()->getParentOfType<scf::ForOp>())
         return;
       extractFixedOuterLoops(op, sizes);
     });
   }
 
   ListOption<int64_t> sizes{
       *this, "test-outer-loop-sizes", llvm::cl::MiscFlags::CommaSeparated,
       llvm::cl::desc(