Revert accidentally updated part.

Author: lialan

compiler/src/iree/compiler/Codegen/LLVMGPU/KernelConfig.cpp

@@ -1446,12 +1446,12 @@ static LogicalResult setContractConfig(IREE::GPU::TargetAttr target,
     return failure();
   }
 
-  auto setMatmulConfig = [&entryPoint, &op,
-                          &target](int64_t tileX, int64_t tileY, int64_t tileK,
-                                   ArrayRef<int64_t> workgroupSize,
-                                   ArrayRef<int32_t> subgroupSizes,
-                                   unsigned softwarePipelineDepth,
-                                   CodeGenPipeline pipeline) {
+  auto setMatmulConfig = [&entryPoint, &op](int64_t tileX, int64_t tileY,
+                                            int64_t tileK,
+                                            ArrayRef<int64_t> workgroupSize,
+                                            ArrayRef<int32_t> subgroupSizes,
+                                            unsigned softwarePipelineDepth,
+                                            CodeGenPipeline pipeline) {
     TileSizesListType tileSizes;
     unsigned numParallelLoops = op.getNumParallelLoops();
     unsigned numReductionLoops = op.getNumReductionLoops();
@@ -1509,8 +1509,7 @@ static LogicalResult setContractConfig(IREE::GPU::TargetAttr target,
       SmallVector<NamedAttribute, 1> pipelineAttrs;
       auto pipelineOptions = IREE::GPU::GPUPipelineOptionsAttr::get(
           context, /*prefetch_num_stages=*/0,
-          /*no_reduce_shared_memory_bank_conflicts=*/
-          IREE::GPU::targetSupportsGlobalLoadDMA(target),
+          /*no_reduce_shared_memory_bank_conflicts=*/true,
           /*use_igemm_convolution=*/false,
           /*reorder_workgroups_strategy=*/std::nullopt);
       pipelineAttrs.emplace_back(

compiler/src/iree/compiler/Codegen/LLVMGPU/test/config_matmul.mlir

@@ -153,7 +153,7 @@ func.func @matmul_DYN_1_4096(%arg0: !TA, %arg1: !TB, %arg2: !TC, %arg3: !DTC, %a
 !DTC = !iree_tensor_ext.dispatch.tensor<readwrite:tensor<32x32xf32>>
 
 //      CHECK:    #translation = #iree_codegen.translation_info<pipeline = LLVMGPUTileAndFuse
-// CHECK-SAME:    workgroup_size = [64, 16, 1] subgroup_size = 64, {gpu_pipeline_options = #iree_gpu.pipeline_options<no_reduce_shared_memory_bank_conflicts = false, use_igemm_convolution = false>}>
+// CHECK-SAME:    workgroup_size = [64, 16, 1] subgroup_size = 64, {gpu_pipeline_options = #iree_gpu.pipeline_options<no_reduce_shared_memory_bank_conflicts = true, use_igemm_convolution = false>}>
 func.func @matmul_32_32_DYN(%arg0: !TA, %arg1: !TB, %arg2: !TC, %arg3: !DTC) {
    // CHECK:     #iree_gpu.lowering_config<{reduction = [0, 0, 1], thread = [1, 8, 0], workgroup = [64, 128, 1]}
   %0 = linalg.matmul ins(%arg0, %arg1 : !TA, !TB) outs(%arg2 : !TC) -> !TC
@@ -168,7 +168,7 @@ func.func @matmul_32_32_DYN(%arg0: !TA, %arg1: !TB, %arg2: !TC, %arg3: !DTC) {
 !TC = tensor<4096x4096xf32>
 !DTC = !iree_tensor_ext.dispatch.tensor<readwrite:tensor<4096x4096xf32>>
 //      CHECK:         #translation = #iree_codegen.translation_info<pipeline = LLVMGPUTileAndFuse
-// CHECK-SAME:         workgroup_size = [64, 16, 1] subgroup_size = 64, {gpu_pipeline_options = #iree_gpu.pipeline_options<no_reduce_shared_memory_bank_conflicts = false, use_igemm_convolution = false>}>
+// CHECK-SAME:         workgroup_size = [64, 16, 1] subgroup_size = 64, {gpu_pipeline_options = #iree_gpu.pipeline_options<no_reduce_shared_memory_bank_conflicts = true, use_igemm_convolution = false>}>
 func.func @matmul_4096_4096_DYN(%arg0: !TA, %arg1: !TB, %arg2: !TC, %arg3: !DTC) {
    // CHECK:      #iree_gpu.lowering_config<{reduction = [0, 0, 1], thread = [1, 8, 0], workgroup = [64, 128, 1]}
   %0 = linalg.matmul ins(%arg0, %arg1 : !TA, !TB) outs(%arg2 : !TC) -> !TC

compiler/src/iree/compiler/Codegen/LLVMGPU/test/config_matvec.mlir

@@ -438,7 +438,7 @@ func.func @not_vmt() {
   return
 }
 
-//   CHECK-DAG: #[[$TRANSLATION:.+]] = #iree_codegen.translation_info<pipeline = LLVMGPUTileAndFuse workgroup_size = [32, 1, 1] subgroup_size = 64, {gpu_pipeline_options = #iree_gpu.pipeline_options<no_reduce_shared_memory_bank_conflicts = false, use_igemm_convolution = false>}>
+//   CHECK-DAG: #[[$TRANSLATION:.+]] = #iree_codegen.translation_info<pipeline = LLVMGPUTileAndFuse workgroup_size = [32, 1, 1] subgroup_size = 64, {gpu_pipeline_options = #iree_gpu.pipeline_options<no_reduce_shared_memory_bank_conflicts = true, use_igemm_convolution = false>}>
 //       CHECK: func.func @not_vmt()
 //  CHECK-SAME:     translation_info = #[[$TRANSLATION]]
 //       CHECK:   linalg.generic