-
Notifications
You must be signed in to change notification settings - Fork 221
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
MIOpenTensile Part2 : Support low precision types #556
Merged
Merged
Changes from all commits
Commits
Show all changes
166 commits
Select commit
Hold shift + click to select a range
3201498
add miopentensile path
ce1adon 46daceb
use miotensile in hip backend
ce1adon bb28471
hip tidy
ce1adon 3c6d5bd
update miotensile
ce1adon a0d8cbc
generalize miotensile stride
ce1adon 93e51f2
disable miopentensile in ocl backend
ce1adon 0aa1a64
remove ocl backend consecutive identical branches
ce1adon e00580c
disable use of miotensile in hip-clang
ce1adon 583dd50
Merge branch 'develop' into miotensile
ce1adon 736940c
update miopentensile version
ce1adon 5e29d30
Merge branch 'develop' into miotensile
ce1adon 33b0fbd
Merge branch 'develop' into miotensile
ce1adon 0998746
Merge branch 'develop' into miotensile
ce1adon 818544f
use const arg in gemm
ce1adon c6e5611
miot introduce transpose feature
ce1adon 7bfe9f7
update miot version
ce1adon 1c02111
Merge branch 'develop' into miotensile
ce1adon 2a73f16
disable miotensile in hip-clang in cmakelists
ce1adon 5d0639b
disable miotensile except in hcc in cmakelists
ce1adon 17f5a59
disable ocl wrw2 solvers
ce1adon 836fbd4
Merge branch 'develop' into miotensile
ce1adon 1beaf51
Merge branch 'disable-oclwrw' into miotensile
ce1adon bd706c5
clang tidy
ce1adon cecd3e9
Merge branch 'disable-oclwrw' into miotensile
ce1adon d691097
enable miopentensile built with hipclang
ce1adon 7d2c634
add miotensile in hip-clang docker
ce1adon 2bd15dc
add python3-venv in hip-clang docker
ce1adon 9580b0a
hip-clang docker update
ce1adon 3460907
update hip-clang docker compiler
93148be
update miopentensile info
3c04ea2
Merge branch 'develop' into miotensile
ce1adon a744b47
disable miotensile with hcc
ce1adon b6ddabe
set miopentensile off & add dedicated test stage
ce1adon e8f2a9d
Revert "clang tidy"
ce1adon c5c18a8
Revert "disable ocl wrw2 solvers"
ce1adon 785fda2
workaround for issue 2534
ce1adon 048799a
Adjust test config
ce1adon 653e6df
Adjust gemm env_var
ce1adon 6ca9951
use miotensile in winograd3x3multipass solver
ce1adon c6d3d7c
revert gemm enforce backend values
ce1adon 3692844
skip igemm tests in test stage iv (miotensile tests)
ce1adon 658f441
skip igemm tests in test stage iv
ce1adon f2e3bf5
[skip ci] add comment for cmake's skip rule
ce1adon 8678cea
use parallel jobs for tensile
215bdff
specify gpu architecture(s) in miotensile build process
1780dbb
fix jenkins docker arg
716b09f
fix jenkins docker arg string
a15cde6
support low-precision datatypes
e17e05e
enable bfloat16 miopentensile path
fc738db
Merge branch 'develop' into miotensile
ce1adon fb452bf
rename gpu target param for miotensile
555aa2c
Merge branch 'miotensile' of https://github.com/ROCmSoftwarePlatform/…
c089443
specify gpu architecture(s) in Jenkinsfile buildJob
83c8082
relocate miotensile gpu targets arg
fd94b50
add symlink to /opt/rocm in Dockerfile
9a7ba7f
revise Dockerfile to pass GPU architecture flag for dependencies
10fa0fd
Merge branch 'develop' into miotensile
ce1adon 67a06c7
add jenkins tests for last/latest version of miopentensile
569bd7b
remove workaround for issue #2534
6b28286
winograd workaround condition update: current miopentensile doesn't s…
8d50d0e
clang format
20b6e44
Merge branch 'develop' into miotensile
ce1adon 29c0524
fix jenkinsfile typo & move GPU architecture flag to cget init in hip…
ce1adon 402b02e
revise conv_multipass_wino3x3WrW GEMM workaround for miopentensile
ce1adon 99581d7
clang format
ce1adon 0e2dd07
miopentensile stage rename
ce1adon b49fee4
reset miopentensile version in hcc
ce1adon 4d7cc64
revert wino3x3WrW datatype constraint in miopentensile
ce1adon ee6b154
Merge branch 'develop' into miotensile
3402052
remove miopentensile from dependency list & update version
ce1adon cc0b6e7
Merge branch 'miotensile' into miotensile-lowprec
c7c17ea
remove datatype limit in multipass wino3x3wrw
bae6b2f
add low precision tests for miopentensile in Jenkins
7678578
update tensile version for int8
ce1adon dcdec86
miopentensile to 3-10-x
ce1adon f6bc06a
revert miopentensile latest version
ce1adon b76f85f
revise int8 conv test validation check based on latest conv host
ce1adon f858aa1
revise format
ce1adon 62156d6
fix bug & clang format
ce1adon 98b2acd
fix merge conflicts
ce1adon 0449ee0
skip not applicable tests
ce1adon 4a151de
skip redundant tests for low precision
ce1adon c6f86ea
Merge branch 'develop' into miotensile-lowprec
ce1adon 4e2f8ec
Merge branch 'develop' into miotensile-lowprec
ce1adon 964b839
add symlink & update miot version
ce1adon 47dbb6c
[skip ci] include miopentensile in MP BD winograd path
ce1adon d62f2de
update target id feature in docker and jenkins files
ce1adon 2460147
Merge branch 'develop' into miotensile-lowprec
ce1adon 0fd4160
skip miopentensile install when target id env not present
ce1adon 65477cf
Merge branch 'develop' into miotensile-lowprec
ce1adon f1cd1aa
targetid-fixes(01) Simplify & improve logging of build commands. Some…
atamazov 4977bde
targetid-fixes(02) Refactor: device_name.hpp -> handle.cpp
atamazov 79a8a11
targetid-fixes(03) Fix OCL build error. Remove dead code: dumpKernel()
atamazov 5e672e0
targetid-fixes(04) Move HIP_PACKAGE_VERSION_FLAT to config.h
atamazov 74a4fc5
targetid-fixes(05) Refactor GetDeviceName()
atamazov df4bd5a
targetid-fixes(06) Rework workaround for issue 1711
atamazov 125010e
targetid-fixes(07) Fix CompileTime logging for generated sources (MIO…
atamazov f6c6442
targetid-fixes(08) Store target properties in the handle (HIP)
atamazov 30b9a12
targetid-fixes(09) Store target properties in the handle (OCL)
atamazov ec11cec
targetid-fixes(10) Fix TargetProperties::Init()
atamazov ed74896
targetid-fixes(11) Simplifications & formatting
atamazov 3280330
targetid-fixes(12) Clang-tidy fix
atamazov fe31c30
targetid-fixes(13) [nfc] Target features: init to defaults, engage in…
atamazov 5edcc27
Merge branch 'develop' into miotensile-lowprec
ce1adon ab4106e
update GEMM macro
ce1adon 18a6581
Merge branch 'miotensile-lowprec' of https://github.com/ROCmSoftwareP…
ce1adon 165189f
targetid-fixes(14) [nfc] Fix: Add missing Handle::GetTargetProperties…
atamazov c8b3557
targetid-fixes(15) [nfc] Refactor & Pass TargetProperties to HipBuild…
atamazov 56a1857
targetid-fixes(16) Clang-tidy fix
atamazov 5b232eb
targetid-fixes(17) Use target features to in perf- and find-databases
atamazov 24e7c79
targetid-fixes(18) Fix bug from PR 511 (comparison of external_tool_v…
atamazov 5d58fcf
targetid-fixes(19) Use xnack feature in HIP builds
atamazov d8b8b8f
targetid-fixes(20) HIP builds: Use --cuda-gpu-arch instead of --offlo…
atamazov f3b88f8
targetid-fixes(21) OCL builds (HIP be): Update CO version selection o…
atamazov 875a38f
targetid-fixes(22) Clang-tidy fix
atamazov ee53f31
targetid-fixes(23) OCL builds (OCL be): Enable logging of build options
atamazov c6904be
targetid-fixes(24) OCL builds (OCL be): Avoid duplication of error me…
atamazov b7b8d99
targetid-fixes(25) Workaround for cppcheck
atamazov 085245c
targetid-fixes(26) Fix COMGR build errors & some 'unused macro' warnings
atamazov ad8d5e7
targetid-fixes(27) [tests] Enable W/A for missing macros in PCH for a…
atamazov 2717fcc
re-group miopentensile tests & rename targetID docker
ce1adon e52e9d7
Merge remote-tracking branch 'origin/develop' into miotensile-lowprec
ce1adon abeac2b
Merge remote-tracking branch 'origin/targetid-fixes' into miotensile-…
ce1adon cc03d67
disable extra implicit gemm test in miopentensile tests
ce1adon 845253a
Merge remote-tracking branch 'origin/develop' into miotensile-lowprec
ce1adon 86e7f7a
rename targetID dockers
ce1adon 09b04fa
Merge remote-tracking branch 'origin/develop' into miotensile-lowprec
ce1adon d248bff
Merge branch 'develop' into miotensile-lowprec
ce1adon 8d4540a
update MIOpenTensile and docker image
ce1adon 76a124f
update Tensile arch string
ce1adon fb7b62a
workaround for hipcc
ce1adon 691113b
restrict hipcc workaround and env var export in targetID env
ce1adon 361249f
Merge branch 'develop' into miotensile-lowprec
ce1adon 83caf61
fix missing property target_id in JenkinsFile
ce1adon 35fc6d9
test miopentensile on latest CI framework
ce1adon cb14a16
Add MIOPEN_DEBUG_HIP_KERNELS=0 to reduce testing time
ce1adon 439042c
rearrange miopentensile tests
ce1adon 2e0346d
regulate stage names
ce1adon fb54873
Revert "test miopentensile on latest CI framework"
ce1adon 301608a
merge develop
ce1adon 60c8e5c
fix conflicts
ce1adon 002e51b
clang format
ce1adon 93a1041
merge tensile stages
ce1adon fe6e75d
replace tensile stages with function
ce1adon 2bb47ef
test miopentensile stages
ce1adon cb4d064
remove image param
ce1adon ca67bcc
accept custom docker image
ce1adon 1153d5d
Revert "test miopentensile stages"
ce1adon a9ac78c
remove MIOPEN_DEBUG_IMPLICIT_GEMM_NON_XDLOPS_INLINE_ASM in jenkins en…
ce1adon 89f94c3
build docker with parallel-jobs=4
ce1adon 23134d9
revert comment
ce1adon 499cb91
update miopentensile version
ce1adon 69d736a
revert build docker with parallel-jobs=4
ce1adon 7f7cda0
update mainline build versoin
ce1adon a688a71
unify image name
ce1adon eea15d8
revert mainline build version
ce1adon a36086d
add comment to distinguish GEMM logic
ce1adon 86629f2
disable MIOpenTensile stage
ce1adon 360fa6b
Merge branch 'develop' into miotensile-lowprec
atamazov 0fc0fd1
Refactor test/CmakeLists.txt. Separate MIOpenTensile related code fro…
atamazov 6ef88d7
Merge branch 'develop' into miotensile-lowprec
atamazov 255122b
Fix stage names. Add MIOPENTENSILE_LATEST parameter.
atamazov f71a778
MIOpenTensile tests OFF by default.
atamazov 0eefeb4
Fix syntax error
atamazov aa49456
Enable testing of LATEST by default
atamazov ef3c289
Fix syntax error
atamazov File filter
Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Large diffs are not rendered by default.
Oops, something went wrong.
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
|
@@ -74,9 +74,9 @@ MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_COMPILE_ONLY) | |
#if MIOPEN_USE_GEMM | ||
#ifdef CPPCHECK | ||
// Keep the value unknown in cppcheck since this can differ between opencl and hip | ||
static bool IsUseRocBlas; | ||
static bool IsBF16PathValid; | ||
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. @atamazov |
||
#else | ||
static const bool IsUseRocBlas = (MIOPEN_USE_ROCBLAS == 1); | ||
static const bool IsBF16PathValid = (MIOPEN_USE_ROCBLAS == 1 || MIOPEN_USE_MIOPENTENSILE == 1); | ||
#endif | ||
|
||
static inline bool IsAnyBufferBF16(const TensorDescriptor& xDesc, | ||
|
@@ -669,7 +669,7 @@ bool ConvolutionDescriptor::IsGemmApplicableWrw(const TensorDescriptor& dyDesc, | |
{ | ||
#if MIOPEN_USE_GEMM | ||
if(!miopen::IsDisabled(MIOPEN_DEBUG_CONV_GEMM{}) && | ||
!(IsAnyBufferBF16(xDesc, dyDesc, dwDesc) && !IsUseRocBlas)) | ||
!(IsAnyBufferBF16(xDesc, dyDesc, dwDesc) && !IsBF16PathValid)) | ||
{ | ||
const std::size_t spatial_dim = GetSpatialDimension(); | ||
const auto wei_spatial = boost::adaptors::slice(dwDesc.GetLengths(), 2, 2 + spatial_dim); | ||
|
@@ -1644,7 +1644,7 @@ void ConvolutionDescriptor::FindConvBwdWeightsAlgorithm(Handle& handle, | |
perf_db = UserFindDbRecord::TryLoad(handle, problem, [&](DbRecord& record) { | ||
#if MIOPEN_USE_GEMM | ||
if(!miopen::IsDisabled(MIOPEN_DEBUG_CONV_GEMM{}) && | ||
!(IsAnyBufferBF16(xDesc, dyDesc, dwDesc) && !IsUseRocBlas)) | ||
!(IsAnyBufferBF16(xDesc, dyDesc, dwDesc) && !IsBF16PathValid)) | ||
{ | ||
const bool time_precision = (!IsDisabled(MIOPEN_CONV_PRECISE_ROCBLAS_TIMING{})); | ||
|
||
|
@@ -1934,7 +1934,7 @@ void ConvolutionDescriptor::BackwardWeightsGemm(Handle& handle, | |
{ | ||
MIOPEN_THROW("GEMM convolution is disabled"); | ||
} | ||
if(IsAnyBufferBF16(tensors.xDesc, tensors.dyDesc, tensors.dwDesc) && !IsUseRocBlas) | ||
if(IsAnyBufferBF16(tensors.xDesc, tensors.dyDesc, tensors.dwDesc) && !IsBF16PathValid) | ||
{ | ||
MIOPEN_THROW("GEMM convolution is unsupported"); | ||
} | ||
|
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Oops, something went wrong.
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
[Q] Why 6416? Shall we promote this number in the future (and when)?
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Remember how long this PR has been staying here? :)
Let me know which mainline# is used currently
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Hmm... Do we want this to advance automatically, to the latest mainline build?
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
What is the latest mainline build number?
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
You can find it here, but I do not know how good or bad is it.