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.

Sign up for GitHub

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

Jump to bottom

selecting the convolution type in MIOpenDriver #57

Closed
aryamazaheri opened this issue Oct 23, 2018 · 11 comments
Closed

selecting the convolution type in MIOpenDriver #57

aryamazaheri opened this issue Oct 23, 2018 · 11 comments

Comments

@aryamazaheri
Copy link

Hi,
How can I use the MIOpen driver to run other types of convolutions like Winograd or FFT? Does the driver select the right convolution based on the given input sizes?
@atamazov
Copy link
Contributor

It selects the fastest algorithm ("type") available for given problem configuration ("input sizes").

@aryamazaheri
Copy link
Author

Thanks for letting me know. Is there any way to know which algorithm is being used?

@dagamayank
Copy link
Contributor

@aryamazaheri you can pass -t 1 on the cmdline, which will output the algorithm being used in numerical notation. You can then lookup here to translate that to the actual algorithm from this table.

@atamazov
Copy link
Contributor

You're welcome. Use -t 1:

atamazov@spb-hlc-fiji1h:~/github/MLOpen2/build/Release$ ./bin/MIOpenDriver convfp16 -x 1 -y 1 -V 0 -t 1
MIOpenDriver: convfp16 -x 1 -y 1 -V 0 -t 1
MIOpen Forward Conv. Algorithm: 1
GPU Kernel Time Forward Conv. Elapsed: 0.116054 ms (average)
MIOpen Backward Data Conv. Algorithm: 1
GPU Kernel Time Backward Data Conv. Elapsed: 0.084463 ms (average)
MIOpen Backward Weights Conv. Algorithm: 1
GPU Kernel Time Backward Weights Conv. Elapsed: 0.419914 ms (average)

Driver prints GPU times and algorithm numbers in this case. Numbers are here:
https://github.com/ROCmSoftwarePlatform/MIOpen/blob/933ccc200489fc6c58512dcf6274577ef6809451/include/miopen/miopen.h#L682-L688

@atamazov
Copy link
Contributor

@dagamayank 👬 ))

@aryamazaheri
Copy link
Author

Great. One more thing, I tried to play around with the input arguments to trigger the MIOpen to choose the Winograd or any other convolution algorithm but it seems that it only chooses the Direct convolution. Can you please show me an example convolution command which triggers Winograd convolution?

@atamazov
Copy link
Contributor

Try -x 3 -y 3 -c 18

@aryamazaheri
Copy link
Author

Hmm. Doesn't use Winograd at all.

./bin/MIOpenDriver conv -F 1 -P 1 -W 32 -H 32 -c 18 -k 5 -x 3 -y 3 -t 1
MIOpenDriver: conv -F 1 -P 1 -W 32 -H 32 -c 18 -k 5 -x 3 -y 3 -t 1
MIOpen Forward Conv. Algorithm: 1
GPU Kernel Time Forward Conv. Elapsed: 0.248961 ms
Forward Convolution Verifies on CPU and GPU (3.75939e-08)

@atamazov
Copy link
Contributor

With -k 5 Winograd is not the fastest one. I use -k 10 and winograd wins.

You can use export MIOPEN_LOG_LEVEL=5 to see the detailed log.

@aryamazaheri
Copy link
Author

Weird! It doesn't choose the Winograd with -k 10 either.

