diff --git a/dnn-providers/integration-tests/SUPPORT_MATRIX_OVERVIEW.md b/dnn-providers/integration-tests/SUPPORT_MATRIX_OVERVIEW.md index 44f2a3c75633..d7dea68ead79 100644 --- a/dnn-providers/integration-tests/SUPPORT_MATRIX_OVERVIEW.md +++ b/dnn-providers/integration-tests/SUPPORT_MATRIX_OVERVIEW.md @@ -2,7 +2,7 @@ -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 (`Attributes` in the flatbuffers schema) that make up eac | **Matmul** | 461 | 35 | [Matmul](#op-matmul) (16), [MoeGroupedMatmul](#op-moegroupedmatmul) (13, **no test graph**), [MoeGroupedMatmulBwd](#op-moegroupedmatmulbwd) (6, **no test graph**) | | **Pointwise** | 6697 | 58 | [Pointwise](#op-pointwise) (58) | | **Reduction** | 114 | 26 | [Reduction](#op-reduction) (26) | -| **Resample** | 48 | 34 | [ResampleBwd](#op-resamplebwd) (15, **no test graph**), [ResampleFwd](#op-resamplefwd) (19) | +| **Resample** | 95 | 34 | [ResampleBwd](#op-resamplebwd) (15, **no test graph**), [ResampleFwd](#op-resamplefwd) (19) | | **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 | **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` | | **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` | | **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` | -| **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` | +| **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` | | **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` | | **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` | | **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` | diff --git a/dnn-providers/integration-tests/integration-test-bundles/comprehensive/ResampleFwd/Default/support.json b/dnn-providers/integration-tests/integration-test-bundles/comprehensive/ResampleFwd/Default/support.json index d4e6a9881ce7..d8c6c18cfa8b 100644 --- a/dnn-providers/integration-tests/integration-test-bundles/comprehensive/ResampleFwd/Default/support.json +++ b/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 diff --git a/dnn-providers/integration-tests/integration-test-bundles/full/ResampleFwd/Default/support.json b/dnn-providers/integration-tests/integration-test-bundles/full/ResampleFwd/Default/support.json index babb41a2cd77..0a9868c231a5 100644 --- a/dnn-providers/integration-tests/integration-test-bundles/full/ResampleFwd/Default/support.json +++ b/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 diff --git a/dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/Default/support.json b/dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/Default/support.json index c80f6f85403b..c7b7cffeb657 100644 --- a/dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/Default/support.json +++ b/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 diff --git a/dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/SymmetricPadding/graph.template.json b/dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/SymmetricPadding/graph.template.json new file mode 100644 index 000000000000..2b94e8877255 --- /dev/null +++ b/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 + } + ] +} diff --git a/dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/SymmetricPadding/support.json b/dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/SymmetricPadding/support.json new file mode 100644 index 000000000000..3792c3e454c7 --- /dev/null +++ b/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 +} diff --git a/dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/SymmetricPadding/sweep.json b/dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/SymmetricPadding/sweep.json new file mode 100644 index 000000000000..98a4f501f214 --- /dev/null +++ b/dnn-providers/integration-tests/integration-test-bundles/quick/ResampleFwd/SymmetricPadding/sweep.json @@ -0,0 +1,3394 @@ +{ + "version": 1, + "cases": [ + { + "id": "2_3_4_3_bfp16_nchw_padding_modezero_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 3 + ], + "strides": [ + 36, + 12, + 3, + 1 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/avg_exclude_2d_nchw_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_avg_exclude_2d_nchw_pad1_win3_stride2" + } + }, + { + "id": "2_8_17_16_bfp16_nhwc_padding_modezero_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 8, + 17, + 16 + ], + "strides": [ + 2176, + 1, + 128, + 8 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 2, + 8, + 33, + 31 + ], + "strides": [ + 8184, + 1, + 248, + 8 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/avg_exclude_2d_nhwc_large_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_avg_exclude_2d_nhwc_large_pad1_win3_stride2" + } + }, + { + "id": "2_3_4_3_bfp16_nhwc_padding_modezero_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 3 + ], + "strides": [ + 36, + 1, + 9, + 3 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 1, + 15, + 3 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/avg_exclude_2d_nhwc_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_avg_exclude_2d_nhwc_pad1_win3_stride2" + } + }, + { + "id": "1_2_2_3_2_bfp16_ncdhw_padding_modezero_pad_postpad1x1x1_prepad1x1x1", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 2, + 3, + 2 + ], + "strides": [ + 24, + 12, + 6, + 2, + 1 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 60, + 15, + 3, + 1 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1, + 1 + ], + "pre_padding": [ + 1, + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2, + 2 + ], + "window": [ + 3, + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/avg_exclude_3d_ncdhw_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_avg_exclude_3d_ncdhw_pad1_win3_stride2" + } + }, + { + "id": "1_2_2_3_2_bfp16_ndhwc_padding_modezero_pad_postpad1x1x1_prepad1x1x1", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 2, + 3, + 2 + ], + "strides": [ + 24, + 1, + 12, + 4, + 2 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 1, + 30, + 6, + 2 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1, + 1 + ], + "pre_padding": [ + 1, + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2, + 2 + ], + "window": [ + 3, + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/avg_exclude_3d_ndhwc_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_avg_exclude_3d_ndhwc_pad1_win3_stride2" + } + }, + { + "id": "2_3_4_4_bfp16_nchw_padding_modezero_pad_postpad1x0_prepad1x0", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 4 + ], + "strides": [ + 48, + 16, + 4, + 1 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 0 + ], + "pre_padding": [ + 1, + 0 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 2, + 1 + ], + "window": [ + 3, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/avg_include_2d_nchw_pad1x0_win3x2_stride2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_avg_include_2d_nchw_pad1x0_win3x2_stride2x1" + } + }, + { + "id": "2_3_4_4_bfp16_nhwc_padding_modezero_pad_postpad1x0_prepad1x0", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 4 + ], + "strides": [ + 48, + 1, + 12, + 3 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 1, + 15, + 3 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 0 + ], + "pre_padding": [ + 1, + 0 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 2, + 1 + ], + "window": [ + 3, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/avg_include_2d_nhwc_pad1x0_win3x2_stride2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_avg_include_2d_nhwc_pad1x0_win3x2_stride2x1" + } + }, + { + "id": "1_2_5_2_4_bfp16_ncdhw_padding_modezero_pad_postpad1x0x1_prepad1x0x1", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 5, + 2, + 4 + ], + "strides": [ + 80, + 40, + 8, + 4, + 1 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 60, + 15, + 3, + 1 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 0, + 1 + ], + "pre_padding": [ + 1, + 0, + 1 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 1, + 2, + 1 + ], + "window": [ + 2, + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/avg_include_3d_ncdhw_pad1x0x1_win2_stride1x2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_avg_include_3d_ncdhw_pad1x0x1_win2_stride1x2x1" + } + }, + { + "id": "1_2_5_2_4_bfp16_ndhwc_padding_modezero_pad_postpad1x0x1_prepad1x0x1", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 5, + 2, + 4 + ], + "strides": [ + 80, + 1, + 16, + 8, + 2 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 1, + 30, + 6, + 2 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 0, + 1 + ], + "pre_padding": [ + 1, + 0, + 1 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 1, + 2, + 1 + ], + "window": [ + 2, + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/avg_include_3d_ndhwc_pad1x0x1_win2_stride1x2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_avg_include_3d_ndhwc_pad1x0x1_win2_stride1x2x1" + } + }, + { + "id": "2_3_3_2_bfp16_nchw_padding_modeneg_inf_pad_postpad0x0_prepad0x0", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 3, + 2 + ], + "strides": [ + 18, + 6, + 2, + 1 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 0, + 0 + ], + "pre_padding": [ + 0, + 0 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 2, + 2 + ], + "window": [ + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/avg_include_neg_inf_pad_2d_nchw_pad0_win2_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_avg_include_neg_inf_pad_2d_nchw_pad0_win2_stride2" + } + }, + { + "id": "2_3_4_3_bfp16_nchw_padding_modeneg_inf_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 3 + ], + "strides": [ + 36, + 12, + 3, + 1 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "maxpool", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/max_2d_nchw_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_max_2d_nchw_pad1_win3_stride2" + } + }, + { + "id": "2_3_4_3_bfp16_nhwc_padding_modeneg_inf_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 3 + ], + "strides": [ + 36, + 1, + 9, + 3 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 1, + 15, + 3 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "maxpool", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/max_2d_nhwc_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_max_2d_nhwc_pad1_win3_stride2" + } + }, + { + "id": "1_2_3_2_2_bfp16_ncdhw_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 3, + 2, + 2 + ], + "strides": [ + 24, + 12, + 4, + 2, + 1 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 60, + 15, + 3, + 1 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 0, + 0, + 0 + ], + "pre_padding": [ + 0, + 0, + 0 + ], + "resample_mode": "maxpool", + "stride": [ + 1, + 2, + 1 + ], + "window": [ + 2, + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/max_3d_ncdhw_pad0_win2_stride1x2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_max_3d_ncdhw_pad0_win2_stride1x2x1" + } + }, + { + "id": "1_2_3_2_2_bfp16_ndhwc_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 3, + 2, + 2 + ], + "strides": [ + 24, + 1, + 8, + 4, + 2 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 1, + 30, + 6, + 2 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 0, + 0, + 0 + ], + "pre_padding": [ + 0, + 0, + 0 + ], + "resample_mode": "maxpool", + "stride": [ + 1, + 2, + 1 + ], + "window": [ + 2, + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/max_3d_ndhwc_pad0_win2_stride1x2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_max_3d_ndhwc_pad0_win2_stride1x2x1" + } + }, + { + "id": "2_3_3_2_bfp16_nchw_padding_modezero_pad_postpad0x0_prepad0x0", + "values": { + "io_data_type": "bfloat16", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 3, + 2 + ], + "strides": [ + 18, + 6, + 2, + 1 + ], + "data_type": "bfloat16" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "bfloat16" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 0, + 0 + ], + "pre_padding": [ + 0, + 0 + ], + "resample_mode": "maxpool", + "stride": [ + 2, + 2 + ], + "window": [ + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness/max_zero_pad_2d_nchw_pad0_win2_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricBfp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricBfp16.Correctness_max_zero_pad_2d_nchw_pad0_win2_stride2" + } + }, + { + "id": "2_3_4_3_fp16_nchw_padding_modezero_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 3 + ], + "strides": [ + 36, + 12, + 3, + 1 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/avg_exclude_2d_nchw_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_avg_exclude_2d_nchw_pad1_win3_stride2" + } + }, + { + "id": "2_8_17_16_fp16_nhwc_padding_modezero_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 8, + 17, + 16 + ], + "strides": [ + 2176, + 1, + 128, + 8 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 2, + 8, + 33, + 31 + ], + "strides": [ + 8184, + 1, + 248, + 8 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/avg_exclude_2d_nhwc_large_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_avg_exclude_2d_nhwc_large_pad1_win3_stride2" + } + }, + { + "id": "2_3_4_3_fp16_nhwc_padding_modezero_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 3 + ], + "strides": [ + 36, + 1, + 9, + 3 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 1, + 15, + 3 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/avg_exclude_2d_nhwc_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_avg_exclude_2d_nhwc_pad1_win3_stride2" + } + }, + { + "id": "1_2_2_3_2_fp16_ncdhw_padding_modezero_pad_postpad1x1x1_prepad1x1x1", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 2, + 3, + 2 + ], + "strides": [ + 24, + 12, + 6, + 2, + 1 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 60, + 15, + 3, + 1 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1, + 1 + ], + "pre_padding": [ + 1, + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2, + 2 + ], + "window": [ + 3, + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/avg_exclude_3d_ncdhw_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_avg_exclude_3d_ncdhw_pad1_win3_stride2" + } + }, + { + "id": "1_2_2_3_2_fp16_ndhwc_padding_modezero_pad_postpad1x1x1_prepad1x1x1", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 2, + 3, + 2 + ], + "strides": [ + 24, + 1, + 12, + 4, + 2 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 1, + 30, + 6, + 2 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1, + 1 + ], + "pre_padding": [ + 1, + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2, + 2 + ], + "window": [ + 3, + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/avg_exclude_3d_ndhwc_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_avg_exclude_3d_ndhwc_pad1_win3_stride2" + } + }, + { + "id": "2_3_4_4_fp16_nchw_padding_modezero_pad_postpad1x0_prepad1x0", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 4 + ], + "strides": [ + 48, + 16, + 4, + 1 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 0 + ], + "pre_padding": [ + 1, + 0 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 2, + 1 + ], + "window": [ + 3, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/avg_include_2d_nchw_pad1x0_win3x2_stride2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_avg_include_2d_nchw_pad1x0_win3x2_stride2x1" + } + }, + { + "id": "2_3_4_4_fp16_nhwc_padding_modezero_pad_postpad1x0_prepad1x0", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 4 + ], + "strides": [ + 48, + 1, + 12, + 3 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 1, + 15, + 3 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 0 + ], + "pre_padding": [ + 1, + 0 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 2, + 1 + ], + "window": [ + 3, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/avg_include_2d_nhwc_pad1x0_win3x2_stride2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_avg_include_2d_nhwc_pad1x0_win3x2_stride2x1" + } + }, + { + "id": "1_2_5_2_4_fp16_ncdhw_padding_modezero_pad_postpad1x0x1_prepad1x0x1", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 5, + 2, + 4 + ], + "strides": [ + 80, + 40, + 8, + 4, + 1 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 60, + 15, + 3, + 1 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 0, + 1 + ], + "pre_padding": [ + 1, + 0, + 1 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 1, + 2, + 1 + ], + "window": [ + 2, + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/avg_include_3d_ncdhw_pad1x0x1_win2_stride1x2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_avg_include_3d_ncdhw_pad1x0x1_win2_stride1x2x1" + } + }, + { + "id": "1_2_5_2_4_fp16_ndhwc_padding_modezero_pad_postpad1x0x1_prepad1x0x1", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 5, + 2, + 4 + ], + "strides": [ + 80, + 1, + 16, + 8, + 2 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 1, + 30, + 6, + 2 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 0, + 1 + ], + "pre_padding": [ + 1, + 0, + 1 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 1, + 2, + 1 + ], + "window": [ + 2, + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/avg_include_3d_ndhwc_pad1x0x1_win2_stride1x2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_avg_include_3d_ndhwc_pad1x0x1_win2_stride1x2x1" + } + }, + { + "id": "2_3_3_2_fp16_nchw_padding_modeneg_inf_pad_postpad0x0_prepad0x0", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 3, + 2 + ], + "strides": [ + 18, + 6, + 2, + 1 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 0, + 0 + ], + "pre_padding": [ + 0, + 0 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 2, + 2 + ], + "window": [ + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/avg_include_neg_inf_pad_2d_nchw_pad0_win2_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_avg_include_neg_inf_pad_2d_nchw_pad0_win2_stride2" + } + }, + { + "id": "2_3_4_3_fp16_nchw_padding_modeneg_inf_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 3 + ], + "strides": [ + 36, + 12, + 3, + 1 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "maxpool", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/max_2d_nchw_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_max_2d_nchw_pad1_win3_stride2" + } + }, + { + "id": "2_3_4_3_fp16_nhwc_padding_modeneg_inf_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 3 + ], + "strides": [ + 36, + 1, + 9, + 3 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 1, + 15, + 3 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "maxpool", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/max_2d_nhwc_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_max_2d_nhwc_pad1_win3_stride2" + } + }, + { + "id": "1_2_3_2_2_fp16_ncdhw_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 3, + 2, + 2 + ], + "strides": [ + 24, + 12, + 4, + 2, + 1 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 60, + 15, + 3, + 1 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 0, + 0, + 0 + ], + "pre_padding": [ + 0, + 0, + 0 + ], + "resample_mode": "maxpool", + "stride": [ + 1, + 2, + 1 + ], + "window": [ + 2, + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/max_3d_ncdhw_pad0_win2_stride1x2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_max_3d_ncdhw_pad0_win2_stride1x2x1" + } + }, + { + "id": "1_2_3_2_2_fp16_ndhwc_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 3, + 2, + 2 + ], + "strides": [ + 24, + 1, + 8, + 4, + 2 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 1, + 30, + 6, + 2 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 0, + 0, + 0 + ], + "pre_padding": [ + 0, + 0, + 0 + ], + "resample_mode": "maxpool", + "stride": [ + 1, + 2, + 1 + ], + "window": [ + 2, + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/max_3d_ndhwc_pad0_win2_stride1x2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_max_3d_ndhwc_pad0_win2_stride1x2x1" + } + }, + { + "id": "2_3_3_2_fp16_nchw_padding_modezero_pad_postpad0x0_prepad0x0", + "values": { + "io_data_type": "half", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 3, + 2 + ], + "strides": [ + 18, + 6, + 2, + 1 + ], + "data_type": "half" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "half" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 0, + 0 + ], + "pre_padding": [ + 0, + 0 + ], + "resample_mode": "maxpool", + "stride": [ + 2, + 2 + ], + "window": [ + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness/max_zero_pad_2d_nchw_pad0_win2_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp16", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp16.Correctness_max_zero_pad_2d_nchw_pad0_win2_stride2" + } + }, + { + "id": "2_3_4_3_fp32_nchw_padding_modezero_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 3 + ], + "strides": [ + 36, + 12, + 3, + 1 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/avg_exclude_2d_nchw_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_avg_exclude_2d_nchw_pad1_win3_stride2" + } + }, + { + "id": "2_8_17_16_fp32_nhwc_padding_modezero_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 8, + 17, + 16 + ], + "strides": [ + 2176, + 1, + 128, + 8 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 2, + 8, + 33, + 31 + ], + "strides": [ + 8184, + 1, + 248, + 8 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/avg_exclude_2d_nhwc_large_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_avg_exclude_2d_nhwc_large_pad1_win3_stride2" + } + }, + { + "id": "2_3_4_3_fp32_nhwc_padding_modezero_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 3 + ], + "strides": [ + 36, + 1, + 9, + 3 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 1, + 15, + 3 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/avg_exclude_2d_nhwc_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_avg_exclude_2d_nhwc_pad1_win3_stride2" + } + }, + { + "id": "1_2_2_3_2_fp32_ncdhw_padding_modezero_pad_postpad1x1x1_prepad1x1x1", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 2, + 3, + 2 + ], + "strides": [ + 24, + 12, + 6, + 2, + 1 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 60, + 15, + 3, + 1 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1, + 1 + ], + "pre_padding": [ + 1, + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2, + 2 + ], + "window": [ + 3, + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/avg_exclude_3d_ncdhw_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_avg_exclude_3d_ncdhw_pad1_win3_stride2" + } + }, + { + "id": "1_2_2_3_2_fp32_ndhwc_padding_modezero_pad_postpad1x1x1_prepad1x1x1", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 2, + 3, + 2 + ], + "strides": [ + 24, + 1, + 12, + 4, + 2 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 1, + 30, + 6, + 2 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 1, + 1 + ], + "pre_padding": [ + 1, + 1, + 1 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 2, + 2, + 2 + ], + "window": [ + 3, + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/avg_exclude_3d_ndhwc_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_avg_exclude_3d_ndhwc_pad1_win3_stride2" + } + }, + { + "id": "2_3_4_4_fp32_nchw_padding_modezero_pad_postpad1x0_prepad1x0", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 4 + ], + "strides": [ + 48, + 16, + 4, + 1 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 0 + ], + "pre_padding": [ + 1, + 0 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 2, + 1 + ], + "window": [ + 3, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/avg_include_2d_nchw_pad1x0_win3x2_stride2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_avg_include_2d_nchw_pad1x0_win3x2_stride2x1" + } + }, + { + "id": "2_3_4_4_fp32_nhwc_padding_modezero_pad_postpad1x0_prepad1x0", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 4 + ], + "strides": [ + 48, + 1, + 12, + 3 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 1, + 15, + 3 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 0 + ], + "pre_padding": [ + 1, + 0 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 2, + 1 + ], + "window": [ + 3, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/avg_include_2d_nhwc_pad1x0_win3x2_stride2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_avg_include_2d_nhwc_pad1x0_win3x2_stride2x1" + } + }, + { + "id": "1_2_5_2_4_fp32_ncdhw_padding_modezero_pad_postpad1x0x1_prepad1x0x1", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 5, + 2, + 4 + ], + "strides": [ + 80, + 40, + 8, + 4, + 1 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 60, + 15, + 3, + 1 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 0, + 1 + ], + "pre_padding": [ + 1, + 0, + 1 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 1, + 2, + 1 + ], + "window": [ + 2, + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/avg_include_3d_ncdhw_pad1x0x1_win2_stride1x2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_avg_include_3d_ncdhw_pad1x0x1_win2_stride1x2x1" + } + }, + { + "id": "1_2_5_2_4_fp32_ndhwc_padding_modezero_pad_postpad1x0x1_prepad1x0x1", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 5, + 2, + 4 + ], + "strides": [ + 80, + 1, + 16, + 8, + 2 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 1, + 30, + 6, + 2 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 1, + 0, + 1 + ], + "pre_padding": [ + 1, + 0, + 1 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 1, + 2, + 1 + ], + "window": [ + 2, + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/avg_include_3d_ndhwc_pad1x0x1_win2_stride1x2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_avg_include_3d_ndhwc_pad1x0x1_win2_stride1x2x1" + } + }, + { + "id": "2_3_3_2_fp32_nchw_padding_modeneg_inf_pad_postpad0x0_prepad0x0", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 3, + 2 + ], + "strides": [ + 18, + 6, + 2, + 1 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 0, + 0 + ], + "pre_padding": [ + 0, + 0 + ], + "resample_mode": "avgpool_include_padding", + "stride": [ + 2, + 2 + ], + "window": [ + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/avg_include_neg_inf_pad_2d_nchw_pad0_win2_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_avg_include_neg_inf_pad_2d_nchw_pad0_win2_stride2" + } + }, + { + "id": "2_3_4_3_fp32_nchw_padding_modeneg_inf_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 3 + ], + "strides": [ + 36, + 12, + 3, + 1 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "maxpool", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/max_2d_nchw_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_max_2d_nchw_pad1_win3_stride2" + } + }, + { + "id": "2_3_4_3_fp32_nhwc_padding_modeneg_inf_pad_postpad1x1_prepad1x1", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 4, + 3 + ], + "strides": [ + 36, + 1, + 9, + 3 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 1, + 15, + 3 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 1, + 1 + ], + "pre_padding": [ + 1, + 1 + ], + "resample_mode": "maxpool", + "stride": [ + 2, + 2 + ], + "window": [ + 3, + 3 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/max_2d_nhwc_pad1_win3_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_max_2d_nhwc_pad1_win3_stride2" + } + }, + { + "id": "1_2_3_2_2_fp32_ncdhw_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 3, + 2, + 2 + ], + "strides": [ + 24, + 12, + 4, + 2, + 1 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 60, + 15, + 3, + 1 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 0, + 0, + 0 + ], + "pre_padding": [ + 0, + 0, + 0 + ], + "resample_mode": "maxpool", + "stride": [ + 1, + 2, + 1 + ], + "window": [ + 2, + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/max_3d_ncdhw_pad0_win2_stride1x2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_max_3d_ncdhw_pad0_win2_stride1x2x1" + } + }, + { + "id": "1_2_3_2_2_fp32_ndhwc_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 3, + 2, + 2 + ], + "strides": [ + 24, + 1, + 8, + 4, + 2 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 4, + 5, + 3 + ], + "strides": [ + 120, + 1, + 30, + 6, + 2 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 0, + 0, + 0 + ], + "pre_padding": [ + 0, + 0, + 0 + ], + "resample_mode": "maxpool", + "stride": [ + 1, + 2, + 1 + ], + "window": [ + 2, + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/max_3d_ndhwc_pad0_win2_stride1x2x1", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_max_3d_ndhwc_pad0_win2_stride1x2x1" + } + }, + { + "id": "2_3_3_2_fp32_nchw_padding_modezero_pad_postpad0x0_prepad0x0", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 2, + 3, + 3, + 2 + ], + "strides": [ + 18, + 6, + 2, + 1 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 2, + 3, + 7, + 5 + ], + "strides": [ + 105, + 35, + 5, + 1 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 0, + 0 + ], + "pre_padding": [ + 0, + 0 + ], + "resample_mode": "maxpool", + "stride": [ + 2, + 2 + ], + "window": [ + 2, + 2 + ] + } + }, + "metadata": { + "format_version": 1, + "generator": "capture-bundles", + "generator_version": "1.0.0", + "notes": "Captured from C++ graph test Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness/max_zero_pad_2d_nchw_pad0_win2_stride2", + "operation": "Smoke/IntegrationGpuResampleForwardSymmetricFp32", + "seed": 1, + "reference_source": "c++ integration suite: Smoke/IntegrationGpuResampleForwardSymmetricFp32.Correctness_max_zero_pad_2d_nchw_pad0_win2_stride2" + } + }, + { + "id": "1_2_64_64_fp32_nhwc_padding_modezero_pad_postpad0x0_prepad0x0", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 1, + 1 + ], + "strides": [ + 2, + 1, + 2, + 2 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 64, + 64 + ], + "strides": [ + 8192, + 1, + 128, + 2 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "zero_pad", + "post_padding": [ + 0, + 0 + ], + "pre_padding": [ + 0, + 0 + ], + "resample_mode": "avgpool_exclude_padding", + "stride": [ + 64, + 64 + ], + "window": [ + 64, + 64 + ] + } + }, + "metadata": { + "format_version": 1, + "reference_source": "graph logged via HIPDNN_LOG_GRAPH_DIR from miopen-provider Smoke/IntegrationGpuResampleFwdPositiveControls.BuildsGraph/NhwcGlobalAvgPool", + "seed": 1, + "generator": "import_graph.py" + } + }, + { + "id": "1_2_16_16_16_fp32_ndhwc_padding_modeneg_inf_pad_postpad0x0x0_prepad0x0x0", + "values": { + "io_data_type": "float", + "tensors": [ + { + "uid": 0, + "dims": [ + 1, + 2, + 1, + 1, + 1 + ], + "strides": [ + 2, + 1, + 2, + 2, + 2 + ], + "data_type": "float" + }, + { + "uid": 1, + "dims": [ + 1, + 2, + 16, + 16, + 16 + ], + "strides": [ + 8192, + 1, + 512, + 32, + 2 + ], + "data_type": "float" + } + ], + "attributes": { + "padding_mode": "neg_inf_pad", + "post_padding": [ + 0, + 0, + 0 + ], + "pre_padding": [ + 0, + 0, + 0 + ], + "resample_mode": "maxpool", + "stride": [ + 16, + 16, + 16 + ], + "window": [ + 16, + 16, + 16 + ] + } + }, + "metadata": { + "format_version": 1, + "reference_source": "graph logged via HIPDNN_LOG_GRAPH_DIR from miopen-provider Smoke/IntegrationGpuResampleFwdPositiveControls.BuildsGraph/NdhwcGlobalMaxPool", + "seed": 1, + "generator": "import_graph.py" + } + } + ] +} diff --git a/dnn-providers/integration-tests/integration-test-bundles/standard/ResampleFwd/Default/support.json b/dnn-providers/integration-tests/integration-test-bundles/standard/ResampleFwd/Default/support.json index 8c8ad3b55983..95be95573405 100644 --- a/dnn-providers/integration-tests/integration-test-bundles/standard/ResampleFwd/Default/support.json +++ b/dnn-providers/integration-tests/integration-test-bundles/standard/ResampleFwd/Default/support.json @@ -36,6 +36,24 @@ ] } } + ], + "MIOPEN_ENGINE": [ + { + "cases": [ + "2_3_3_2_fp32_nchw_padding_modeneg_inf_pad_postpad0x0_prepad0x0" + ], + "support": { + "gfx90a": [ + "linux" + ], + "gfx942": [ + "linux" + ], + "gfx950": [ + "linux" + ] + } + } ] }, "version": 1 diff --git a/dnn-providers/miopen-provider/CMakeLists.txt b/dnn-providers/miopen-provider/CMakeLists.txt index e3946c853cec..9a2668bcc3cc 100644 --- a/dnn-providers/miopen-provider/CMakeLists.txt +++ b/dnn-providers/miopen-provider/CMakeLists.txt @@ -221,6 +221,9 @@ add_library( engines/plans/MiopenConvPlanBuilder.cpp engines/plans/MiopenConvWrwPlan.cpp engines/plans/MiopenPointwiseTensorChecks.cpp + engines/plans/MiopenResampleFwdChecks.cpp + engines/plans/MiopenResampleFwdPlan.cpp + engines/plans/MiopenResampleFwdPlanBuilder.cpp engines/plans/MiopenUnaryActivationChecks.cpp engines/plans/MiopenUnaryActivationPlan.cpp engines/plans/MiopenUnaryActivationPlanBuilder.cpp @@ -229,6 +232,7 @@ add_library( MiopenContainer.cpp MiopenConvDescriptor.cpp MiopenPluginPublic.cpp + MiopenPoolingDescriptor.cpp MiopenTensor.cpp MiopenUtils.cpp ) diff --git a/dnn-providers/miopen-provider/MiopenContainer.cpp b/dnn-providers/miopen-provider/MiopenContainer.cpp index 75679ca1c3c2..cc06d6710628 100644 --- a/dnn-providers/miopen-provider/MiopenContainer.cpp +++ b/dnn-providers/miopen-provider/MiopenContainer.cpp @@ -8,6 +8,7 @@ #include "engines/plans/MiopenBinaryPointwisePlanBuilder.hpp" #include "engines/plans/MiopenConvFwdBiasActivPlanBuilder.hpp" #include "engines/plans/MiopenConvPlanBuilder.hpp" +#include "engines/plans/MiopenResampleFwdPlanBuilder.hpp" #include "engines/plans/MiopenUnaryActivationPlanBuilder.hpp" #include @@ -52,6 +53,7 @@ const std::vector& MiopenContainer::getEngine engine->addPlanBuilder(std::make_unique(false)); engine->addPlanBuilder(std::make_unique()); engine->addPlanBuilder(std::make_unique()); + engine->addPlanBuilder(std::make_unique()); return engine; }}, @@ -68,6 +70,7 @@ const std::vector& MiopenContainer::getEngine engine->addPlanBuilder(std::make_unique(true)); engine->addPlanBuilder(std::make_unique()); engine->addPlanBuilder(std::make_unique()); + engine->addPlanBuilder(std::make_unique()); return engine; }} diff --git a/dnn-providers/miopen-provider/MiopenPoolingDescriptor.cpp b/dnn-providers/miopen-provider/MiopenPoolingDescriptor.cpp new file mode 100644 index 000000000000..4841acf2d01c --- /dev/null +++ b/dnn-providers/miopen-provider/MiopenPoolingDescriptor.cpp @@ -0,0 +1,100 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include + +#include + +#include "MiopenPoolingDescriptor.hpp" +#include "MiopenUtils.hpp" + +namespace miopen_plugin +{ + +MiopenPoolingDescriptor::MiopenPoolingDescriptor(miopenPoolingMode_t mode, + const std::vector& window, + const std::vector& padding, + const std::vector& stride, + miopenIndexType_t indexType) +{ + if(window.size() != padding.size() || window.size() != stride.size()) + { + throw hipdnn_plugin_sdk::HipdnnPluginException( + HIPDNN_PLUGIN_STATUS_BAD_PARAM, + "MiopenPoolingDescriptor: window, padding and stride must have the same length"); + } + + // MIOpen pooling handles 4D and 5D tensors only, i.e. 2 or 3 spatial dims. + if(window.size() != 2 && window.size() != 3) + { + throw hipdnn_plugin_sdk::HipdnnPluginException( + HIPDNN_PLUGIN_STATUS_BAD_PARAM, + "MiopenPoolingDescriptor: only 2 or 3 spatial dimensions are supported, got " + + std::to_string(window.size())); + } + + THROW_ON_MIOPEN_FAILURE(miopenCreatePoolingDescriptor(&_descriptor)); + + try + { + THROW_ON_MIOPEN_FAILURE(miopenSetNdPoolingDescriptor(_descriptor, + mode, + static_cast(window.size()), + window.data(), + padding.data(), + stride.data())); + + // Setting the N-d descriptor resets the index type and the workspace index mode, so both + // must be applied afterwards. MIOpen defaults to mask mode (an offset inside the window), + // which its N-d max pooling solver rejects when indices are saved. + THROW_ON_MIOPEN_FAILURE(miopenSetPoolingIndexType(_descriptor, indexType)); + THROW_ON_MIOPEN_FAILURE( + miopenSetPoolingWorkSpaceIndexMode(_descriptor, miopenPoolingWorkspaceIndexImage)); + } + catch(...) + { + LOG_ON_MIOPEN_FAILURE(miopenDestroyPoolingDescriptor(_descriptor)); + _descriptor = nullptr; + throw; + } +} + +MiopenPoolingDescriptor::MiopenPoolingDescriptor(MiopenPoolingDescriptor&& other) noexcept + : _descriptor(other._descriptor) +{ + other._descriptor = nullptr; +} + +MiopenPoolingDescriptor& + MiopenPoolingDescriptor::operator=(MiopenPoolingDescriptor&& other) noexcept +{ + if(this == &other) + { + return *this; + } + + if(_descriptor != nullptr) + { + LOG_ON_MIOPEN_FAILURE(miopenDestroyPoolingDescriptor(_descriptor)); + } + + _descriptor = other._descriptor; + other._descriptor = nullptr; + + return *this; +} + +MiopenPoolingDescriptor::~MiopenPoolingDescriptor() +{ + if(_descriptor != nullptr) + { + LOG_ON_MIOPEN_FAILURE(miopenDestroyPoolingDescriptor(_descriptor)); + } +} + +miopenPoolingDescriptor_t MiopenPoolingDescriptor::poolingDescriptor() const +{ + return _descriptor; +} + +} // namespace miopen_plugin diff --git a/dnn-providers/miopen-provider/MiopenPoolingDescriptor.hpp b/dnn-providers/miopen-provider/MiopenPoolingDescriptor.hpp new file mode 100644 index 000000000000..d76541e10a45 --- /dev/null +++ b/dnn-providers/miopen-provider/MiopenPoolingDescriptor.hpp @@ -0,0 +1,39 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#pragma once + +#include + +#include "MiopenApi.hpp" + +namespace miopen_plugin +{ + +// RAII owner of an N-d miopenPoolingDescriptor_t. Indices, when requested, are always recorded +// in image mode: a flattened spatial offset into the unpadded input, which is hipDNN's index +// convention. +class MiopenPoolingDescriptor +{ +public: + MiopenPoolingDescriptor(miopenPoolingMode_t mode, + const std::vector& window, + const std::vector& padding, + const std::vector& stride, + miopenIndexType_t indexType); + + MiopenPoolingDescriptor(const MiopenPoolingDescriptor&) = delete; + MiopenPoolingDescriptor& operator=(const MiopenPoolingDescriptor&) = delete; + + MiopenPoolingDescriptor(MiopenPoolingDescriptor&& other) noexcept; + MiopenPoolingDescriptor& operator=(MiopenPoolingDescriptor&& other) noexcept; + + ~MiopenPoolingDescriptor(); + + miopenPoolingDescriptor_t poolingDescriptor() const; + +private: + miopenPoolingDescriptor_t _descriptor{nullptr}; +}; + +} // namespace miopen_plugin diff --git a/dnn-providers/miopen-provider/docs/OperationSupport.md b/dnn-providers/miopen-provider/docs/OperationSupport.md index 70b3cce52da7..a617c984af0f 100644 --- a/dnn-providers/miopen-provider/docs/OperationSupport.md +++ b/dnn-providers/miopen-provider/docs/OperationSupport.md @@ -24,6 +24,7 @@ The following table lists all operations currently supported in hipDNN: | Convolution Wgrad | FP16, BFP16, FP32 | NCL, NLC, NCHW, NHWC, NCDHW, NDHWC | Cross-correlation only2,6, Deterministic5 | | Pointwise Activation (standalone) | FP16, FP32 | NCHW, NHWC | Single-node graph7 | | Pointwise Binary Elementwise (standalone) | FP16, FP32 | NCHW, NCDHW | Single-node graph8 | +| Resample Forward (pooling, standalone) | FP16, BFP16, FP32 | NCHW, NHWC, NCDHW, NDHWC | Single-node graph9, Deterministic5 | ¹ See Batchnorm Operations note below ² See Convolution Operations note below @@ -33,6 +34,7 @@ The following table lists all operations currently supported in hipDNN: ⁶ 3D tensors are internally padded to 4D for MIOpen compatibility. For convolution, the padding also extends the padding, stride and dilation vectors with a trailing spatial dimension of length 1. ⁷ See Standalone Activation note below ⁸ See Standalone Binary Elementwise note below +⁹ See Standalone Resample Forward note below ## Detailed Requirements @@ -93,6 +95,15 @@ The following table lists all operations currently supported in hipDNN: > must share one datatype. SUB is executed internally as Add with `in_1` scaled by −1. In-place > execution is not supported. +> [!NOTE] +> **Standalone Resample Forward:** A single-node graph supports MAXPOOL, AVGPOOL_EXCLUDE_PADDING +> and AVGPOOL_INCLUDE_PADDING across ranks 4-5, with packed X and Y sharing one datatype and +> layout, and an FP32 compute type. Padding must be symmetric and smaller than the window; unless +> it is all zero, MAXPOOL needs NEG_INF_PAD and the average modes need ZERO_PAD. Y must have the +> floor-mode output size. The index output (`generate_index`) is supported for MAXPOOL in +> NCHW/NCDHW only, as INT8 or INT32; with INT8 the input spatial size may not exceed 128. Very +> large windows on channels-last tensors are not supported. Max pooling ignores NaN inputs. + > [!NOTE] > **Sparse Support:** All operations currently work with dense tensors only. @@ -101,7 +112,7 @@ The following table lists all operations currently supported in hipDNN: ## Deterministic Engine Support -The MIOpen provider offers a deterministic execution engine (`MIOPEN_ENGINE_DETERMINISTIC`) for convolution, unary activation, and binary pointwise operations. This engine guarantees bit-reproducible results across multiple executions with the same inputs. +The MIOpen provider offers a deterministic execution engine (`MIOPEN_ENGINE_DETERMINISTIC`) for convolution, unary activation, binary pointwise and resample forward (pooling) operations. This engine guarantees bit-reproducible results across multiple executions with the same inputs. To use the deterministic engine, set it as the preferred engine on your graph before building: diff --git a/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdChecks.cpp b/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdChecks.cpp new file mode 100644 index 000000000000..e9f9a623ae93 --- /dev/null +++ b/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdChecks.cpp @@ -0,0 +1,468 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include +#include +#include +#include + +#include +#include +#include +#include +#include + +#include "engines/plans/MiopenResampleFwdChecks.hpp" + +namespace miopen_plugin::resample_fwd_applicability +{ + +using hipdnn_flatbuffers_sdk::data_objects::DataType; +using hipdnn_flatbuffers_sdk::data_objects::NodeAttributes; +using hipdnn_flatbuffers_sdk::data_objects::PaddingMode; +using hipdnn_flatbuffers_sdk::data_objects::ResampleFwdAttributes; +using hipdnn_flatbuffers_sdk::data_objects::ResampleMode; +using hipdnn_flatbuffers_sdk::data_objects::TensorAttributes; +using hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph; + +namespace +{ + +// The prefix isSupported() logs. Kept out of every throw below so it never doubles. +constexpr auto LOG_PREFIX = "Resample forward plan builder: "; + +[[noreturn]] void decline(const std::string& reason) +{ + throw hipdnn_plugin_sdk::HipdnnPluginException(HIPDNN_PLUGIN_STATUS_BAD_PARAM, reason); +} + +enum class Layout +{ + CHANNELS_FIRST, // NCHW / NCDHW + CHANNELS_LAST, // NHWC / NDHWC +}; + +struct TensorInfo +{ + const TensorAttributes* attr; + int64_t uid; + std::vector dims; + std::vector strides; +}; + +TensorInfo describeTensor(const IGraph& opGraph, int64_t uid, const std::string& name) +{ + const auto& tensorMap = opGraph.getTensorMap(); + const auto it = tensorMap.find(uid); + if(it == tensorMap.end() || it->second == nullptr) + { + decline(name + " tensor (uid " + std::to_string(uid) + + ") does not resolve in the graph's tensor map"); + } + + const auto* attr = it->second; + if(attr->virtual_()) + { + decline(name + " tensor must not be virtual"); + } + if(hipdnn_flatbuffers_sdk::utilities::isPassByValueTensor(attr)) + { + decline(name + " tensor must not be pass-by-value"); + } + if(attr->ragged_offset_tensor_uid().has_value()) + { + decline(name + " tensor must not be ragged"); + } + if(attr->dims() == nullptr || attr->strides() == nullptr + || attr->dims()->size() != attr->strides()->size()) + { + decline(name + " tensor must have one stride per dimension"); + } + + TensorInfo info{attr, + uid, + std::vector(attr->dims()->begin(), attr->dims()->end()), + std::vector(attr->strides()->begin(), attr->strides()->end())}; + + // MIOpen's pooling kernels address tensors with 32-bit offsets. + int64_t elementCount = 1; + for(const auto dim : info.dims) + { + if(dim <= 0) + { + decline(name + " tensor has a non-positive dimension"); + } + if(elementCount > std::numeric_limits::max() / dim) + { + decline(name + " tensor element count exceeds INT32_MAX"); + } + elementCount *= dim; + } + + return info; +} + +std::vector packedStrides(const std::vector& dims, Layout layout) +{ + using namespace hipdnn_data_sdk::utilities; + return layout == Layout::CHANNELS_FIRST ? generateStrides(dims) + : generateStrides(dims, strideOrderNhwc(dims.size())); +} + +// True when the tensor puts every element where a packed tensor of `layout` would. Length-1 +// axes are never stepped along, so their strides do not matter. +bool addressesLike(const TensorInfo& tensor, Layout layout) +{ + const auto expected = packedStrides(tensor.dims, layout); + for(size_t i = 0; i < tensor.dims.size(); ++i) + { + if(tensor.dims[i] > 1 && tensor.strides[i] != expected[i]) + { + return false; + } + } + return true; +} + +std::vector toSpatialVector(const flatbuffers::Vector* values, + size_t spatialRank, + const std::string& name, + int64_t minValue) +{ + if(values == nullptr || values->size() != spatialRank) + { + decline(name + " must have one entry per spatial dimension (" + std::to_string(spatialRank) + + ")"); + } + + std::vector result; + result.reserve(spatialRank); + for(const auto value : *values) + { + if(value < minValue || value > std::numeric_limits::max()) + { + decline(name + " values must be in [" + std::to_string(minValue) + ", INT_MAX], got " + + std::to_string(value)); + } + result.push_back(static_cast(value)); + } + return result; +} + +miopenPoolingMode_t toMiopenPoolingMode(ResampleMode mode) +{ + switch(mode) + { + case ResampleMode::MAXPOOL: + return miopenPoolingMax; + case ResampleMode::AVGPOOL_EXCLUDE_PADDING: + return miopenPoolingAverage; + case ResampleMode::AVGPOOL_INCLUDE_PADDING: + return miopenPoolingAverageInclusive; + default: + decline("unsupported resample mode " + + std::string(hipdnn_flatbuffers_sdk::data_objects::EnumNameResampleMode(mode))); + } +} + +miopenIndexType_t toMiopenIndexType(DataType dataType) +{ + // hipDNN's index is signed, MIOpen's unsigned; the range check in resolve() keeps every + // index MIOpen can write representable in the signed type. + switch(dataType) + { + case DataType::INT8: + return miopenIndexUint8; + case DataType::INT32: + return miopenIndexUint32; + default: + decline("index tensor data type must be INT8 or INT32, got " + + std::string(hipdnn_flatbuffers_sdk::data_objects::EnumNameDataType(dataType))); + } +} + +// MIOpen pools channels-last tensors only by transposing them around its tiled NCHW kernels +// (TransposedPoolingFwd2d/Nd); unlike channels-first there is no naive fallback. A tiled kernel is +// applicable only while its per-lane tile fits in scratch, sized as sizeof_private_memory() in +// MIOpen's src/solver/pooling/forward2d.cpp and forwardNd.cpp; past that the forward call fails +// with "No solver found". The limit assumes wave64, the smaller of MIOpen's two. +bool channelsLastTileFits(const ResampleFwdConfig& config) +{ + constexpr uint64_t K_MAX_WAVE_SCRATCH_SIZE = uint64_t{1024} * 8191; + constexpr uint64_t K_LANE_SCRATCH_LIMIT = K_MAX_WAVE_SCRATCH_SIZE / 64; + + auto roundUpTo16 = [](uint64_t bytes) { return (bytes + 15) / 16 * 16; }; + + // Every factor is at least 1, so stopping as soon as the tile outgrows the limit also keeps + // the product from overflowing. + uint64_t tileBytes = config.dataType == DataType::FLOAT ? 4 : 2; + auto grow = [&tileBytes](uint64_t factor) { + if(factor > K_LANE_SCRATCH_LIMIT || tileBytes * factor > K_LANE_SCRATCH_LIMIT) + { + return false; + } + tileBytes *= factor; + return true; + }; + + const auto inputsSpanned = [&config](size_t dim, uint64_t outputsPerLane) { + return (outputsPerLane - 1) * static_cast(config.stride[dim]) + + static_cast(config.window[dim]); + }; + + if(config.window.size() == 2) + { + // 16 output rows of one output column per lane, plus their 16 fp32 accumulators. + return grow(inputsSpanned(0, 16)) && grow(inputsSpanned(1, 1)) + && roundUpTo16(tileBytes) + roundUpTo16(sizeof(float) * 16) <= K_LANE_SCRATCH_LIMIT; + } + + // A 4 x 4 x 4 block of outputs per lane. + return grow(inputsSpanned(0, 4)) && grow(inputsSpanned(1, 4)) && grow(inputsSpanned(2, 4)) + && roundUpTo16(tileBytes) <= K_LANE_SCRATCH_LIMIT; +} + +} // namespace + +ResampleFwdConfig resolve(const IGraph& opGraph) +{ + if(opGraph.nodeCount() != 1) + { + decline("applicable only for single-node graphs. Graph has " + + std::to_string(opGraph.nodeCount()) + " nodes"); + } + + if(!opGraph.hasOnlySupportedAttributes( + std::set{NodeAttributes::ResampleFwdAttributes})) + { + decline("graph contains unsupported node attributes"); + } + + if(opGraph.getNode(0).compute_data_type() != DataType::FLOAT) + { + decline("only supports nodes with an fp32 compute_data_type"); + } + + const auto& attrs = opGraph.getNodeWrapper(0).attributesAs(); + + ResampleFwdConfig config; + config.mode = toMiopenPoolingMode(attrs.resample_mode()); + const bool isMaxPool = attrs.resample_mode() == ResampleMode::MAXPOOL; + + if(attrs.x_tensor_uid() == attrs.y_tensor_uid()) + { + decline("does not support in-place execution; X and Y must be distinct tensors"); + } + + const auto x = describeTensor(opGraph, attrs.x_tensor_uid(), "X"); + const auto y = describeTensor(opGraph, attrs.y_tensor_uid(), "Y"); + + const auto rank = x.dims.size(); + if(rank != 4 && rank != 5) + { + decline("X must be 4D or 5D, got rank " + std::to_string(rank) + + "; MIOpen pooling has no 1D path"); + } + if(y.dims.size() != rank) + { + decline("X and Y must have the same rank"); + } + + const auto dataType = x.attr->data_type(); + if(dataType != DataType::FLOAT && dataType != DataType::HALF && dataType != DataType::BFLOAT16) + { + decline("X data type must be FLOAT, HALF or BFLOAT16, got " + + std::string(hipdnn_flatbuffers_sdk::data_objects::EnumNameDataType(dataType))); + } + if(y.attr->data_type() != dataType) + { + decline("Y data type must match X"); + } + + const auto spatialRank = rank - 2; + config.window = toSpatialVector(attrs.window(), spatialRank, "window", 1); + config.stride = toSpatialVector(attrs.stride(), spatialRank, "stride", 1); + config.padding = toSpatialVector(attrs.pre_padding(), spatialRank, "pre_padding", 0); + + if(toSpatialVector(attrs.post_padding(), spatialRank, "post_padding", 0) != config.padding) + { + decline("asymmetric padding is not supported; MIOpen pools with one symmetric pad per " + "spatial dimension, so pre_padding must equal post_padding"); + } + + for(size_t i = 0; i < spatialRank; ++i) + { + if(config.padding[i] >= config.window[i]) + { + // A window that lies entirely in the padding is where hipDNN and MIOpen disagree: + // hipDNN's index sentinel is -1, MIOpen writes 0. + decline("padding must be smaller than the window in every spatial dimension, got " + "padding " + + std::to_string(config.padding[i]) + " and window " + + std::to_string(config.window[i]) + " at spatial dimension " + + std::to_string(i)); + } + } + + // MIOpen's max pooling can only select a padded position when no input in the window is above + // the lowest finite value of the type (all -inf, NaN or lowest), and its average pooling counts + // padded positions as zero. Any other padding_mode, PADDING_NOT_SET included, only agrees with + // that when there are no padded positions. + const auto implicitPaddingMode = isMaxPool ? PaddingMode::NEG_INF_PAD : PaddingMode::ZERO_PAD; + const bool hasPadding + = std::any_of(config.padding.begin(), config.padding.end(), [](int p) { return p > 0; }); + if(attrs.padding_mode() != implicitPaddingMode && hasPadding) + { + decline("padding_mode " + + std::string( + hipdnn_flatbuffers_sdk::data_objects::EnumNamePaddingMode(attrs.padding_mode())) + + " with resample mode " + + hipdnn_flatbuffers_sdk::data_objects::EnumNameResampleMode(attrs.resample_mode()) + + " requires all padding to be zero; MIOpen pads this mode as " + + hipdnn_flatbuffers_sdk::data_objects::EnumNamePaddingMode(implicitPaddingMode)); + } + + if(x.dims[0] != y.dims[0] || x.dims[1] != y.dims[1]) + { + decline("X and Y must have the same batch and channel dimensions"); + } + + for(size_t i = 0; i < spatialRank; ++i) + { + const auto paddedInput = x.dims[i + 2] + 2 * static_cast(config.padding[i]); + if(paddedInput < config.window[i]) + { + decline("the window must fit inside the padded input at spatial dimension " + + std::to_string(i)); + } + + // MIOpen does not check Y against its own (floor mode) output size formula. + const auto expected = (paddedInput - config.window[i]) / config.stride[i] + 1; + if(y.dims[i + 2] != expected) + { + decline("Y spatial dimension " + std::to_string(i) + " is " + + std::to_string(y.dims[i + 2]) + ", expected " + std::to_string(expected) + + " = floor((x + pre_padding + post_padding - window) / stride) + 1"); + } + + // With a single output the stride never steps, so the smallest stride that still gives + // one output changes nothing. MIOpen's estimates do depend on it: V2 sizes channels-last + // scratch for an input of y * stride - 2 * pad + window - 1, and the tiled kernels size + // their tile from the stride. Global pooling would otherwise overshoot both. + if(expected == 1) + { + config.stride[i] = static_cast(paddedInput - config.window[i] + 1); + } + } + + Layout layout = Layout::CHANNELS_FIRST; + if(addressesLike(x, Layout::CHANNELS_FIRST) && addressesLike(y, Layout::CHANNELS_FIRST)) + { + layout = Layout::CHANNELS_FIRST; + } + else if(addressesLike(x, Layout::CHANNELS_LAST) && addressesLike(y, Layout::CHANNELS_LAST)) + { + layout = Layout::CHANNELS_LAST; + } + else + { + decline("X and Y must both be packed channels-first (NCHW/NCDHW) or both packed " + "channels-last (NHWC/NDHWC)"); + } + + config.dataType = dataType; + if(layout == Layout::CHANNELS_LAST && !channelsLastTileFits(config)) + { + decline("the window and stride are too large for MIOpen's channels-last (NHWC/NDHWC) " + "pooling kernels"); + } + + config.xUid = x.uid; + config.xDims = x.dims; + config.xStrides = packedStrides(x.dims, layout); + config.yUid = y.uid; + config.yDims = y.dims; + config.yStrides = packedStrides(y.dims, layout); + + const bool generateIndex = attrs.generate_index().value_or(false); + const auto indexUid = attrs.index_tensor_uid(); + // Average pooling produces no index, so for it generate_index without an index tensor asks + // for nothing and is ignored; the frontend builds such graphs. + if(generateIndex && !indexUid.has_value() && isMaxPool) + { + decline("generate_index is set but the node has no index tensor"); + } + if(!generateIndex && indexUid.has_value()) + { + decline("the node has an index tensor but generate_index is not set"); + } + + if(indexUid.has_value()) + { + if(!isMaxPool) + { + decline("only max pooling produces indices"); + } + + // MIOpen's channels-last path transposes X and Y around a channels-first kernel but + // leaves the indices in packed NC(D)HW order. + if(layout != Layout::CHANNELS_FIRST) + { + decline("max pooling with generate_index requires channels-first (NCHW/NCDHW) X and " + "Y"); + } + + if(*indexUid == x.uid || *indexUid == y.uid) + { + decline("the index tensor must be distinct from X and Y"); + } + + const auto index = describeTensor(opGraph, *indexUid, "index"); + config.indexType = toMiopenIndexType(index.attr->data_type()); + + if(index.dims != y.dims) + { + decline("index tensor dimensions must equal Y dimensions"); + } + if(!addressesLike(index, Layout::CHANNELS_FIRST)) + { + decline("index tensor must be packed channels-first, the order MIOpen writes indices " + "in"); + } + + // Image-mode indices are flattened spatial offsets into X, from 0 to the X spatial size + // minus one. X's element count is already known to fit in INT32, so only an INT8 index + // can overflow. + int64_t spatialSize = 1; + for(size_t i = 2; i < rank; ++i) + { + spatialSize *= x.dims[i]; + } + if(index.attr->data_type() == DataType::INT8 + && spatialSize - 1 > std::numeric_limits::max()) + { + decline("X spatial size " + std::to_string(spatialSize) + + " exceeds the range of an INT8 index"); + } + + config.indexUid = indexUid; + } + + return config; +} + +bool isSupported(const IGraph& opGraph) +{ + try + { + resolve(opGraph); + return true; + } + catch(const std::exception& e) + { + HIPDNN_PLUGIN_LOG_INFO(LOG_PREFIX << e.what()); + return false; + } +} + +} // namespace miopen_plugin::resample_fwd_applicability diff --git a/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdChecks.hpp b/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdChecks.hpp new file mode 100644 index 000000000000..01baca10f752 --- /dev/null +++ b/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdChecks.hpp @@ -0,0 +1,58 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#pragma once + +#include +#include +#include + +#include +#include + +#include "MiopenApi.hpp" + +namespace miopen_plugin::resample_fwd_applicability +{ + +// Everything MiopenResampleFwdPlan needs, resolved from a graph the checks accepted. +// +// The X and Y strides are the canonical packed strides of the one layout both tensors share, +// not the graph's own strides. They address the same elements (only the strides of length-1 +// axes can differ), but MIOpen infers a tensor's layout from its strides and picks its native or +// transposing pooling path per tensor, while it sizes the transpose workspace from Y alone. +// Canonical strides make MIOpen classify X and Y identically. +struct ResampleFwdConfig +{ + miopenPoolingMode_t mode{miopenPoolingMax}; + std::vector window; + std::vector padding; // pre_padding == post_padding, MIOpen takes one symmetric pad + std::vector stride; + + hipdnn_flatbuffers_sdk::data_objects::DataType dataType{ + hipdnn_flatbuffers_sdk::data_objects::DataType::UNSET}; + + int64_t xUid{0}; + std::vector xDims; + std::vector xStrides; + + int64_t yUid{0}; + std::vector yDims; + std::vector yStrides; + + // Set when the node has generate_index. MIOpen writes the indices, packed in NC(D)HW order, + // to the front of its workspace; the plan copies them into this tensor. + std::optional indexUid; + miopenIndexType_t indexType{miopenIndexUint8}; +}; + +// Resolves a single-node ResampleFwd graph to a MIOpen pooling forward configuration. Throws +// HipdnnPluginException(HIPDNN_PLUGIN_STATUS_BAD_PARAM) naming the reason when the graph is not +// supported. Does not look at is_override_shape_enabled; the plan builder declines that first. +ResampleFwdConfig resolve(const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph& opGraph); + +// Returns whether resolve() accepts the graph, logging the reason at INFO with prefix +// "Resample forward plan builder: " if it does not. +bool isSupported(const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph& opGraph); + +} // namespace miopen_plugin::resample_fwd_applicability diff --git a/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlan.cpp b/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlan.cpp new file mode 100644 index 000000000000..94c93746a734 --- /dev/null +++ b/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlan.cpp @@ -0,0 +1,142 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include +#include +#include + +#include + +#include "MiopenUtils.hpp" +#include "engines/plans/MiopenResampleFwdPlan.hpp" + +namespace +{ + +size_t indexTypeSize(miopenIndexType_t indexType) +{ + switch(indexType) + { + case miopenIndexUint8: + return sizeof(uint8_t); + case miopenIndexUint16: + return sizeof(uint16_t); + case miopenIndexUint32: + return sizeof(uint32_t); + case miopenIndexUint64: + return sizeof(uint64_t); + default: + throw hipdnn_plugin_sdk::HipdnnPluginException(HIPDNN_PLUGIN_STATUS_INTERNAL_ERROR, + "MiopenResampleFwdPlan: unknown index type"); + } +} + +size_t elementCount(const std::vector& dims) +{ + return static_cast( + std::accumulate(dims.begin(), dims.end(), int64_t{1}, std::multiplies<>())); +} + +} // namespace + +namespace miopen_plugin +{ + +MiopenResampleFwdPlan::MiopenResampleFwdPlan( + const resample_fwd_applicability::ResampleFwdConfig& config) + : _x(config.xUid, config.dataType, config.xDims, config.xStrides) + , _y(config.yUid, config.dataType, config.yDims, config.yStrides) + , _pooling(config.mode, config.window, config.padding, config.stride, config.indexType) + , _indexUid(config.indexUid) +{ + if(_indexUid.has_value()) + { + _indexBytes = elementCount(config.yDims) * indexTypeSize(config.indexType); + } + + THROW_ON_MIOPEN_FAILURE(miopenPoolingGetWorkSpaceSizeV2( + _pooling.poolingDescriptor(), _y.tensorDescriptor(), &_workspaceSize)); + + if(_workspaceSize < _indexBytes) + { + throw hipdnn_plugin_sdk::HipdnnPluginException( + HIPDNN_PLUGIN_STATUS_INTERNAL_ERROR, + "MiopenResampleFwdPlan: MIOpen's pooling workspace (" + std::to_string(_workspaceSize) + + " bytes) cannot hold the " + std::to_string(_indexBytes) + " bytes of indices"); + } +} + +size_t + MiopenResampleFwdPlan::getWorkspaceSize([[maybe_unused]] const HipdnnMiopenHandle& handle) const +{ + return _workspaceSize; +} + +void MiopenResampleFwdPlan::execute(const HipdnnMiopenHandle& handle, + const hipdnnPluginDeviceBuffer_t* deviceBuffers, + uint32_t numDeviceBuffers, + void* workspace) const +{ + const auto xBuf + = hipdnn_plugin_sdk::findDeviceBuffer(_x.uid(), deviceBuffers, numDeviceBuffers); + const auto yBuf + = hipdnn_plugin_sdk::findDeviceBuffer(_y.uid(), deviceBuffers, numDeviceBuffers); + + void* indexPtr = nullptr; + if(_indexUid.has_value()) + { + indexPtr + = hipdnn_plugin_sdk::findDeviceBuffer(*_indexUid, deviceBuffers, numDeviceBuffers).ptr; + if(indexPtr == nullptr) + { + throw hipdnn_plugin_sdk::HipdnnPluginException( + HIPDNN_PLUGIN_STATUS_INVALID_VALUE, + "MiopenResampleFwdPlan: the index tensor has no device buffer"); + } + } + + // MIOpen never checks the size of the workspace it is given, and rejects a null one only for + // max pooling that saves indices: a missing or short workspace is otherwise an out-of-bounds + // write on the device, not an error. + if(_workspaceSize > 0 && workspace == nullptr) + { + throw hipdnn_plugin_sdk::HipdnnPluginException( + HIPDNN_PLUGIN_STATUS_INVALID_VALUE, + "MiopenResampleFwdPlan: requires " + std::to_string(_workspaceSize) + + " workspace bytes but none was provided"); + } + + const float alpha = 1.0f; + const float beta = 0.0f; + + THROW_ON_MIOPEN_FAILURE(miopenPoolingForward(handle.miopenHandle, + _pooling.poolingDescriptor(), + &alpha, + _x.tensorDescriptor(), + xBuf.ptr, + &beta, + _y.tensorDescriptor(), + yBuf.ptr, + _indexUid.has_value(), + workspace, + _workspaceSize)); + + if(_indexUid.has_value()) + { + // Enqueue on MIOpen's stream so the copy is ordered after the pooling kernel. + hipStream_t stream = nullptr; + THROW_ON_MIOPEN_FAILURE(miopenGetStream(handle.miopenHandle, &stream)); + + const auto status + = hipMemcpyAsync(indexPtr, workspace, _indexBytes, hipMemcpyDeviceToDevice, stream); + if(status != hipSuccess) + { + throw hipdnn_plugin_sdk::HipdnnPluginException( + HIPDNN_PLUGIN_STATUS_INTERNAL_ERROR, + std::string("MiopenResampleFwdPlan: copying the pooling indices failed: ") + + hipGetErrorString(status)); + } + } +} + +} // namespace miopen_plugin diff --git a/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlan.hpp b/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlan.hpp new file mode 100644 index 000000000000..a9991bc27f82 --- /dev/null +++ b/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlan.hpp @@ -0,0 +1,56 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#pragma once + +#include +#include +#include + +#include + +#include "HipdnnMiopenHandle.hpp" +#include "MiopenPoolingDescriptor.hpp" +#include "MiopenTensor.hpp" +#include "engines/plans/MiopenResampleFwdChecks.hpp" + +namespace miopen_plugin +{ + +// Execution plan for a standalone resample forward (pooling) node, backed by +// miopenPoolingForward. +// +// The workspace is MIOpen's own (miopenPoolingGetWorkSpaceSizeV2): room for the indices plus, +// for channels-last tensors, the transposed copies of X and Y. When the node generates an +// index, MIOpen leaves it in the first yElements * sizeof(index) bytes of the workspace and the +// plan copies it into the index tensor. +class MiopenResampleFwdPlan : public hipdnn_plugin_sdk::IPlan +{ +public: + explicit MiopenResampleFwdPlan(const resample_fwd_applicability::ResampleFwdConfig& config); + + MiopenResampleFwdPlan(const MiopenResampleFwdPlan&) = delete; + MiopenResampleFwdPlan& operator=(const MiopenResampleFwdPlan&) = delete; + + MiopenResampleFwdPlan(MiopenResampleFwdPlan&&) = delete; + MiopenResampleFwdPlan& operator=(MiopenResampleFwdPlan&&) = delete; + + ~MiopenResampleFwdPlan() override = default; + + size_t getWorkspaceSize(const HipdnnMiopenHandle& handle) const override; + + void execute(const HipdnnMiopenHandle& handle, + const hipdnnPluginDeviceBuffer_t* deviceBuffers, + uint32_t numDeviceBuffers, + void* workspace = nullptr) const override; + +private: + MiopenTensor _x; + MiopenTensor _y; + MiopenPoolingDescriptor _pooling; + std::optional _indexUid; + size_t _indexBytes{0}; + size_t _workspaceSize{0}; +}; + +} // namespace miopen_plugin diff --git a/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlanBuilder.cpp b/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlanBuilder.cpp new file mode 100644 index 000000000000..678d09070094 --- /dev/null +++ b/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlanBuilder.cpp @@ -0,0 +1,76 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include + +#include +#include + +#include "engines/plans/MiopenResampleFwdChecks.hpp" +#include "engines/plans/MiopenResampleFwdPlan.hpp" +#include "engines/plans/MiopenResampleFwdPlanBuilder.hpp" + +namespace miopen_plugin +{ + +bool MiopenResampleFwdPlanBuilder::isApplicable( + [[maybe_unused]] const HipdnnMiopenHandle& handle, + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph& opGraph) const +{ + if(opGraph.getGraph().is_override_shape_enabled()) + { + HIPDNN_PLUGIN_LOG_INFO("Resample forward plan builder does not support override shapes"); + return false; + } + + return resample_fwd_applicability::isSupported(opGraph); +} + +size_t MiopenResampleFwdPlanBuilder::getMaxWorkspaceSize( + const HipdnnMiopenHandle& handle, + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph& opGraph, + [[maybe_unused]] const HipdnnMiopenSettings& executionSettings) const +{ + const MiopenResampleFwdPlan plan(resample_fwd_applicability::resolve(opGraph)); + return plan.getWorkspaceSize(handle); +} + +void MiopenResampleFwdPlanBuilder::initializeExecutionSettings( + [[maybe_unused]] const HipdnnMiopenHandle& handle, + [[maybe_unused]] const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph& opGraph, + [[maybe_unused]] const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IEngineConfig& + engineConfig, + [[maybe_unused]] HipdnnMiopenSettings& executionSettings) const +{ +} + +void MiopenResampleFwdPlanBuilder::buildPlan( + [[maybe_unused]] const HipdnnMiopenHandle& handle, + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph& opGraph, + [[maybe_unused]] const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IEngineConfig& + engineConfig, + HipdnnMiopenContext& executionContext) const +{ + const auto& nodeWrapper = opGraph.getNodeWrapper(0); + const auto& attrs + = nodeWrapper.attributesAs(); + + HIPDNN_PLUGIN_LOG_INFO( + "Building resample forward " + << hipdnn_flatbuffers_sdk::data_objects::EnumNameResampleMode(attrs.resample_mode()) + << " plan for node: " << nodeWrapper.name()); + + auto plan + = std::make_unique(resample_fwd_applicability::resolve(opGraph)); + executionContext.setPlan(std::move(plan)); +} + +std::vector + MiopenResampleFwdPlanBuilder::getCustomKnobs( + [[maybe_unused]] const HipdnnMiopenHandle& handle, + [[maybe_unused]] const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph& opGraph) const +{ + return {}; +} + +} // namespace miopen_plugin diff --git a/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlanBuilder.hpp b/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlanBuilder.hpp new file mode 100644 index 000000000000..bec442d7a82b --- /dev/null +++ b/dnn-providers/miopen-provider/engines/plans/MiopenResampleFwdPlanBuilder.hpp @@ -0,0 +1,55 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#pragma once + +#include + +#include +#include + +#include "HipdnnMiopenContext.hpp" +#include "HipdnnMiopenHandle.hpp" +#include "HipdnnMiopenSettings.hpp" + +namespace miopen_plugin +{ + +// PlanBuilder for a standalone resample forward (pooling) node: MAXPOOL, +// AVGPOOL_EXCLUDE_PADDING and AVGPOOL_INCLUDE_PADDING, routed to miopenPoolingForward. +class MiopenResampleFwdPlanBuilder : public hipdnn_plugin_sdk::IPlanBuilder +{ +public: + MiopenResampleFwdPlanBuilder() = default; + ~MiopenResampleFwdPlanBuilder() override = default; + + MiopenResampleFwdPlanBuilder(const MiopenResampleFwdPlanBuilder&) = delete; + MiopenResampleFwdPlanBuilder& operator=(const MiopenResampleFwdPlanBuilder&) = delete; + + bool isApplicable( + const HipdnnMiopenHandle& handle, + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph& opGraph) const override; + + size_t getMaxWorkspaceSize(const HipdnnMiopenHandle& handle, + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph& opGraph, + const HipdnnMiopenSettings& executionSettings) const override; + + void initializeExecutionSettings( + const HipdnnMiopenHandle& handle, + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph& opGraph, + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IEngineConfig& engineConfig, + HipdnnMiopenSettings& executionSettings) const override; + + void buildPlan(const HipdnnMiopenHandle& handle, + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph& opGraph, + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IEngineConfig& engineConfig, + HipdnnMiopenContext& executionContext) const override; + + std::vector getCustomKnobs( + const HipdnnMiopenHandle& handle, + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph& opGraph) const override; +}; + +} // namespace miopen_plugin diff --git a/dnn-providers/miopen-provider/tests/CMakeLists.txt b/dnn-providers/miopen-provider/tests/CMakeLists.txt index 1c75613d4609..1aa0d3b103b5 100644 --- a/dnn-providers/miopen-provider/tests/CMakeLists.txt +++ b/dnn-providers/miopen-provider/tests/CMakeLists.txt @@ -36,6 +36,9 @@ add_executable( engines/plans/TestMiopenConvPlanBuilder.cpp engines/plans/TestMiopenConvWrwPlan.cpp engines/plans/TestMiopenPointwiseTensorChecks.cpp + engines/plans/TestMiopenResampleFwdChecks.cpp + engines/plans/TestMiopenResampleFwdPlan.cpp + engines/plans/TestMiopenResampleFwdPlanBuilder.cpp engines/plans/TestMiopenUnaryActivationPlanBuilder.cpp engines/TestMiopenEngine.cpp main.cpp @@ -45,6 +48,7 @@ add_executable( TestMiopenHipdnnEnginePluginHandle.cpp TestMiopenHipdnnPluginExecutionContext.cpp TestMiopenPluginPublic.cpp + TestMiopenPoolingDescriptor.cpp TestMiopenTensor.cpp TestMiopenUtils.cpp ) diff --git a/dnn-providers/miopen-provider/tests/TestMiopenEngineManager.cpp b/dnn-providers/miopen-provider/tests/TestMiopenEngineManager.cpp index 4d2960c7cd41..c48f7fed4df4 100644 --- a/dnn-providers/miopen-provider/tests/TestMiopenEngineManager.cpp +++ b/dnn-providers/miopen-provider/tests/TestMiopenEngineManager.cpp @@ -4,16 +4,21 @@ #include #include #include +#include #include +#include +#include #include #include +#include #include #include #include #include "HipdnnMiopenContext.hpp" #include "HipdnnMiopenHandle.hpp" +#include "MiopenContainer.hpp" #include "mocks/MockEngine.hpp" #include "mocks/MockHipdnnMiopenContext.hpp" @@ -243,3 +248,34 @@ TEST(TestMiopenEngineManager, InitializeExecutionContextThrowsOnInvalidEngineId) manager.initializeExecutionContext(dummyHandle, mockGraph, mockEngineConfig, execCtx), hipdnn_plugin_sdk::HipdnnPluginException); } + +// The real engine set: MIOpen's pooling forward is bit-reproducible (one thread per output, no +// atomics), so both the default and the deterministic engine take a ResampleFwd graph. +TEST(TestGpuMiopenContainer, DefaultAndDeterministicEnginesAcceptResampleFwd) +{ + SKIP_IF_NO_DEVICES(); + + using hipdnn_data_sdk::utilities::MIOPEN_ENGINE_DETERMINISTIC_ID; + using hipdnn_data_sdk::utilities::MIOPEN_ENGINE_ID; + using hipdnn_flatbuffers_sdk::data_objects::PaddingMode; + using hipdnn_flatbuffers_sdk::data_objects::ResampleMode; + using hipdnn_flatbuffers_sdk::flatbuffer_utilities::GraphWrapper; + + MiopenContainer container; + HipdnnMiopenHandle handle; + auto& manager = container.getEngineManager(); + + for(const auto mode : {ResampleMode::MAXPOOL, ResampleMode::AVGPOOL_EXCLUDE_PADDING}) + { + SCOPED_TRACE(hipdnn_flatbuffers_sdk::data_objects::EnumNameResampleMode(mode)); + auto fbb = createValidResampleFwdGraph( + mode == ResampleMode::MAXPOOL ? std::optional(true) : std::nullopt, + mode, + mode == ResampleMode::MAXPOOL ? PaddingMode::NEG_INF_PAD : PaddingMode::ZERO_PAD); + const GraphWrapper graph(fbb.GetBufferPointer(), fbb.GetSize()); + + const auto applicable = manager.getApplicableEngineIds(handle, graph); + EXPECT_THAT(applicable, ::testing::Contains(MIOPEN_ENGINE_ID)); + EXPECT_THAT(applicable, ::testing::Contains(MIOPEN_ENGINE_DETERMINISTIC_ID)); + } +} diff --git a/dnn-providers/miopen-provider/tests/TestMiopenPoolingDescriptor.cpp b/dnn-providers/miopen-provider/tests/TestMiopenPoolingDescriptor.cpp new file mode 100644 index 000000000000..0672be919ed3 --- /dev/null +++ b/dnn-providers/miopen-provider/tests/TestMiopenPoolingDescriptor.cpp @@ -0,0 +1,223 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include +#include +#include + +#include "MiopenApi.hpp" +#include +#include + +#include "MiopenPoolingDescriptor.hpp" + +using namespace miopen_plugin; + +struct PoolingDescriptorCase +{ + std::string name; + miopenPoolingMode_t mode; + std::vector window; + std::vector padding; + std::vector stride; + miopenIndexType_t indexType; +}; + +namespace +{ + +std::vector getRoundTripCases() +{ + return { + {"Max2dUint8", miopenPoolingMax, {2, 3}, {1, 0}, {2, 1}, miopenIndexUint8}, + {"Max2dUint32", miopenPoolingMax, {3, 2}, {0, 1}, {1, 2}, miopenIndexUint32}, + {"Average2d", miopenPoolingAverage, {3, 3}, {1, 1}, {2, 2}, miopenIndexUint8}, + {"AverageInclusive2d", + miopenPoolingAverageInclusive, + {2, 2}, + {0, 0}, + {2, 2}, + miopenIndexUint16}, + {"Max3dUint32", miopenPoolingMax, {2, 3, 4}, {1, 0, 2}, {2, 1, 3}, miopenIndexUint32}, + {"Max3dUint64", miopenPoolingMax, {1, 1, 1}, {0, 0, 0}, {1, 1, 1}, miopenIndexUint64}, + {"Average3d", miopenPoolingAverage, {3, 3, 3}, {1, 1, 1}, {1, 2, 3}, miopenIndexUint8}, + {"AverageInclusive3d", + miopenPoolingAverageInclusive, + {2, 4, 2}, + {1, 2, 1}, + {2, 2, 2}, + miopenIndexUint32}, + }; +} + +void expectBadParam(const std::vector& window, + const std::vector& padding, + const std::vector& stride) +{ + try + { + const MiopenPoolingDescriptor descriptor( + miopenPoolingMax, window, padding, stride, miopenIndexUint8); + ADD_FAILURE() << "expected HipdnnPluginException"; + } + catch(const hipdnn_plugin_sdk::HipdnnPluginException& e) + { + EXPECT_EQ(e.getStatus(), HIPDNN_PLUGIN_STATUS_BAD_PARAM); + } +} + +} // namespace + +class TestMiopenPoolingDescriptorRoundTrip : public ::testing::TestWithParam +{ +}; + +TEST_P(TestMiopenPoolingDescriptorRoundTrip, DescriptorReportsConfiguredValues) +{ + const auto& param = GetParam(); + const MiopenPoolingDescriptor descriptor( + param.mode, param.window, param.padding, param.stride, param.indexType); + ASSERT_NE(descriptor.poolingDescriptor(), nullptr); + + const auto spatialRank = static_cast(param.window.size()); + miopenPoolingMode_t mode{}; + int nbDims = 0; + std::vector window(param.window.size(), -1); + std::vector padding(param.window.size(), -1); + std::vector stride(param.window.size(), -1); + ASSERT_EQ(miopenGetNdPoolingDescriptor(descriptor.poolingDescriptor(), + spatialRank, + &mode, + &nbDims, + window.data(), + padding.data(), + stride.data()), + miopenStatusSuccess); + + EXPECT_EQ(mode, param.mode); + EXPECT_EQ(nbDims, spatialRank); + EXPECT_EQ(window, param.window); + EXPECT_EQ(padding, param.padding); + EXPECT_EQ(stride, param.stride); + + // miopenSetNdPoolingDescriptor resets both of these, so they only survive if the descriptor + // applies them afterwards. + miopenIndexType_t indexType{}; + ASSERT_EQ(miopenGetPoolingIndexType(descriptor.poolingDescriptor(), &indexType), + miopenStatusSuccess); + EXPECT_EQ(indexType, param.indexType); + + // Image mode for 2D too, where MIOpen's own default is mask mode. + miopenPoolingWorkspaceIndexMode_t indexMode{}; + ASSERT_EQ(miopenGetPoolingWorkSpaceIndexMode(descriptor.poolingDescriptor(), &indexMode), + miopenStatusSuccess); + EXPECT_EQ(indexMode, miopenPoolingWorkspaceIndexImage); +} + +INSTANTIATE_TEST_SUITE_P(AllCases, + TestMiopenPoolingDescriptorRoundTrip, + ::testing::ValuesIn(getRoundTripCases()), + [](const ::testing::TestParamInfo& info) { + return info.param.name; + }); + +TEST(TestMiopenPoolingDescriptor, ThrowsBadParamForPaddingLengthMismatch) +{ + expectBadParam({2, 2}, {0}, {2, 2}); +} + +TEST(TestMiopenPoolingDescriptor, ThrowsBadParamForStrideLengthMismatch) +{ + expectBadParam({2, 2}, {0, 0}, {2, 2, 2}); +} + +TEST(TestMiopenPoolingDescriptor, ThrowsBadParamForWindowLengthMismatch) +{ + expectBadParam({2, 2, 2}, {0, 0}, {2, 2}); +} + +TEST(TestMiopenPoolingDescriptor, ThrowsBadParamForNoSpatialDims) +{ + expectBadParam({}, {}, {}); +} + +TEST(TestMiopenPoolingDescriptor, ThrowsBadParamForOneSpatialDim) +{ + expectBadParam({2}, {0}, {2}); +} + +TEST(TestMiopenPoolingDescriptor, ThrowsBadParamForFourSpatialDims) +{ + expectBadParam({2, 2, 2, 2}, {0, 0, 0, 0}, {1, 1, 1, 1}); +} + +TEST(TestMiopenPoolingDescriptor, MoveConstructorTransfersOwnershipAndNullsSource) +{ + MiopenPoolingDescriptor source(miopenPoolingMax, {2, 2}, {0, 0}, {2, 2}, miopenIndexUint32); + auto* const handle = source.poolingDescriptor(); + ASSERT_NE(handle, nullptr); + + const MiopenPoolingDescriptor destination(std::move(source)); + + EXPECT_EQ(destination.poolingDescriptor(), handle); + // NOLINTNEXTLINE(bugprone-use-after-move) + EXPECT_EQ(source.poolingDescriptor(), nullptr); + + miopenIndexType_t indexType{}; + ASSERT_EQ(miopenGetPoolingIndexType(destination.poolingDescriptor(), &indexType), + miopenStatusSuccess); + EXPECT_EQ(indexType, miopenIndexUint32); +} + +TEST(TestMiopenPoolingDescriptor, MoveAssignmentTransfersOwnershipAndNullsSource) +{ + MiopenPoolingDescriptor source( + miopenPoolingAverage, {3, 3, 3}, {1, 1, 1}, {1, 1, 1}, miopenIndexUint8); + MiopenPoolingDescriptor destination( + miopenPoolingMax, {2, 2}, {0, 0}, {2, 2}, miopenIndexUint32); + auto* const handle = source.poolingDescriptor(); + ASSERT_NE(handle, nullptr); + ASSERT_NE(destination.poolingDescriptor(), handle); + + destination = std::move(source); + + EXPECT_EQ(destination.poolingDescriptor(), handle); + // NOLINTNEXTLINE(bugprone-use-after-move) + EXPECT_EQ(source.poolingDescriptor(), nullptr); + + miopenPoolingMode_t mode{}; + int nbDims = 0; + ASSERT_EQ(miopenGetNdPoolingDescriptor( + destination.poolingDescriptor(), 0, &mode, &nbDims, nullptr, nullptr, nullptr), + miopenStatusSuccess); + EXPECT_EQ(mode, miopenPoolingAverage); + EXPECT_EQ(nbDims, 3); +} + +TEST(TestMiopenPoolingDescriptor, MoveAssignmentFromMovedFromSourceLeavesBothNull) +{ + MiopenPoolingDescriptor first(miopenPoolingMax, {2, 2}, {0, 0}, {2, 2}, miopenIndexUint8); + const MiopenPoolingDescriptor keeper(std::move(first)); + MiopenPoolingDescriptor second(miopenPoolingMax, {2, 2}, {0, 0}, {2, 2}, miopenIndexUint8); + + // NOLINTNEXTLINE(bugprone-use-after-move) + second = std::move(first); + + EXPECT_EQ(second.poolingDescriptor(), nullptr); + // NOLINTNEXTLINE(bugprone-use-after-move) + EXPECT_EQ(first.poolingDescriptor(), nullptr); + EXPECT_NE(keeper.poolingDescriptor(), nullptr); +} + +TEST(TestMiopenPoolingDescriptor, SelfMoveAssignmentKeepsDescriptor) +{ + MiopenPoolingDescriptor descriptor(miopenPoolingMax, {2, 2}, {0, 0}, {2, 2}, miopenIndexUint8); + auto* const handle = descriptor.poolingDescriptor(); + + // Through a reference so the compiler does not reject the self-move outright. + auto& alias = descriptor; + descriptor = std::move(alias); + + // NOLINTNEXTLINE(bugprone-use-after-move) + EXPECT_EQ(descriptor.poolingDescriptor(), handle); +} diff --git a/dnn-providers/miopen-provider/tests/engines/plans/TestMiopenResampleFwdChecks.cpp b/dnn-providers/miopen-provider/tests/engines/plans/TestMiopenResampleFwdChecks.cpp new file mode 100644 index 000000000000..8b1a4d39ecfa --- /dev/null +++ b/dnn-providers/miopen-provider/tests/engines/plans/TestMiopenResampleFwdChecks.cpp @@ -0,0 +1,684 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +// Applicability rules of resample_fwd_applicability::resolve(), one test per rule, each matching +// its decline message, plus the boundaries and transformations the provider computes itself. None +// of these tests needs a GPU: resolve() only reads the graph. + +#include +#include +#include +#include +#include + +#include +#include + +#include +#include +#include + +#include "engines/plans/MiopenResampleFwdChecks.hpp" + +using namespace miopen_plugin; + +using hipdnn_flatbuffers_sdk::data_objects::DataType; +using hipdnn_flatbuffers_sdk::data_objects::PaddingMode; +using hipdnn_flatbuffers_sdk::data_objects::ResampleMode; +using hipdnn_flatbuffers_sdk::flatbuffer_utilities::GraphWrapper; +using hipdnn_test_sdk::utilities::createResampleFwdGraph; +using hipdnn_test_sdk::utilities::ResampleFwdGraphSpec; +using resample_fwd_applicability::ResampleFwdConfig; + +namespace +{ + +constexpr int64_t INT_MAX_VALUE = std::numeric_limits::max(); + +// 2x3x6x6 -> 2x3x3x3 with a 3x3 window, stride 2 and padding 1, so the padding rules have +// something to bite on. +ResampleFwdGraphSpec defaultSpec(ResampleMode mode = ResampleMode::MAXPOOL, + std::optional indexType = std::nullopt) +{ + return ResampleFwdGraphSpec::packed( + mode, {2, 3, 6, 6}, {3, 3}, {2, 2}, {1, 1}, DataType::FLOAT, indexType); +} + +ResampleFwdConfig resolveSpec(const ResampleFwdGraphSpec& spec) +{ + auto fbb = createResampleFwdGraph(spec); + const GraphWrapper graph(fbb.GetBufferPointer(), fbb.GetSize()); + return resample_fwd_applicability::resolve(graph); +} + +void expectAccepted(const ResampleFwdGraphSpec& spec) +{ + auto fbb = createResampleFwdGraph(spec); + const GraphWrapper graph(fbb.GetBufferPointer(), fbb.GetSize()); + ASSERT_TRUE(graph.isValid()); + + EXPECT_TRUE(resample_fwd_applicability::isSupported(graph)); + try + { + resample_fwd_applicability::resolve(graph); + } + catch(const hipdnn_plugin_sdk::HipdnnPluginException& e) + { + ADD_FAILURE() << "resolve() declined: " << e.what(); + } +} + +// Declined with BAD_PARAM, for the reason the message fragment names. The fragment guards against +// a decline test passing because some other rule happened to fire first. +void expectDeclined(const flatbuffers::FlatBufferBuilder& fbb, const std::string& reasonFragment) +{ + const GraphWrapper graph(fbb.GetBufferPointer(), fbb.GetSize()); + ASSERT_TRUE(graph.isValid()); + + EXPECT_FALSE(resample_fwd_applicability::isSupported(graph)); + try + { + resample_fwd_applicability::resolve(graph); + ADD_FAILURE() << "resolve() accepted a graph it should decline (" << reasonFragment << ")"; + } + catch(const hipdnn_plugin_sdk::HipdnnPluginException& e) + { + EXPECT_EQ(e.getStatus(), HIPDNN_PLUGIN_STATUS_BAD_PARAM); + EXPECT_THAT(e.what(), ::testing::HasSubstr(reasonFragment)); + } +} + +void expectDeclined(const ResampleFwdGraphSpec& spec, const std::string& reasonFragment) +{ + expectDeclined(createResampleFwdGraph(spec), reasonFragment); +} + +// Max pooling with a 1x1 window and stride, so Y matches X and only the X spatial size matters. +ResampleFwdGraphSpec identityWindowSpec(const std::vector& xDims, DataType indexType) +{ + const std::vector ones(xDims.size() - 2, 1); + const std::vector zeros(xDims.size() - 2, 0); + return ResampleFwdGraphSpec::packed( + ResampleMode::MAXPOOL, xDims, ones, ones, zeros, DataType::FLOAT, indexType); +} + +} // namespace + +// ============================================================================ +// Graph-level rules. +// ============================================================================ + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForMultiNodeGraph) +{ + auto spec = defaultSpec(); + spec.nodeCount = 2; + expectDeclined(spec, "single-node"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForOtherNodeType) +{ + expectDeclined(hipdnn_test_sdk::utilities::createValidResampleBwdGraph(), + "unsupported node attributes"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForNonFloatComputeType) +{ + auto spec = defaultSpec(); + spec.computeDataType = DataType::HALF; + expectDeclined(spec, "compute_data_type"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForModeNotSet) +{ + auto spec = defaultSpec(); + spec.mode = ResampleMode::NOT_SET; + expectDeclined(spec, "unsupported resample mode"); +} + +// ============================================================================ +// X and Y tensor rules. Every per-tensor rule is checked on X; Y goes through the same check. +// ============================================================================ + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForYAliasingX) +{ + auto spec = defaultSpec(); + spec.y.uid = spec.x.uid; + expectDeclined(spec, "in-place"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForUnresolvableX) +{ + auto spec = defaultSpec(); + spec.x.inGraph = false; + expectDeclined(spec, "X tensor (uid 1) does not resolve"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForUnresolvableY) +{ + auto spec = defaultSpec(); + spec.y.inGraph = false; + expectDeclined(spec, "Y tensor (uid 2) does not resolve"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForVirtualX) +{ + auto spec = defaultSpec(); + spec.x.isVirtual = true; + expectDeclined(spec, "X tensor must not be virtual"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForPassByValueX) +{ + auto spec = defaultSpec(); + spec.x.isRuntimePassByValue = true; + expectDeclined(spec, "X tensor must not be pass-by-value"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForRaggedX) +{ + auto spec = defaultSpec(); + spec.x.raggedOffsetTensorUid = 7; + expectDeclined(spec, "X tensor must not be ragged"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForMissingXStrides) +{ + auto spec = defaultSpec(); + spec.x.strides = std::nullopt; + expectDeclined(spec, "X tensor must have one stride per dimension"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForZeroDimInX) +{ + auto spec = defaultSpec(); + spec.x.dims[0] = 0; + expectDeclined(spec, "X tensor has a non-positive dimension"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForXElementCountAboveInt32Max) +{ + // 46341^2 > INT32_MAX; the stride collapses Y to a single element. + expectDeclined(ResampleFwdGraphSpec::packed( + ResampleMode::MAXPOOL, {1, 1, 46341, 46341}, {1, 1}, {46341, 46341}, {0, 0}), + "X tensor element count exceeds INT32_MAX"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForThreeDimensionalTensors) +{ + expectDeclined(ResampleFwdGraphSpec::packed(ResampleMode::MAXPOOL, {2, 3, 6}, {2}, {2}, {0}), + "X must be 4D or 5D"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForRankMismatch) +{ + auto spec = defaultSpec(); + spec.y.dims = {2, 3, 3, 3, 1}; + spec.y.strides = std::vector{27, 9, 3, 1, 1}; + expectDeclined(spec, "X and Y must have the same rank"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForUnsupportedDataType) +{ + auto spec = defaultSpec(); + spec.x.dataType = DataType::DOUBLE; + spec.y.dataType = DataType::DOUBLE; + expectDeclined(spec, "X data type must be FLOAT, HALF or BFLOAT16"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForMismatchedYDataType) +{ + auto spec = defaultSpec(); + spec.y.dataType = DataType::HALF; + expectDeclined(spec, "Y data type must match X"); +} + +// ============================================================================ +// Pooling parameter rules. +// ============================================================================ + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForWindowLengthMismatch) +{ + auto spec = defaultSpec(); + spec.window = std::vector{3, 3, 3}; + expectDeclined(spec, "window must have one entry per spatial dimension"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForZeroStride) +{ + auto spec = defaultSpec(); + spec.stride = std::vector{0, 2}; + expectDeclined(spec, "stride values must be in [1, INT_MAX]"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForWindowAboveIntMax) +{ + auto spec = defaultSpec(); + spec.window = std::vector{INT_MAX_VALUE + 1, 3}; + expectDeclined(spec, "window values must be in [1, INT_MAX]"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForAsymmetricPadding) +{ + auto spec = defaultSpec(); + spec.postPadding = std::vector{0, 0}; + expectDeclined(spec, "asymmetric padding"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForPaddingEqualToWindow) +{ + expectDeclined( + ResampleFwdGraphSpec::packed(ResampleMode::MAXPOOL, {2, 3, 6, 6}, {2, 2}, {2, 2}, {0, 2}), + "padding must be smaller than the window"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedTrueForPaddingOneBelowWindow) +{ + expectAccepted( + ResampleFwdGraphSpec::packed(ResampleMode::MAXPOOL, {1, 1, 4, 4}, {3, 3}, {1, 1}, {2, 2})); +} + +// MIOpen pads max pooling with -inf and average pooling with zeros. Any other padding_mode is only +// equivalent when no position is padded. +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForMismatchedPaddingModeWithPadding) +{ + auto spec = defaultSpec(ResampleMode::MAXPOOL); + spec.paddingMode = PaddingMode::ZERO_PAD; + expectDeclined(spec, "requires all padding to be zero"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedTrueForMismatchedPaddingModeWithoutPadding) +{ + auto spec = ResampleFwdGraphSpec::packed( + ResampleMode::AVGPOOL_INCLUDE_PADDING, {2, 3, 6, 6}, {3, 3}, {2, 2}, {0, 0}); + spec.paddingMode = PaddingMode::NEG_INF_PAD; + expectAccepted(spec); +} + +// ============================================================================ +// Shape rules. +// ============================================================================ + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForBatchMismatch) +{ + auto spec = defaultSpec(); + spec.y.dims = {1, 3, 3, 3}; + spec.y.strides = std::vector{27, 9, 3, 1}; + expectDeclined(spec, "same batch and channel dimensions"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForWindowLargerThanPaddedInput) +{ + // 2 + 2 * 1 < 5 in both spatial dimensions. + auto spec + = ResampleFwdGraphSpec::packed(ResampleMode::MAXPOOL, {1, 1, 4, 4}, {5, 5}, {1, 1}, {1, 1}); + spec.x.dims = {1, 1, 2, 2}; + spec.x.strides = std::vector{4, 4, 2, 1}; + spec.y.dims = {1, 1, 1, 1}; + spec.y.strides = std::vector{1, 1, 1, 1}; + expectDeclined(spec, "window must fit inside the padded input"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedTrueForWindowCoveringPaddedInputExactly) +{ + expectAccepted(ResampleFwdGraphSpec::packed( + ResampleMode::AVGPOOL_INCLUDE_PADDING, {1, 1, 2, 2}, {4, 4}, {1, 1}, {1, 1})); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForOutputOfOtherThanFloorModeSize) +{ + auto spec = defaultSpec(); + spec.y.dims = {2, 3, 4, 3}; + spec.y.strides = std::vector{36, 12, 3, 1}; + expectDeclined(spec, "Y spatial dimension 0 is 4, expected 3"); +} + +// ============================================================================ +// Layout rules. +// ============================================================================ + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForChannelsFirstXWithChannelsLastY) +{ + auto spec = defaultSpec(); + spec.y.strides = std::vector{27, 1, 9, 3}; + expectDeclined(spec, "both be packed channels-first"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedTrueForArbitraryStridesOnLengthOneAxes) +{ + auto spec + = ResampleFwdGraphSpec::packed(ResampleMode::MAXPOOL, {1, 1, 6, 6}, {3, 3}, {2, 2}, {1, 1}); + spec.x.strides = std::vector{7, 999, 6, 1}; + spec.y.strides = std::vector{0, 1, 3, 1}; + expectAccepted(spec); +} + +// ============================================================================ +// Index rules. +// ============================================================================ + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForMaxPoolGenerateIndexWithoutIndexTensor) +{ + auto spec = defaultSpec(ResampleMode::MAXPOOL); + spec.generateIndex = true; + expectDeclined(spec, "generate_index is set but the node has no index tensor"); +} + +// Average pooling produces no index: the flag without an index tensor asks for nothing and is +// ignored. Graph::resample() builds this graph for average pooling. +TEST(TestMiopenResampleFwdChecks, IsSupportedTrueForAveragePoolingGenerateIndexWithoutIndexTensor) +{ + auto spec = defaultSpec(ResampleMode::AVGPOOL_EXCLUDE_PADDING); + spec.generateIndex = true; + expectAccepted(spec); + EXPECT_FALSE(resolveSpec(spec).indexUid.has_value()); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForIndexTensorWithoutGenerateIndex) +{ + auto spec = defaultSpec(ResampleMode::MAXPOOL, DataType::INT32); + spec.generateIndex = std::nullopt; + expectDeclined(spec, "generate_index is not set"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForIndexWithAveragePooling) +{ + expectDeclined(defaultSpec(ResampleMode::AVGPOOL_EXCLUDE_PADDING, DataType::INT32), + "only max pooling produces indices"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForIndexWithChannelsLastTensors) +{ + auto spec = defaultSpec(ResampleMode::MAXPOOL, DataType::INT32); + spec.makeChannelsLast(); + expectDeclined(spec, "requires channels-first"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForIndexAliasingX) +{ + auto spec = defaultSpec(ResampleMode::MAXPOOL, DataType::INT32); + spec.index->uid = spec.x.uid; + expectDeclined(spec, "index tensor must be distinct from X and Y"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForUnresolvableIndex) +{ + auto spec = defaultSpec(ResampleMode::MAXPOOL, DataType::INT32); + spec.index->inGraph = false; + expectDeclined(spec, "index tensor (uid 3) does not resolve"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForUnsupportedIndexDataType) +{ + expectDeclined(defaultSpec(ResampleMode::MAXPOOL, DataType::INT64), + "index tensor data type must be INT8 or INT32"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForIndexDimsMismatch) +{ + auto spec = defaultSpec(ResampleMode::MAXPOOL, DataType::INT32); + spec.index->dims = {2, 3, 3, 4}; + spec.index->strides = std::vector{36, 12, 4, 1}; + expectDeclined(spec, "index tensor dimensions must equal Y dimensions"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForChannelsLastIndex) +{ + auto spec = defaultSpec(ResampleMode::MAXPOOL, DataType::INT32); + spec.index->strides = std::vector{27, 1, 9, 3}; + expectDeclined(spec, "index tensor must be packed channels-first"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedTrueForInt8IndexAtSpatialSize128) +{ + // Indices run up to 8 * 16 - 1 = 127, the largest INT8 value. N and C do not count. + expectAccepted(identityWindowSpec({1, 2, 8, 16}, DataType::INT8)); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForInt8IndexAtSpatialSize129) +{ + expectDeclined(identityWindowSpec({1, 2, 3, 43}, DataType::INT8), + "X spatial size 129 exceeds the range of"); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedTrueForInt32IndexAtSpatialSize129) +{ + expectAccepted(identityWindowSpec({1, 2, 3, 43}, DataType::INT32)); +} + +// ============================================================================ +// resolve() output. +// ============================================================================ + +TEST(TestMiopenResampleFwdChecks, ResolveMapsResampleModesToMiopenModes) +{ + EXPECT_EQ(resolveSpec(defaultSpec(ResampleMode::MAXPOOL)).mode, miopenPoolingMax); + EXPECT_EQ(resolveSpec(defaultSpec(ResampleMode::AVGPOOL_EXCLUDE_PADDING)).mode, + miopenPoolingAverage); + EXPECT_EQ(resolveSpec(defaultSpec(ResampleMode::AVGPOOL_INCLUDE_PADDING)).mode, + miopenPoolingAverageInclusive); +} + +TEST(TestMiopenResampleFwdChecks, ResolveReportsPoolingParameters) +{ + // Y: H = (7 + 2 - 3) / 2 + 1 = 4, W = (5 + 0 - 2) / 1 + 1 = 4. + const auto config = resolveSpec(ResampleFwdGraphSpec::packed( + ResampleMode::AVGPOOL_EXCLUDE_PADDING, {1, 2, 7, 5}, {3, 2}, {2, 1}, {1, 0})); + EXPECT_EQ(config.window, (std::vector{3, 2})); + EXPECT_EQ(config.padding, (std::vector{1, 0})); + EXPECT_EQ(config.stride, (std::vector{2, 1})); +} + +TEST(TestMiopenResampleFwdChecks, ResolveReportsDataTypeUidsAndDims) +{ + auto spec = ResampleFwdGraphSpec::packed(ResampleMode::MAXPOOL, + {2, 3, 6, 6}, + {3, 3}, + {2, 2}, + {1, 1}, + DataType::BFLOAT16, + DataType::INT32); + spec.x.uid = 10; + spec.y.uid = 20; + spec.index->uid = 30; + + const auto config = resolveSpec(spec); + EXPECT_EQ(config.dataType, DataType::BFLOAT16); + EXPECT_EQ(config.xUid, 10); + EXPECT_EQ(config.yUid, 20); + EXPECT_EQ(config.indexUid, std::optional(30)); + EXPECT_EQ(config.xDims, (std::vector{2, 3, 6, 6})); + EXPECT_EQ(config.yDims, (std::vector{2, 3, 3, 3})); +} + +TEST(TestMiopenResampleFwdChecks, ResolveReportsIndexTypes) +{ + const auto int32Config = resolveSpec(defaultSpec(ResampleMode::MAXPOOL, DataType::INT32)); + EXPECT_EQ(int32Config.indexUid, std::optional(ResampleFwdGraphSpec::INDEX_UID)); + EXPECT_EQ(int32Config.indexType, miopenIndexUint32); + + const auto int8Config = resolveSpec(defaultSpec(ResampleMode::MAXPOOL, DataType::INT8)); + EXPECT_EQ(int8Config.indexType, miopenIndexUint8); + + EXPECT_FALSE(resolveSpec(defaultSpec(ResampleMode::MAXPOOL)).indexUid.has_value()); +} + +TEST(TestMiopenResampleFwdChecks, ResolveReportsChannelsLastStrides) +{ + auto spec = defaultSpec(); + spec.makeChannelsLast(); + const auto config = resolveSpec(spec); + EXPECT_EQ(config.xStrides, (std::vector{108, 1, 18, 3})); + EXPECT_EQ(config.yStrides, (std::vector{27, 1, 9, 3})); +} + +TEST(TestMiopenResampleFwdChecks, ResolveReplacesStridesOfLengthOneAxes) +{ + auto spec + = ResampleFwdGraphSpec::packed(ResampleMode::MAXPOOL, {1, 1, 6, 6}, {3, 3}, {2, 2}, {1, 1}); + spec.x.strides = std::vector{7, 999, 6, 1}; + spec.y.strides = std::vector{0, 1, 3, 1}; + + const auto config = resolveSpec(spec); + EXPECT_EQ(config.xStrides, (std::vector{36, 36, 6, 1})); + EXPECT_EQ(config.yStrides, (std::vector{9, 9, 3, 1})); +} + +TEST(TestMiopenResampleFwdChecks, ResolvePrefersChannelsFirstWhenBothLayoutsMatch) +{ + // With C == 1, NHWC strides address like NCHW for both tensors. + auto spec = ResampleFwdGraphSpec::packed( + ResampleMode::AVGPOOL_EXCLUDE_PADDING, {2, 1, 4, 4}, {2, 2}, {2, 2}, {0, 0}); + spec.makeChannelsLast(); + + const auto config = resolveSpec(spec); + EXPECT_EQ(config.xStrides, (std::vector{16, 16, 4, 1})); + EXPECT_EQ(config.yStrides, (std::vector{4, 4, 2, 1})); +} + +TEST(TestMiopenResampleFwdChecks, ResolveFollowsChannelsLastInputForAmbiguousOutput) +{ + // 2x4x3x3 -> 2x4x1x1 with channels-last X. Y has only length-1 spatial axes, so its + // channels-first-style strides match both layouts; X settles it as channels-last. + auto spec + = ResampleFwdGraphSpec::packed(ResampleMode::MAXPOOL, {2, 4, 3, 3}, {3, 3}, {1, 1}, {0, 0}); + spec.makeChannelsLast(); + spec.y.strides = std::vector{4, 1, 1, 1}; + + const auto config = resolveSpec(spec); + EXPECT_EQ(config.xStrides, (std::vector{36, 1, 12, 4})); + EXPECT_EQ(config.yStrides, (std::vector{4, 1, 4, 4})); +} + +// ============================================================================ +// Stride normalization: a spatial dimension with a single output never steps its window, so +// resolve() replaces its stride by the smallest one that still yields one output, +// x + 2 * pad - window + 1. Dimensions with more outputs keep their stride. +// ============================================================================ + +TEST(TestMiopenResampleFwdChecks, ResolveNormalizesStrideOfGlobalPooling) +{ + const auto config = resolveSpec(ResampleFwdGraphSpec::packed( + ResampleMode::MAXPOOL, {1, 2, 64, 64}, {64, 64}, {64, 64}, {0, 0})); + EXPECT_EQ(config.yDims, (std::vector{1, 2, 1, 1})); + EXPECT_EQ(config.stride, (std::vector{1, 1})); +} + +TEST(TestMiopenResampleFwdChecks, ResolveNormalizesStrideAgainstPaddedInput) +{ + // 5 + 2 * 1 - 3 + 1 = 5; the padding itself is left alone. + const auto config = resolveSpec(ResampleFwdGraphSpec::packed( + ResampleMode::MAXPOOL, {1, 1, 5, 5}, {3, 3}, {10, 10}, {1, 1})); + EXPECT_EQ(config.stride, (std::vector{5, 5})); + EXPECT_EQ(config.padding, (std::vector{1, 1})); +} + +TEST(TestMiopenResampleFwdChecks, ResolveNormalizesOnlySingleOutputDims) +{ + // H: one output, 7 - 3 + 1 = 5. W: (12 - 3) / 4 + 1 = 3 outputs, stride kept. + const auto config = resolveSpec(ResampleFwdGraphSpec::packed( + ResampleMode::MAXPOOL, {1, 1, 7, 12}, {3, 3}, {100, 4}, {0, 0})); + EXPECT_EQ(config.yDims, (std::vector{1, 1, 1, 3})); + EXPECT_EQ(config.stride, (std::vector{5, 4})); +} + +// ============================================================================ +// Channels-last tile limit. MIOpen pools NHWC/NDHWC only through its tiled Forward2d/ForwardNd +// kernels, which are applicable while their per-lane tile fits 1024 * 8191 / 64 = 131056 bytes: +// 2D: roundUp16(esize * ((16 - 1) * stride_h + window_h) * window_w) + 64 +// 3D: roundUp16(esize * prod((4 - 1) * stride + window)) +// with the strides after normalization. Channels-first graphs fall back to the naive kernel. +// ============================================================================ + +namespace +{ + +constexpr const char* TILE_LIMIT_REASON = "too large for MIOpen's channels-last"; + +ResampleFwdGraphSpec tileSpec(const std::vector& xDims, + const std::vector& window, + const std::vector& stride, + bool channelsLast, + DataType dataType, + const std::optional>& padding = std::nullopt) +{ + auto spec + = ResampleFwdGraphSpec::packed(ResampleMode::MAXPOOL, + xDims, + window, + stride, + padding.value_or(std::vector(window.size(), 0)), + dataType); + if(channelsLast) + { + spec.makeChannelsLast(); + } + return spec; +} + +} // namespace + +TEST(TestMiopenResampleFwdChecks, IsSupportedTrueForPlanarFloatTileAtLimit) +{ + // 4 * (15 * 545 + 12) * 4 = 130992, + 64 = 131056. H keeps two outputs so its stride is not + // normalized. + expectAccepted(tileSpec({1, 2, 557, 4}, {12, 4}, {545, 1}, true, DataType::FLOAT)); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForPlanarFloatTileOverLimit) +{ + // One window row more: 4 * 8188 * 4 = 131008, + 64 = 131072. + expectDeclined(tileSpec({1, 2, 558, 4}, {13, 4}, {545, 1}, true, DataType::FLOAT), + TILE_LIMIT_REASON); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedTrueForPlanarFloatTileOverLimitInChannelsFirst) +{ + expectAccepted(tileSpec({1, 2, 558, 4}, {13, 4}, {545, 1}, false, DataType::FLOAT)); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedTrueForPlanar16BitTileAtLimit) +{ + // 2 * (15 * 545 + 12) * 8 = 130992, + 64 = 131056. + expectAccepted(tileSpec({1, 2, 557, 8}, {12, 8}, {545, 1}, true, DataType::HALF)); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForFloatTileThatFitsAs16Bit) +{ + // The 16-bit at-limit shape needs twice the bytes in fp32. + expectDeclined(tileSpec({1, 2, 557, 8}, {12, 8}, {545, 1}, true, DataType::FLOAT), + TILE_LIMIT_REASON); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedTrueForVolumetricFloatTileUnderLimit) +{ + // Window 7, stride 8: 4 * 31^3 = 119164. + expectAccepted(tileSpec({1, 2, 16, 16, 16}, {7, 7, 7}, {8, 8, 8}, true, DataType::FLOAT)); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForVolumetricFloatTileOverLimit) +{ + // Window 8, stride 8: 4 * 32^3 = 131072. + expectDeclined(tileSpec({1, 2, 16, 16, 16}, {8, 8, 8}, {8, 8, 8}, true, DataType::FLOAT), + TILE_LIMIT_REASON); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedTrueForChannelsLastGlobalPoolingThanksToNormalization) +{ + // Unnormalized, stride 64 would need 4 * (15 * 64 + 64) * 64 = 262144 bytes; normalized to 1, + // 4 * (15 + 64) * 64 = 20224. + expectAccepted(tileSpec({1, 2, 64, 64}, {64, 64}, {64, 64}, true, DataType::FLOAT)); +} + +TEST(TestMiopenResampleFwdChecks, IsSupportedFalseForChannelsLastIntMaxWindowWithoutOverflow) +{ + // Padding INT_MAX - 1 lets an INT_MAX window and stride produce two outputs, so nothing is + // normalized and the tile estimate would overflow without the early stop. + const auto big = INT_MAX_VALUE; + expectDeclined(tileSpec({1, 2, 2, 2}, + {big, 1}, + {big, 1}, + true, + DataType::FLOAT, + std::vector{big - 1, 0}), + TILE_LIMIT_REASON); +} diff --git a/dnn-providers/miopen-provider/tests/engines/plans/TestMiopenResampleFwdPlan.cpp b/dnn-providers/miopen-provider/tests/engines/plans/TestMiopenResampleFwdPlan.cpp new file mode 100644 index 000000000000..e4316440d75c --- /dev/null +++ b/dnn-providers/miopen-provider/tests/engines/plans/TestMiopenResampleFwdPlan.cpp @@ -0,0 +1,452 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +// Executes MiopenResampleFwdPlan on small inputs with values worked out by hand. Each test fails +// when the plan's own code is wrong: the descriptor or tensors it builds from the configuration, +// the workspace it reports and checks, the alpha and beta it passes, or the index it copies out. +// How MIOpen pools is MIOpen's to test. + +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +#include "MiopenApi.hpp" +#include +#include +#include +#include +#include +#include +#include +#include + +#include "HipdnnMiopenHandle.hpp" +#include "engines/plans/MiopenResampleFwdChecks.hpp" +#include "engines/plans/MiopenResampleFwdPlan.hpp" + +using namespace miopen_plugin; + +using hipdnn_data_sdk::types::half; +using hipdnn_data_sdk::utilities::iterateAlongDimensions; +using hipdnn_data_sdk::utilities::Tensor; +using hipdnn_flatbuffers_sdk::data_objects::DataType; +using hipdnn_flatbuffers_sdk::data_objects::ResampleMode; +using hipdnn_flatbuffers_sdk::flatbuffer_utilities::GraphWrapper; +using hipdnn_test_sdk::utilities::createResampleFwdGraph; +using hipdnn_test_sdk::utilities::ResampleFwdGraphSpec; + +namespace +{ + +constexpr float Y_SENTINEL = -12345.0f; +constexpr int INDEX_SENTINEL = -7; + +std::vector ascending(int64_t count, float first = 0.0f) +{ + std::vector values(static_cast(count)); + std::iota(values.begin(), values.end(), first); + return values; +} + +// Y and the indices in row-major logical order over the Y dims; `index` is empty when the node +// has no index tensor. +struct PoolingResult +{ + std::vector y; + std::vector index; +}; + +void expectYNear(const std::vector& actual, + const std::vector& expected, + double relativeTolerance) +{ + ASSERT_EQ(actual.size(), expected.size()); + for(size_t i = 0; i < actual.size(); ++i) + { + const auto want = static_cast(expected[i]); + EXPECT_NEAR( + static_cast(actual[i]), want, relativeTolerance * std::max(1.0, std::abs(want))) + << "Y element " << i; + } +} + +void expectMatches(const PoolingResult& actual, + const PoolingResult& expected, + double relativeTolerance = 1e-5) +{ + expectYNear(actual.y, expected.y, relativeTolerance); + EXPECT_EQ(actual.index, expected.index); +} + +// Host/device tensors for one execution: X filled from logical (row-major) values, Y and the +// index pre-filled with sentinels so unwritten elements show up. +template +class PoolingTensors +{ +public: + PoolingTensors(const ResampleFwdGraphSpec& spec, const std::vector& xValues) + : _spec(spec) + , _x(spec.x.dims, spec.x.strides.value()) + , _y(spec.y.dims, spec.y.strides.value()) + { + size_t next = 0; + iterateAlongDimensions(_spec.x.dims, [&](const std::vector& index) { + _x.setHostValue(static_cast(xValues.at(next++)), index); + }); + _x.markHostModified(); + + _y.fillWithValue(static_cast(Y_SENTINEL)); + + if(_spec.index.has_value()) + { + _index.emplace(_spec.index->dims, _spec.index->strides.value()); + _index->fillWithValue(static_cast(INDEX_SENTINEL)); + } + } + + std::vector deviceBuffers(bool includeIndex = true) + { + std::vector buffers{{_spec.x.uid, _x.rawDeviceData()}, + {_spec.y.uid, _y.rawDeviceData()}}; + if(includeIndex && _index.has_value()) + { + buffers.push_back({_spec.index->uid, _index->rawDeviceData()}); + } + return buffers; + } + + // Reads Y and the index back in logical order. The device work must have completed. + PoolingResult read() + { + PoolingResult result; + + _y.markDeviceModified(); + _y.rawHostData(); + iterateAlongDimensions(_spec.y.dims, [&](const std::vector& index) { + result.y.push_back(static_cast(_y.getHostValue(index))); + }); + + if(_index.has_value()) + { + _index->markDeviceModified(); + _index->rawHostData(); + iterateAlongDimensions(_spec.index->dims, [&](const std::vector& index) { + result.index.push_back(static_cast(_index->getHostValue(index))); + }); + } + return result; + } + +private: + ResampleFwdGraphSpec _spec; + Tensor _x; + Tensor _y; + std::optional> _index; +}; + +std::unique_ptr buildPlan(const ResampleFwdGraphSpec& spec) +{ + auto fbb = createResampleFwdGraph(spec); + const GraphWrapper graph(fbb.GetBufferPointer(), fbb.GetSize()); + return std::make_unique(resample_fwd_applicability::resolve(graph)); +} + +// 4x4 input pooled 2x2 / 2 into 2x2. Window maxima (row, col): (0,1)=5, (1,2)=8, (3,1)=11, +// (2,3)=7. +const std::vector& planarMaxInput() +{ + // clang-format off + static const std::vector s_values{ + 1, 5, 2, 0, + 3, 4, 8, 6, + 9, 2, 1, 7, + 0, 11, 3, 2, + }; + // clang-format on + return s_values; +} + +const PoolingResult& planarMaxExpected() +{ + static const PoolingResult s_expected{{5, 8, 11, 7}, {1, 6, 13, 11}}; + return s_expected; +} + +ResampleFwdGraphSpec planarMaxSpec(std::optional indexType, + DataType dataType = DataType::FLOAT) +{ + return ResampleFwdGraphSpec::packed( + ResampleMode::MAXPOOL, {1, 1, 4, 4}, {2, 2}, {2, 2}, {0, 0}, dataType, indexType); +} + +// `channels` slices of 3x3, 3x3 window, stride 1, padding 1: corner windows see 4 valid inputs, +// edge windows 6 and the centre window all 9. +ResampleFwdGraphSpec paddedAverageSpec(int64_t channels = 1) +{ + return ResampleFwdGraphSpec::packed( + ResampleMode::AVGPOOL_EXCLUDE_PADDING, {1, channels, 3, 3}, {3, 3}, {1, 1}, {1, 1}); +} + +// paddedAverageSpec over 1..9: window sums divided by 4 (corners), 6 (edges) and 9 (centre). +const std::vector& paddedAverageExpected() +{ + static const std::vector s_expected{ + 3.0f, 3.5f, 4.0f, 4.5f, 5.0f, 5.5f, 6.0f, 6.5f, 7.0f}; + return s_expected; +} + +class ResampleFwdPlanFixture : public ::testing::Test +{ +protected: + void SetUp() override + { + SKIP_IF_NO_DEVICES(); + _handle = std::make_unique(); + } + + void synchronize() const + { + hipStream_t stream = nullptr; + ASSERT_EQ(miopenGetStream(_handle->miopenHandle, &stream), miopenStatusSuccess); + ASSERT_EQ(hipStreamSynchronize(stream), hipSuccess); + } + + // Builds the plan for `spec`, executes it with a workspace of the size the plan asks for and + // returns Y (and the index) in logical order. + template + PoolingResult run(const ResampleFwdGraphSpec& spec, const std::vector& xValues) + { + const auto plan = buildPlan(spec); + PoolingTensors tensors(spec, xValues); + auto buffers = tensors.deviceBuffers(); + + const hipdnn_data_sdk::utilities::Workspace workspace(plan->getWorkspaceSize(*_handle)); + plan->execute( + *_handle, buffers.data(), static_cast(buffers.size()), workspace.get()); + synchronize(); + + return tensors.read(); + } + + std::unique_ptr _handle; +}; + +} // namespace + +class TestGpuMiopenResampleFwdPlan : public ResampleFwdPlanFixture +{ +}; + +// ============================================================================ +// Workspace: MIOpen's size for the index type the plan passes. +// ============================================================================ + +TEST_F(TestGpuMiopenResampleFwdPlan, WorkspaceHoldsInt32Indices) +{ + const auto plan = buildPlan(planarMaxSpec(DataType::INT32)); + EXPECT_EQ(plan->getWorkspaceSize(*_handle), 4 * sizeof(int32_t)); +} + +TEST_F(TestGpuMiopenResampleFwdPlan, WorkspaceHoldsInt8Indices) +{ + const auto plan = buildPlan(planarMaxSpec(DataType::INT8)); + EXPECT_EQ(plan->getWorkspaceSize(*_handle), 4 * sizeof(int8_t)); +} + +// ============================================================================ +// Execution. Y and the index start as sentinels, so a wrong alpha or beta, or an index that is +// not copied out, shows in the values. +// ============================================================================ + +TEST_F(TestGpuMiopenResampleFwdPlan, MaxPoolWritesValuesAndInt32Indices) +{ + expectMatches(run(planarMaxSpec(DataType::INT32), planarMaxInput()), planarMaxExpected(), 0.0); +} + +TEST_F(TestGpuMiopenResampleFwdPlan, MaxPoolWritesInt8Indices) +{ + expectMatches(run(planarMaxSpec(DataType::INT8), planarMaxInput()), + planarMaxExpected(), + 0.0); +} + +TEST_F(TestGpuMiopenResampleFwdPlan, MaxPoolWritesHalfValues) +{ + // The inputs are small integers, exact in half precision. + expectMatches( + run(planarMaxSpec(DataType::INT32, DataType::HALF), planarMaxInput()), + planarMaxExpected(), + 0.0); +} + +TEST_F(TestGpuMiopenResampleFwdPlan, MaxPoolWithoutIndexWritesValues) +{ + const auto result = run(planarMaxSpec(std::nullopt), planarMaxInput()); + expectYNear(result.y, planarMaxExpected().y, 0.0); + EXPECT_TRUE(result.index.empty()); +} + +TEST_F(TestGpuMiopenResampleFwdPlan, AvgExcludePaddingDividesByValidCount) +{ + const auto result = run(paddedAverageSpec(), ascending(9, 1.0f)); + expectYNear(result.y, paddedAverageExpected(), 1e-5); + EXPECT_TRUE(result.index.empty()); +} + +TEST_F(TestGpuMiopenResampleFwdPlan, AvgPoolingNhwcWritesValues) +{ + // Two channels: 1..9 and 10..18, whose averages are those of channel 0 plus 9. The plan hands + // MIOpen channels-last strides, so MIOpen transposes X and Y through the workspace. + auto spec = paddedAverageSpec(/*channels=*/2); + spec.makeChannelsLast(); + ASSERT_GT(buildPlan(spec)->getWorkspaceSize(*_handle), 0u); + + auto expected = paddedAverageExpected(); + for(const auto value : paddedAverageExpected()) + { + expected.push_back(value + 9.0f); + } + expectYNear(run(spec, ascending(18, 1.0f)).y, expected, 1e-5); +} + +TEST_F(TestGpuMiopenResampleFwdPlan, PoolingChannelsLastStyleInputIntoNchwOutput) +{ + // X {2, 3, 1, 1} carries NHWC strides {3, 1, 3, 3}, which address like packed NCHW; Y is + // packed NCHW. Taken at face value MIOpen would read X as NHWC and Y as NCHW; the plan hands + // it canonical NCHW strides for both. Every 2x2 window holds the single valid pixel. + auto spec = ResampleFwdGraphSpec::packed(ResampleMode::MAXPOOL, + {2, 3, 1, 1}, + {2, 2}, + {1, 1}, + {1, 1}, + DataType::FLOAT, + DataType::INT32); + spec.x.strides = std::vector{3, 1, 3, 3}; + ASSERT_EQ(spec.y.strides.value(), (std::vector{12, 4, 2, 1})); + + const auto x = ascending(6, 1.0f); + PoolingResult expected; + for(const auto value : x) + { + expected.y.insert(expected.y.end(), 4, value); + expected.index.insert(expected.index.end(), 4, 0); + } + expectMatches(run(spec, x), expected, 0.0); +} + +// ============================================================================ +// Execution errors +// ============================================================================ + +TEST_F(TestGpuMiopenResampleFwdPlan, ExecuteThrowsInvalidValueWithoutRequiredWorkspace) +{ + const auto spec = planarMaxSpec(DataType::INT32); + const auto plan = buildPlan(spec); + ASSERT_GT(plan->getWorkspaceSize(*_handle), 0u); + + PoolingTensors tensors(spec, planarMaxInput()); + auto buffers = tensors.deviceBuffers(); + try + { + plan->execute(*_handle, buffers.data(), static_cast(buffers.size()), nullptr); + FAIL() << "expected HipdnnPluginException"; + } + catch(const hipdnn_plugin_sdk::HipdnnPluginException& e) + { + EXPECT_EQ(e.getStatus(), HIPDNN_PLUGIN_STATUS_INVALID_VALUE); + } +} + +TEST_F(TestGpuMiopenResampleFwdPlan, ExecuteThrowsInvalidValueWithoutWorkspaceForNhwcAveragePooling) +{ + // No index, but channels-last tensors still need MIOpen's transpose buffers. + auto spec = ResampleFwdGraphSpec::packed( + ResampleMode::AVGPOOL_INCLUDE_PADDING, {2, 3, 5, 5}, {3, 3}, {2, 2}, {1, 1}); + spec.makeChannelsLast(); + const auto plan = buildPlan(spec); + ASSERT_GT(plan->getWorkspaceSize(*_handle), 0u); + + PoolingTensors tensors(spec, ascending(150)); + auto buffers = tensors.deviceBuffers(); + try + { + plan->execute(*_handle, buffers.data(), static_cast(buffers.size()), nullptr); + FAIL() << "expected HipdnnPluginException"; + } + catch(const hipdnn_plugin_sdk::HipdnnPluginException& e) + { + EXPECT_EQ(e.getStatus(), HIPDNN_PLUGIN_STATUS_INVALID_VALUE); + } +} + +TEST_F(TestGpuMiopenResampleFwdPlan, ExecuteSucceedsWithoutWorkspaceWhenNoneIsRequired) +{ + const auto spec = paddedAverageSpec(); + const auto plan = buildPlan(spec); + ASSERT_EQ(plan->getWorkspaceSize(*_handle), 0u); + + PoolingTensors tensors(spec, ascending(9, 1.0f)); + auto buffers = tensors.deviceBuffers(); + ASSERT_NO_THROW( + plan->execute(*_handle, buffers.data(), static_cast(buffers.size()), nullptr)); + synchronize(); + + expectYNear(tensors.read().y, paddedAverageExpected(), 1e-5); +} + +TEST_F(TestGpuMiopenResampleFwdPlan, ExecuteThrowsInvalidValueForMissingIndexBuffer) +{ + const auto spec = planarMaxSpec(DataType::INT32); + const auto plan = buildPlan(spec); + + PoolingTensors tensors(spec, planarMaxInput()); + auto buffers = tensors.deviceBuffers(/*includeIndex=*/false); + const hipdnn_data_sdk::utilities::Workspace workspace(plan->getWorkspaceSize(*_handle)); + try + { + plan->execute( + *_handle, buffers.data(), static_cast(buffers.size()), workspace.get()); + FAIL() << "expected HipdnnPluginException"; + } + catch(const hipdnn_plugin_sdk::HipdnnPluginException& e) + { + EXPECT_EQ(e.getStatus(), HIPDNN_PLUGIN_STATUS_INVALID_VALUE); + } +} + +TEST_F(TestGpuMiopenResampleFwdPlan, ExecuteThrowsInvalidValueForNullIndexPointer) +{ + const auto spec = planarMaxSpec(DataType::INT32); + const auto plan = buildPlan(spec); + + PoolingTensors tensors(spec, planarMaxInput()); + auto buffers = tensors.deviceBuffers(); + ASSERT_EQ(buffers.size(), 3u); + ASSERT_EQ(buffers[2].uid, ResampleFwdGraphSpec::INDEX_UID); + buffers[2].ptr = nullptr; + + const hipdnn_data_sdk::utilities::Workspace workspace(plan->getWorkspaceSize(*_handle)); + try + { + plan->execute( + *_handle, buffers.data(), static_cast(buffers.size()), workspace.get()); + FAIL() << "expected HipdnnPluginException"; + } + catch(const hipdnn_plugin_sdk::HipdnnPluginException& e) + { + EXPECT_EQ(e.getStatus(), HIPDNN_PLUGIN_STATUS_INVALID_VALUE); + } + synchronize(); + + // The plan threw before MIOpen ran: Y and the index still hold their sentinels. + const auto result = tensors.read(); + EXPECT_EQ(result.y, std::vector(4, Y_SENTINEL)); + EXPECT_EQ(result.index, std::vector(4, INDEX_SENTINEL)); +} diff --git a/dnn-providers/miopen-provider/tests/engines/plans/TestMiopenResampleFwdPlanBuilder.cpp b/dnn-providers/miopen-provider/tests/engines/plans/TestMiopenResampleFwdPlanBuilder.cpp new file mode 100644 index 000000000000..85e4c3b33819 --- /dev/null +++ b/dnn-providers/miopen-provider/tests/engines/plans/TestMiopenResampleFwdPlanBuilder.cpp @@ -0,0 +1,159 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include +#include +#include +#include + +#include + +#include +#include +#include +#include +#include + +#include "HipdnnMiopenContext.hpp" +#include "HipdnnMiopenHandle.hpp" +#include "HipdnnMiopenSettings.hpp" +#include "engines/plans/MiopenResampleFwdPlanBuilder.hpp" + +using namespace miopen_plugin; + +using hipdnn_flatbuffers_sdk::flatbuffer_utilities::GraphWrapper; +using hipdnn_test_sdk::utilities::createResampleFwdGraph; +using hipdnn_test_sdk::utilities::MockEngineConfig; +using hipdnn_test_sdk::utilities::ResampleFwdGraphSpec; + +namespace +{ + +// A spec the checks decline: asymmetric padding. +ResampleFwdGraphSpec unsupportedSpec() +{ + ResampleFwdGraphSpec spec; + spec.postPadding = std::vector{0, 0}; + return spec; +} + +} // namespace + +class TestGpuMiopenResampleFwdPlanBuilder : public ::testing::Test +{ +protected: + void SetUp() override + { + SKIP_IF_NO_DEVICES(); + _handle = std::make_unique(); + } + + bool isApplicable(const ResampleFwdGraphSpec& spec) const + { + auto fbb = createResampleFwdGraph(spec); + const GraphWrapper graph(fbb.GetBufferPointer(), fbb.GetSize()); + return _planBuilder.isApplicable(*_handle, graph); + } + + size_t maxWorkspaceSize(const ResampleFwdGraphSpec& spec) const + { + auto fbb = createResampleFwdGraph(spec); + const GraphWrapper graph(fbb.GetBufferPointer(), fbb.GetSize()); + const HipdnnMiopenSettings settings; + return _planBuilder.getMaxWorkspaceSize(*_handle, graph, settings); + } + + MiopenResampleFwdPlanBuilder _planBuilder; + MockEngineConfig _engineConfig; + std::unique_ptr _handle; +}; + +// ============================================================================ +// isApplicable +// ============================================================================ + +TEST_F(TestGpuMiopenResampleFwdPlanBuilder, IsApplicableDefersToTheChecks) +{ + EXPECT_TRUE(isApplicable(ResampleFwdGraphSpec{})); + EXPECT_FALSE(isApplicable(unsupportedSpec())); +} + +TEST_F(TestGpuMiopenResampleFwdPlanBuilder, IsApplicableReturnsFalseForOverrideShapeEnabledGraph) +{ + ResampleFwdGraphSpec spec; + spec.overrideShapeEnabled = true; + EXPECT_FALSE(isApplicable(spec)); +} + +// ============================================================================ +// getMaxWorkspaceSize and buildPlan +// ============================================================================ + +TEST_F(TestGpuMiopenResampleFwdPlanBuilder, BuildPlanWorkspaceMatchesMaxWorkspaceSize) +{ + // Channels-last max pooling: MIOpen's workspace holds the index region and the transposed X + // and Y, so the builder and the plan must size it from the same resolved configuration. + ResampleFwdGraphSpec spec; + spec.makeChannelsLast(); + auto fbb = createResampleFwdGraph(spec); + const GraphWrapper graph(fbb.GetBufferPointer(), fbb.GetSize()); + const auto expected = maxWorkspaceSize(spec); + ASSERT_GT(expected, 0u); + + HipdnnMiopenContext context; + ASSERT_NO_THROW(_planBuilder.buildPlan(*_handle, graph, _engineConfig, context)); + ASSERT_TRUE(context.hasValidPlan()); + EXPECT_EQ(context.plan().getWorkspaceSize(*_handle), expected); +} + +TEST_F(TestGpuMiopenResampleFwdPlanBuilder, MaxWorkspaceSizeThrowsBadParamForUnsupportedGraph) +{ + try + { + maxWorkspaceSize(unsupportedSpec()); + FAIL() << "expected HipdnnPluginException"; + } + catch(const hipdnn_plugin_sdk::HipdnnPluginException& e) + { + EXPECT_EQ(e.getStatus(), HIPDNN_PLUGIN_STATUS_BAD_PARAM); + } +} + +TEST_F(TestGpuMiopenResampleFwdPlanBuilder, BuildPlanThrowsBadParamForUnsupportedGraph) +{ + auto fbb = createResampleFwdGraph(unsupportedSpec()); + const GraphWrapper graph(fbb.GetBufferPointer(), fbb.GetSize()); + + HipdnnMiopenContext context; + try + { + _planBuilder.buildPlan(*_handle, graph, _engineConfig, context); + FAIL() << "expected HipdnnPluginException"; + } + catch(const hipdnn_plugin_sdk::HipdnnPluginException& e) + { + EXPECT_EQ(e.getStatus(), HIPDNN_PLUGIN_STATUS_BAD_PARAM); + } + EXPECT_FALSE(context.hasValidPlan()); +} + +TEST_F(TestGpuMiopenResampleFwdPlanBuilder, InitializeExecutionSettingsLeavesSettingsUntouched) +{ + auto fbb = createResampleFwdGraph(ResampleFwdGraphSpec{}); + const GraphWrapper graph(fbb.GetBufferPointer(), fbb.GetSize()); + + HipdnnMiopenSettings settings; + EXPECT_NO_THROW( + _planBuilder.initializeExecutionSettings(*_handle, graph, _engineConfig, settings)); + EXPECT_FALSE(settings.workspaceSizeLimit().has_value()); + EXPECT_FALSE(settings.defaultWorkspaceSize().has_value()); + EXPECT_FALSE(settings.benchmarkingEnabled()); +} + +TEST_F(TestGpuMiopenResampleFwdPlanBuilder, GetCustomKnobsReturnsEmpty) +{ + auto fbb = createResampleFwdGraph(ResampleFwdGraphSpec{}); + const GraphWrapper graph(fbb.GetBufferPointer(), fbb.GetSize()); + + EXPECT_TRUE(_planBuilder.getCustomKnobs(*_handle, graph).empty()); +} diff --git a/projects/hipdnn/docs/OperationSupport.md b/projects/hipdnn/docs/OperationSupport.md index afe2019fa33b..befaa000bbb9 100644 --- a/projects/hipdnn/docs/OperationSupport.md +++ b/projects/hipdnn/docs/OperationSupport.md @@ -21,6 +21,7 @@ hipDNN operations are implemented through plugins. Each plugin provides its own - Fused operation graphs - Standalone pointwise activation - Standalone pointwise binary elementwise (ADD, SUB, MUL, MAX, MIN) + - Standalone resample forward - **[hipBLASLt Provider Plugin](../../../dnn-providers/hipblaslt-provider/docs/OperationSupport.md)** - Integration with AMD's hipBLASLt library that provides optimized GEMM operations. - Stand-alone Matmul (FP32/FP16/BF16 inputs and outputs, with optional fused bias and activation) diff --git a/projects/hipdnn/docs/user-guides/reference/miopen-plugin.rst b/projects/hipdnn/docs/user-guides/reference/miopen-plugin.rst index eda6367fec01..2ff47ab94b71 100644 --- a/projects/hipdnn/docs/user-guides/reference/miopen-plugin.rst +++ b/projects/hipdnn/docs/user-guides/reference/miopen-plugin.rst @@ -85,10 +85,14 @@ This table lists all operations supported in hipDNN: - ``FP16``, ``FP32`` - NCHW, NCDHW - Single-node graph⁷ + * - Resample Forward (pooling, standalone) + - ``FP16``, ``BFP16``, ``FP32`` + - NCHW, NHWC, NCDHW, NDHWC + - Single-node graph⁸ .. note:: - - For annotations ¹ through ⁴, ⁶ and ⁷, see :ref:`operations`. + - For annotations ¹ through ⁴ and ⁶ through ⁸, see :ref:`operations`. - For annotation ⁵, see :ref:`detailed` for more information. .. _detailed: @@ -148,6 +152,7 @@ Operation notes - **Standalone Activation**: A single-node pointwise graph supports ReLU, Clipped ReLU (configurable upper clip), CLAMP (configurable lower/upper clips), Leaky ReLU (configurable negative slope), Sigmoid and Tanh. - **ReLU Parameters**: An unset ``relu_lower_clip`` is treated as 0, an unset ``relu_upper_clip`` as no upper bound, and an unset ``relu_lower_clip_slope`` or a slope of exactly 0 as no leaky slope, so equivalent encodings select the same MIOpen activation. A non-zero ``relu_lower_clip`` without an upper clip is not supported. - ⁷ **Standalone Binary Elementwise**: A single-node pointwise graph supports ADD, SUB, MUL, MAX and MIN across ranks 3-5, with packed channels-first strides and ``FP16``/``FP32`` datatypes only. ``in_0`` must be the full-shape operand; only ``in_1`` may broadcast per-axis at matching rank, and all three tensors must share one datatype. In-place execution (``out_0`` aliasing an input) isn't supported. +- ⁸ **Standalone Resample Forward**: A single-node graph supports ``MAXPOOL``, ``AVGPOOL_EXCLUDE_PADDING`` and ``AVGPOOL_INCLUDE_PADDING`` across ranks 4-5, with packed X and Y sharing one datatype and layout, and an ``FP32`` compute type. Padding must be symmetric and smaller than the window; unless it is all zero, ``MAXPOOL`` needs ``NEG_INF_PAD`` and the average modes need ``ZERO_PAD``. Y must have the floor-mode output size. The index output (``generate_index``) is supported for ``MAXPOOL`` in NCHW/NCDHW only, as ``INT8`` or ``INT32``; with ``INT8`` the input spatial size may not exceed 128. Very large windows on channels-last tensors aren't supported. Max pooling ignores NaN inputs. - **Sparse Support**: All operations only work with dense tensors. Knobs diff --git a/projects/hipdnn/docs/user-guides/reference/plugins.rst b/projects/hipdnn/docs/user-guides/reference/plugins.rst index ebadc73ef636..c1771fa2432f 100644 --- a/projects/hipdnn/docs/user-guides/reference/plugins.rst +++ b/projects/hipdnn/docs/user-guides/reference/plugins.rst @@ -19,6 +19,7 @@ Plugins - Batchnorm operations (Training, Backward, Inference) - Fused operation graphs - Standalone pointwise activation and binary elementwise operations + - Standalone resample forward - :ref:`hipblaslt`: Provides integration with AMD's `hipBLASLt library `_ that provides optimized GEMM operations. diff --git a/projects/hipdnn/test_sdk/include/hipdnn_test_sdk/utilities/FlatbufferGraphTestUtils.hpp b/projects/hipdnn/test_sdk/include/hipdnn_test_sdk/utilities/FlatbufferGraphTestUtils.hpp index a4d918550ef0..b037c1f4fa1b 100644 --- a/projects/hipdnn/test_sdk/include/hipdnn_test_sdk/utilities/FlatbufferGraphTestUtils.hpp +++ b/projects/hipdnn/test_sdk/include/hipdnn_test_sdk/utilities/FlatbufferGraphTestUtils.hpp @@ -6,6 +6,8 @@ #include #include #include +#include +#include #include #include @@ -3498,6 +3500,208 @@ inline flatbuffers::FlatBufferBuilder return builder; } +// A tensor of a ResampleFwdGraphSpec. +struct ResampleFwdTensorSpec +{ + // A dense, non-virtual tensor that is part of the graph. + static ResampleFwdTensorSpec dense(int64_t uid, + std::vector dims, + std::vector strides, + hipdnn_flatbuffers_sdk::data_objects::DataType dataType + = hipdnn_flatbuffers_sdk::data_objects::DataType::FLOAT) + { + ResampleFwdTensorSpec tensor; + tensor.uid = uid; + tensor.dims = std::move(dims); + tensor.strides = std::move(strides); + tensor.dataType = dataType; + return tensor; + } + + int64_t uid = 0; + std::vector dims; + // nullopt leaves the strides field out of the flatbuffer. + std::optional> strides; + hipdnn_flatbuffers_sdk::data_objects::DataType dataType + = hipdnn_flatbuffers_sdk::data_objects::DataType::FLOAT; + bool isVirtual = false; + bool isRuntimePassByValue = false; + std::optional raggedOffsetTensorUid; + // false leaves the tensor out of the graph, so the node's reference to `uid` dangles. + bool inGraph = true; +}; + +// Configuration for a single-node ResampleFwdAttributes graph. The defaults describe a valid graph: +// fp32 max pooling of a packed NCHW 2x3x6x6 tensor with a 3x3 window, stride 2 and padding 1 +// (NEG_INF_PAD) into 2x3x3x3, without an index. +struct ResampleFwdGraphSpec +{ + static constexpr int64_t X_UID = 1; + static constexpr int64_t Y_UID = 2; + static constexpr int64_t INDEX_UID = 3; + + // Pools `xDims` with `window`, `stride` and the same `padding` before and after every spatial + // dimension: packed channels-first X and Y of `dataType`, Y sized by the floor-mode formula, and + // NEG_INF_PAD for MAXPOOL or ZERO_PAD for the average modes. With `indexType`, the node + // generates an index into a packed channels-first tensor shaped like Y. + static ResampleFwdGraphSpec + packed(hipdnn_flatbuffers_sdk::data_objects::ResampleMode mode, + const std::vector& xDims, + const std::vector& window, + const std::vector& stride, + const std::vector& padding, + hipdnn_flatbuffers_sdk::data_objects::DataType dataType + = hipdnn_flatbuffers_sdk::data_objects::DataType::FLOAT, + std::optional indexType + = std::nullopt) + { + using hipdnn_data_sdk::utilities::generateStrides; + + std::vector yDims{xDims.at(0), xDims.at(1)}; + for(size_t i = 0; i + 2 < xDims.size(); ++i) + { + yDims.push_back((xDims[i + 2] + 2 * padding.at(i) - window.at(i)) / stride.at(i) + 1); + } + + ResampleFwdGraphSpec spec; + spec.mode = mode; + spec.paddingMode = mode == hipdnn_flatbuffers_sdk::data_objects::ResampleMode::MAXPOOL + ? hipdnn_flatbuffers_sdk::data_objects::PaddingMode::NEG_INF_PAD + : hipdnn_flatbuffers_sdk::data_objects::PaddingMode::ZERO_PAD; + spec.x = ResampleFwdTensorSpec::dense(X_UID, xDims, generateStrides(xDims), dataType); + spec.y = ResampleFwdTensorSpec::dense(Y_UID, yDims, generateStrides(yDims), dataType); + if(indexType.has_value()) + { + spec.index = ResampleFwdTensorSpec::dense( + INDEX_UID, yDims, generateStrides(yDims), indexType.value()); + spec.generateIndex = true; + } + spec.window = window; + spec.stride = stride; + spec.setPadding(padding); + return spec; + } + + // Gives X and Y packed channels-last strides (NHWC, NDHWC). The index keeps its strides. + void makeChannelsLast() + { + using hipdnn_data_sdk::utilities::generateStrides; + using hipdnn_data_sdk::utilities::strideOrderNhwc; + x.strides = generateStrides(x.dims, strideOrderNhwc(x.dims.size())); + y.strides = generateStrides(y.dims, strideOrderNhwc(y.dims.size())); + } + + // Sets the same padding before and after every spatial dimension. + void setPadding(const std::vector& padding) + { + prePadding = padding; + postPadding = padding; + } + + hipdnn_flatbuffers_sdk::data_objects::ResampleMode mode + = hipdnn_flatbuffers_sdk::data_objects::ResampleMode::MAXPOOL; + hipdnn_flatbuffers_sdk::data_objects::PaddingMode paddingMode + = hipdnn_flatbuffers_sdk::data_objects::PaddingMode::NEG_INF_PAD; + ResampleFwdTensorSpec x = ResampleFwdTensorSpec::dense(X_UID, {2, 3, 6, 6}, {108, 36, 6, 1}); + ResampleFwdTensorSpec y = ResampleFwdTensorSpec::dense(Y_UID, {2, 3, 3, 3}, {27, 9, 3, 1}); + // Sets the node's index_tensor_uid; nullopt leaves it unset. + std::optional index; + // nullopt leaves the vector out of the node. + std::optional> window = std::vector{3, 3}; + std::optional> stride = std::vector{2, 2}; + std::optional> prePadding = std::vector{1, 1}; + std::optional> postPadding = std::vector{1, 1}; + std::optional generateIndex; + hipdnn_flatbuffers_sdk::data_objects::DataType computeDataType + = hipdnn_flatbuffers_sdk::data_objects::DataType::FLOAT; + bool overrideShapeEnabled = false; + // Number of identical resample nodes in the graph. + uint32_t nodeCount = 1; +}; + +// Builds a ResampleFwdAttributes graph from `spec`. Graph tensor uids are unique, so a tensor +// whose uid an earlier one already has is emitted once, and the node references it twice. +inline flatbuffers::FlatBufferBuilder createResampleFwdGraph(const ResampleFwdGraphSpec& spec) +{ + namespace data_objects = hipdnn_flatbuffers_sdk::data_objects; + + flatbuffers::FlatBufferBuilder builder; + std::vector> tensorAttributes; + std::vector emittedUids; + + const auto addTensor = [&](const ResampleFwdTensorSpec& tensor, const char* name) { + if(!tensor.inGraph + || std::find(emittedUids.begin(), emittedUids.end(), tensor.uid) != emittedUids.end()) + { + return; + } + emittedUids.push_back(tensor.uid); + tensorAttributes.push_back(data_objects::CreateTensorAttributesDirect( + builder, + tensor.uid, + name, + tensor.dataType, + tensor.strides.has_value() ? &tensor.strides.value() : nullptr, + &tensor.dims, + tensor.isVirtual, + data_objects::TensorValue::NONE, + 0, + tensor.isRuntimePassByValue, + tensor.raggedOffsetTensorUid)); + }; + + addTensor(spec.x, "x"); + addTensor(spec.y, "y"); + if(spec.index.has_value()) + { + addTensor(spec.index.value(), "index"); + } + + const auto optionalVector = [](const std::optional>& values) { + return values.has_value() ? &values.value() : nullptr; + }; + const auto indexUid = spec.index.has_value() + ? ::flatbuffers::Optional(spec.index->uid) + : ::flatbuffers::Optional(::flatbuffers::nullopt); + + std::vector> nodes; + for(uint32_t i = 0; i < spec.nodeCount; ++i) + { + const auto attributes + = data_objects::CreateResampleFwdAttributesDirect(builder, + spec.x.uid, + spec.y.uid, + indexUid, + optionalVector(spec.prePadding), + optionalVector(spec.postPadding), + optionalVector(spec.stride), + optionalVector(spec.window), + spec.mode, + spec.paddingMode, + spec.generateIndex); + + const auto name = "resample_fwd_" + std::to_string(i); + nodes.push_back( + data_objects::CreateNodeDirect(builder, + name.c_str(), + spec.computeDataType, + data_objects::NodeAttributes::ResampleFwdAttributes, + attributes.Union())); + } + + const auto graph = data_objects::CreateGraphDirect(builder, + "test", + data_objects::DataType::FLOAT, + data_objects::DataType::FLOAT, + data_objects::DataType::FLOAT, + &tensorAttributes, + &nodes, + flatbuffers::nullopt, + spec.overrideShapeEnabled); + builder.Finish(graph); + return builder; +} + inline flatbuffers::FlatBufferBuilder createValidResampleBwdGraph(bool useIndex = true, hipdnn_flatbuffers_sdk::data_objects::ResampleMode resampleMode diff --git a/projects/hipdnn/test_sdk/tests/CMakeLists.txt b/projects/hipdnn/test_sdk/tests/CMakeLists.txt index b7397b7a28dc..76e1472e2eeb 100644 --- a/projects/hipdnn/test_sdk/tests/CMakeLists.txt +++ b/projects/hipdnn/test_sdk/tests/CMakeLists.txt @@ -54,6 +54,7 @@ add_executable( utilities/TestGraphTestHelpers.cpp utilities/TestFlatbufferDatatypeMapping.cpp utilities/TestFlatbufferGraphMutators.cpp + utilities/TestFlatbufferGraphTestUtils.cpp utilities/TestFlatbufferTensorAttributesUtils.cpp utilities/TestLoadGraphAndTensors.cpp utilities/TestLoggingUtils.cpp diff --git a/projects/hipdnn/test_sdk/tests/utilities/TestFlatbufferGraphTestUtils.cpp b/projects/hipdnn/test_sdk/tests/utilities/TestFlatbufferGraphTestUtils.cpp new file mode 100644 index 000000000000..93cee0b88e6c --- /dev/null +++ b/projects/hipdnn/test_sdk/tests/utilities/TestFlatbufferGraphTestUtils.cpp @@ -0,0 +1,196 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include +#include +#include + +#include + +#include +#include + +using hipdnn_flatbuffers_sdk::data_objects::DataType; +using hipdnn_flatbuffers_sdk::data_objects::GetGraph; +using hipdnn_flatbuffers_sdk::data_objects::PaddingMode; +using hipdnn_flatbuffers_sdk::data_objects::ResampleMode; +using hipdnn_test_sdk::utilities::createResampleFwdGraph; +using hipdnn_test_sdk::utilities::ResampleFwdGraphSpec; + +namespace +{ + +std::vector toVector(const flatbuffers::Vector& values) +{ + return {values.begin(), values.end()}; +} + +} // namespace + +// ============================================================================ +// ResampleFwdGraphSpec and createResampleFwdGraph +// ============================================================================ + +TEST(TestResampleFwdGraphSpec, PackedSizesFloorModeOutputAndMatchesPaddingMode) +{ + // Y: H = (7 + 2 - 3) / 2 + 1 = 4, W = (5 - 2) / 1 + 1 = 4. + const auto spec = ResampleFwdGraphSpec::packed(ResampleMode::AVGPOOL_INCLUDE_PADDING, + {1, 2, 7, 5}, + {3, 2}, + {2, 1}, + {1, 0}, + DataType::HALF); + + EXPECT_EQ(spec.mode, ResampleMode::AVGPOOL_INCLUDE_PADDING); + EXPECT_EQ(spec.paddingMode, PaddingMode::ZERO_PAD); + EXPECT_EQ(spec.x.dims, (std::vector{1, 2, 7, 5})); + EXPECT_EQ(spec.x.strides, (std::vector{70, 35, 5, 1})); + EXPECT_EQ(spec.x.dataType, DataType::HALF); + EXPECT_EQ(spec.y.dims, (std::vector{1, 2, 4, 4})); + EXPECT_EQ(spec.y.strides, (std::vector{32, 16, 4, 1})); + EXPECT_EQ(spec.y.dataType, DataType::HALF); + EXPECT_EQ(spec.window, (std::vector{3, 2})); + EXPECT_EQ(spec.stride, (std::vector{2, 1})); + EXPECT_EQ(spec.prePadding, (std::vector{1, 0})); + EXPECT_EQ(spec.postPadding, (std::vector{1, 0})); + EXPECT_FALSE(spec.index.has_value()); + EXPECT_FALSE(spec.generateIndex.has_value()); +} + +TEST(TestResampleFwdGraphSpec, PackedWithIndexTypeAddsChannelsFirstIndexShapedLikeY) +{ + const auto spec = ResampleFwdGraphSpec::packed(ResampleMode::MAXPOOL, + {2, 3, 6, 6}, + {3, 3}, + {2, 2}, + {1, 1}, + DataType::FLOAT, + DataType::INT8); + + EXPECT_EQ(spec.paddingMode, PaddingMode::NEG_INF_PAD); + ASSERT_TRUE(spec.index.has_value()); + EXPECT_EQ(spec.index->uid, ResampleFwdGraphSpec::INDEX_UID); + EXPECT_EQ(spec.index->dims, spec.y.dims); + EXPECT_EQ(spec.index->strides, (std::vector{27, 9, 3, 1})); + EXPECT_EQ(spec.index->dataType, DataType::INT8); + EXPECT_EQ(spec.generateIndex, std::optional(true)); +} + +TEST(TestResampleFwdGraphSpec, MakeChannelsLastRestridesXAndYButNotTheIndex) +{ + auto spec = ResampleFwdGraphSpec::packed(ResampleMode::MAXPOOL, + {2, 3, 6, 6}, + {3, 3}, + {2, 2}, + {1, 1}, + DataType::FLOAT, + DataType::INT32); + spec.makeChannelsLast(); + + EXPECT_EQ(spec.x.strides, (std::vector{108, 1, 18, 3})); + EXPECT_EQ(spec.y.strides, (std::vector{27, 1, 9, 3})); + ASSERT_TRUE(spec.index.has_value()); + EXPECT_EQ(spec.index->strides, (std::vector{27, 9, 3, 1})); +} + +TEST(TestCreateResampleFwdGraph, WritesTensorsAndNodeAttributes) +{ + auto spec = ResampleFwdGraphSpec::packed(ResampleMode::MAXPOOL, + {2, 3, 6, 6}, + {3, 3}, + {2, 2}, + {1, 1}, + DataType::FLOAT, + DataType::INT32); + spec.computeDataType = DataType::HALF; + + const auto builder = createResampleFwdGraph(spec); + const auto* graph = GetGraph(builder.GetBufferPointer()); + + ASSERT_EQ(graph->tensors()->size(), 3u); + EXPECT_EQ(graph->tensors()->Get(2)->uid(), ResampleFwdGraphSpec::INDEX_UID); + EXPECT_EQ(graph->tensors()->Get(2)->data_type(), DataType::INT32); + ASSERT_EQ(graph->nodes()->size(), 1u); + EXPECT_EQ(graph->nodes()->Get(0)->compute_data_type(), DataType::HALF); + EXPECT_FALSE(graph->is_override_shape_enabled()); + + const auto* attributes = graph->nodes()->Get(0)->attributes_as_ResampleFwdAttributes(); + ASSERT_NE(attributes, nullptr); + EXPECT_EQ(attributes->x_tensor_uid(), ResampleFwdGraphSpec::X_UID); + EXPECT_EQ(attributes->y_tensor_uid(), ResampleFwdGraphSpec::Y_UID); + ASSERT_TRUE(attributes->index_tensor_uid().has_value()); + EXPECT_EQ(attributes->index_tensor_uid().value(), ResampleFwdGraphSpec::INDEX_UID); + EXPECT_EQ(toVector(*attributes->window()), (std::vector{3, 3})); + EXPECT_EQ(toVector(*attributes->stride()), (std::vector{2, 2})); + EXPECT_EQ(toVector(*attributes->pre_padding()), (std::vector{1, 1})); + EXPECT_EQ(toVector(*attributes->post_padding()), (std::vector{1, 1})); + EXPECT_EQ(attributes->resample_mode(), ResampleMode::MAXPOOL); + EXPECT_EQ(attributes->padding_mode(), PaddingMode::NEG_INF_PAD); + ASSERT_TRUE(attributes->generate_index().has_value()); + EXPECT_TRUE(attributes->generate_index().value()); +} + +TEST(TestCreateResampleFwdGraph, LeavesOutTensorNotInGraph) +{ + ResampleFwdGraphSpec spec; + spec.y.inGraph = false; + + const auto builder = createResampleFwdGraph(spec); + const auto* graph = GetGraph(builder.GetBufferPointer()); + + ASSERT_EQ(graph->tensors()->size(), 1u); + EXPECT_EQ(graph->tensors()->Get(0)->uid(), ResampleFwdGraphSpec::X_UID); + // The node still names y, so its output uid dangles. + const auto* attributes = graph->nodes()->Get(0)->attributes_as_ResampleFwdAttributes(); + ASSERT_NE(attributes, nullptr); + EXPECT_EQ(attributes->y_tensor_uid(), ResampleFwdGraphSpec::Y_UID); + EXPECT_FALSE(attributes->index_tensor_uid().has_value()); +} + +TEST(TestCreateResampleFwdGraph, WritesNoStridesOrPoolingVectorsWhenUnset) +{ + ResampleFwdGraphSpec spec; + spec.x.strides = std::nullopt; + spec.window = std::nullopt; + spec.generateIndex = false; + + const auto builder = createResampleFwdGraph(spec); + const auto* graph = GetGraph(builder.GetBufferPointer()); + + ASSERT_EQ(graph->tensors()->size(), 2u); + EXPECT_EQ(graph->tensors()->Get(0)->strides(), nullptr); + ASSERT_NE(graph->tensors()->Get(1)->strides(), nullptr); + const auto* attributes = graph->nodes()->Get(0)->attributes_as_ResampleFwdAttributes(); + ASSERT_NE(attributes, nullptr); + EXPECT_EQ(attributes->window(), nullptr); + ASSERT_NE(attributes->stride(), nullptr); + ASSERT_TRUE(attributes->generate_index().has_value()); + EXPECT_FALSE(attributes->generate_index().value()); +} + +TEST(TestCreateResampleFwdGraph, EmitsAliasedTensorOnce) +{ + ResampleFwdGraphSpec spec; + spec.y.uid = spec.x.uid; + + const auto builder = createResampleFwdGraph(spec); + const auto* graph = GetGraph(builder.GetBufferPointer()); + + ASSERT_EQ(graph->tensors()->size(), 1u); + const auto* attributes = graph->nodes()->Get(0)->attributes_as_ResampleFwdAttributes(); + ASSERT_NE(attributes, nullptr); + EXPECT_EQ(attributes->x_tensor_uid(), attributes->y_tensor_uid()); +} + +TEST(TestCreateResampleFwdGraph, WritesNodeCountAndOverrideShapeFlag) +{ + ResampleFwdGraphSpec spec; + spec.nodeCount = 2; + spec.overrideShapeEnabled = true; + + const auto builder = createResampleFwdGraph(spec); + const auto* graph = GetGraph(builder.GetBufferPointer()); + + EXPECT_EQ(graph->nodes()->size(), 2u); + EXPECT_TRUE(graph->is_override_shape_enabled()); +}