MP_bidirect_winograd with xdlops [FP32] by shurale-nkn · Pull Request #454 · ROCm/MIOpen

shurale-nkn · 2020-09-22T13:46:04Z

Added new MP_bidirect_winograd solver with xdlops support.
This solver is tunable.

gfx908 time improvements without new solver tuning

N-C-K-HW-xy-pq-uv	direction	time reduce	acceleration	develop, best algo	this PR
256-128-128-2828-33-11-11	FWD	15%	18%	ConvMPBidirectWinograd<6-3>	ConvMPBidirectWinograd_xdlops<6-3>
256-128-128-2828-33-11-11	BWD	15%	18%	ConvMPBidirectWinograd<6-3>	ConvMPBidirectWinograd_xdlops<6-3>
256-256-256-1414-33-11-11	FWD	39%	64%	ConvBinWinogradRxSf2x3	ConvMPBidirectWinograd_xdlops<5-3>
256-256-256-1414-33-11-11	BWD	38%	61%	ConvBinWinograd3x3U	ConvMPBidirectWinograd_xdlops<5-3>
256-512-512-77-33-11-11	FWD	54%	117%	ConvMPBidirectWinograd<6-3>	ConvMPBidirectWinograd_xdlops<4-3>
256-512-512-77-33-11-11	BWD	51%	104%	ConvMPBidirectWinograd<6-3>	ConvMPBidirectWinograd_xdlops<4-3>

daniellowell · 2020-09-24T15:10:00Z

@shurale-nkn Can you share target configs and performance data?

shurale-nkn · 2020-09-24T16:18:08Z

@shurale-nkn Can you share target configs and performance data?

@daniellowell
If you need more data, I can create a private thread.

atamazov · 2020-09-25T22:28:58Z

@shurale-nkn

If you need more data, I can create a private thread.

Better open an ticket for discussions at https://github.com/AMDComputeLibraries/MLOpen

removed RunAndMeasureSolution from ConvMPBidirectWinograd_xdlops solver

ConvMPBidirectWinograd_xdlops<2-3> returned

asroy · 2020-09-28T22:00:14Z

@shurale-nkn Has rocblas been tuned for these configs?

atamazov

Minor issues

shurale-nkn · 2020-09-29T11:44:58Z

@shurale-nkn Has rocblas been tuned for these configs?

@asroy
If we speak about ConvMPBidirectWinograd algo in develop, then no, because ticket for tuning created a week ago. They wouldn't have had time to do it.

If about ConvMPBidirectWinograd_xdlops, than no, ConvHipImplicitGemmForwardV4R4Xdlops not tuned, used config from EuristicInit .

Kirpich30000

minor changes

Kirpich30000

LGTM

src/solver.cpp conflict resolved

codecov · 2020-10-02T15:44:29Z

Codecov Report

Merging #454 into develop will decrease coverage by 0.07%.
The diff coverage is 25.86%.

@@             Coverage Diff             @@
##           develop     #454      +/-   ##
===========================================
- Coverage    53.24%   53.16%   -0.08%     
===========================================
  Files          296      296              
  Lines        44986    45096     +110     
===========================================
+ Hits         23953    23976      +23     
- Misses       21033    21120      +87

Impacted Files	Coverage Δ
src/include/miopen/solver.hpp	`15.26% <0.00%> (-0.62%)`	⬇️
src/solver/conv_MP_bidirectional_winograd.cpp	`18.86% <20.65%> (+8.34%)`	⬆️
src/mlo_dir_conv.cpp	`88.52% <100.00%> (ø)`
src/solver.cpp	`91.62% <100.00%> (+0.43%)`	⬆️
src/md5.cpp	`90.57% <0.00%> (-2.90%)`	⬇️
src/ocl/clhelper.cpp	`86.66% <0.00%> (-1.12%)`	⬇️

Continue to review full report at Codecov.

Legend - Click here to learn more
Δ = absolute <relative> (impact), ø = not affected, ? = missing data
Powered by Codecov. Last update 26b60d3...4eef1a1. Read the comment docs.

daniellowell · 2020-10-05T14:27:10Z

@shurale-nkn Merge conflicts and blocking review. Please update this ASAP, otherwise it will not get into ROCm 3.10.

Kirpich30000

LGTM

atamazov

Good job of reusing existing kernels/solvers

atamazov · 2020-10-05T21:33:53Z

Please do not forget to document MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_WORKSPACE_MAX in separate PR. Also it would be nice to get rid of // cppcheck-suppress unreadVariable for WinoOffsets.

shurale-nkn · 2020-10-06T12:25:00Z

@daniellowell
Update! Git doesn't display this, but the testing was completed successfully.

