Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

implicit gemm backward weight with xdlops and atomic add #413

Closed
wants to merge 26 commits into from

Conversation

ltqin
Copy link
Contributor

@ltqin ltqin commented Sep 2, 2020

Realize backward weight function, some optimization:

  1. use v2 version of gridwise-GEMM
  2. use buffer_atomic_fadd

src/mlo_dir_conv.cpp Outdated Show resolved Hide resolved
src/solver.cpp Outdated Show resolved Hide resolved
src/solver/conv_hip_implicit_gemm_wrw_v4r4_xdlops.cpp Outdated Show resolved Hide resolved
src/solver/conv_hip_implicit_gemm_wrw_v4r4_xdlops.cpp Outdated Show resolved Hide resolved
src/hip/hip_build_utils.cpp Outdated Show resolved Hide resolved
@whchung whchung self-requested a review September 3, 2020 17:19
@daniellowell
Copy link
Contributor

@ltqin Please get this PR into shape.

@shaojiewang
Copy link
Contributor

If this PR is ready to review, please remove the Draft tag to start the Jenkins test.

@shaojiewang
Copy link
Contributor

shaojiewang commented Sep 10, 2020

@ltqin Please resolve conflicts.

@asroy asroy requested a review from shaojiewang September 10, 2020 15:05
@asroy asroy self-requested a review September 10, 2020 15:24
@atamazov
Copy link
Contributor

@shaojiewang

...please remove the Draft tag to start the Jenkins test.

Our CI does not observe Draft status. A new testing round is started after each commit into branch.

@shaojiewang
Copy link
Contributor

@shaojiewang

...please remove the Draft tag to start the Jenkins test.

Our CI does not observe Draft status. A new testing round is started after each commit into branch.

Yes.
And why does this PR never trigger a CI test?

@daniellowell
Copy link
Contributor

@asroy @zjing14 Final reviews.

asroy
asroy previously approved these changes Sep 24, 2020
zjing14
zjing14 previously approved these changes Sep 24, 2020
shaojiewang
shaojiewang previously approved these changes Sep 26, 2020
@ltqin ltqin dismissed stale reviews from shaojiewang, zjing14, and asroy via e7e3dfc September 28, 2020 10:00
@daniellowell
Copy link
Contributor

@atamazov Please unblock this PR is all change requests are resolved.

@asroy asroy self-requested a review September 28, 2020 21:46
asroy
asroy previously approved these changes Sep 28, 2020
Copy link
Contributor

@asleepzzz asleepzzz left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

invokerparams already changed in develop branch
please merge it and change

@daniellowell
Copy link
Contributor

This PR keeps failing in the exact same way:

MIOpen(HIP): Info2 [FindRecordUnsafe] Looking for key 32-26-26-5x5-192-28-28-16-1x1-1x1-1x1-0-NCHW-FP32-W in file /tmp/miopen-tmp-a4ab-21b7-b1b9-7524/miopen.test.find_db
MIOpen(HIP): Info2 [FindRecordUnsafe] Key match: 32-26-26-5x5-192-28-28-16-1x1-1x1-1x1-0-NCHW-FP32-W
MIOpen(HIP): Info2 [FindRecordUnsafe] Contents found: miopenConvolutionBwdWeightsAlgoImplicitGEMM:ConvHipImplicitGemmWrwV4R4Xdlops,1.45585,0,miopenConvolutionBwdWeightsAlgoImplicitGEMM,<unused>;miopenConvolutionBwdWeightsAlgoWinograd:ConvBinWinogradRxSf2x3,0.443042,0,miopenConvolutionBwdWeightsAlgoWinograd,<unused>;miopenConvolutionBwdWeightsAlgoGEMM:gemm,17.9022,12979200,rocBlas,<unused>;miopenConvolutionBwdWeightsAlgoDirect:ConvOclBwdWrW53,0.430721,9830400,miopenConvolutionBwdWeightsAlgoDirect,<unused>
MIOpen(HIP): Info2 [Measure] Db::FindRecord time: 0.082628 ms
MIOpen(HIP): Info [EuristicInit] 32,64,8,32,32,4,0,1
MIOpen(HIP): Info [CompileWrwSolution] solver_id = ConvOclBwdWrW53
MIOpen(HIP): Info2 [GetInvoker] Returning an invoker for problem 32x26x26x5x5x192x28x28x16xNCHWxFP32x1x1x1x1x1x1x1xW and solver ConvOclBwdWrW53
MIOpen(HIP): Info [FindConvBwdWeightsAlgorithm] miopenConvolutionBwdWeightsAlgoDirect	0.430721	9830400
MIOpen(HIP): Info [FindConvBwdWeightsAlgorithm] BWrW Chosen Algorithm: ConvOclBwdWrW53 , 9830400, 0.430721
MIOpen(HIP): Info [Test] Find(), find-db enabled: 61.1031
MIOpen(HIP): Info [Test] Speedup: 2.29189
FAILED: find_db_speedup(2.29189) >= 3(3): /var/jenkins/workspace/MLLibs_MIOpen_PR-413/test/find_db.cpp:198
CMake Error at test_test_find_db.cmake:7 (message):
  Test failed

@atamazov can you assist on this? Looks like the DB access is timing out.

@codecov
Copy link

codecov bot commented Sep 30, 2020

Codecov Report

Merging #413 into develop will decrease coverage by 0.67%.
The diff coverage is 0.89%.

Impacted file tree graph

@@             Coverage Diff             @@
##           develop     #413      +/-   ##
===========================================
- Coverage    54.26%   53.58%   -0.68%     
===========================================
  Files          293      294       +1     
  Lines        44048    44608     +560     
===========================================
+ Hits         23904    23905       +1     
- Misses       20144    20703     +559     
Impacted Files Coverage Δ
src/include/miopen/solver.hpp 17.34% <0.00%> (-0.62%) ⬇️
src/mlo_dir_conv.cpp 88.52% <ø> (ø)
.../solver/conv_hip_implicit_gemm_wrw_v4r4_xdlops.cpp 0.54% <0.54%> (ø)
src/solver.cpp 91.00% <100.00%> (+0.09%) ⬆️
src/include/miopen/find_controls.hpp 80.00% <0.00%> (-4.00%) ⬇️
src/md5.cpp 90.57% <0.00%> (-2.90%) ⬇️
src/include/miopen/sqlite_db.hpp 86.08% <0.00%> (+0.43%) ⬆️

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 49f56c3...3f9badc. Read the comment docs.

@daniellowell
Copy link
Contributor

This PR moved to #472

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

None yet

7 participants