Skip to content

[NFC][MLIR] Add {} for else when if body has {} #139422

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom
Merged
merged 1 commit into from
May 12, 2025
Merged

[NFC][MLIR] Add {} for else when if body has {} #139422

merged 1 commit into from
May 12, 2025

Conversation

jurahul
Copy link
Contributor

@jurahul jurahul commented May 11, 2025

No description provided.

@jurahul jurahul marked this pull request as ready for review May 11, 2025 19:33
@llvmbot
Copy link
Member

llvmbot commented May 11, 2025 •
edited
Loading

@llvm/pr-subscribers-mlir-affine
@llvm/pr-subscribers-mlir-scf
@llvm/pr-subscribers-mlir-execution-engine
@llvm/pr-subscribers-mlir-presburger
@llvm/pr-subscribers-mlir-core

@llvm/pr-subscribers-mlir-llvm

Author: Rahul Joshi (jurahul)

Changes

Full diff: https://github.com/llvm/llvm-project/pull/139422.diff

16 Files Affected:

  • (modified) mlir/lib/Analysis/Presburger/IntegerRelation.cpp (+2-1)
  • (modified) mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp (+2-1)
  • (modified) mlir/lib/Dialect/Affine/Analysis/LoopAnalysis.cpp (+2-1)
  • (modified) mlir/lib/Dialect/Linalg/IR/LinalgOps.cpp (+2-1)
  • (modified) mlir/lib/Dialect/Mesh/Interfaces/ShardingInterface.cpp (+2-1)
  • (modified) mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp (+2-1)
  • (modified) mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp (+2-1)
  • (modified) mlir/lib/Dialect/SCF/Transforms/LoopPipelining.cpp (+2-1)
  • (modified) mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp (+2-1)
  • (modified) mlir/lib/Dialect/XeGPU/Transforms/XeGPUSubgroupDistribute.cpp (+2-1)
  • (modified) mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp (+2-1)
  • (modified) mlir/lib/Target/LLVM/NVVM/Target.cpp (+2-1)
  • (modified) mlir/lib/Target/LLVMIR/ModuleImport.cpp (+2-1)
  • (modified) mlir/lib/Transforms/OpStats.cpp (+2-2)
  • (modified) mlir/test/Dialect/SCF/canonicalize.mlir (+4-4)
  • (modified) mlir/test/lib/Dialect/Test/TestOpDefs.cpp (+2-1)
diff --git a/mlir/lib/Analysis/Presburger/IntegerRelation.cpp b/mlir/lib/Analysis/Presburger/IntegerRelation.cpp
index 097cb9c2201aa..e1fd20fe48009 100644
--- a/mlir/lib/Analysis/Presburger/IntegerRelation.cpp
+++ b/mlir/lib/Analysis/Presburger/IntegerRelation.cpp
@@ -2398,8 +2398,9 @@ bool IntegerRelation::removeDuplicateConstraints() {
       addEquality(getInequality(k));
       removeInequality(k);
       removeInequality(l);
-    } else
+    } else {
       *this = getEmpty(getSpace());
+    }
     break;
   }
 
diff --git a/mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp b/mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp
index 5296013189b9e..400003d37bf20 100644
--- a/mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp
+++ b/mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp
@@ -896,8 +896,9 @@ class VectorReductionOpConversion
     } else if (kind == vector::CombiningKind::MAXNUMF) {
       result = createFPReductionComparisonOpLowering<LLVM::vector_reduce_fmax>(
           rewriter, loc, llvmType, operand, acc, fmf);
-    } else
+    } else {
       return failure();
+    }
 
     rewriter.replaceOp(reductionOp, result);
     return success();
diff --git a/mlir/lib/Dialect/Affine/Analysis/LoopAnalysis.cpp b/mlir/lib/Dialect/Affine/Analysis/LoopAnalysis.cpp
index 411e79171b898..fe53d03249369 100644
--- a/mlir/lib/Dialect/Affine/Analysis/LoopAnalysis.cpp
+++ b/mlir/lib/Dialect/Affine/Analysis/LoopAnalysis.cpp
@@ -233,8 +233,9 @@ std::optional<uint64_t> mlir::affine::getConstantTripCount(AffineForOp forOp) {
             std::min(*tripCount, static_cast<uint64_t>(constExpr.getValue()));
       else
         tripCount = constExpr.getValue();
-    } else
+    } else {
       return std::nullopt;
+    }
   }
   return tripCount;
 }
diff --git a/mlir/lib/Dialect/Linalg/IR/LinalgOps.cpp b/mlir/lib/Dialect/Linalg/IR/LinalgOps.cpp
index fce0751430305..33bc1fb05f25e 100644
--- a/mlir/lib/Dialect/Linalg/IR/LinalgOps.cpp
+++ b/mlir/lib/Dialect/Linalg/IR/LinalgOps.cpp
@@ -4266,8 +4266,9 @@ void ElementwiseOp::regionBuilder(ImplicitLocOpBuilder &b, Block &block,
     result = helper.buildTernaryFn(kind.ternaryFn, block.getArgument(0),
                                    block.getArgument(1), block.getArgument(2));
 
-  } else
+  } else {
     assert(false && "found unhandled category in elemwise");
+  }
 
   yields.push_back(result);
   helper.yieldOutputs(yields);
diff --git a/mlir/lib/Dialect/Mesh/Interfaces/ShardingInterface.cpp b/mlir/lib/Dialect/Mesh/Interfaces/ShardingInterface.cpp
index 80f4c5cd5afca..dca2b1a52166a 100644
--- a/mlir/lib/Dialect/Mesh/Interfaces/ShardingInterface.cpp
+++ b/mlir/lib/Dialect/Mesh/Interfaces/ShardingInterface.cpp
@@ -57,8 +57,9 @@ checkOperandAffineExprRecursively(AffineExpr expr,
     } else if (rhs.getKind() == AffineExprKind::DimId &&
                lhs.getKind() == AffineExprKind::Constant) {
       dimExpr = rhs;
-    } else
+    } else {
       return failure();
+    }
     unsigned position = cast<AffineDimExpr>(dimExpr).getPosition();
     if ((size_t)position >= seenIds.size() || seenIds[position])
       return failure();
diff --git a/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp b/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp
index eb6e710cae69e..2bf7aaa46db11 100644
--- a/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp
+++ b/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp
@@ -662,8 +662,9 @@ static ParseResult parseClauseWithRegionArgs(
                 parser.parseInteger(mapIndicesVec.emplace_back()) ||
                 parser.parseRSquare())
               return failure();
-          } else
+          } else {
             mapIndicesVec.push_back(-1);
+          }
         }
 
         return success();
diff --git a/mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp b/mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp
index 551411bb14765..57c27231f2144 100644
--- a/mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp
+++ b/mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp
@@ -594,9 +594,10 @@ transform::LoopFuseSiblingOp::apply(transform::TransformRewriter &rewriter,
   } else if (isForallWithIdenticalConfiguration(target, source)) {
     fusedLoop = fuseIndependentSiblingForallLoops(
         cast<scf::ForallOp>(target), cast<scf::ForallOp>(source), rewriter);
-  } else
+  } else {
     return emitSilenceableFailure(target->getLoc())
            << "operations cannot be fused";
+  }
 
   assert(fusedLoop && "failed to fuse operations");
 
