Add support for the dilations attribute to Pooling ops #2105

mirza-halilcevic · 2023-08-22T10:02:36Z

Add parsing support for the dilations attribute
Modify reference implementation
Resolves Pool dilations attribute migraphx-benchmark/AMDMIGraphX#70

umangyadav · 2023-08-22T12:58:40Z

For such changes, it would require tests inside
ref_ops_test.cpp which tests whether reference implementation inside MIGraphX is correct or not.
op_shape_test.cpp. to test for the compute_shape() part
onnx_test.cpp to test whether MIGraphX is parsing operator correctly or not.

and verify tests. which would run the operator/program on GPU and compare it with the reference implementation. Looks like MIOpen doesn't support dilations so we won't have GPU implementation.

src/include/migraphx/op/pooling.hpp

pfultz2 · 2023-08-22T18:41:39Z

Looks like MIOpen doesn't support dilations so we won't have GPU implementation.

@umangyadav Is there a way we can insert some other operators to get the same dilation result on the GPU, similar to how we do for non-symmetrical padding? Like using step or reshape/broadcast/reshape?

attila-dusnoki-htec · 2023-08-28T15:53:31Z

We updated the test part where dilations where missing, and added new tests as well.
Note: this version does not support non-default dtilations on GPU yet.

implementation.

codecov · 2023-09-04T16:01:33Z

Codecov Report

Attention: 4 lines in your changes are missing coverage. Please review.

❗ No coverage uploaded for pull request base (develop@c7bae54). Click here to learn what that means.

❗ Current head 39491c4 differs from pull request most recent head 1e83d83. Consider uploading reports for the commit 1e83d83 to get more accurate results

Files	Patch %	Lines
src/include/migraphx/op/pooling.hpp	85.71%	3 Missing ⚠️
src/rewrite_pooling.cpp	98.44%	1 Missing ⚠️

Additional details and impacted files

@@            Coverage Diff             @@
##             develop    #2105   +/-   ##
==========================================
  Coverage           ?   91.39%           
==========================================
  Files              ?      440           
  Lines              ?    16591           
  Branches           ?        0           
==========================================
  Hits               ?    15162           
  Misses             ?     1429           
  Partials           ?        0

☔ View full report in Codecov by Sentry.
📢 Have feedback on the report? Share it here.

attila-dusnoki-htec · 2023-09-07T18:39:54Z

cc @pfultz2
I looked at "step" op, but still did not managed to figure out what you meant. 😞
Instead I ended up collecting the axes, and creating the input with gather.
It works, but has a serious limitation: the axes is a vector and can become too large, which results in bad alloc.

pfultz2 · 2023-09-08T16:10:31Z

I looked at "step" op, but still did not managed to figure out what you meant.

