Another update

Author: lialan

mlir/include/mlir/Conversion/GPUToROCDL/GPUToROCDLPass.h

@@ -10,7 +10,6 @@
 
 #include "mlir/Conversion/GPUToROCDL/Runtimes.h"
 #include "mlir/Conversion/LLVMCommon/LoweringOptions.h"
-#include "mlir/Dialect/AMDGPU/Utils/Chipset.h"
 #include <memory>
 
 namespace mlir {

mlir/include/mlir/Conversion/Passes.td

@@ -608,10 +608,6 @@ def ConvertGpuOpsToROCDLOps : Pass<"convert-gpu-to-rocdl", "gpu::GPUModuleOp"> {
                clEnumValN(::mlir::gpu::amd::Runtime::HIP, "HIP", "HIP"),
                clEnumValN(::mlir::gpu::amd::Runtime::OpenCL, "OpenCL",
                           "OpenCL"))}]>,
-    Option<"subgroupSize", "subgroup-size", "unsigned",
-           "0",
-           "specify subgroup size for the kernel, if left empty, the default "
-           "value will be decided by the target chipset.">,
     ListOption<"allowedDialects", "allowed-dialects", "std::string",
                "Run conversion patterns of only the specified dialects">,
   ];

mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp

@@ -90,7 +90,6 @@ static Value truncOrExtToLLVMType(ConversionPatternRewriter &rewriter,
   int64_t indexBitwidth = converter.getIndexTypeBitwidth();
   auto indexBitwidthType =
       IntegerType::get(rewriter.getContext(), converter.getIndexTypeBitwidth());
-  // TODO: use <=> in C++20.
   if (indexBitwidth > intWidth) {
     return rewriter.create<LLVM::SExtOp>(loc, indexBitwidthType, value);
   }
@@ -281,7 +280,7 @@ struct GPUSubgroupIdOpToROCDL final
         rewriter.create<LLVM::AddOp>(loc, int32Type, workitemIdX,
                                      dimYxIdZPlusIdYTimesDimX, flags);
     Value subgroupSize = rewriter.create<ROCDL::WavefrontSizeOp>(
-        loc, rewriter.getI32Type(), nullptr);
+        loc, rewriter.getI32Type(), /*upper_bound = */ nullptr);
     Value waveIdOp = rewriter.create<LLVM::UDivOp>(
         loc, workitemIdXPlusDimYxIdZPlusIdYTimesDimX, subgroupSize);
     rewriter.replaceOp(op, {truncOrExtToLLVMType(rewriter, loc, waveIdOp,