diff --git a/mlir/lib/Dialect/SCF/Transforms/LoopPipelining.cpp b/mlir/lib/Dialect/SCF/Transforms/LoopPipelining.cpp
index 3bff1488bde3d..4aacbe739ca5d 100644
--- a/mlir/lib/Dialect/SCF/Transforms/LoopPipelining.cpp
+++ b/mlir/lib/Dialect/SCF/Transforms/LoopPipelining.cpp
@@ -417,8 +417,9 @@ scf::ForOp LoopPipelinerInternal::createKernelLoop(
                       [maxStage - defStage->second];
       assert(valueVersion);
       newLoopArg.push_back(valueVersion);
-    } else
+    } else {
       newLoopArg.push_back(forOp.getInitArgs()[retVal.index()]);
+    }
   }
   for (auto escape : crossStageValues) {
     LiverangeInfo &info = escape.second;
diff --git a/mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp b/mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp
index e172f02c00f51..ad1267381c4f2 100644
--- a/mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp
+++ b/mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp
@@ -124,8 +124,9 @@ static bool haveNoReadsAfterWriteExceptSameIndex(
                   OperationEquivalence::Flags::IgnoreLocations)) {
             return WalkResult::interrupt();
           }
-        } else
+        } else {
           return WalkResult::interrupt();
+        }
       }
     }
     return WalkResult::advance();
diff --git a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUSubgroupDistribute.cpp b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUSubgroupDistribute.cpp
index 2300d9e3bd43f..acdbe7be1358b 100644
--- a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUSubgroupDistribute.cpp
+++ b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUSubgroupDistribute.cpp
@@ -157,8 +157,9 @@ void LayoutInfo::print(raw_ostream &os) const {
     laneLayout.print(os);
     os << ", lane_data: ";
     laneData.print(os);
-  } else
+  } else {
     os << "Not assigned.";
+  }
 }
 
 LayoutInfo LayoutInfo::meet(const LayoutInfo &lhs, const LayoutInfo &rhs) {
diff --git a/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp b/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp
index c250340c38fc7..acb5d9de07510 100644
--- a/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp
+++ b/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp
@@ -65,8 +65,9 @@ static sycl::device getDefaultDevice() {
       return syclDevice;
     }
     throw std::runtime_error("getDefaultDevice failed");
-  } else
+  } else {
     return syclDevice;
+  }
 }
 
 static sycl::context getDefaultContext() {
diff --git a/mlir/lib/Target/LLVM/NVVM/Target.cpp b/mlir/lib/Target/LLVM/NVVM/Target.cpp
index 914a349696617..565f153b50c4f 100644
--- a/mlir/lib/Target/LLVM/NVVM/Target.cpp
+++ b/mlir/lib/Target/LLVM/NVVM/Target.cpp
@@ -606,9 +606,10 @@ NVPTXSerializer::compileToBinaryNVPTX(const std::string &ptxCode) {
           nvPTXCompilerGetErrorLog(compiler, log.data()));
       emitError(loc) << "NVPTX compiler invocation failed, error log: "
                      << log.data();
-    } else
+    } else {
       emitError(loc) << "NVPTX compiler invocation failed with error code: "
                      << status;
+    }
     return std::nullopt;
   }
 
diff --git a/mlir/lib/Target/LLVMIR/ModuleImport.cpp b/mlir/lib/Target/LLVMIR/ModuleImport.cpp
index 77094d4b75f38..7d7d0bb02a3d7 100644
--- a/mlir/lib/Target/LLVMIR/ModuleImport.cpp
+++ b/mlir/lib/Target/LLVMIR/ModuleImport.cpp
@@ -2680,8 +2680,9 @@ ModuleImport::convertParameterAttribute(llvm::AttributeSet llvmParamAttrs,
       const llvm::ConstantRange &value = llvmAttr.getValueAsConstantRange();
       mlirAttr = builder.getAttr<LLVM::ConstantRangeAttr>(value.getLower(),
                                                           value.getUpper());
-    } else
+    } else {
       llvm_unreachable("unexpected parameter attribute kind");
+    }
     paramAttrs.push_back(builder.getNamedAttr(mlirName, mlirAttr));
   }
 
diff --git a/mlir/lib/Transforms/OpStats.cpp b/mlir/lib/Transforms/OpStats.cpp
index 6746ed52396af..0dc3fe9569437 100644
--- a/mlir/lib/Transforms/OpStats.cpp
+++ b/mlir/lib/Transforms/OpStats.cpp
@@ -51,9 +51,9 @@ void PrintOpStatsPass::runOnOperation() {
   // Compute the operation statistics for the currently visited operation.
   getOperation()->walk(
       [&](Operation *op) { ++opCount[op->getName().getStringRef()]; });
-  if (printAsJSON) {
+  if (printAsJSON)
     printSummaryInJSON();
-  } else
+  else
     printSummary();
   markAllAnalysesPreserved();
 }
diff --git a/mlir/test/Dialect/SCF/canonicalize.mlir b/mlir/test/Dialect/SCF/canonicalize.mlir
index c18bd617216f1..5e32a3a78c032 100644
--- a/mlir/test/Dialect/SCF/canonicalize.mlir
+++ b/mlir/test/Dialect/SCF/canonicalize.mlir
@@ -149,7 +149,7 @@ func.func @one_unused(%cond: i1) -> (index) {
 // CHECK:             call @side_effect() : () -> ()
 // CHECK:             [[C1:%.*]] = "test.value1"
 // CHECK:             scf.yield [[C1]] : index
-// CHECK:           } else
+// CHECK:           } else {
 // CHECK:             [[C3:%.*]] = "test.value3"
 // CHECK:             scf.yield [[C3]] : index
 // CHECK:           }
@@ -185,12 +185,12 @@ func.func @nested_unused(%cond1: i1, %cond2: i1) -> (index) {
 // CHECK:               call @side_effect() : () -> ()
 // CHECK:               [[C1:%.*]] = "test.value1"
 // CHECK:               scf.yield [[C1]] : index
-// CHECK:             } else
+// CHECK:             } else {
 // CHECK:               [[C3:%.*]] = "test.value3"
 // CHECK:               scf.yield [[C3]] : index
 // CHECK:             }
 // CHECK:             scf.yield [[V1]] : index
-// CHECK:           } else
+// CHECK:           } else {
 // CHECK:             [[C1_2:%.*]] = "test.value1_2"
 // CHECK:             scf.yield [[C1_2]] : index
 // CHECK:           }
@@ -215,7 +215,7 @@ func.func @all_unused(%cond: i1) {
 // CHECK-LABEL:   func @all_unused
 // CHECK:           scf.if %{{.*}} {
 // CHECK:             call @side_effect() : () -> ()
-// CHECK:           } else
+// CHECK:           } else {
 // CHECK:             call @side_effect() : () -> ()
 // CHECK:           }
 // CHECK:           return
diff --git a/mlir/test/lib/Dialect/Test/TestOpDefs.cpp b/mlir/test/lib/Dialect/Test/TestOpDefs.cpp
index 454a12bac9ab3..b5a8bd10d6b68 100644
--- a/mlir/test/lib/Dialect/Test/TestOpDefs.cpp
+++ b/mlir/test/lib/Dialect/Test/TestOpDefs.cpp
@@ -801,8 +801,9 @@ void TestReflectBoundsOp::inferResultRanges(
     unsigned bitwidth = intTy.getWidth();
     sIntTy = b.getIntegerType(bitwidth, /*isSigned=*/true);
     uIntTy = b.getIntegerType(bitwidth, /*isSigned=*/false);
-  } else
+  } else {
     sIntTy = uIntTy = type;
+  }
 
   setUminAttr(b.getIntegerAttr(uIntTy, range.umin()));
   setUmaxAttr(b.getIntegerAttr(uIntTy, range.umax()));

@llvmbot
Copy link
Member

llvmbot commented May 11, 2025

@llvm/pr-subscribers-flang-openmp

Author: Rahul Joshi (jurahul)

Changes

Full diff: https://github.com/llvm/llvm-project/pull/139422.diff

16 Files Affected:

  • (modified) mlir/lib/Analysis/Presburger/IntegerRelation.cpp (+2-1)
  • (modified) mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp (+2-1)
  • (modified) mlir/lib/Dialect/Affine/Analysis/LoopAnalysis.cpp (+2-1)
  • (modified) mlir/lib/Dialect/Linalg/IR/LinalgOps.cpp (+2-1)
  • (modified) mlir/lib/Dialect/Mesh/Interfaces/ShardingInterface.cpp (+2-1)
  • (modified) mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp (+2-1)
  • (modified) mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp (+2-1)
  • (modified) mlir/lib/Dialect/SCF/Transforms/LoopPipelining.cpp (+2-1)
  • (modified) mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp (+2-1)
  • (modified) mlir/lib/Dialect/XeGPU/Transforms/XeGPUSubgroupDistribute.cpp (+2-1)
  • (modified) mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp (+2-1)
  • (modified) mlir/lib/Target/LLVM/NVVM/Target.cpp (+2-1)
  • (modified) mlir/lib/Target/LLVMIR/ModuleImport.cpp (+2-1)
  • (modified) mlir/lib/Transforms/OpStats.cpp (+2-2)
  • (modified) mlir/test/Dialect/SCF/canonicalize.mlir (+4-4)
  • (modified) mlir/test/lib/Dialect/Test/TestOpDefs.cpp (+2-1)
diff --git a/mlir/lib/Analysis/Presburger/IntegerRelation.cpp b/mlir/lib/Analysis/Presburger/IntegerRelation.cpp
index 097cb9c2201aa..e1fd20fe48009 100644
--- a/mlir/lib/Analysis/Presburger/IntegerRelation.cpp
+++ b/mlir/lib/Analysis/Presburger/IntegerRelation.cpp
@@ -2398,8 +2398,9 @@ bool IntegerRelation::removeDuplicateConstraints() {
       addEquality(getInequality(k));
       removeInequality(k);
       removeInequality(l);
-    } else
+    } else {
       *this = getEmpty(getSpace());
+    }
     break;
   }
 
