Forward and backward weights padding kernel

mlir/lib/Dialect/MIOpen/Transforms/AffineTransforms.cpp

@@ -65,50 +65,12 @@ AffineMap AffineTransforms::buildIndexAffineMap(miopen::TransformOp op) {
           // leftPad = 1 rightPad = 2
           auto leftPad =
               parameters.getValue()[j * 2].cast<IntegerAttr>().getInt();
-          auto rightPad =
-              parameters.getValue()[j * 2 + 1].cast<IntegerAttr>().getInt();
 
           auto srcDim = srcDimAttr.getValue()[j].cast<IntegerAttr>().getInt();
           auto destDim = destDimAttr.getValue()[j].cast<IntegerAttr>().getInt();
 
           auto expr = getAffineDimExpr(destDim, op.getContext()) + getAffineConstantExpr(-leftPad, op.getContext());
-          if (leftPad == 0 && rightPad != 0) {
-            // when leftPad == 0 , your original expr is just minus leftpad, but
-            // leftpad is zero, affinemap do not have minus out of boundary
-            // check depends on minus symbol , it will not do out of boundary
-            // check even rightpad part is oob example of leftPad == 0 &&
-            // rightPad != 0:
-            //
-            // srcIndex0 srcIndex1 ... srcIndex[src_size - 1]
-            // dstIndex0 dstIndex1 ... dstIndex[src_size - 1] dstIndex[rightpad]
-            // index0    index1        index[src_size -1]     index[src_size]
-            // can't find index[src_size] in src
-            // so we need to force it to  do out of boundary check ,
-            //
-            // the idea :
-            // dst index :
-            // dstIndex0 dstIndex1 ... dstIndex[src_size -1]  dstIndex[rightpad]
-            // src index computed:
-            // srcIndex0 srcIndex1 ... srcIndex[src_size - 1] src_size+1
-            //
-            // how to achieve it:
-            // dstIndex + (dstIndex/srcsize) + 1 - 1
-            //
-            // the expr is :
-            // dstIndex + ceildiv(dstIndex+1/srcsize) - 1
-            // the same with above but the
-            // minus symbol exist after optimization
-            //
-            // but if we use the equation above, when srcsize = 1
-            // affinemap will optimized  and no minus symbol
-            // just add more 1 can generate minus symbol
-            // the final expr is :
-            // dstIndex + ceildiv((dstIndex+2)/(srcsize+1)) - 1
-            expr = ((getAffineDimExpr(destDim, op.getContext()) + 2)
-                        .ceilDiv(inputShape[srcDim] + 1)) +
-                   getAffineDimExpr(destDim, op.getContext()) -
-                   getAffineConstantExpr(1, op.getContext());
-          }
+
           affExprsMap.insert({srcDim, expr});
         }
       } else if (transformAttr.getValue() == "Merge" ||

mlir/lib/Dialect/MIOpen/Tuning/GridwiseGemmParams.cpp

@@ -177,8 +177,7 @@ LogicalResult PopulateParams::paramsFromCtx(
                             << " PARAMETERS!\n");
 
     InitParams paddingParam = getUniversalParameters();
-    if ((gemmSize.gemmN % paddingParam.gemmNPerBlock == 0) &&
-        (gemmSize.gemmM % paddingParam.gemmMPerBlock == 0)) {
+    if (ctx.opType != miopen::ConvOpType::Conv2DBwdDataOpType) {
       LLVM_DEBUG(llvm::dbgs() << "BUT PADDING KERNEL CAN EXECUTE IT\n");
 
       for (auto &params : initParameters) {
@@ -194,7 +193,9 @@ LogicalResult PopulateParams::paramsFromCtx(
         break;
       }
     } else {
-      LLVM_DEBUG(llvm::dbgs() << "PADDING KERNEL only support gemmK now\n");
+      LLVM_DEBUG(
+          llvm::dbgs()
+          << "PADDING KERNEL only support forward, backward weights now\n");
     }
   } else {
     LLVM_DEBUG(llvm::dbgs() << "Successfully picked tuning params from backup"
@@ -380,8 +381,7 @@ LogicalResult PopulateParamsXDL::paramsFromCtx(
                             << " PARAMETERS!\n");
 
     InitParams paddingParam = getUniversalParameters();
-    if ((gemmSize.gemmN % paddingParam.gemmNPerBlock == 0) &&
-        (gemmSize.gemmM % paddingParam.gemmMPerBlock == 0)) {
+    if (ctx.opType != miopen::ConvOpType::Conv2DBwdDataOpType) {
       LLVM_DEBUG(llvm::dbgs() << "BUT PADDING KERNEL CAN EXECUTE IT\n");
       for (auto &params : initParameters) {
         res = populatePaddingKernelDerived(ctx, params, gemmSize,
@@ -395,7 +395,9 @@ LogicalResult PopulateParamsXDL::paramsFromCtx(
         break;
       }
     } else {
-      LLVM_DEBUG(llvm::dbgs() << "PADDING KERNEL only support gemmK now\n");
+      LLVM_DEBUG(
+          llvm::dbgs()
+          << "PADDING KERNEL only support forward, backward weights now\n");
     }
   } else {
     LLVM_DEBUG(llvm::dbgs() << "Successfully picked tuning params from backup"

mlir/test/Dialect/MIOpen/lowering_filter_tensor_ckyx_cnhw_knhw.mlir

@@ -74,4 +74,4 @@ func @miopen_conv2d_bwd_weight_ckyx_cnhw_knhw(%filter : memref<1x8x128x3x3xf32>,
 // CHECK:       upper_layer_names = ["gemmM"]
 // CHECK:       lower_layer_names = ["c", "y", "x"]
 // CHECK:       upper_layer_names = ["gemmN"]
-// CHECK-NEXT:  miopen.transform(%arg1)
+// CHECK-NEXT:  miopen.transform

mlir/test/Dialect/MIOpen/lowering_gridwise_gemm_position_cyxk_chwn_khwn.mlir

@@ -71,7 +71,9 @@ func @miopen_conv2d_bwd_weight_cyxk_chwn_khwn(%filter : memref<1x8x3x3x128xf32>,
 // CHECK-NEXT:  miopen.transform
 // CHECK-NEXT:  miopen.transform
 // CHECK-NEXT:  miopen.transform
+// CHECK-NEXT:  miopen.transform
 // CHECK:       gridwise_gemm_argument_position = 1
 // CHECK-NEXT:  miopen.transform
+// CHECK-NEXT:  miopen.transform
 // CHECK:       gridwise_gemm_argument_position = 0
-// CHECK-NEXT:  miopen.gridwise_gemm(%4, %3, %0)
+// CHECK-NEXT:  miopen.gridwise_gemm(%6, %5, %1)

mlir/test/Dialect/MIOpen/lowering_input_tensor_cyxk_cnhw_knhw.mlir

@@ -66,10 +66,14 @@ func @miopen_conv2d_bwd_weight_cyxk_cnhw_knhw(%filter : memref<1x8x3x3x128xf32>,
 }
 // CHECK-LABEL: func @miopen_conv2d_bwd_weight
 // CHECK-NEXT:  miopen.transform(%arg0)
+// CHECK-NEXT:  miopen.transform
+// CHECK:       upper_layer_layout = ["gemmG", "gemmM", "gemmNPad"]
 // CHECK-NEXT:  miopen.transform(%arg1)
 // CHECK:       upper_layer_layout = ["gi", "ci", "ni", "hipad", "wipad"]
 // CHECK-NEXT:  miopen.transform
 // CHECK:       upper_layer_layout = ["gi", "ci", "ni", "y", "ho", "x", "wo"]
 // CHECK-NEXT:  miopen.transform
 // CHECK:       upper_layer_layout = ["gemmG", "gemmK", "gemmN"]
+// CHECK-NEXT:  miopen.transform
+// CHECK:       upper_layer_layout = ["gemmG", "gemmK", "gemmNPad"]
 // CHECK-NEXT:  miopen.transform(%arg2)

mlir/test/Dialect/MIOpen/lowering_memref_kcyx_nchw_nkhw.mlir

@@ -66,9 +66,11 @@ func @miopen_conv2d_bwd_weight_kcyx_nchw_nkhw(%filter : memref<1x128x8x3x3xf32>,
   return
 }
 // CHECK-LABEL: func @miopen_conv2d_bwd_weight
-// CHECK-NEXT:  {{miopen.transform.*{.*}.*memref.*memref}}
-// CHECK-NEXT:  {{miopen.transform.*{.*}.*memref.*memref}}
-// CHECK-NEXT:  {{miopen.transform.*{.*}.*memref.*memref}}
-// CHECK-NEXT:  {{miopen.transform.*{.*}.*memref.*memref}}
-// CHECK-NEXT:  {{miopen.transform.*{.*}.*memref.*memref}}
+// CHECK-NEXT:  {{miopen.transform.*{.*"g", "k", "c", "y", "x".*}.*memref.*memref}}
+// CHECK-NEXT:  {{miopen.transform.*{.*"gemmG", "gemmM", "gemmNPad".*}.*memref.*memref}}
+// CHECK-NEXT:  {{miopen.transform.*{.*"ni", "gi", "ci", "hipad", "wipad".*}.*memref.*memref}}
+// CHECK-NEXT:  {{miopen.transform.*{.*"ni", "gi", "ci", "y", "ho", "x", "wo".*}.*memref.*memref}}
+// CHECK-NEXT:  {{miopen.transform.*{.*"gemmG", "gemmK", "gemmN".*}.*memref.*memref}}
+// CHECK-NEXT:  {{miopen.transform.*{.*"gemmG", "gemmK", "gemmNPad".*}.*memref.*memref}}
+// CHECK-NEXT:  {{miopen.transform.*{.*"gemmG", "gemmK", "gemmM".*}.*memref.*memref}}
 // CHECK-NEXT:  {{miopen.gridwise_gemm.*{.*}.*memref.*memref.*memref}}

mlir/test/Dialect/MIOpen/lowering_padding_kernel.mlir

@@ -18,9 +18,9 @@ func @miopen_conv2d_kcyx_nchw_nkhw_padding_kernel(%filter : memref<32x128x2x3x3x
   return
 }
 // CHECK-LABEL: func @miopen_conv2d
-// CHECK-NEXT:  {{miopen.transform.*{.*extraPad = true, gemmKExtra = 14 : i32, gemmMExtra = 0 : i32,.*}.*memref.*memref}}
+// CHECK-NEXT:  {{miopen.transform.*{.*extraPad = true, gemmKExtra = 14 : i32, gemmMExtra = 0 : i32, gemmNExtra = 0 : i32,.*}.*memref.*memref}}
 // CHECK-NEXT:  {{miopen.transform.*{.*}.*memref.*memref}}
-// CHECK-NEXT:  {{miopen.transform.*{.*extraPad = true, gemmKExtra = 14 : i32, gemmNExtra = 0 : i32,.*}.*memref.*memref}}
+// CHECK-NEXT:  {{miopen.transform.*{.*extraPad = true, gemmKExtra = 14 : i32, gemmMExtra = 0 : i32, gemmNExtra = 0 : i32,.*}.*memref.*memref}}
 // CHECK-NEXT:  {{miopen.transform.*{.*}.*memref.*memref}}
 // CHECK-NEXT:  {{miopen.transform.*{.*}.*memref.*memref}}
 // CHECK-NEXT:  {{miopen.transform.*{.*}.*memref.*memref}}
@@ -41,8 +41,8 @@ func @miopen_conv2d_kcyx_nchw_nkhw_no_extra_padding(%filter : memref<1x128x64x3x
   return
 }
 // CHECK-LABEL: func @miopen_conv2d
-// CHECK-NEXT:  {{miopen.transform.*{.*extraPad = false, gemmKExtra = 0 : i32, gemmMExtra = 0 : i32,.*}.*memref.*memref}}
-// CHECK-NEXT:  {{miopen.transform.*{.*extraPad = false, gemmKExtra = 0 : i32, gemmNExtra = 0 : i32,.*}.*memref.*memref}}
+// CHECK-NEXT:  {{miopen.transform.*{.*extraPad = false, gemmKExtra = 0 : i32, gemmMExtra = 0 : i32, gemmNExtra = 0 : i32,.*}.*memref.*memref}}
+// CHECK-NEXT:  {{miopen.transform.*{.*extraPad = false, gemmKExtra = 0 : i32, gemmMExtra = 0 : i32, gemmNExtra = 0 : i32,.*}.*memref.*memref}}
 // CHECK-NEXT:  {{miopen.transform.*{.*}.*memref.*memref}}
 // CHECK-NEXT:  {{miopen.transform.*{.*}.*memref.*memref}}
 // CHECK-NEXT:  {{miopen.transform.*{.*extraPad = false, gemmMExtra = 0 : i32, gemmNExtra = 0 : i32,.*}.*memref.*memref}}
@@ -62,9 +62,9 @@ func @miopen_conv2d_kcyx_nchw_nkhw_partial_padding_kernel(%filter : memref<32x12
   return
 }
 // CHECK-LABEL: func @miopen_conv2d
-// CHECK-NEXT:  {{miopen.transform.*{.*extraPad = true, gemmKExtra = 14 : i32, gemmMExtra = 0 : i32,.*}.*memref.*memref}}
+// CHECK-NEXT:  {{miopen.transform.*{.*extraPad = true, gemmKExtra = 14 : i32, gemmMExtra = 0 : i32, gemmNExtra = 0 : i32,.*}.*memref.*memref}}
 // CHECK-NEXT:  {{miopen.transform.*{.*}.*memref.*memref}}
-// CHECK-NEXT:  {{miopen.transform.*{.*extraPad = true, gemmKExtra = 14 : i32, gemmNExtra = 0 : i32,.*}.*memref.*memref}}
+// CHECK-NEXT:  {{miopen.transform.*{.*extraPad = true, gemmKExtra = 14 : i32, gemmMExtra = 0 : i32, gemmNExtra = 0 : i32,.*}.*memref.*memref}}
 // CHECK-NEXT:  {{miopen.transform.*{.*}.*memref.*memref}}
 // CHECK-NEXT:  {{miopen.transform.*{.*}.*memref.*memref}}
 // CHECK-NEXT:  {{miopen.transform.*{.*}.*memref.*memref}}

mlir/test/Dialect/MIOpen/lowering_top_level.mlir

@@ -121,12 +121,14 @@ func @miopen_conv2d_bwd_weight(%filter : memref<1x128x8x3x3xf32>, %input : memre
   return
 }
 // CHECK-LABEL: func {{@miopen_conv2d_bwd_weight.*%arg0.*%arg1.*%arg2}}
-// CHECK-NOT:   miopen.conv2d_bwd_data
-// CHECK-NEXT:  miopen.transform(%arg0)
+// CHECK-NOT:   miopen.conv2d_bwd_weight
+// CHECK-NEXT:  {{miopen.transform\(%arg0\).* upper_layer_layout = \["gemmG", "gemmM", "gemmN"\].*}}
+// CHECK-NEXT:  {{miopen.transform.* upper_layer_layout = \["gemmG", "gemmM", "gemmNPad"\].*}}
 // CHECK-NEXT:  miopen.transform(%arg1)
 // CHECK-NEXT:  miopen.transform
-// CHECK-NEXT:  miopen.transform
-// CHECK-NEXT:  miopen.transform(%arg2)
+// CHECK-NEXT:  {{miopen.transform.* upper_layer_layout = \["gemmG", "gemmK", "gemmN"\].*}}
+// CHECK-NEXT:  {{miopen.transform.* upper_layer_layout = \["gemmG", "gemmK", "gemmNPad"\].*}}
+// CHECK-NEXT:  {{miopen.transform\(%arg2\).* upper_layer_layout = \["gemmG", "gemmK", "gemmM"\].*}}
 // CHECK-NEXT:  miopen.gridwise_gemm
 
 func @miopen_conv2d_bwd_weight_f16(%filter : memref<1x128x8x3x3xf16>, %input : memref<128x1x8x32x32xf16>, %output : memref<128x1x128x30x30xf16>) {
@@ -142,11 +144,63 @@ func @miopen_conv2d_bwd_weight_f16(%filter : memref<1x128x8x3x3xf16>, %input : m
   } : memref<1x128x8x3x3xf16>, memref<128x1x8x32x32xf16>, memref<128x1x128x30x30xf16>
   return
 }
-// CHECK-LABEL: func {{@miopen_conv2d_bwd_weight.*%arg0.*%arg1.*%arg2}}
-// CHECK-NOT:   miopen.conv2d_bwd_data
-// CHECK-NEXT:  miopen.transform(%arg0)
+// CHECK-LABEL: func {{@miopen_conv2d_bwd_weight_f16.*%arg0.*%arg1.*%arg2}}
+// CHECK-NOT:   miopen.conv2d_bwd_weight
+// CHECK-NEXT:  {{miopen.transform\(%arg0\).* upper_layer_layout = \["gemmG", "gemmM", "gemmN"\].*}}
+// CHECK-NEXT:  {{miopen.transform.* upper_layer_layout = \["gemmG", "gemmM", "gemmNPad"\].*}}
 // CHECK-NEXT:  miopen.transform(%arg1)
 // CHECK-NEXT:  miopen.transform
+// CHECK-NEXT:  {{miopen.transform.* upper_layer_layout = \["gemmG", "gemmK", "gemmN"\].*}}
+// CHECK-NEXT:  {{miopen.transform.* upper_layer_layout = \["gemmG", "gemmK", "gemmNPad"\].*}}
+// CHECK-NEXT:  {{miopen.transform\(%arg2\).* upper_layer_layout = \["gemmG", "gemmK", "gemmM"\].*}}
+// CHECK-NEXT:  miopen.gridwise_gemm
+
+func @miopen_conv2d_bwd_weight_padALL(%filter : memref<1x20x8x3x3xf32>, %input : memref<7x1x8x32x32xf32>, %output : memref<7x1x20x30x30xf32>) {
+  miopen.conv2d_bwd_weight(%filter, %input, %output) {
+    arch = "gfx906",
+    num_cu = 64,
+    filter_layout = ["g", "k", "c", "y", "x"],
+    input_layout = ["ni", "gi", "ci", "hi", "wi"],
+    output_layout = ["no", "go", "ko", "ho", "wo"],
+    dilations = [1, 1],
+    strides = [1, 1],
+    padding = [0, 0, 0 ,0]
+  } : memref<1x20x8x3x3xf32>, memref<7x1x8x32x32xf32>, memref<7x1x20x30x30xf32>
+  return
+}
+// CHECK-LABEL: func {{@miopen_conv2d_bwd_weight_padALL.*%arg0.*%arg1.*%arg2}}
+// CHECK-NOT:   miopen.conv2d_bwd_weight
+// CHECK-NEXT:  {{miopen.transform\(%arg0\).* upper_layer_layout = \["gemmG", "gemmM", "gemmN"\].*}}
+// CHECK-NEXT:  {{miopen.transform.* upper_layer_layout = \["gemmG", "gemmMPad", "gemmNPad"\].*}}
+// CHECK-NEXT:  miopen.transform(%arg1)
+// CHECK-NEXT:  miopen.transform
+// CHECK-NEXT:  {{miopen.transform.* upper_layer_layout = \["gemmG", "gemmK", "gemmN"\].*}}
+// CHECK-NEXT:  {{miopen.transform.* upper_layer_layout = \["gemmG", "gemmKPad", "gemmNPad"\].*}}
+// CHECK-NEXT:  {{miopen.transform\(%arg2\).* upper_layer_layout = \["gemmG", "gemmK", "gemmM"\].*}}
+// CHECK-NEXT:  {{miopen.transform.* upper_layer_layout = \["gemmG", "gemmKPad", "gemmMPad"\].*}}
+// CHECK-NEXT:  miopen.gridwise_gemm
+
+func @miopen_conv2d_bwd_weight_padALL_f16(%filter : memref<1x20x8x3x3xf16>, %input : memref<7x1x8x32x32xf16>, %output : memref<7x1x20x30x30xf16>) {
+  miopen.conv2d_bwd_weight(%filter, %input, %output) {
+    arch = "gfx906",
+    num_cu = 64,
+    filter_layout = ["g", "k", "c", "y", "x"],
+    input_layout = ["ni", "gi", "ci", "hi", "wi"],
+    output_layout = ["no", "go", "ko", "ho", "wo"],
+    dilations = [1, 1],
+    strides = [1, 1],
+    padding = [0, 0, 0 ,0]
+  } : memref<1x20x8x3x3xf16>, memref<7x1x8x32x32xf16>, memref<7x1x20x30x30xf16>
+  return
+}
+// CHECK-LABEL: func {{@miopen_conv2d_bwd_weight_padALL_f16.*%arg0.*%arg1.*%arg2}}
+// CHECK-NOT:   miopen.conv2d_bwd_weight
+// CHECK-NEXT:  {{miopen.transform\(%arg0\).* upper_layer_layout = \["gemmG", "gemmM", "gemmN"\].*}}
+// CHECK-NEXT:  {{miopen.transform.* upper_layer_layout = \["gemmG", "gemmMPad", "gemmNPad"\].*}}
+// CHECK-NEXT:  miopen.transform(%arg1)
 // CHECK-NEXT:  miopen.transform
-// CHECK-NEXT:  miopen.transform(%arg2)
+// CHECK-NEXT:  {{miopen.transform.* upper_layer_layout = \["gemmG", "gemmK", "gemmN"\].*}}
+// CHECK-NEXT:  {{miopen.transform.* upper_layer_layout = \["gemmG", "gemmKPad", "gemmNPad"\].*}}
+// CHECK-NEXT:  {{miopen.transform\(%arg2\).* upper_layer_layout = \["gemmG", "gemmK", "gemmM"\].*}}
+// CHECK-NEXT:  {{miopen.transform.* upper_layer_layout = \["gemmG", "gemmKPad", "gemmMPad"\].*}}
 // CHECK-NEXT:  miopen.gridwise_gemm

mlir/test/mlir-miopen-driver/auto_e2e/padding_kernel_all_fwd.mlir

@@ -0,0 +1,21 @@
+// RUN: mlir-miopen-driver --operation=conv2d -t f32 -p=false  -fil_layout=gkcyx -in_layout=ngchw -out_layout=ngkhw -batchsize=20 -groupsize=1 -in_channels=3 -out_channels=6 -in_h=32 -in_w=32 -fil_h=7 -fil_w=7 --dilation_h=1 --dilation_w=1 --padding_h=3 --padding_w=3 --conv_stride_h=2 --conv_stride_w=2 -pv %random_data %xdlops -c | mlir-rocm-runner --shared-libs=%rocm_wrapper_library_dir/librocm-runtime-wrappers%shlibext,%linalg_test_lib_dir/libmlir_runner_utils%shlibext --entry-point-result=void | FileCheck %s --check-prefix=CHECK_RESNET50_CONFIG1
+
+// CHECK_RESNET50_CONFIG1: Unranked Memref base@ = 0x{{.*}} rank = 1 offset = 0 sizes = [1] strides = [1] data =
+// CHECK_RESNET50_CONFIG1: [1]
+
+// RUN: mlir-miopen-driver --operation=conv2d -t f16 -p=false  -fil_layout=gkcyx -in_layout=ngchw -out_layout=ngkhw -batchsize=20 -groupsize=1 -in_channels=3 -out_channels=6 -in_h=32 -in_w=32 -fil_h=7 -fil_w=7 --dilation_h=1 --dilation_w=1 --padding_h=3 --padding_w=3 --conv_stride_h=2 --conv_stride_w=2 -pv %random_data %xdlops -c | mlir-rocm-runner --shared-libs=%rocm_wrapper_library_dir/librocm-runtime-wrappers%shlibext,%linalg_test_lib_dir/libmlir_runner_utils%shlibext --entry-point-result=void | FileCheck %s --check-prefix=CHECK_RESNET50_CONFIG2
+
+// CHECK_RESNET50_CONFIG2: Unranked Memref base@ = 0x{{.*}} rank = 1 offset = 0 sizes = [1] strides = [1] data =
+// CHECK_RESNET50_CONFIG2: [1]
+
+// RUN: mlir-miopen-driver --operation=conv2d -t f32 -p=false  -fil_layout=gkyxc -in_layout=nhwgc -out_layout=nhwgk -batchsize=20 -groupsize=1 -in_channels=3 -out_channels=6 -in_h=32 -in_w=32 -fil_h=7 -fil_w=7 --dilation_h=1 --dilation_w=1 --padding_h=3 --padding_w=3 --conv_stride_h=2 --conv_stride_w=2 -pv %random_data %xdlops -c | mlir-rocm-runner --shared-libs=%rocm_wrapper_library_dir/librocm-runtime-wrappers%shlibext,%linalg_test_lib_dir/libmlir_runner_utils%shlibext --entry-point-result=void | FileCheck %s --check-prefix=CHECK_RESNET50_CONFIG3
+
+// CHECK_RESNET50_CONFIG3: Unranked Memref base@ = 0x{{.*}} rank = 1 offset = 0 sizes = [1] strides = [1] data =
+// CHECK_RESNET50_CONFIG3: [1]
+
+// RUN: mlir-miopen-driver --operation=conv2d -t f16 -p=false  -fil_layout=gkyxc -in_layout=nhwgc -out_layout=nhwgk -batchsize=20 -groupsize=1 -in_channels=3 -out_channels=6 -in_h=32 -in_w=32 -fil_h=7 -fil_w=7 --dilation_h=1 --dilation_w=1 --padding_h=3 --padding_w=3 --conv_stride_h=2 --conv_stride_w=2 -pv %random_data %xdlops -c | mlir-rocm-runner --shared-libs=%rocm_wrapper_library_dir/librocm-runtime-wrappers%shlibext,%linalg_test_lib_dir/libmlir_runner_utils%shlibext --entry-point-result=void | FileCheck %s --check-prefix=CHECK_RESNET50_CONFIG4
+
+// CHECK_RESNET50_CONFIG4: Unranked Memref base@ = 0x{{.*}} rank = 1 offset = 0 sizes = [1] strides = [1] data =
+// CHECK_RESNET50_CONFIG4: [1]
+
+