[GPU] Allow vectorization for dynamic shapes with inner static dims (#…

…19850)

For the Tile and Fuse vectorize pipeline config setup, needing all dims
static for vectorization seems unnecessary as dynamaic dims will get
tiled to 1 anyway due to other checks that skip dynamic dims. Now, if the
innermost dim is static we will end up vectorizing.
Fixes : #19843

---------

Signed-off-by: Nirvedh Meshram <[email protected]>

compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/ConfigUtils.cpp

@@ -559,11 +559,10 @@ LogicalResult setTileAndFuseLoweringConfig(IREE::GPU::TargetAttr target,
                    [](AffineMap map) { return map.isProjectedPermutation(); });
   bool powTwo =
       llvm::all_of(linalgOp->getOperands(), elementHasPowerOfTwoBitwidth);
-  bool staticShape = llvm::none_of(loopBounds, ShapedType::isDynamic);
 
   // Require all affine maps to be projected permutation so that we can
   // generate vector transfer ops.
-  bool vectorizable = projPerm && powTwo && staticShape;
+  bool vectorizable = projPerm && powTwo;
 
   const unsigned minBitwidth = getMinElementBitwidth(linalgOp);
   // Make sure we use a tile size that results in some integral number of bytes.

compiler/src/iree/compiler/Codegen/LLVMGPU/test/ROCDL/config_tile_and_fuse.mlir

@@ -179,8 +179,8 @@ module {
 //       CHECK:   linalg.matmul {{.*}}lowering_config = #iree_gpu.lowering_config
 //  CHECK-SAME:     promote_operands = [0, 1]
 //  CHECK-SAME:     reduction = [0, 0, 4]
-//  CHECK-SAME:     thread = [1, 1, 0]
-//  CHECK-SAME:     workgroup = [1, 64, 0]
+//  CHECK-SAME:     thread = [1, 4, 0]
+//  CHECK-SAME:     workgroup = [1, 256, 0]
 
 // -----
 
@@ -198,8 +198,8 @@ module {
 // CHECK-LABEL: func.func @elementwise_dynamic_dim
 //  CHECK-SAME:   #iree_codegen.translation_info<pipeline = LLVMGPUTileAndFuse workgroup_size = [64, 1, 1] subgroup_size = 64>
 //       CHECK:   linalg.add {{.*}}lowering_config = #iree_gpu.lowering_config
-//  CHECK-SAME:     thread = [1, 1]
-//  CHECK-SAME:     workgroup = [1, 64]
+//  CHECK-SAME:     thread = [1, 4]
+//  CHECK-SAME:     workgroup = [1, 256]
 
 // -----