You would insert a step operator with the same value as the dilation, and using the axes for the image. Like `make_op("step", {{"axes", {2, 3}}, {"steps", {dilation_for_h, dilation_for_w}}}).

pfultz2 · 2023-09-18T17:20:02Z

src/include/migraphx/op/pooling.hpp

+                {
+                    if(idx_w[axis] % dilations[axis])
+                    {
+                        pool_size -= 1;


pool_size should be adjusted before calling shape_for_each.

pfultz2 · 2023-09-18T17:21:13Z

src/include/migraphx/op/pooling.hpp

@@ -390,7 +411,15 @@ struct pooling
                    // this is a padding element.  Padding locations
                    // don't contribute to average or max pooling total but can play in
                    // lpnorm pooling.
-                    output_val = op(output_val, 0);


This should just be output_val = op(output_val, op.template init<Type>());. Using 0 is a bug.

pfultz2 · 2023-09-18T17:22:01Z

src/include/migraphx/rewrite_pooling.hpp

+
+    private:
+    void replace_with_reduce(module& m, instruction_ref ins) const;
+    void replace_dilations_with_gather_pooling(module& m, instruction_ref ins) const;


These should be static functions defined in rewrite_pooling.cpp file.

pfultz2 · 2023-09-18T17:30:02Z

src/targets/gpu/include/migraphx/gpu/miopen.hpp

+        for(const auto& d : op.dilations)
+            ss << d << ",";
+        ss << "]";
+        MIGRAPHX_THROW(ss.str());


Just do MIGRAPHX_THROW("Unsupported dilations for pooling: [" + to_string_range(op.dilations) + "]") instead.

attila-dusnoki-htec · 2023-09-19T16:15:43Z

Addressed requested changes, except this:

pool_size should be adjusted before calling shape_for_each.

I tried to calculate it beforehand, but did not managed to resolve the dyn pad cases properly 😞 (avgpool_dyn_auto_pad_test, avgpool_dyn_pad_ceil_test)
The current approach is the most straighforward. But i can try again if that is a blocker.

CharlieL7 · 2023-11-21T16:08:11Z

test/rewrite_pooling_test.cpp

@@ -80,6 +81,483 @@ TEST_CASE(rewrite_pooling_test)
                 migraphx::make_op("reduce_max", {{"axes", {2, 3, 4}}}));
 }

+TEST_CASE(rewrite_pooling_dialtions_test)


Suggested change

TEST_CASE(rewrite_pooling_dialtions_test)

TEST_CASE(rewrite_pooling_dilations_test)

CharlieL7 · 2023-11-21T16:15:12Z

test/rewrite_pooling_test.cpp

+    auto result = p.eval({}).back();
+    std::vector<float> results_vector;
+    result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); });
+    std::vector<float> gold{0.7812, 1.0449, 2.7666, 2.6859};


How are you getting these gold values? Want to make sure that we know these are correct.

causten requested review from CharlieL7 and umangyadav August 22, 2023 12:35

umangyadav reviewed Aug 22, 2023

View reviewed changes

src/include/migraphx/op/pooling.hpp Show resolved Hide resolved

attila-dusnoki-htec mentioned this pull request Aug 24, 2023

Bump onnx from 1.10.2 to 1.14.0 #2121

Closed

mirza-halilcevic and others added 11 commits September 1, 2023 11:51

Introduce dilations attribute to pooling operators reference

96971e3

implementation.

Make some code styling improvements.

f666861

Add dilations to existing pooling tests

ceb0325

Fix missing dilations from pooling throw message

31111b7

Fix dilation calc logic for pooling

59a7768

Add pooling dilation shape tests

6f097e8

Add pooling dilation ref tests

5d92ec1

Add pooling dilation onnx tests

e3aaae0

Add throw for GPU non-default dilations

e5f06f5

Merge branch 'develop' into pooling_dilations

5e8cdce

format.py fixes

53a10bc

gyulaz-htec force-pushed the pooling_dilations branch from 8250a47 to 53a10bc Compare September 1, 2023 13:30

Add missed dilations

8ddbfa5

attila-dusnoki-htec added 2 commits September 8, 2023 07:32

Add dilated pooling rewrite

f361d5c

Merge remote-tracking branch 'upstream/develop' into pooling_dilations

d639171

attila-dusnoki-htec force-pushed the pooling_dilations branch from 734ebf1 to d639171 Compare September 8, 2023 07:56

mirza-halilcevic marked this pull request as ready for review September 11, 2023 08:05

Merge remote-tracking branch 'upstream/develop' into pooling_dilations

3426ebf

pfultz2 reviewed Sep 18, 2023

View reviewed changes

Address review comments

6e43d0c

Merge remote-tracking branch 'upstream/develop' into pooling_dilations

d8de555

attila-dusnoki-htec force-pushed the pooling_dilations branch from 84b4676 to d8de555 Compare October 3, 2023 08:19

attila-dusnoki-htec added 2 commits October 11, 2023 10:39

Merge remote-tracking branch 'upstream/develop' into pooling_dilations

1e76547

Merge remote-tracking branch 'upstream/develop' into pooling_dilations

195cebd

causten requested review from umangyadav and pfultz2 October 25, 2023 13:07

Merge branch 'develop' into pooling_dilations

e3fcc5a

CharlieL7 reviewed Nov 21, 2023

View reviewed changes

pfultz2 approved these changes Nov 21, 2023

View reviewed changes

Merge branch 'develop' into pooling_dilations

1e83d83

causten merged commit 19bd9c4 into ROCm:develop Nov 22, 2023
11 of 12 checks passed

causten deleted the pooling_dilations branch November 22, 2023 18:59

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Add support for the dilations attribute to Pooling ops #2105

Add support for the dilations attribute to Pooling ops #2105

mirza-halilcevic commented Aug 22, 2023 •

edited

Loading

umangyadav commented Aug 22, 2023 •

edited

Loading

pfultz2 commented Aug 22, 2023 •

edited

Loading

attila-dusnoki-htec commented Aug 28, 2023

codecov bot commented Sep 4, 2023 •

edited

Loading

attila-dusnoki-htec commented Sep 7, 2023

pfultz2 commented Sep 8, 2023

pfultz2 Sep 18, 2023

pfultz2 Sep 18, 2023

pfultz2 Sep 18, 2023

pfultz2 Sep 18, 2023

attila-dusnoki-htec commented Sep 19, 2023

CharlieL7 Nov 21, 2023

CharlieL7 Nov 21, 2023

	TEST_CASE(rewrite_pooling_dialtions_test)
	TEST_CASE(rewrite_pooling_dilations_test)

Add support for the dilations attribute to Pooling ops #2105

Add support for the dilations attribute to Pooling ops #2105

Conversation

mirza-halilcevic commented Aug 22, 2023 • edited Loading

umangyadav commented Aug 22, 2023 • edited Loading

pfultz2 commented Aug 22, 2023 • edited Loading

attila-dusnoki-htec commented Aug 28, 2023

codecov bot commented Sep 4, 2023 • edited Loading

Codecov Report

attila-dusnoki-htec commented Sep 7, 2023

pfultz2 commented Sep 8, 2023

pfultz2 Sep 18, 2023

Choose a reason for hiding this comment

pfultz2 Sep 18, 2023

Choose a reason for hiding this comment

pfultz2 Sep 18, 2023

Choose a reason for hiding this comment

pfultz2 Sep 18, 2023

Choose a reason for hiding this comment

attila-dusnoki-htec commented Sep 19, 2023

CharlieL7 Nov 21, 2023

Choose a reason for hiding this comment

CharlieL7 Nov 21, 2023

Choose a reason for hiding this comment

mirza-halilcevic commented Aug 22, 2023 •

edited

Loading

umangyadav commented Aug 22, 2023 •

edited

Loading

pfultz2 commented Aug 22, 2023 •

edited

Loading

codecov bot commented Sep 4, 2023 •

edited

Loading