diff --git a/mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp b/mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp
index 5296013189b9e..400003d37bf20 100644
--- a/mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp
+++ b/mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp
@@ -896,8 +896,9 @@ class VectorReductionOpConversion
     } else if (kind == vector::CombiningKind::MAXNUMF) {
       result = createFPReductionComparisonOpLowering<LLVM::vector_reduce_fmax>(
           rewriter, loc, llvmType, operand, acc, fmf);
-    } else
+    } else {
       return failure();
+    }
 
     rewriter.replaceOp(reductionOp, result);
     return success();
diff --git a/mlir/lib/Dialect/Affine/Analysis/LoopAnalysis.cpp b/mlir/lib/Dialect/Affine/Analysis/LoopAnalysis.cpp
index 411e79171b898..fe53d03249369 100644
--- a/mlir/lib/Dialect/Affine/Analysis/LoopAnalysis.cpp
+++ b/mlir/lib/Dialect/Affine/Analysis/LoopAnalysis.cpp
@@ -233,8 +233,9 @@ std::optional<uint64_t> mlir::affine::getConstantTripCount(AffineForOp forOp) {
             std::min(*tripCount, static_cast<uint64_t>(constExpr.getValue()));
       else
         tripCount = constExpr.getValue();
-    } else
+    } else {
       return std::nullopt;
+    }
   }
   return tripCount;
 }
diff --git a/mlir/lib/Dialect/Linalg/IR/LinalgOps.cpp b/mlir/lib/Dialect/Linalg/IR/LinalgOps.cpp
index fce0751430305..33bc1fb05f25e 100644
--- a/mlir/lib/Dialect/Linalg/IR/LinalgOps.cpp
+++ b/mlir/lib/Dialect/Linalg/IR/LinalgOps.cpp
@@ -4266,8 +4266,9 @@ void ElementwiseOp::regionBuilder(ImplicitLocOpBuilder &b, Block &block,
     result = helper.buildTernaryFn(kind.ternaryFn, block.getArgument(0),
                                    block.getArgument(1), block.getArgument(2));
 
-  } else
+  } else {
     assert(false && "found unhandled category in elemwise");
+  }
 
   yields.push_back(result);
   helper.yieldOutputs(yields);
diff --git a/mlir/lib/Dialect/Mesh/Interfaces/ShardingInterface.cpp b/mlir/lib/Dialect/Mesh/Interfaces/ShardingInterface.cpp
index 80f4c5cd5afca..dca2b1a52166a 100644
--- a/mlir/lib/Dialect/Mesh/Interfaces/ShardingInterface.cpp
+++ b/mlir/lib/Dialect/Mesh/Interfaces/ShardingInterface.cpp
@@ -57,8 +57,9 @@ checkOperandAffineExprRecursively(AffineExpr expr,
     } else if (rhs.getKind() == AffineExprKind::DimId &&
                lhs.getKind() == AffineExprKind::Constant) {
       dimExpr = rhs;
-    } else
+    } else {
       return failure();
+    }
     unsigned position = cast<AffineDimExpr>(dimExpr).getPosition();
     if ((size_t)position >= seenIds.size() || seenIds[position])
       return failure();
diff --git a/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp b/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp
index eb6e710cae69e..2bf7aaa46db11 100644
--- a/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp
+++ b/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp
@@ -662,8 +662,9 @@ static ParseResult parseClauseWithRegionArgs(
                 parser.parseInteger(mapIndicesVec.emplace_back()) ||
                 parser.parseRSquare())
               return failure();
-          } else
+          } else {
             mapIndicesVec.push_back(-1);
+          }
         }
 
         return success();
diff --git a/mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp b/mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp
index 551411bb14765..57c27231f2144 100644
--- a/mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp
+++ b/mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp
@@ -594,9 +594,10 @@ transform::LoopFuseSiblingOp::apply(transform::TransformRewriter &rewriter,
   } else if (isForallWithIdenticalConfiguration(target, source)) {
     fusedLoop = fuseIndependentSiblingForallLoops(
         cast<scf::ForallOp>(target), cast<scf::ForallOp>(source), rewriter);
-  } else
+  } else {
     return emitSilenceableFailure(target->getLoc())
            << "operations cannot be fused";
+  }
 
   assert(fusedLoop && "failed to fuse operations");
 
diff --git a/mlir/lib/Dialect/SCF/Transforms/LoopPipelining.cpp b/mlir/lib/Dialect/SCF/Transforms/LoopPipelining.cpp
index 3bff1488bde3d..4aacbe739ca5d 100644
--- a/mlir/lib/Dialect/SCF/Transforms/LoopPipelining.cpp
+++ b/mlir/lib/Dialect/SCF/Transforms/LoopPipelining.cpp
@@ -417,8 +417,9 @@ scf::ForOp LoopPipelinerInternal::createKernelLoop(
                       [maxStage - defStage->second];
       assert(valueVersion);
       newLoopArg.push_back(valueVersion);
-    } else
+    } else {
       newLoopArg.push_back(forOp.getInitArgs()[retVal.index()]);
+    }
   }
   for (auto escape : crossStageValues) {
     LiverangeInfo &info = escape.second;
diff --git a/mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp b/mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp
index e172f02c00f51..ad1267381c4f2 100644
--- a/mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp
+++ b/mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp
@@ -124,8 +124,9 @@ static bool haveNoReadsAfterWriteExceptSameIndex(
                   OperationEquivalence::Flags::IgnoreLocations)) {
             return WalkResult::interrupt();
           }
