feat(miopen-provider): support standalone resample forward (pooling)

dnn-providers/integration-tests/SUPPORT_MATRIX_OVERVIEW.md

@@ -2,7 +2,7 @@
 
 <!-- Generated by python3 dnn-providers/integration-tests/scripts/support_features/feature_matrix.py; edit the script, not this file. -->
 
-Generated from the committed support claims (12353 graphs, 490 catalogued options). Percentages are shares of each op's API options, not of its graphs. Options count singly, not in combination.
+Generated from the committed support claims (12400 graphs, 490 catalogued options). Percentages are shares of each op's API options, not of its graphs. Options count singly, not in combination.
 
 **Jump to:** [RDNA (consumer)](#rdna-consumer) · [CDNA (data center)](#cdna-data-center) · [By engine](#by-engine) · [Op groups](#op-groups) · [Options per op](#options-per-op) · [GPU targets](#gpu-targets)
 
@@ -43,7 +43,7 @@ Generated from the committed support claims (12353 graphs, 490 catalogued option
 | **[Matmul](#group-matmul)** (35) | 🟡 37% | 🟡 37% | 🟡 37% |
 | **[Pointwise](#group-pointwise)** (58) | 🟡 20% | 🟡 20% | 🟡 20% |
 | **[Reduction](#group-reduction)** (26) | ❌ | ❌ | ❌ |
-| **[Resample](#group-resample)** (34) | 🟡 47% | 🟡 47% | 🟡 47% |
+| **[Resample](#group-resample)** (34) | 🟡 52% | 🟡 52% | 🟡 52% |
 | **[RMSNorm](#group-rmsnorm)** (38) | 🟡 86% | 🟡 86% | 🟡 86% |
 
 ## By engine
@@ -96,6 +96,7 @@ An op appears under an engine that claims at least one of its graphs somewhere.
 | **[ConvolutionFwd](#op-convolutionfwd)** (17) | 🟡 76% | 🟡 76% | 🟡 76% | 🟡 76% | ❌ | 🟡 76% | ❌ | 🟡 76% | 🟡 76% | 🟡 76% | ❌ | 🟡 76% |
 | **[ConvolutionWrw](#op-convolutionwrw)** (17) | 🟡 76% | 🟡 76% | 🟡 76% | 🟡 76% | ❌ | 🟡 76% | ❌ | 🟡 76% | 🟡 76% | 🟡 76% | ❌ | 🟡 76% |
 | **[Pointwise](#op-pointwise)** (58) | 🟡 13% | 🟡 17% | 🟡 17% | 🟡 17% | ❌ | 🟡 17% | ❌ | 🟡 17% | 🟡 17% | 🟡 17% | ❌ | 🟡 17% |
+| **[ResampleFwd](#op-resamplefwd)** (19) | 🟡 84% | 🟡 84% | 🟡 84% | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ | ❌ |
 
 ### MIOPEN_ENGINE_DETERMINISTIC
 
@@ -170,7 +171,7 @@ The API node types (`<Op>Attributes` in the flatbuffers schema) that make up eac
 | <a id="group-matmul"></a>**Matmul** | 461 | 35 | [Matmul](#op-matmul) (16), [MoeGroupedMatmul](#op-moegroupedmatmul) (13, **no test graph**), [MoeGroupedMatmulBwd](#op-moegroupedmatmulbwd) (6, **no test graph**) |
 | <a id="group-pointwise"></a>**Pointwise** | 6697 | 58 | [Pointwise](#op-pointwise) (58) |
 | <a id="group-reduction"></a>**Reduction** | 114 | 26 | [Reduction](#op-reduction) (26) |
-| <a id="group-resample"></a>**Resample** | 48 | 34 | [ResampleBwd](#op-resamplebwd) (15, **no test graph**), [ResampleFwd](#op-resamplefwd) (19) |
+| <a id="group-resample"></a>**Resample** | 95 | 34 | [ResampleBwd](#op-resamplebwd) (15, **no test graph**), [ResampleFwd](#op-resamplefwd) (19) |
 | <a id="group-rmsnorm"></a>**RMSNorm** | 5008 | 38 | [RMSNorm](#op-rmsnorm) (21), [RMSNormBackward](#op-rmsnormbackward) (17) |
 
 ## Options per op
@@ -199,7 +200,7 @@ Every API option the percentages count, catalogued per op in `dnn-providers/inte
 | <a id="op-pointwise"></a>**Pointwise** (58) | `mode:gelu_approx_tanh_fwd`, `mode:relu_fwd`, `mode:sigmoid_fwd`, `mode:swish_fwd`, `mode:tanh_fwd`, `mode:add`, `mode:max_op`, `mode:min_op`, `mode:mul`, `mode:relu_bwd`, `mode:sub`, `dtype:fp32`, `dtype:fp16`, `dtype:bf16`, `precision:mixed`, `broadcast:in`, `relu:lower_clip`, `relu:upper_clip`, `relu:leaky` | `mode:abs`, `mode:ceil`, `mode:elu_fwd`, `mode:erf`, `mode:exp`, `mode:floor`, `mode:gelu_fwd`, `mode:gen_index`, `mode:identity`, `mode:log`, `mode:logical_not`, `mode:neg`, `mode:reciprocal`, `mode:rsqrt`, `mode:sin`, `mode:softplus_fwd`, `mode:sqrt`, `mode:tan`, `mode:add_square`, `mode:cmp_eq`, `mode:cmp_ge`, `mode:cmp_gt`, `mode:cmp_le`, `mode:cmp_lt`, `mode:cmp_neq`, `mode:div`, `mode:elu_bwd`, `mode:gelu_approx_tanh_bwd`, `mode:gelu_bwd`, `mode:logical_and`, `mode:logical_or`, `mode:sigmoid_bwd`, `mode:softplus_bwd`, `mode:swish_bwd`, `mode:tanh_bwd`, `mode:binary_select`, `swish:beta`, `elu:alpha`, `softplus:beta` |
 | <a id="op-reduction"></a>**Reduction** (26) | `mode:add`, `mode:mul`, `mode:min`, `mode:max`, `mode:amax`, `mode:avg`, `mode:norm1`, `mode:norm2`, `mode:mul_no_zeros`, `dtype:fp32`, `dtype:fp16`, `dtype:bf16`, `dims:4`, `reduce:all`, `reduce:trailing`, `reduce:leading`, `layout:nchw`, `layout:nhwc`, `deterministic:off` | `precision:mixed_output`, `dims:1`, `dims:2`, `dims:3`, `dims:5`, `reduce:interleaved`, `deterministic:on` |
 | <a id="op-resamplebwd"></a>**ResampleBwd** (15) | — | `mode:maxpool`, `mode:avgpool_exclude_padding`, `mode:avgpool_include_padding`, `dtype:fp32`, `dtype:fp16`, `dtype:bf16`, `spatial:2d`, `spatial:3d`, `layout:nchw`, `layout:nhwc`, `padding:symmetric`, `padding:asymmetric`, `stride:strided`, `window:overlapping`, `window:gapped` |
-| <a id="op-resamplefwd"></a>**ResampleFwd** (19) | `mode:maxpool`, `mode:avgpool_exclude_padding`, `mode:avgpool_include_padding`, `padding_mode:zero_pad`, `dtype:fp32`, `dtype:fp16`, `dtype:bf16`, `spatial:2d`, `spatial:3d`, `layout:nchw`, `layout:nhwc`, `padding:symmetric`, `padding:asymmetric`, `stride:strided`, `window:overlapping`, `index:generated` | `padding_mode:neg_inf_pad`, `window:gapped`, `index:none` |
+| <a id="op-resamplefwd"></a>**ResampleFwd** (19) | `mode:maxpool`, `mode:avgpool_exclude_padding`, `mode:avgpool_include_padding`, `padding_mode:neg_inf_pad`, `padding_mode:zero_pad`, `dtype:fp32`, `dtype:fp16`, `dtype:bf16`, `spatial:2d`, `spatial:3d`, `layout:nchw`, `layout:nhwc`, `padding:symmetric`, `padding:asymmetric`, `stride:strided`, `window:overlapping`, `index:generated`, `index:none` | `window:gapped` |
 | <a id="op-rmsnorm"></a>**RMSNorm** (21) | `dtype:fp32`, `dtype:fp16`, `dtype:bf16`, `precision:mixed_io`, `precision:mixed_affine`, `precision:mixed_stats`, `dims:4`, `dims:5`, `layout:nchw`, `layout:nhwc`, `norm:last_dim`, `norm:multi_dim`, `norm:all_but_batch`, `bias:none`, `bias:bound`, `phase:inference`, `phase:training`, `epsilon:fixed`, `epsilon:runtime` | `dims:2`, `dims:3` |
 | <a id="op-rmsnormbackward"></a>**RMSNormBackward** (17) | `dtype:fp32`, `dtype:fp16`, `dtype:bf16`, `precision:mixed_io`, `precision:mixed_affine`, `precision:mixed_stats`, `dims:4`, `dims:5`, `layout:nchw`, `layout:nhwc`, `norm:last_dim`, `norm:multi_dim`, `norm:all_but_batch`, `dbias:none`, `dbias:bound` | `dims:2`, `dims:3` |
 | <a id="op-sdpa"></a>**Sdpa** (39) | `dtype:bf16`, `dtype:fp16`, `dtype:fp8e4m3`, `dtype:fp8e5m2`, `head_dim:64`, `head_dim:128`, `head_dim:192x128`, `head_dim:256`, `heads:mha`, `heads:gqa`, `heads:mqa`, `layout:bhsd`, `layout:bshd`, `layout:ragged`, `mask:none`, `mask:causal_top_left`, `mask:causal_bottom_right`, `mask:sliding_window`, `seq:q_eq_kv`, `seq:q_ne_kv`, `seq:decode`, `varlen:seq_lens`, `scale:fixed`, `scale:runtime`, `out:stats`, `out:max_sum_exp`, `fp8:amax`, `precision:mixed_output`, `bias:additive`, `bias:alibi`, `dropout:on`, `kv:paged`, `input:block_mask`, `input:sink_token` | `dtype:fp8e4m3_fnuz`, `dtype:fp8e5m2_fnuz`, `head_dim:192`, `scale:default`, `out:rng_dump` |

dnn-providers/integration-tests/integration-test-bundles/comprehensive/ResampleFwd/Default/support.json

@@ -38,6 +38,26 @@
           ]
         }
       }
+    ],
+    "MIOPEN_ENGINE": [
+      {
+        "cases": [
+          "1_2_3_2_2_bfp16_ncdhw_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0",
+          "1_2_4_3_bfp16_nchw_padding_modeneg_inf_pad_postpad0x0_prepad0x0",
+          "2_3_3_2_fp16_nchw_padding_modeneg_inf_pad_postpad0x0_prepad0x0"
+        ],
+        "support": {
+          "gfx90a": [
+            "linux"
+          ],
+          "gfx942": [
+            "linux"
+          ],
+          "gfx950": [
+            "linux"
+          ]
+        }
+      }
     ]
   },
   "version": 1

dnn-providers/integration-tests/integration-test-bundles/full/ResampleFwd/Default/support.json

@@ -40,6 +40,25 @@
           ]
         }
       }
+    ],
+    "MIOPEN_ENGINE": [
+      {
+        "cases": [
+          "1_2_3_2_2_fp16_ncdhw_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0",
+          "2_3_3_2_bfp16_nchw_padding_modeneg_inf_pad_postpad0x0_prepad0x0"
+        ],
+        "support": {
+          "gfx90a": [
+            "linux"
+          ],
+          "gfx942": [
+            "linux"
+          ],
+          "gfx950": [
+            "linux"
+          ]
+        }
+      }
     ]
   },
   "version": 1

dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/Default/support.json

@@ -38,6 +38,26 @@
           ]
         }
       }
+    ],
+    "MIOPEN_ENGINE": [
+      {
+        "cases": [
+          "1_2_3_2_2_fp32_ncdhw_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0",
+          "1_2_4_3_fp16_nchw_padding_modeneg_inf_pad_postpad0x0_prepad0x0",
+          "1_2_4_3_fp32_nchw_padding_modeneg_inf_pad_postpad0x0_prepad0x0"
+        ],
+        "support": {
+          "gfx90a": [
+            "linux"
+          ],
+          "gfx942": [
+            "linux"
+          ],
+          "gfx950": [
+            "linux"
+          ]
+        }
+      }
     ]
   },
   "version": 1

dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/SymmetricPadding/graph.template.json

@@ -0,0 +1,52 @@
+{
+  "compute_data_type": "float",
+  "intermediate_data_type": "float",
+  "io_data_type": "${case.io_data_type}",
+  "is_override_shape_enabled": false,
+  "min_required_engine_api_version": {
+    "major": 1,
+    "minor": 0,
+    "patch": 0
+  },
+  "name": "ResampleFwdTest",
+  "nodes": [
+    {
+      "compute_data_type": "float",
+      "generate_index": false,
+      "name": "ResampleFwd_0",
+      "padding_mode": "${case.attributes.padding_mode}",
+      "post_padding": "${case.attributes.post_padding}",
+      "pre_padding": "${case.attributes.pre_padding}",
+      "resample_mode": "${case.attributes.resample_mode}",
+      "stride": "${case.attributes.stride}",
+      "type": "ResampleFwdAttributes",
+      "window": "${case.attributes.window}",
+      "x_tensor_uid": 1,
+      "y_tensor_uid": 0
+    }
+  ],
+  "tensors": [
+    {
+      "alignment": 16,
+      "data_type": "${case.data_type}",
+      "dims": "${case.dims}",
+      "is_runtime_pass_by_value": false,
+      "name": "ResampleFwd_0::Y",
+      "ragged_offset_multiplier": 1,
+      "strides": "${case.strides}",
+      "uid": 0,
+      "virtual": false
+    },
+    {
+      "alignment": 16,
+      "data_type": "${case.data_type}",
+      "dims": "${case.dims}",
+      "is_runtime_pass_by_value": false,
+      "name": "X",
+      "ragged_offset_multiplier": 1,
+      "strides": "${case.strides}",
+      "uid": 1,
+      "virtual": false
+    }
+  ]
+}

dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/SymmetricPadding/support.json

@@ -0,0 +1,69 @@
+{
+  "claims": {
+    "MIOPEN_ENGINE": [
+      {
+        "cases": [
+          "1_2_16_16_16_fp32_ndhwc_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0",
+          "1_2_2_3_2_bfp16_ncdhw_padding_modezero_pad_postpad1x1x1_prepad1x1x1",
+          "1_2_2_3_2_bfp16_ndhwc_padding_modezero_pad_postpad1x1x1_prepad1x1x1",
+          "1_2_2_3_2_fp16_ncdhw_padding_modezero_pad_postpad1x1x1_prepad1x1x1",
+          "1_2_2_3_2_fp16_ndhwc_padding_modezero_pad_postpad1x1x1_prepad1x1x1",
+          "1_2_2_3_2_fp32_ncdhw_padding_modezero_pad_postpad1x1x1_prepad1x1x1",
+          "1_2_2_3_2_fp32_ndhwc_padding_modezero_pad_postpad1x1x1_prepad1x1x1",
+          "1_2_3_2_2_bfp16_ncdhw_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0",
+          "1_2_3_2_2_bfp16_ndhwc_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0",
+          "1_2_3_2_2_fp16_ncdhw_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0",
+          "1_2_3_2_2_fp16_ndhwc_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0",
+          "1_2_3_2_2_fp32_ncdhw_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0",
+          "1_2_3_2_2_fp32_ndhwc_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0",
+          "1_2_5_2_4_bfp16_ncdhw_padding_modezero_pad_postpad1x0x1_prepad1x0x1",
+          "1_2_5_2_4_bfp16_ndhwc_padding_modezero_pad_postpad1x0x1_prepad1x0x1",
+          "1_2_5_2_4_fp16_ncdhw_padding_modezero_pad_postpad1x0x1_prepad1x0x1",
+          "1_2_5_2_4_fp16_ndhwc_padding_modezero_pad_postpad1x0x1_prepad1x0x1",
+          "1_2_5_2_4_fp32_ncdhw_padding_modezero_pad_postpad1x0x1_prepad1x0x1",
+          "1_2_5_2_4_fp32_ndhwc_padding_modezero_pad_postpad1x0x1_prepad1x0x1",
+          "1_2_64_64_fp32_nhwc_padding_modezero_pad_postpad0x0_prepad0x0",
+          "2_3_3_2_bfp16_nchw_padding_modeneg_inf_pad_postpad0x0_prepad0x0",
+          "2_3_3_2_bfp16_nchw_padding_modezero_pad_postpad0x0_prepad0x0",
+          "2_3_3_2_fp16_nchw_padding_modeneg_inf_pad_postpad0x0_prepad0x0",
+          "2_3_3_2_fp16_nchw_padding_modezero_pad_postpad0x0_prepad0x0",
+          "2_3_3_2_fp32_nchw_padding_modeneg_inf_pad_postpad0x0_prepad0x0",
+          "2_3_3_2_fp32_nchw_padding_modezero_pad_postpad0x0_prepad0x0",
+          "2_3_4_3_bfp16_nchw_padding_modeneg_inf_pad_postpad1x1_prepad1x1",
+          "2_3_4_3_bfp16_nchw_padding_modezero_pad_postpad1x1_prepad1x1",
+          "2_3_4_3_bfp16_nhwc_padding_modeneg_inf_pad_postpad1x1_prepad1x1",
+          "2_3_4_3_bfp16_nhwc_padding_modezero_pad_postpad1x1_prepad1x1",
+          "2_3_4_3_fp16_nchw_padding_modeneg_inf_pad_postpad1x1_prepad1x1",
+          "2_3_4_3_fp16_nchw_padding_modezero_pad_postpad1x1_prepad1x1",
+          "2_3_4_3_fp16_nhwc_padding_modeneg_inf_pad_postpad1x1_prepad1x1",
+          "2_3_4_3_fp16_nhwc_padding_modezero_pad_postpad1x1_prepad1x1",
+          "2_3_4_3_fp32_nchw_padding_modeneg_inf_pad_postpad1x1_prepad1x1",
+          "2_3_4_3_fp32_nchw_padding_modezero_pad_postpad1x1_prepad1x1",
+          "2_3_4_3_fp32_nhwc_padding_modeneg_inf_pad_postpad1x1_prepad1x1",
+          "2_3_4_3_fp32_nhwc_padding_modezero_pad_postpad1x1_prepad1x1",
+          "2_3_4_4_bfp16_nchw_padding_modezero_pad_postpad1x0_prepad1x0",
+          "2_3_4_4_bfp16_nhwc_padding_modezero_pad_postpad1x0_prepad1x0",
+          "2_3_4_4_fp16_nchw_padding_modezero_pad_postpad1x0_prepad1x0",
+          "2_3_4_4_fp16_nhwc_padding_modezero_pad_postpad1x0_prepad1x0",
+          "2_3_4_4_fp32_nchw_padding_modezero_pad_postpad1x0_prepad1x0",
+          "2_3_4_4_fp32_nhwc_padding_modezero_pad_postpad1x0_prepad1x0",
+          "2_8_17_16_bfp16_nhwc_padding_modezero_pad_postpad1x1_prepad1x1",
+          "2_8_17_16_fp16_nhwc_padding_modezero_pad_postpad1x1_prepad1x1",
+          "2_8_17_16_fp32_nhwc_padding_modezero_pad_postpad1x1_prepad1x1"
+        ],
+        "support": {
+          "gfx90a": [
+            "linux"
+          ],
+          "gfx942": [
+            "linux"
+          ],
+          "gfx950": [
+            "linux"
+          ]
+        }
+      }
+    ]
+  },
+  "version": 1
+}