More cleanup to satisfy clang werror rules

lib/Conversion/NDArrayToLinalg/NDArrayToLinalg.cpp

@@ -1242,7 +1242,7 @@ struct ConvertNDArrayToLinalgPass
     // Convert unknown types to itself
     auto convT2T = [](::mlir::Type type) { return type; };
     // Convert NDArrayType to (tensorType)
-    auto convNDArray2RankedTensor = [&ctxt](::imex::ndarray::NDArrayType type)
+    auto convNDArray2RankedTensor = [](::imex::ndarray::NDArrayType type)
         -> std::optional<::mlir::Type> { return type.getTensorType(); };
 
     typeConverter.addConversion(convT2T);

lib/Conversion/XeGPUToSPIRV/XeGPUToSPIRV.cpp

@@ -426,6 +426,7 @@ class LoadStorePrefetchNdToLsc : public OpConversionPattern<OpType> {
   }
 };
 
+#if 0
 std::optional<xegpu::CreateNdDescOp> findDescOp(mlir::Value val) {
   if (auto op = val.getDefiningOp()) {
     if (auto descOp = dyn_cast<xegpu::CreateNdDescOp>(op)) {
@@ -442,6 +443,7 @@ std::optional<xegpu::CreateNdDescOp> findDescOp(mlir::Value val) {
   // Add more support
   return std::nullopt;
 }
+#endif
 
 template <typename OpType>
 class LoadStorePrefetchNdToRawSend : public OpConversionPattern<OpType> {
@@ -1138,9 +1140,6 @@ class FenceToVCPattern : public OpConversionPattern<FenceOp> {
     case mlir::xegpu::MemoryScope::SLM:
       sfid = lscSFID::TGM;
       break;
-    default:
-      llvm_unreachable("unsupported value for memory_kind attribute");
-      break;
     }
 
     switch (op.getFenceScope()) {
@@ -1150,9 +1149,6 @@ class FenceToVCPattern : public OpConversionPattern<FenceOp> {
     case mlir::xegpu::FenceScope::GPU:
       fence_scope = lscFenceScope::GPU;
       break;
-    default:
-      llvm_unreachable("unsupported value for fence_scope attribute");
-      break;
     }
 
     SmallVector<Value> args{pred, i8_val(sfid), i8_val(fence_op),
@@ -1252,7 +1248,7 @@ struct VectorExtract final : public OpConversionPattern<vector::ExtractOp> {
 
 static uint64_t getFirstIntValue(mlir::ArrayAttr attr) {
   return (*attr.getAsValueRange<IntegerAttr>().begin()).getZExtValue();
-};
+}
 
 struct VectorExtractStridedSlice final
     : public OpConversionPattern<vector::ExtractStridedSliceOp> {

lib/Conversion/XeGPUToVC/XeGPUToVC.cpp

@@ -1091,9 +1091,6 @@ class FenceToVCPattern : public OpConversionPattern<::mlir::xegpu::FenceOp> {
     case mlir::xegpu::MemoryScope::SLM:
       sfid = lscSFID::TGM;
       break;
-    default:
-      llvm_unreachable("unsupported value for memory_kind attribute");
-      break;
     }
 
     switch (op.getFenceScope()) {
@@ -1103,9 +1100,6 @@ class FenceToVCPattern : public OpConversionPattern<::mlir::xegpu::FenceOp> {
     case mlir::xegpu::FenceScope::GPU:
       fence_scope = lscFenceScope::GPU;
       break;
-    default:
-      llvm_unreachable("unsupported value for fence_scope attribute");
-      break;
     }
 
     SmallVector<Value> args{pred, i8_val(sfid), i8_val(fence_op),
@@ -1212,7 +1206,7 @@ struct VectorExtractVC final
 
 static uint64_t getFirstIntValue(mlir::ArrayAttr attr) {
   return (*attr.getAsValueRange<IntegerAttr>().begin()).getZExtValue();
-};
+}
 
 struct VectorExtractStridedSliceVC final
     : public OpConversionPattern<vector::ExtractStridedSliceOp> {

lib/Conversion/XeTileToXeGPU/XeTileOpConversion.cpp

@@ -135,7 +135,7 @@ mlir::Value mergeVectorsWrapper(mlir::ValueRange ins,
     }
   }
   return shuffleOps[0];
-};
+}
 
 // a unified function lowering Unpack and Pack ops.
 static llvm::SmallVector<mlir::Value>

lib/Dialect/Dist/IR/DistOps.cpp

@@ -194,8 +194,9 @@ void InitDistArrayOp::build(::mlir::OpBuilder &odsBuilder,
     shapes.emplace_back(
         mlir::cast<::imex::ndarray::NDArrayType>(p.getType()).getShape());
   }
+  auto resShape = getShapeFromValues(l_offset);
   ::mlir::ArrayRef<int64_t> lOffs =
-      s_Offs.size() ? s_Offs : getShapeFromValues(l_offset);
+      s_Offs.size() ? s_Offs : resShape;
   ::mlir::SmallVector<::mlir::Attribute> nEnvs(environments);
   nEnvs.emplace_back(::imex::dist::DistEnvAttr::get(team, lOffs, shapes));
   auto arType = ::imex::ndarray::NDArrayType::get(g_shape, elTyp, nEnvs);

lib/Dialect/NDArray/IR/EWOp.h

@@ -33,4 +33,4 @@ bool replaceOperandInplaceWithCast(::mlir::PatternRewriter &rewriter,
     }
   }
   return false;
-};
+}

lib/Dialect/NDArray/IR/NDArrayOps.cpp

@@ -121,7 +121,7 @@ imex::ndarray::NDArrayBase::cloneWithEnv(::mlir::Attribute env) const {
   envs.emplace_back(env);
   return NDArrayType::get(t.getShape(), t.getElementType(), envs,
                           t.getLayout());
-};
+}
 
 bool imex::ndarray::NDArrayBase::isValidElementType(Type type) {
   return type.isIntOrIndexOrFloat();

lib/Dialect/NDArray/Transforms/AddGPURegions.cpp

@@ -80,7 +80,7 @@ matchAndRewritePTOP(::mlir::Operation *op, ::mlir::PatternRewriter &rewriter,
   // create a region with given env and clone creator op within and yield it
   auto rOp = rewriter.create<::imex::region::EnvironmentRegionOp>(
       op->getLoc(), env, std::nullopt, op->getResultTypes(),
-      [op, &env](::mlir::OpBuilder &builder, ::mlir::Location loc) {
+      [op](::mlir::OpBuilder &builder, ::mlir::Location loc) {
         auto cOp = builder.clone(*op);
         (void)builder.create<::imex::region::EnvironmentRegionYieldOp>(
             loc, cOp->getResults());

lib/ExecutionEngine/LEVELZERORUNTIME/LevelZeroRuntimeWrappers.cpp

@@ -135,8 +135,8 @@ getDriverAndDevice(ze_device_type_t deviceType = ZE_DEVICE_TYPE_GPU) {
 namespace imex {
 namespace profiling {
 // defining two types representing kernel start and kernel end
-_IMEX_PROFILING_TRAITS_SPEC(command_start);
-_IMEX_PROFILING_TRAITS_SPEC(command_end);
+_IMEX_PROFILING_TRAITS_SPEC(command_start)
+_IMEX_PROFILING_TRAITS_SPEC(command_end)
 } // namespace profiling
 } // namespace imex
 

lib/Transforms/BF16ToGPU.cpp

@@ -39,7 +39,6 @@ struct BF16ToGPUPass : public BF16ToGPUBase<BF16ToGPUPass> {
       // 1-1: Create new FunctionType and replace old FunctionType
       auto oftype = op.getFunctionType();
       llvm::SmallVector<mlir::Type, 4> argTypes;
-      ArrayRef<Type> inputTypes;
       ArrayRef<Type> resultTypes;
       for (Type t : oftype.getInputs()) {
         MemRefType m = mlir::dyn_cast<MemRefType>(t);

lib/Utils/XeCommon.cpp

@@ -28,15 +28,15 @@ int getOperandIndex(mlir::Operation *op, mlir::Value operand) {
       return i;
   }
   return -1;
-};
+}
 
 mlir::BlockArgument getArgForOperand(mlir::scf::ForOp &op,
                                      mlir::Value operand) {
   auto idx = getOperandIndex(op, operand);
   auto numControls = op.getNumControlOperands();
   assert(idx >= (int)numControls);
   return op.getRegionIterArg(idx - numControls);
-};
+}
 
 bool isSupportedModule(mlir::gpu::GPUModuleOp mod) {
   bool hasTileTyInFuncTy = false;