-        } else
+        } else {
           return WalkResult::interrupt();
+        }
       }
     }
     return WalkResult::advance();
diff --git a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUSubgroupDistribute.cpp b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUSubgroupDistribute.cpp
index 2300d9e3bd43f..acdbe7be1358b 100644
--- a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUSubgroupDistribute.cpp
+++ b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUSubgroupDistribute.cpp
@@ -157,8 +157,9 @@ void LayoutInfo::print(raw_ostream &os) const {
     laneLayout.print(os);
     os << ", lane_data: ";
     laneData.print(os);
-  } else
+  } else {
     os << "Not assigned.";
+  }
 }
 
 LayoutInfo LayoutInfo::meet(const LayoutInfo &lhs, const LayoutInfo &rhs) {
diff --git a/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp b/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp
index c250340c38fc7..acb5d9de07510 100644
--- a/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp
+++ b/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp
@@ -65,8 +65,9 @@ static sycl::device getDefaultDevice() {
       return syclDevice;
     }
     throw std::runtime_error("getDefaultDevice failed");
-  } else
+  } else {
     return syclDevice;
+  }
 }
 
 static sycl::context getDefaultContext() {
diff --git a/mlir/lib/Target/LLVM/NVVM/Target.cpp b/mlir/lib/Target/LLVM/NVVM/Target.cpp
index 914a349696617..565f153b50c4f 100644
--- a/mlir/lib/Target/LLVM/NVVM/Target.cpp
+++ b/mlir/lib/Target/LLVM/NVVM/Target.cpp
@@ -606,9 +606,10 @@ NVPTXSerializer::compileToBinaryNVPTX(const std::string &ptxCode) {
           nvPTXCompilerGetErrorLog(compiler, log.data()));
       emitError(loc) << "NVPTX compiler invocation failed, error log: "
                      << log.data();
-    } else
+    } else {
       emitError(loc) << "NVPTX compiler invocation failed with error code: "
                      << status;
+    }
     return std::nullopt;
   }
 
diff --git a/mlir/lib/Target/LLVMIR/ModuleImport.cpp b/mlir/lib/Target/LLVMIR/ModuleImport.cpp
index 77094d4b75f38..7d7d0bb02a3d7 100644
--- a/mlir/lib/Target/LLVMIR/ModuleImport.cpp
+++ b/mlir/lib/Target/LLVMIR/ModuleImport.cpp
@@ -2680,8 +2680,9 @@ ModuleImport::convertParameterAttribute(llvm::AttributeSet llvmParamAttrs,
       const llvm::ConstantRange &value = llvmAttr.getValueAsConstantRange();
       mlirAttr = builder.getAttr<LLVM::ConstantRangeAttr>(value.getLower(),
                                                           value.getUpper());
-    } else
+    } else {
       llvm_unreachable("unexpected parameter attribute kind");
+    }
     paramAttrs.push_back(builder.getNamedAttr(mlirName, mlirAttr));
   }
 
diff --git a/mlir/lib/Transforms/OpStats.cpp b/mlir/lib/Transforms/OpStats.cpp
index 6746ed52396af..0dc3fe9569437 100644
--- a/mlir/lib/Transforms/OpStats.cpp
+++ b/mlir/lib/Transforms/OpStats.cpp
@@ -51,9 +51,9 @@ void PrintOpStatsPass::runOnOperation() {
   // Compute the operation statistics for the currently visited operation.
   getOperation()->walk(
       [&](Operation *op) { ++opCount[op->getName().getStringRef()]; });
-  if (printAsJSON) {
+  if (printAsJSON)
     printSummaryInJSON();
-  } else
+  else
     printSummary();
   markAllAnalysesPreserved();
 }
diff --git a/mlir/test/Dialect/SCF/canonicalize.mlir b/mlir/test/Dialect/SCF/canonicalize.mlir
index c18bd617216f1..5e32a3a78c032 100644
--- a/mlir/test/Dialect/SCF/canonicalize.mlir
+++ b/mlir/test/Dialect/SCF/canonicalize.mlir
@@ -149,7 +149,7 @@ func.func @one_unused(%cond: i1) -> (index) {
 // CHECK:             call @side_effect() : () -> ()
 // CHECK:             [[C1:%.*]] = "test.value1"
 // CHECK:             scf.yield [[C1]] : index
-// CHECK:           } else
+// CHECK:           } else {
 // CHECK:             [[C3:%.*]] = "test.value3"
 // CHECK:             scf.yield [[C3]] : index
 // CHECK:           }
@@ -185,12 +185,12 @@ func.func @nested_unused(%cond1: i1, %cond2: i1) -> (index) {
 // CHECK:               call @side_effect() : () -> ()
 // CHECK:               [[C1:%.*]] = "test.value1"
 // CHECK:               scf.yield [[C1]] : index
-// CHECK:             } else
+// CHECK:             } else {
 // CHECK:               [[C3:%.*]] = "test.value3"
 // CHECK:               scf.yield [[C3]] : index
 // CHECK:             }
 // CHECK:             scf.yield [[V1]] : index
-// CHECK:           } else
+// CHECK:           } else {
 // CHECK:             [[C1_2:%.*]] = "test.value1_2"
 // CHECK:             scf.yield [[C1_2]] : index
 // CHECK:           }
@@ -215,7 +215,7 @@ func.func @all_unused(%cond: i1) {
 // CHECK-LABEL:   func @all_unused
 // CHECK:           scf.if %{{.*}} {
 // CHECK:             call @side_effect() : () -> ()
-// CHECK:           } else
+// CHECK:           } else {
 // CHECK:             call @side_effect() : () -> ()
 // CHECK:           }
 // CHECK:           return
diff --git a/mlir/test/lib/Dialect/Test/TestOpDefs.cpp b/mlir/test/lib/Dialect/Test/TestOpDefs.cpp
index 454a12bac9ab3..b5a8bd10d6b68 100644
--- a/mlir/test/lib/Dialect/Test/TestOpDefs.cpp
+++ b/mlir/test/lib/Dialect/Test/TestOpDefs.cpp
@@ -801,8 +801,9 @@ void TestReflectBoundsOp::inferResultRanges(
     unsigned bitwidth = intTy.getWidth();
     sIntTy = b.getIntegerType(bitwidth, /*isSigned=*/true);
     uIntTy = b.getIntegerType(bitwidth, /*isSigned=*/false);
-  } else
+  } else {
     sIntTy = uIntTy = type;
+  }
 
   setUminAttr(b.getIntegerAttr(uIntTy, range.umin()));
   setUmaxAttr(b.getIntegerAttr(uIntTy, range.umax()));

@llvmbot
Copy link
Member

llvmbot commented May 11, 2025

@llvm/pr-subscribers-mlir-gpu

Author: Rahul Joshi (jurahul)

Changes

Full diff: https://github.com/llvm/llvm-project/pull/139422.diff

16 Files Affected:

  • (modified) mlir/lib/Analysis/Presburger/IntegerRelation.cpp (+2-1)
  • (modified) mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp (+2-1)
  • (modified) mlir/lib/Dialect/Affine/Analysis/LoopAnalysis.cpp (+2-1)
  • (modified) mlir/lib/Dialect/Linalg/IR/LinalgOps.cpp (+2-1)
  • (modified) mlir/lib/Dialect/Mesh/Interfaces/ShardingInterface.cpp (+2-1)
  • (modified) mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp (+2-1)
  • (modified) mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp (+2-1)
  • (modified) mlir/lib/Dialect/SCF/Transforms/LoopPipelining.cpp (+2-1)
  • (modified) mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp (+2-1)
  • (modified) mlir/lib/Dialect/XeGPU/Transforms/XeGPUSubgroupDistribute.cpp (+2-1)
  • (modified) mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp (+2-1)
  • (modified) mlir/lib/Target/LLVM/NVVM/Target.cpp (+2-1)
  • (modified) mlir/lib/Target/LLVMIR/ModuleImport.cpp (+2-1)
  • (modified) mlir/lib/Transforms/OpStats.cpp (+2-2)
  • (modified) mlir/test/Dialect/SCF/canonicalize.mlir (+4-4)
  • (modified) mlir/test/lib/Dialect/Test/TestOpDefs.cpp (+2-1)
diff --git a/mlir/lib/Analysis/Presburger/IntegerRelation.cpp b/mlir/lib/Analysis/Presburger/IntegerRelation.cpp
index 097cb9c2201aa..e1fd20fe48009 100644
--- a/mlir/lib/Analysis/Presburger/IntegerRelation.cpp
+++ b/mlir/lib/Analysis/Presburger/IntegerRelation.cpp
@@ -2398,8 +2398,9 @@ bool IntegerRelation::removeDuplicateConstraints() {
       addEquality(getInequality(k));
       removeInequality(k);
       removeInequality(l);
-    } else
+    } else {
       *this = getEmpty(getSpace());
+    }
     break;
   }
 