MIOpenDriver: conv -F 1 -P 1 -W 32 -H 32 -c 18 -k 10 -x 3 -y 3 -t 1
MIOpen(HIP): Info [DetectAmdRocmMetadataVersion] AMDHSA_1_0
MIOpen(HIP): Info [FindSolutionImpl] ConvOclBwdWrW2 (not searchable)
MIOpen(HIP): Info [FindSolutionImpl] ConvOclBwdWrW53 (not searchable)
MIOpen(HIP): Info [FindSolutionImpl] ConvOclDirectFwd
MIOpen(HIP): Info [FindRecordUnsafe] Looking for key: 10-30-30-3x3-18-32-32-100-0x0-1x1-1x1-0-NCHW-FP32-B
MIOpen(HIP): Info [FindRecordUnsafe] File is unreadable: /home/mazaheri/.config/miopen//gfx803_36.cd.updb.txt
MIOpen(HIP): Info [FindRecordUnsafe] Looking for key: 10-30-30-3x3-18-32-32-100-0x0-1x1-1x1-0-NCHW-FP32-B
MIOpen(HIP): Info [FindSolutionImpl] ConvOclDirectFwd
MIOpen(HIP): Info [FindRecordUnsafe] Looking for key: 18-32-32-3x3-10-30-30-100-0x0-1x1-1x1-0-NCHW-FP32-F
MIOpen(HIP): Info [FindRecordUnsafe] File is unreadable: /home/mazaheri/.config/miopen//gfx803_36.cd.updb.txt
MIOpen(HIP): Info [FindRecordUnsafe] Looking for key: 18-32-32-3x3-10-30-30-100-0x0-1x1-1x1-0-NCHW-FP32-F
MIOpen(HIP): Info [SetValues] Record under key: 18-32-32-3x3-10-30-30-100-0x0-1x1-1x1-0-NCHW-FP32-F, content inserted: miopenConvolutionFwdAlgoGEMM:gemm,4.544,583200,0_0_900_162_900_900_10_162
MIOpen(HIP): Info [FindSolutionImpl] ConvBinWinogradRxS (not searchable)
MIOpen(HIP): Info [SetValues] Record under key: 18-32-32-3x3-10-30-30-100-0x0-1x1-1x1-0-NCHW-FP32-F, content inserted: miopenConvolutionFwdAlgoWinograd:ConvBinWinogradRxS,0.40112,0,18x32x32x3x3x10x30x30x100xNCHWxFP32x1
MIOpen(HIP): Info [FindSolutionImpl] ConvOclDirectFwd
MIOpen(HIP): Info [FindRecordUnsafe] Looking for key: 18-32-32-3x3-10-30-30-100-0x0-1x1-1x1-0-NCHW-FP32-F
MIOpen(HIP): Info [FindRecordUnsafe] File is unreadable: /home/mazaheri/.config/miopen//gfx803_36.cd.updb.txt
MIOpen(HIP): Info [FindRecordUnsafe] Looking for key: 18-32-32-3x3-10-30-30-100-0x0-1x1-1x1-0-NCHW-FP32-F
MIOpen(HIP): Info [DirConvFindCore] ConvOclDirectFwd: MIOpenConvUni: 0.3776 < 3.40282e+38
MIOpen(HIP): Info [DirConvFindCore] Selected: ConvOclDirectFwd: MIOpenConvUni: 0.3776, workspce_sz = 0
MIOpen(HIP): Info [SetValues] Record under key: 18-32-32-3x3-10-30-30-100-0x0-1x1-1x1-0-NCHW-FP32-F, content inserted: miopenConvolutionFwdAlgoDirect:ConvOclDirectFwd,0.3776,0,18x32x32x3x3x10x30x30x100xNCHWxFP32x1
MIOpen(HIP): Info [FindConvFwdAlgorithm] miopenConvolutionFwdAlgoDirect	0.3776	0
MIOpen(HIP): Info [FindConvFwdAlgorithm] miopenConvolutionFwdAlgoWinograd	0.40112	0
MIOpen(HIP): Info [FindConvFwdAlgorithm] miopenConvolutionFwdAlgoGEMM	4.544	583200
MIOpen(HIP): Info [ConvolutionForward] workspace = 0
MIOpen(HIP): Info [ConvolutionForward] workspace = 0
MIOpen(HIP): Info [ConvolutionForward] workspace = 0
MIOpen(HIP): Info [ConvolutionForward] workspace = 0
MIOpen(HIP): Info [ConvolutionForward] workspace = 0
MIOpen(HIP): Info [ConvolutionForward] workspace = 0
MIOpen(HIP): Info [ConvolutionForward] workspace = 0
MIOpen(HIP): Info [ConvolutionForward] workspace = 0
MIOpen(HIP): Info [ConvolutionForward] workspace = 0
MIOpen(HIP): Info [ConvolutionForward] workspace = 0
MIOpen Forward Conv. Algorithm: 1
GPU Kernel Time Forward Conv. Elapsed: 0.384800 ms
Forward Convolution Verifies on CPU and GPU (4.17085e-08)