* add xdlops_winograd_solver * format * tidy fix * small code fix removed RunAndMeasureSolution from ConvMPBidirectWinograd_xdlops solver * Multi-pass Winograd documentation updated ConvMPBidirectWinograd_xdlops<2-3> returned * WinoOffsets cppcheck-suppress unreadVariable

shurale-nkn added 3 commits September 18, 2020 13:26

add xdlops_winograd_solver

572b294

format

a15d75b

tidy fix

be4e0ad

shurale-nkn added enhancement value_middle performance complexity_high tuning labels Sep 22, 2020

shurale-nkn self-assigned this Sep 22, 2020

shurale-nkn requested review from Kirpich30000 and atamazov and removed request for atamazov September 22, 2020 13:49

atamazov suggested changes Sep 25, 2020

View reviewed changes

Comment thread src/solver/conv_MP_bidirectional_winograd.cpp

shurale-nkn added 5 commits September 26, 2020 21:18

small code fix

33dd659

removed RunAndMeasureSolution from ConvMPBidirectWinograd_xdlops solver

Multi-pass Winograd documentation updated

cdeeddd

ConvMPBidirectWinograd_xdlops<2-3> returned

Merge remote-tracking branch 'develop' into HEAD

f2e1fe5

post merge Update

5c0edd2

format

d9ad409

atamazov reviewed Sep 28, 2020

View reviewed changes

Comment thread src/solver/conv_MP_bidirectional_winograd.cpp Outdated

Comment thread src/solver/conv_MP_bidirectional_winograd.cpp Outdated

atamazov reviewed Sep 28, 2020

View reviewed changes

Comment thread src/solver/conv_MP_bidirectional_winograd.cpp Outdated

Comment thread src/include/miopen/solver.hpp

atamazov suggested changes Sep 28, 2020

View reviewed changes

Kirpich30000 suggested changes Sep 29, 2020

View reviewed changes

Comment thread src/kernels/xform_bidirect_winograd_code.inc Outdated

code clearing

5f3bbf0

tidy fix

a698519

daniellowell changed the title ~~MP_bidirect_winograd with xdlops~~ MP_bidirect_winograd with xdlops [FP32] Sep 29, 2020

WinoOffsets cppcheck-suppress unreadVariable

e696ac8

Kirpich30000 previously approved these changes Oct 1, 2020

View reviewed changes

Merge remote-tracking branch 'develop' into MP_BD_winograd_xdlops_PR

b03dc4a

src/solver.cpp conflict resolved

shurale-nkn dismissed Kirpich30000’s stale review via b03dc4a October 2, 2020 13:23

Merge remote-tracking branch 'develop' into MP_BD_winograd_xdlops_PR

4eef1a1

shurale-nkn requested a review from atamazov October 5, 2020 16:57

Kirpich30000 approved these changes Oct 5, 2020

View reviewed changes

daniellowell added this to the ROCm 3.10 milestone Oct 5, 2020

atamazov reviewed Oct 5, 2020

View reviewed changes

Comment thread src/solver/conv_MP_bidirectional_winograd.cpp

Comment thread src/solver/conv_MP_bidirectional_winograd.cpp

atamazov approved these changes Oct 5, 2020

View reviewed changes

atamazov added IMPACTS_FIND_DB IMPACTS_PERFORMANCE_DB and removed tuning labels Oct 5, 2020

daniellowell merged commit fcf4ed8 into develop Oct 6, 2020

shurale-nkn deleted the MP_BD_winograd_xdlops_PR branch February 2, 2021 16:00

Conversation

shurale-nkn commented Sep 22, 2020 • edited by atamazov Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

daniellowell commented Sep 24, 2020

Uh oh!

shurale-nkn commented Sep 24, 2020 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

atamazov commented Sep 25, 2020

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

asroy commented Sep 28, 2020

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

atamazov left a comment

Choose a reason for hiding this comment

Uh oh!

shurale-nkn commented Sep 29, 2020

Uh oh!

Kirpich30000 left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Kirpich30000 left a comment

Choose a reason for hiding this comment

Uh oh!

codecov bot commented Oct 2, 2020 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Codecov Report

Uh oh!

daniellowell commented Oct 5, 2020

Uh oh!

Kirpich30000 left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

atamazov left a comment

Choose a reason for hiding this comment

Uh oh!

atamazov commented Oct 5, 2020

Uh oh!

shurale-nkn commented Oct 6, 2020

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

shurale-nkn commented Sep 22, 2020 •

edited by atamazov

Loading

shurale-nkn commented Sep 24, 2020 •

edited

Loading

codecov bot commented Oct 2, 2020 •

edited

Loading