diff --git a/mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp b/mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp
index 5296013189b9e..400003d37bf20 100644
--- a/mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp
+++ b/mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp
@@ -896,8 +896,9 @@ class VectorReductionOpConversion
     } else if (kind == vector::CombiningKind::MAXNUMF) {
       result = createFPReductionComparisonOpLowering<LLVM::vector_reduce_fmax>(
           rewriter, loc, llvmType, operand, acc, fmf);
-    } else
+    } else {
       return failure();
+    }
 
     rewriter.replaceOp(reductionOp, result);
     return success();
diff --git a/mlir/lib/Dialect/Affine/Analysis/LoopAnalysis.cpp b/mlir/lib/Dialect/Affine/Analysis/LoopAnalysis.cpp
index 411e79171b898..fe53d03249369 100644
--- a/mlir/lib/Dialect/Affine/Analysis/LoopAnalysis.cpp
+++ b/mlir/lib/Dialect/Affine/Analysis/LoopAnalysis.cpp
@@ -233,8 +233,9 @@ std::optional<uint64_t> mlir::affine::getConstantTripCount(AffineForOp forOp) {
             std::min(*tripCount, static_cast<uint64_t>(constExpr.getValue()));
       else
         tripCount = constExpr.getValue();
-    } else
+    } else {
       return std::nullopt;
+    }
   }
   return tripCount;
 }
diff --git a/mlir/lib/Dialect/Linalg/IR/LinalgOps.cpp b/mlir/lib/Dialect/Linalg/IR/LinalgOps.cpp
index fce0751430305..33bc1fb05f25e 100644
--- a/mlir/lib/Dialect/Linalg/IR/LinalgOps.cpp
+++ b/mlir/lib/Dialect/Linalg/IR/LinalgOps.cpp
@@ -4266,8 +4266,9 @@ void ElementwiseOp::regionBuilder(ImplicitLocOpBuilder &b, Block &block,
     result = helper.buildTernaryFn(kind.ternaryFn, block.getArgument(0),
                                    block.getArgument(1), block.getArgument(2));
 
-  } else
+  } else {
     assert(false && "found unhandled category in elemwise");
+  }
 
   yields.push_back(result);
   helper.yieldOutputs(yields);
diff --git a/mlir/lib/Dialect/Mesh/Interfaces/ShardingInterface.cpp b/mlir/lib/Dialect/Mesh/Interfaces/ShardingInterface.cpp
index 80f4c5cd5afca..dca2b1a52166a 100644
--- a/mlir/lib/Dialect/Mesh/Interfaces/ShardingInterface.cpp
+++ b/mlir/lib/Dialect/Mesh/Interfaces/ShardingInterface.cpp
@@ -57,8 +57,9 @@ checkOperandAffineExprRecursively(AffineExpr expr,
     } else if (rhs.getKind() == AffineExprKind::DimId &&
                lhs.getKind() == AffineExprKind::Constant) {
       dimExpr = rhs;
-    } else
+    } else {
       return failure();
+    }
     unsigned position = cast<AffineDimExpr>(dimExpr).getPosition();
     if ((size_t)position >= seenIds.size() || seenIds[position])
       return failure();
diff --git a/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp b/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp
index eb6e710cae69e..2bf7aaa46db11 100644
--- a/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp
+++ b/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp
@@ -662,8 +662,9 @@ static ParseResult parseClauseWithRegionArgs(
                 parser.parseInteger(mapIndicesVec.emplace_back()) ||
                 parser.parseRSquare())
               return failure();
-          } else
+          } else {
             mapIndicesVec.push_back(-1);
+          }
         }
 
         return success();
diff --git a/mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp b/mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp
index 551411bb14765..57c27231f2144 100644
--- a/mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp
+++ b/mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp
@@ -594,9 +594,10 @@ transform::LoopFuseSiblingOp::apply(transform::TransformRewriter &rewriter,
   } else if (isForallWithIdenticalConfiguration(target, source)) {
     fusedLoop = fuseIndependentSiblingForallLoops(
         cast<scf::ForallOp>(target), cast<scf::ForallOp>(source), rewriter);
-  } else
+  } else {
     return emitSilenceableFailure(target->getLoc())
            << "operations cannot be fused";
+  }
 
   assert(fusedLoop && "failed to fuse operations");
 
diff --git a/mlir/lib/Dialect/SCF/Transforms/LoopPipelining.cpp b/mlir/lib/Dialect/SCF/Transforms/LoopPipelining.cpp
index 3bff1488bde3d..4aacbe739ca5d 100644
--- a/mlir/lib/Dialect/SCF/Transforms/LoopPipelining.cpp
+++ b/mlir/lib/Dialect/SCF/Transforms/LoopPipelining.cpp
@@ -417,8 +417,9 @@ scf::ForOp LoopPipelinerInternal::createKernelLoop(
                       [maxStage - defStage->second];
       assert(valueVersion);
       newLoopArg.push_back(valueVersion);
-    } else
+    } else {
       newLoopArg.push_back(forOp.getInitArgs()[retVal.index()]);
+    }
   }
   for (auto escape : crossStageValues) {
     LiverangeInfo &info = escape.second;
diff --git a/mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp b/mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp
index e172f02c00f51..ad1267381c4f2 100644
--- a/mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp
+++ b/mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp
@@ -124,8 +124,9 @@ static bool haveNoReadsAfterWriteExceptSameIndex(
                   OperationEquivalence::Flags::IgnoreLocations)) {
             return WalkResult::interrupt();
           }
-        } else
+        } else {
           return WalkResult::interrupt();
+        }
       }
     }
     return WalkResult::advance();