@aryamazaheri
Copy link
Author

OK. With -k 11 it chooses Winograd!

odellus pushed a commit to synpon/MIOpen that referenced this issue Nov 15, 2018
@ce1adon
Merge pull request ROCm#57 from AMDComputeLibraries/develop
19f97f3 
up2dev
@alexandraBara alexandraBara mentioned this issue Sep 11, 2020
Closed
cderb added a commit that referenced this issue Aug 3, 2022
@cderb
Squashed 'fin/' changes from 53d2563fe..30d699b9e
41b34da 
30d699b9e Perf Eval Update (#60)
3535b948c PerfCompile and PerfEval changes (#59)
de79468d2 remove unneccessary solution check, add check for previously modified kernel names (#56)
6924286a2 miopen hash update (#55)
530399575 Refactor googletest infra to align with MIOpen (#53)
71c50d146 Datatype fix for BN (#57)
8abe2f5c6 Perf Eval updates, Add find info (#51)
e1c1ef0f5 filter find compile by solver input (#54)
722feea66 sp/chk precomp kernel 264 (#41)
b9aba2034 Batch norm find compile (#50)
359f3da80 Fix missing link directives in fin binary (#48)
a4020c1ba Cache Miss Fixes (#46)
2ec7ef44d Enable google test and compiling fin in the CI (#47)
8b6b453bc Applicability support for batch norm (#45)
44323aae9 Perf compile/eval for fin (#42)
ebd9aa6bd update member name (#43)
d6d798efe add cu count (#39)
8e1989a9f Add find option for selecting only dynamic solvers (#38)
0e164bf66 setting json version (#37)
f3f7fed18 Remove function redefinition (#36)
e1de51a58 Performance DB de-serialize test (#34)
043cdcdaa Layout support in Fin (#33)
3a1d58236 Hotfix (#32)
ee3f0d543 4.4 Tuning Bugfixes (#31)
832dbe234 Tunability Reporting (#27)
a564a229f include gfx90a_110 (#28)

git-subtree-dir: fin
git-subtree-split: 30d699b9edc014c6076a9649f849bd3c4588d4ab
averinevg pushed a commit that referenced this issue Aug 19, 2022
@cderb @junliume
[Enhancement][Fin] Enhance test for pdb records (#1669)
84d4399 
* add perf cfg validity test to TestSysDbRecord

* remove debug prints

* removing invalid entries from all perf dbs

* VACUUM sqlite

* Squashed 'fin/' changes from 53d2563fe..30d699b9e

30d699b9e Perf Eval Update (#60)
3535b948c PerfCompile and PerfEval changes (#59)
de79468d2 remove unneccessary solution check, add check for previously modified kernel names (#56)
6924286a2 miopen hash update (#55)
530399575 Refactor googletest infra to align with MIOpen (#53)
71c50d146 Datatype fix for BN (#57)
8abe2f5c6 Perf Eval updates, Add find info (#51)
e1c1ef0f5 filter find compile by solver input (#54)
722feea66 sp/chk precomp kernel 264 (#41)
b9aba2034 Batch norm find compile (#50)
359f3da80 Fix missing link directives in fin binary (#48)
a4020c1ba Cache Miss Fixes (#46)
2ec7ef44d Enable google test and compiling fin in the CI (#47)
8b6b453bc Applicability support for batch norm (#45)
44323aae9 Perf compile/eval for fin (#42)
ebd9aa6bd update member name (#43)
d6d798efe add cu count (#39)
8e1989a9f Add find option for selecting only dynamic solvers (#38)
0e164bf66 setting json version (#37)
f3f7fed18 Remove function redefinition (#36)
e1de51a58 Performance DB de-serialize test (#34)
043cdcdaa Layout support in Fin (#33)
3a1d58236 Hotfix (#32)
ee3f0d543 4.4 Tuning Bugfixes (#31)
832dbe234 Tunability Reporting (#27)
a564a229f include gfx90a_110 (#28)

git-subtree-dir: fin
git-subtree-split: 30d699b9edc014c6076a9649f849bd3c4588d4ab

* Squashed 'fin/' changes from 30d699b9e..ea5c844af

ea5c844af fix direction test
3aa412ee1 Update to use revised testSysDbRecord miopen function

git-subtree-dir: fin
git-subtree-split: ea5c844aff8b5d46537aa59034a596fd15cd9e1e

* rename pipe step

* Squashed 'fin/' changes from ea5c844af..c702cb968

c702cb968 format

git-subtree-dir: fin
git-subtree-split: c702cb96800a03b17ee17d03a015dfa38e3883b9

* Squashed 'fin/' changes from c702cb968..d5397abd3

d5397abd3 rename targets

git-subtree-dir: fin
git-subtree-split: d5397abd37b6908bcd96ef750ea5a3ace04cdf3c

* rename archive

Co-authored-by: Jun Liu <Liu.Jun@amd.com>
cderb added a commit that referenced this issue Oct 5, 2022
@cderb
Squashed 'fin/' changes from 53d2563fe..e05dcb421
8875272 
e05dcb421 perf db validation fix (#68)
260d9465d Add INT8 as a data_type v2 (#67)
b6a5b2a77 sync with fin folder in miopen (#62)
0e03399ec prep for Palamida scan (#63)
e6bd05c33 Performance db testing (#61)
30d699b9e Perf Eval Update (#60)
3535b948c PerfCompile and PerfEval changes (#59)
de79468d2 remove unneccessary solution check, add check for previously modified kernel names (#56)
6924286a2 miopen hash update (#55)
530399575 Refactor googletest infra to align with MIOpen (#53)
71c50d146 Datatype fix for BN (#57)
8abe2f5c6 Perf Eval updates, Add find info (#51)
e1c1ef0f5 filter find compile by solver input (#54)
722feea66 sp/chk precomp kernel 264 (#41)
b9aba2034 Batch norm find compile (#50)
359f3da80 Fix missing link directives in fin binary (#48)
a4020c1ba Cache Miss Fixes (#46)
2ec7ef44d Enable google test and compiling fin in the CI (#47)
8b6b453bc Applicability support for batch norm (#45)
44323aae9 Perf compile/eval for fin (#42)
ebd9aa6bd update member name (#43)
d6d798efe add cu count (#39)
8e1989a9f Add find option for selecting only dynamic solvers (#38)
0e164bf66 setting json version (#37)
f3f7fed18 Remove function redefinition (#36)
e1de51a58 Performance DB de-serialize test (#34)
043cdcdaa Layout support in Fin (#33)
3a1d58236 Hotfix (#32)
ee3f0d543 4.4 Tuning Bugfixes (#31)
832dbe234 Tunability Reporting (#27)
a564a229f include gfx90a_110 (#28)

git-subtree-dir: fin
git-subtree-split: e05dcb42187f05fe0d0d1b05b822dc4b750f199e
junliume added a commit that referenced this issue Oct 6, 2022
@cderb @junliume
Perf DB cleanup + Fin pdb verification fix (#1810)
e7d3c60 
* remove datatype 0,1 from perf_db

* rm invalid fp16 entries from pdb

* Squashed 'fin/' changes from 53d2563fe..e05dcb421

e05dcb421 perf db validation fix (#68)
260d9465d Add INT8 as a data_type v2 (#67)
b6a5b2a77 sync with fin folder in miopen (#62)
0e03399ec prep for Palamida scan (#63)
e6bd05c33 Performance db testing (#61)
30d699b9e Perf Eval Update (#60)
3535b948c PerfCompile and PerfEval changes (#59)
de79468d2 remove unneccessary solution check, add check for previously modified kernel names (#56)
6924286a2 miopen hash update (#55)
530399575 Refactor googletest infra to align with MIOpen (#53)
71c50d146 Datatype fix for BN (#57)
8abe2f5c6 Perf Eval updates, Add find info (#51)
e1c1ef0f5 filter find compile by solver input (#54)
722feea66 sp/chk precomp kernel 264 (#41)
b9aba2034 Batch norm find compile (#50)
359f3da80 Fix missing link directives in fin binary (#48)
a4020c1ba Cache Miss Fixes (#46)
2ec7ef44d Enable google test and compiling fin in the CI (#47)
8b6b453bc Applicability support for batch norm (#45)
44323aae9 Perf compile/eval for fin (#42)
ebd9aa6bd update member name (#43)
d6d798efe add cu count (#39)
8e1989a9f Add find option for selecting only dynamic solvers (#38)
0e164bf66 setting json version (#37)
f3f7fed18 Remove function redefinition (#36)
e1de51a58 Performance DB de-serialize test (#34)
043cdcdaa Layout support in Fin (#33)
3a1d58236 Hotfix (#32)
ee3f0d543 4.4 Tuning Bugfixes (#31)
832dbe234 Tunability Reporting (#27)
a564a229f include gfx90a_110 (#28)

git-subtree-dir: fin
git-subtree-split: e05dcb42187f05fe0d0d1b05b822dc4b750f199e

* fix clang-format issue

Co-authored-by: Jun Liu <Liu.Jun@amd.com>
cderb added a commit that referenced this issue Nov 21, 2022
@cderb
Squashed 'fin/' changes from 53d2563fe..49e3e3a62
4042531 
49e3e3a62 clang format
db80b1777 update to using TestPerfCfgParams for pdb validity checks
e48a4fd3a format
a4f85842c exception for non-tunable solvers in params check
d58c42bbd Check params at end of perf tuning (#70)
1a3b47c7b Return status for failed compile commands (#69)
d59962752 out_layout -> in_layout
6ba7a8f3f Rename conv_mode to mode (#64)
513a3da1b [bg/LWPTUNA-173] (#65)
e05dcb421 perf db validation fix (#68)
260d9465d Add INT8 as a data_type v2 (#67)
b6a5b2a77 sync with fin folder in miopen (#62)
0e03399ec prep for Palamida scan (#63)
e6bd05c33 Performance db testing (#61)
30d699b9e Perf Eval Update (#60)
3535b948c PerfCompile and PerfEval changes (#59)
de79468d2 remove unneccessary solution check, add check for previously modified kernel names (#56)
6924286a2 miopen hash update (#55)
530399575 Refactor googletest infra to align with MIOpen (#53)
71c50d146 Datatype fix for BN (#57)
8abe2f5c6 Perf Eval updates, Add find info (#51)
e1c1ef0f5 filter find compile by solver input (#54)
722feea66 sp/chk precomp kernel 264 (#41)
b9aba2034 Batch norm find compile (#50)
359f3da80 Fix missing link directives in fin binary (#48)
a4020c1ba Cache Miss Fixes (#46)
2ec7ef44d Enable google test and compiling fin in the CI (#47)
8b6b453bc Applicability support for batch norm (#45)
44323aae9 Perf compile/eval for fin (#42)
ebd9aa6bd update member name (#43)
d6d798efe add cu count (#39)
8e1989a9f Add find option for selecting only dynamic solvers (#38)
0e164bf66 setting json version (#37)
f3f7fed18 Remove function redefinition (#36)
e1de51a58 Performance DB de-serialize test (#34)
043cdcdaa Layout support in Fin (#33)
3a1d58236 Hotfix (#32)
ee3f0d543 4.4 Tuning Bugfixes (#31)
832dbe234 Tunability Reporting (#27)
a564a229f include gfx90a_110 (#28)

git-subtree-dir: fin
git-subtree-split: 49e3e3a62a7cc54adacbeea95680d35f9a4685de
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Assignees
No one assigned
Labels
None yet
Projects
None yet
Milestone
No milestone
Development

No branches or pull requests
3 participants
@aryamazaheri @dagamayank @atamazov