diff --git a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUSubgroupDistribute.cpp b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUSubgroupDistribute.cpp
index 2300d9e3bd43f..acdbe7be1358b 100644
--- a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUSubgroupDistribute.cpp
+++ b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUSubgroupDistribute.cpp
@@ -157,8 +157,9 @@ void LayoutInfo::print(raw_ostream &os) const {
     laneLayout.print(os);
     os << ", lane_data: ";
     laneData.print(os);
-  } else
+  } else {
     os << "Not assigned.";
+  }
 }
 
 LayoutInfo LayoutInfo::meet(const LayoutInfo &lhs, const LayoutInfo &rhs) {
diff --git a/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp b/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp
index c250340c38fc7..acb5d9de07510 100644
--- a/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp
+++ b/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp
@@ -65,8 +65,9 @@ static sycl::device getDefaultDevice() {
       return syclDevice;
     }
     throw std::runtime_error("getDefaultDevice failed");
-  } else
+  } else {
     return syclDevice;
+  }
 }
 
 static sycl::context getDefaultContext() {
diff --git a/mlir/lib/Target/LLVM/NVVM/Target.cpp b/mlir/lib/Target/LLVM/NVVM/Target.cpp
index 914a349696617..565f153b50c4f 100644
--- a/mlir/lib/Target/LLVM/NVVM/Target.cpp
+++ b/mlir/lib/Target/LLVM/NVVM/Target.cpp
@@ -606,9 +606,10 @@ NVPTXSerializer::compileToBinaryNVPTX(const std::string &ptxCode) {
           nvPTXCompilerGetErrorLog(compiler, log.data()));
       emitError(loc) << "NVPTX compiler invocation failed, error log: "
                      << log.data();
-    } else
+    } else {
       emitError(loc) << "NVPTX compiler invocation failed with error code: "
                      << status;
+    }
     return std::nullopt;
   }
 
diff --git a/mlir/lib/Target/LLVMIR/ModuleImport.cpp b/mlir/lib/Target/LLVMIR/ModuleImport.cpp
index 77094d4b75f38..7d7d0bb02a3d7 100644
--- a/mlir/lib/Target/LLVMIR/ModuleImport.cpp
+++ b/mlir/lib/Target/LLVMIR/ModuleImport.cpp
@@ -2680,8 +2680,9 @@ ModuleImport::convertParameterAttribute(llvm::AttributeSet llvmParamAttrs,
       const llvm::ConstantRange &value = llvmAttr.getValueAsConstantRange();
       mlirAttr = builder.getAttr<LLVM::ConstantRangeAttr>(value.getLower(),
                                                           value.getUpper());
-    } else
+    } else {
       llvm_unreachable("unexpected parameter attribute kind");
+    }
     paramAttrs.push_back(builder.getNamedAttr(mlirName, mlirAttr));
   }
 
diff --git a/mlir/lib/Transforms/OpStats.cpp b/mlir/lib/Transforms/OpStats.cpp
index 6746ed52396af..0dc3fe9569437 100644
--- a/mlir/lib/Transforms/OpStats.cpp
+++ b/mlir/lib/Transforms/OpStats.cpp
@@ -51,9 +51,9 @@ void PrintOpStatsPass::runOnOperation() {
   // Compute the operation statistics for the currently visited operation.
   getOperation()->walk(
       [&](Operation *op) { ++opCount[op->getName().getStringRef()]; });
-  if (printAsJSON) {
+  if (printAsJSON)
     printSummaryInJSON();
-  } else
+  else
     printSummary();
   markAllAnalysesPreserved();
 }
diff --git a/mlir/test/Dialect/SCF/canonicalize.mlir b/mlir/test/Dialect/SCF/canonicalize.mlir
index c18bd617216f1..5e32a3a78c032 100644
--- a/mlir/test/Dialect/SCF/canonicalize.mlir
+++ b/mlir/test/Dialect/SCF/canonicalize.mlir
@@ -149,7 +149,7 @@ func.func @one_unused(%cond: i1) -> (index) {
 // CHECK:             call @side_effect() : () -> ()
 // CHECK:             [[C1:%.*]] = "test.value1"
 // CHECK:             scf.yield [[C1]] : index
-// CHECK:           } else
+// CHECK:           } else {
 // CHECK:             [[C3:%.*]] = "test.value3"
 // CHECK:             scf.yield [[C3]] : index
 // CHECK:           }
@@ -185,12 +185,12 @@ func.func @nested_unused(%cond1: i1, %cond2: i1) -> (index) {
 // CHECK:               call @side_effect() : () -> ()
 // CHECK:               [[C1:%.*]] = "test.value1"
 // CHECK:               scf.yield [[C1]] : index
-// CHECK:             } else
+// CHECK:             } else {
 // CHECK:               [[C3:%.*]] = "test.value3"
 // CHECK:               scf.yield [[C3]] : index
 // CHECK:             }
 // CHECK:             scf.yield [[V1]] : index
-// CHECK:           } else
+// CHECK:           } else {
 // CHECK:             [[C1_2:%.*]] = "test.value1_2"
 // CHECK:             scf.yield [[C1_2]] : index
 // CHECK:           }
@@ -215,7 +215,7 @@ func.func @all_unused(%cond: i1) {
 // CHECK-LABEL:   func @all_unused
 // CHECK:           scf.if %{{.*}} {
 // CHECK:             call @side_effect() : () -> ()
-// CHECK:           } else
+// CHECK:           } else {
 // CHECK:             call @side_effect() : () -> ()
 // CHECK:           }
 // CHECK:           return
diff --git a/mlir/test/lib/Dialect/Test/TestOpDefs.cpp b/mlir/test/lib/Dialect/Test/TestOpDefs.cpp
index 454a12bac9ab3..b5a8bd10d6b68 100644
--- a/mlir/test/lib/Dialect/Test/TestOpDefs.cpp
+++ b/mlir/test/lib/Dialect/Test/TestOpDefs.cpp
@@ -801,8 +801,9 @@ void TestReflectBoundsOp::inferResultRanges(
     unsigned bitwidth = intTy.getWidth();
     sIntTy = b.getIntegerType(bitwidth, /*isSigned=*/true);
     uIntTy = b.getIntegerType(bitwidth, /*isSigned=*/false);
-  } else
+  } else {
     sIntTy = uIntTy = type;
+  }
 
   setUminAttr(b.getIntegerAttr(uIntTy, range.umin()));
   setUmaxAttr(b.getIntegerAttr(uIntTy, range.umax()));

@llvmbot
Copy link
Member

llvmbot commented May 11, 2025

@llvm/pr-subscribers-mlir-linalg

Author: Rahul Joshi (jurahul)

Changes

Full diff: https://github.com/llvm/llvm-project/pull/139422.diff

16 Files Affected:

  • (modified) mlir/lib/Analysis/Presburger/IntegerRelation.cpp (+2-1)
  • (modified) mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp (+2-1)
  • (modified) mlir/lib/Dialect/Affine/Analysis/LoopAnalysis.cpp (+2-1)
  • (modified) mlir/lib/Dialect/Linalg/IR/LinalgOps.cpp (+2-1)
  • (modified) mlir/lib/Dialect/Mesh/Interfaces/ShardingInterface.cpp (+2-1)
  • (modified) mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp (+2-1)
  • (modified) mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp (+2-1)
  • (modified) mlir/lib/Dialect/SCF/Transforms/LoopPipelining.cpp (+2-1)
  • (modified) mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp (+2-1)
  • (modified) mlir/lib/Dialect/XeGPU/Transforms/XeGPUSubgroupDistribute.cpp (+2-1)
  • (modified) mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp (+2-1)
  • (modified) mlir/lib/Target/LLVM/NVVM/Target.cpp (+2-1)
  • (modified) mlir/lib/Target/LLVMIR/ModuleImport.cpp (+2-1)
  • (modified) mlir/lib/Transforms/OpStats.cpp (+2-2)
  • (modified) mlir/test/Dialect/SCF/canonicalize.mlir (+4-4)
  • (modified) mlir/test/lib/Dialect/Test/TestOpDefs.cpp (+2-1)
diff --git a/mlir/lib/Analysis/Presburger/IntegerRelation.cpp b/mlir/lib/Analysis/Presburger/IntegerRelation.cpp
index 097cb9c2201aa..e1fd20fe48009 100644
--- a/mlir/lib/Analysis/Presburger/IntegerRelation.cpp
+++ b/mlir/lib/Analysis/Presburger/IntegerRelation.cpp
@@ -2398,8 +2398,9 @@ bool IntegerRelation::removeDuplicateConstraints() {
       addEquality(getInequality(k));
       removeInequality(k);
       removeInequality(l);
-    } else
+    } else {
       *this = getEmpty(getSpace());
+    }
     break;
   }
 
diff --git a/mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp b/mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp
index 5296013189b9e..400003d37bf20 100644
--- a/mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp
+++ b/mlir/lib/Conversion/VectorToLLVM/ConvertVectorToLLVM.cpp
@@ -896,8 +896,9 @@ class VectorReductionOpConversion
     } else if (kind == vector::CombiningKind::MAXNUMF) {
       result = createFPReductionComparisonOpLowering<LLVM::vector_reduce_fmax>(
           rewriter, loc, llvmType, operand, acc, fmf);
-    } else
+    } else {
       return failure();
+    }
 
     rewriter.replaceOp(reductionOp, result);
     return success();
diff --git a/mlir/lib/Dialect/Affine/Analysis/LoopAnalysis.cpp b/mlir/lib/Dialect/Affine/Analysis/LoopAnalysis.cpp
index 411e79171b898..fe53d03249369 100644
--- a/mlir/lib/Dialect/Affine/Analysis/LoopAnalysis.cpp
+++ b/mlir/lib/Dialect/Affine/Analysis/LoopAnalysis.cpp
@@ -233,8 +233,9 @@ std::optional<uint64_t> mlir::affine::getConstantTripCount(AffineForOp forOp) {
             std::min(*tripCount, static_cast<uint64_t>(constExpr.getValue()));
       else
         tripCount = constExpr.getValue();
-    } else
+    } else {
       return std::nullopt;
+    }
   }
   return tripCount;
 }
diff --git a/mlir/lib/Dialect/Linalg/IR/LinalgOps.cpp b/mlir/lib/Dialect/Linalg/IR/LinalgOps.cpp
index fce0751430305..33bc1fb05f25e 100644
--- a/mlir/lib/Dialect/Linalg/IR/LinalgOps.cpp
+++ b/mlir/lib/Dialect/Linalg/IR/LinalgOps.cpp
@@ -4266,8 +4266,9 @@ void ElementwiseOp::regionBuilder(ImplicitLocOpBuilder &b, Block &block,
     result = helper.buildTernaryFn(kind.ternaryFn, block.getArgument(0),
                                    block.getArgument(1), block.getArgument(2));
 
-  } else
+  } else {
     assert(false && "found unhandled category in elemwise");
+  }
 
   yields.push_back(result);
   helper.yieldOutputs(yields);
diff --git a/mlir/lib/Dialect/Mesh/Interfaces/ShardingInterface.cpp b/mlir/lib/Dialect/Mesh/Interfaces/ShardingInterface.cpp
index 80f4c5cd5afca..dca2b1a52166a 100644
--- a/mlir/lib/Dialect/Mesh/Interfaces/ShardingInterface.cpp
+++ b/mlir/lib/Dialect/Mesh/Interfaces/ShardingInterface.cpp
@@ -57,8 +57,9 @@ checkOperandAffineExprRecursively(AffineExpr expr,
     } else if (rhs.getKind() == AffineExprKind::DimId &&
                lhs.getKind() == AffineExprKind::Constant) {
       dimExpr = rhs;
-    } else
+    } else {
       return failure();
+    }
     unsigned position = cast<AffineDimExpr>(dimExpr).getPosition();
     if ((size_t)position >= seenIds.size() || seenIds[position])
       return failure();
diff --git a/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp b/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp
index eb6e710cae69e..2bf7aaa46db11 100644
--- a/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp
+++ b/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp
@@ -662,8 +662,9 @@ static ParseResult parseClauseWithRegionArgs(
                 parser.parseInteger(mapIndicesVec.emplace_back()) ||
                 parser.parseRSquare())
               return failure();
-          } else
+          } else {
             mapIndicesVec.push_back(-1);
+          }
         }
 
         return success();
diff --git a/mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp b/mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp
index 551411bb14765..57c27231f2144 100644
--- a/mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp
+++ b/mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp
@@ -594,9 +594,10 @@ transform::LoopFuseSiblingOp::apply(transform::TransformRewriter &rewriter,
   } else if (isForallWithIdenticalConfiguration(target, source)) {
     fusedLoop = fuseIndependentSiblingForallLoops(
         cast<scf::ForallOp>(target), cast<scf::ForallOp>(source), rewriter);
-  } else
+  } else {
     return emitSilenceableFailure(target->getLoc())
            << "operations cannot be fused";
+  }
 
   assert(fusedLoop && "failed to fuse operations");
 
diff --git a/mlir/lib/Dialect/SCF/Transforms/LoopPipelining.cpp b/mlir/lib/Dialect/SCF/Transforms/LoopPipelining.cpp
index 3bff1488bde3d..4aacbe739ca5d 100644
--- a/mlir/lib/Dialect/SCF/Transforms/LoopPipelining.cpp
+++ b/mlir/lib/Dialect/SCF/Transforms/LoopPipelining.cpp
@@ -417,8 +417,9 @@ scf::ForOp LoopPipelinerInternal::createKernelLoop(
                       [maxStage - defStage->second];
       assert(valueVersion);
       newLoopArg.push_back(valueVersion);
-    } else
+    } else {
       newLoopArg.push_back(forOp.getInitArgs()[retVal.index()]);
+    }
   }
   for (auto escape : crossStageValues) {
     LiverangeInfo &info = escape.second;
diff --git a/mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp b/mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp
index e172f02c00f51..ad1267381c4f2 100644
--- a/mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp
+++ b/mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp
@@ -124,8 +124,9 @@ static bool haveNoReadsAfterWriteExceptSameIndex(
                   OperationEquivalence::Flags::IgnoreLocations)) {
             return WalkResult::interrupt();
           }
-        } else
+        } else {
           return WalkResult::interrupt();
+        }
       }
     }
     return WalkResult::advance();
diff --git a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUSubgroupDistribute.cpp b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUSubgroupDistribute.cpp
index 2300d9e3bd43f..acdbe7be1358b 100644
--- a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUSubgroupDistribute.cpp
+++ b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUSubgroupDistribute.cpp
@@ -157,8 +157,9 @@ void LayoutInfo::print(raw_ostream &os) const {
     laneLayout.print(os);
     os << ", lane_data: ";
     laneData.print(os);
-  } else
+  } else {
     os << "Not assigned.";
+  }
 }
 
 LayoutInfo LayoutInfo::meet(const LayoutInfo &lhs, const LayoutInfo &rhs) {
diff --git a/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp b/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp
index c250340c38fc7..acb5d9de07510 100644
--- a/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp
+++ b/mlir/lib/ExecutionEngine/SyclRuntimeWrappers.cpp
@@ -65,8 +65,9 @@ static sycl::device getDefaultDevice() {
       return syclDevice;
     }
     throw std::runtime_error("getDefaultDevice failed");
-  } else
+  } else {
     return syclDevice;
+  }
 }
 
 static sycl::context getDefaultContext() {
diff --git a/mlir/lib/Target/LLVM/NVVM/Target.cpp b/mlir/lib/Target/LLVM/NVVM/Target.cpp
index 914a349696617..565f153b50c4f 100644
--- a/mlir/lib/Target/LLVM/NVVM/Target.cpp
+++ b/mlir/lib/Target/LLVM/NVVM/Target.cpp
@@ -606,9 +606,10 @@ NVPTXSerializer::compileToBinaryNVPTX(const std::string &ptxCode) {
           nvPTXCompilerGetErrorLog(compiler, log.data()));
       emitError(loc) << "NVPTX compiler invocation failed, error log: "
                      << log.data();
-    } else
+    } else {
       emitError(loc) << "NVPTX compiler invocation failed with error code: "
                      << status;
+    }
     return std::nullopt;
   }
 
diff --git a/mlir/lib/Target/LLVMIR/ModuleImport.cpp b/mlir/lib/Target/LLVMIR/ModuleImport.cpp
index 77094d4b75f38..7d7d0bb02a3d7 100644
--- a/mlir/lib/Target/LLVMIR/ModuleImport.cpp
+++ b/mlir/lib/Target/LLVMIR/ModuleImport.cpp
@@ -2680,8 +2680,9 @@ ModuleImport::convertParameterAttribute(llvm::AttributeSet llvmParamAttrs,
       const llvm::ConstantRange &value = llvmAttr.getValueAsConstantRange();
       mlirAttr = builder.getAttr<LLVM::ConstantRangeAttr>(value.getLower(),
                                                           value.getUpper());
-    } else
+    } else {
       llvm_unreachable("unexpected parameter attribute kind");
+    }
     paramAttrs.push_back(builder.getNamedAttr(mlirName, mlirAttr));
   }
 
diff --git a/mlir/lib/Transforms/OpStats.cpp b/mlir/lib/Transforms/OpStats.cpp
index 6746ed52396af..0dc3fe9569437 100644
--- a/mlir/lib/Transforms/OpStats.cpp
+++ b/mlir/lib/Transforms/OpStats.cpp
@@ -51,9 +51,9 @@ void PrintOpStatsPass::runOnOperation() {
   // Compute the operation statistics for the currently visited operation.
   getOperation()->walk(
       [&](Operation *op) { ++opCount[op->getName().getStringRef()]; });
-  if (printAsJSON) {
+  if (printAsJSON)
     printSummaryInJSON();
-  } else
+  else
     printSummary();
   markAllAnalysesPreserved();
 }
diff --git a/mlir/test/Dialect/SCF/canonicalize.mlir b/mlir/test/Dialect/SCF/canonicalize.mlir
index c18bd617216f1..5e32a3a78c032 100644
--- a/mlir/test/Dialect/SCF/canonicalize.mlir
+++ b/mlir/test/Dialect/SCF/canonicalize.mlir
@@ -149,7 +149,7 @@ func.func @one_unused(%cond: i1) -> (index) {
 // CHECK:             call @side_effect() : () -> ()
 // CHECK:             [[C1:%.*]] = "test.value1"
 // CHECK:             scf.yield [[C1]] : index
-// CHECK:           } else
+// CHECK:           } else {
 // CHECK:             [[C3:%.*]] = "test.value3"
 // CHECK:             scf.yield [[C3]] : index
 // CHECK:           }
@@ -185,12 +185,12 @@ func.func @nested_unused(%cond1: i1, %cond2: i1) -> (index) {
 // CHECK:               call @side_effect() : () -> ()
 // CHECK:               [[C1:%.*]] = "test.value1"
 // CHECK:               scf.yield [[C1]] : index
-// CHECK:             } else
+// CHECK:             } else {
 // CHECK:               [[C3:%.*]] = "test.value3"
 // CHECK:               scf.yield [[C3]] : index
 // CHECK:             }
 // CHECK:             scf.yield [[V1]] : index
-// CHECK:           } else
+// CHECK:           } else {
 // CHECK:             [[C1_2:%.*]] = "test.value1_2"
 // CHECK:             scf.yield [[C1_2]] : index
 // CHECK:           }
@@ -215,7 +215,7 @@ func.func @all_unused(%cond: i1) {
 // CHECK-LABEL:   func @all_unused
 // CHECK:           scf.if %{{.*}} {
 // CHECK:             call @side_effect() : () -> ()
-// CHECK:           } else
+// CHECK:           } else {
 // CHECK:             call @side_effect() : () -> ()
 // CHECK:           }
 // CHECK:           return
diff --git a/mlir/test/lib/Dialect/Test/TestOpDefs.cpp b/mlir/test/lib/Dialect/Test/TestOpDefs.cpp
index 454a12bac9ab3..b5a8bd10d6b68 100644
--- a/mlir/test/lib/Dialect/Test/TestOpDefs.cpp
+++ b/mlir/test/lib/Dialect/Test/TestOpDefs.cpp
@@ -801,8 +801,9 @@ void TestReflectBoundsOp::inferResultRanges(
     unsigned bitwidth = intTy.getWidth();
     sIntTy = b.getIntegerType(bitwidth, /*isSigned=*/true);
     uIntTy = b.getIntegerType(bitwidth, /*isSigned=*/false);
-  } else
+  } else {
     sIntTy = uIntTy = type;
+  }
 
   setUminAttr(b.getIntegerAttr(uIntTy, range.umin()));
   setUmaxAttr(b.getIntegerAttr(uIntTy, range.umax()));

kazutakahirata
kazutakahirata approved these changes May 11, 2025
View reviewed changes
Copy link
Contributor

@kazutakahirata kazutakahirata left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM. Thanks!

Dinistro
Dinistro approved these changes May 12, 2025
View reviewed changes
Copy link
Contributor

@Dinistro Dinistro left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM % the change in the .mlir test file.

mlir/test/Dialect/SCF/canonicalize.mlir
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I'm not sure if you really want to do these modifications in the tests. I suggest to abstain from touching these

joker-eph
joker-eph approved these changes May 12, 2025
View reviewed changes
Copy link
Collaborator

@joker-eph joker-eph left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LG, but not sure about the tests: in general we should test the minimum amount of things and it's not clear to me whether these CHECKs do?

@jurahul
Copy link
Contributor Author

jurahul commented May 12, 2025

Right, I changed them as they showed up in my grep for this. Note that SCF ops.mlir test checks for these braces as well, and some existing checks in that same file check for opening { after else. Also, all the sub-tests I am modifying check for the opening { of the scf.if so checking the { for else is consistent with that as well. So I propose we accept the test changes here.

@jurahul jurahul merged commit b17f3c6 into llvm:main May 12, 2025
25 checks passed
@jurahul jurahul deleted the mlir_else_brace branch May 12, 2025 17:29
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers

@joker-eph joker-eph joker-eph approved these changes

@Dinistro Dinistro Dinistro approved these changes

@kazutakahirata kazutakahirata kazutakahirata approved these changes

Assignees
No one assigned
Labels
flang:openmp mlir:affine mlir:core MLIR Core Infrastructure mlir:execution-engine mlir:gpu mlir:linalg mlir:llvm mlir:openmp mlir:presburger mlir:scf mlir
Projects
None yet
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.
5 participants
@jurahul @llvmbot @joker-eph @Dinistro @kazutakahirata