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

Added missing weight sorting as WTI calculation in fusion #3105

Merged
merged 9 commits into from
Jul 24, 2024
Merged

Added missing weight sorting as WTI calculation in fusion #3105

merged 9 commits into from
Jul 24, 2024

Conversation

DrizztDoUrden
Copy link
Contributor

This was missing since #2388. Reimplemented as WTI to use in the fallback. Numbers are preserved mostly as were before #2388 with some changes due to how WTI is used: they have been normalized to range of (0;1] and 0 values replaced by -2.

Thanks to @atamazov for noticing.

averinevg
averinevg reviewed Jul 10, 2024
View reviewed changes
src/fusion.cpp Show resolved Hide resolved
src/solver/conv_bin_winoRxS_fused.cpp Outdated Show resolved Hide resolved
src/solver/conv_ocl_dir2Dfwd_fused.cpp Outdated Show resolved Hide resolved
src/solver/conv_winoRxS_fused.cpp Outdated Show resolved Hide resolved
src/solver/conv_asm_1x1u_bias_activ_fused.cpp Outdated Show resolved Hide resolved
atamazov
atamazov suggested changes Jul 11, 2024
View reviewed changes
Copy link
Contributor

@atamazov atamazov left a comment

Choose a reason for hiding this comment

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

Good, but I have to ask for some quality-related changes.

src/solver/conv_asm_1x1u_bias_activ_fused.cpp Outdated
// Solvers with positive values are considered first priority
// If none found solvers with negative are used, -2.0f meaning a completely unknown value
// So lower absolute values mean higher priority
return -.5f;
Copy link
Contributor

Choose a reason for hiding this comment

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

[Q] Are you pretty sure that ConvBiasActivAsm1x1U is better than all other applicable ones that return wti_unknown?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

No, but they had these priorities before #2388

src/solver/conv_asm_1x1u_bias_activ_fused.cpp Outdated Show resolved Hide resolved
src/solver/conv_bin_winoRxS_fused.cpp Outdated Show resolved Hide resolved
src/solver/conv_bin_winoRxS_fused.cpp Outdated Show resolved Hide resolved
src/include/miopen/solver.hpp Outdated
@@ -108,6 +108,8 @@ struct SolverBase
/// run-time parameters.
virtual bool IsDynamic() const { return false; }

static constexpr float wti_unknown = -2;
Copy link
Contributor

Choose a reason for hiding this comment

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

[notice] google C++ style recommends camelCase with prefix "k", i.e. kWtiUnknown (which is very like a kind of hungarian notation which I personally hate). I would prefer WTI_UNKNOWN for this specific case. Decide yourself, no feedback required.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Caps is generally reserved for macros to indicate caution required when interacting with them. But I change it to any style if you want.

Copy link
Contributor

Choose a reason for hiding this comment

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

Up to you, my comment is just a notice

src/solver/conv_asm_1x1u_bias_activ_fused.cpp Outdated Show resolved Hide resolved
src/solver/conv_ocl_dir2Dfwd_fused.cpp Outdated Show resolved Hide resolved
src/solver/conv_ocl_dir2Dfwd_fused.cpp Outdated Show resolved Hide resolved
src/solver/conv_winoRxS_fused.cpp
const FusionDescription& problem) const
{
const auto conv_problem = problem.GetConvProblem(0, miopen::conv::Direction::Forward);
return conv::ConvBinWinogradRxSf2x3g1().GetWti(ctx, conv_problem);
Copy link
Contributor

Choose a reason for hiding this comment

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

[notice] Maybe multiply with 1.001f just to notice some small overhead?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

I don't think it would be useful because WTI is not meant to be comparable between primitives. I think multiplication by 1.001 would mean that the rest of the kernel is .1% less optimal than non-conv part of its competitors.

src/include/miopen/solver.hpp Outdated
@@ -108,6 +108,8 @@ struct SolverBase
/// run-time parameters.
virtual bool IsDynamic() const { return false; }

static constexpr float wti_unknown = -2;
Copy link
Contributor

Choose a reason for hiding this comment

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

[notice] Introducing a symbol here logically leads us to introduce symbols in other similar places.

DrizztDoUrden added a commit that referenced this pull request Jul 15, 2024
@DrizztDoUrden
Added missing weight sorting as WTI calculation in fusion #3105
53beafa
averinevg
averinevg approved these changes Jul 17, 2024
View reviewed changes
Copy link
Collaborator

@averinevg averinevg left a comment

Choose a reason for hiding this comment

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

LGTM

@DrizztDoUrden
Copy link
Contributor Author

@junliume @JehandadKhan ping

@junliume junliume merged commit 9685f42 into develop Jul 24, 2024
140 of 141 checks passed
@junliume junliume deleted the ddu/missing-weight-sorting branch July 24, 2024 00:21
DrizztDoUrden added a commit that referenced this pull request Jul 26, 2024
@DrizztDoUrden
Support for invoker caching and find modes to fusion find (#3095)
d4e63ee 
* implemented support for invoker caching and find modes to fusion find

* Fixed some errors

[Fixes CI] Fixed tidy hang (#3114)

* Fixed tidy hang by disabling a check on bugged versions

* Added a comment on why this is necessary
Added missing weight sorting as WTI calculation in fusion #3105

Fixed build

fixed fusion to force fallback mode even in hybrid mode

Changed how find results are registered as find 1.0 results for fusion

Added support for having multiple find mode variables for different primitives with different default values.

Changed fusion to use MIOPEN_FIND_MODE_FUSION with Fast mode as default

Reverted fusion defaulting to hybrid mode

Fixed build

Fixed redundant find on failed fallback

tidy fix
This was referenced Jul 26, 2024
Merged
Merged
junliume pushed a commit that referenced this pull request Aug 1, 2024
@DrizztDoUrden
Implemented find modes and fallback for FusionPlanDescriptor::Compile (…
6a240ed 
…#3158)

* Support for invoker caching and find modes to fusion find (#3095)

* implemented support for invoker caching and find modes to fusion find

* Fixed some errors

* [Fixes CI] Fixed tidy hang (#3114)

* Fixed tidy hang by disabling a check on bugged versions

* Added a comment on why this is necessary

* Added missing weight sorting as WTI calculation in fusion #3105

* Fixed build

* fixed fusion to force fallback mode even in hybrid mode

* Changed how find results are registered as find 1.0 results for fusion

* Added support for having multiple find mode variables for different primitives with different default values.

* Changed fusion to use MIOPEN_FIND_MODE_FUSION with Fast mode as default

* Reverted fusion defaulting to hybrid mode

* Fixed build

* Fixed redundant find on failed fallback

* tidy fix
jangdan added a commit to jangdan/MIOpen that referenced this pull request Aug 1, 2024
@jangdan
Squashed commit of the following:
7c864cf 
commit da4ba7f
Author: Artur Wojcik <[email protected]>
Date:   Tue Jul 30 07:37:57 2024 +0200

    [Windows] Temporary workaround on rocMLIR issue (ROCm#3164)

commit bd993eb
Author: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>
Date:   Fri Jul 26 23:38:36 2024 -0700

    Bump rocm-docs-core[api_reference] from 1.5.0 to 1.6.1 in /docs/sphinx (ROCm#3160)

    Bumps [rocm-docs-core[api_reference]](https://github.com/ROCm/rocm-docs-core) from 1.5.0 to 1.6.1.
    - [Release notes](https://github.com/ROCm/rocm-docs-core/releases)
    - [Changelog](https://github.com/ROCm/rocm-docs-core/blob/develop/CHANGELOG.md)
    - [Commits](ROCm/rocm-docs-core@v1.5.0...v1.6.1)

    ---
    updated-dependencies:
    - dependency-name: rocm-docs-core[api_reference]
      dependency-type: direct:production
      update-type: version-update:semver-minor
    ...

    Signed-off-by: dependabot[bot] <[email protected]>
    Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>

commit e158084
Author: BrianHarrisonAMD <[email protected]>
Date:   Wed Jul 24 16:29:22 2024 -0600

    [gtest] Fix log test error (ROCm#3150)

commit 3e983ed
Author: Vasilii Filippov <[email protected]>
Date:   Wed Jul 24 19:31:00 2024 +0200

    Fixed redundant find on failed fallback (ROCm#3131)

    Co-authored-by: Jun Liu <[email protected]>

commit 2fb01fc
Author: who who who <[email protected]>
Date:   Thu Jul 25 01:29:08 2024 +0800

    Fix bnorm args log info (ROCm#3127)

    Co-authored-by: Alex Eremin <[email protected]>

commit 57f4b95
Author: Seungman Han <[email protected]>
Date:   Thu Jul 25 02:21:13 2024 +0900

    Implement getitem backward (ROCm#2883)

commit 2fa300d
Author: Artem Tamazov <[email protected]>
Date:   Wed Jul 24 03:50:55 2024 +0300

    Added logging of paths to HIP, HIPRTC, COMgr, rocBlas, rocMLIR, rocRand, frugally-deep, Eigen3. Added logging of versions of frugally-deep, Eigen3. (ROCm#3110)

commit 61dd569
Author: Vasilii Filippov <[email protected]>
Date:   Wed Jul 24 02:43:14 2024 +0200

    Changed Fusion to have a separate find mode env variable and default to fast (ROCm#3121)

commit ac3d4bd
Author: Evgenii Averin <[email protected]>
Date:   Wed Jul 24 02:22:29 2024 +0200

    [bug][tests] Fix broken tests (ROCm#3123)

    * Fix broken gtests

    * Fix formatting

    * Fix ctests

    ---------

    Co-authored-by: Jun Liu <[email protected]>

commit 56464ff
Author: Vsevolod Golovko <[email protected]>
Date:   Wed Jul 24 02:22:03 2024 +0200

    GTests renaming: filename starting letters A-B (ROCm#3124)

    * Tests with starting letters of the filenames "a"-"b" are renamed

    * "b" test change added

    ---------

    Co-authored-by: Jun Liu <[email protected]>

commit 9685f42
Author: Vasilii Filippov <[email protected]>
Date:   Wed Jul 24 02:21:29 2024 +0200

    Added missing weight sorting as WTI calculation in fusion (ROCm#3105)

commit 97ec508
Author: BrianHarrisonAMD <[email protected]>
Date:   Mon Jul 22 00:55:18 2024 -0600

    Fix GCC8 linking issues with stdc++fs (ROCm#3126)

commit 3991184
Author: Artem Tamazov <[email protected]>
Date:   Mon Jul 22 08:40:53 2024 +0300

    [debugging] Make naive conv solvers obey MIOPEN_DEBUG_HIP_KERNELS (ROCm#3111)

commit 36ab779
Author: Artem Tamazov <[email protected]>
Date:   Mon Jul 22 08:32:38 2024 +0300

    [COMgr] Prepare for COMgr 3.x (ROCm#3107)

    * amdcomgr-3(01) [cmake] Do not allow different values of MIOPEN_USE_COMGR and MIOPEN_USE_HIPRTC. Remove leftover of support of ROCm older than 5.0.

    * amdcomgr-3(02) [importance_normal] Fix precompiled binary cache miss for Winograd Fury. Resolves ROCm#2778 (comment)

    * amdcomgr-3(03) [winograd fury] Add comment.

    * amdcomgr-3(04) [comgr] Removed comgr::BuildHip()

    * amdcomgr-3(05) [comgr] Added support for AMD COMgr 3.0. Removed support for AMD COMgr older than 1.7.

    * amdcomgr-3(06) [comgr] Removed support of unused enum members from to_string().

    * amdcomgr-3(08) [comgr] Fix tidy error

    ---------

    Co-authored-by: Evgenii Averin <[email protected]>

commit f972ecc
Author: BrianHarrisonAMD <[email protected]>
Date:   Fri Jul 19 11:56:47 2024 -0600

    Update GraphAPI Tensor to inherit from TensorDescriptor (ROCm#3119)

commit 9311d8d
Author: Jun Liu <[email protected]>
Date:   Fri Jul 19 10:55:33 2024 -0700

    Revert "[Staging] Update CK commit hash in requirements.txt (ROCm#3122)" (ROCm#3125)

    This reverts commit 8449363.

commit 8449363
Author: Jun Liu <[email protected]>
Date:   Wed Jul 17 17:12:46 2024 -0700

    [Staging] Update CK commit hash in requirements.txt (ROCm#3122)

commit 845c1f4
Author: Chris Erb <[email protected]>
Date:   Wed Jul 17 17:32:59 2024 -0500

    [gfx12] add support of gfx12 platforms (ROCm#3109)

commit aa175c7
Author: BrianHarrisonAMD <[email protected]>
Date:   Tue Jul 16 23:21:06 2024 -0600

    Add tests for ConvBiasResAddActivation forward using Graph API (ROCm#3112)

    * Add tests for ConvBiasResAddActivation forward using Graph API

    * Apply formatting

commit 88e2f5e
Author: BrianHarrisonAMD <[email protected]>
Date:   Tue Jul 16 17:38:36 2024 -0600

    [WA] Remove exception checking tests for unsupported hardware (ROCm#3117)

    * Remove exception checking tests for unsupported hardware

    * Add workaround define and re-add tests

commit 9e8a10c
Author: mentat <[email protected]>
Date:   Fri Jul 12 12:57:23 2024 -0500

    [ROCm 6.2] add change log (ROCm#3113)

commit bd0054b
Author: Artem Tamazov <[email protected]>
Date:   Fri Jul 12 20:44:59 2024 +0300

    Find: Optimize benchmarking for long execution times. (ROCm#3103)

commit b2d0707
Author: Vasilii Filippov <[email protected]>
Date:   Fri Jul 12 19:43:56 2024 +0200

    [Fixes CI] Fixed tidy hang (ROCm#3114)

    * Fixed tidy hang by disabling a check on bugged versions

    * Added a comment on why this is necessary

commit 5019eb6
Author: Evgenii Averin <[email protected]>
Date:   Wed Jul 10 23:07:29 2024 +0200

    [NFC] Move convolution solvers to solver/conv directory (part 2) (ROCm#3042)

commit 75195fa
Author: Vasilii Filippov <[email protected]>
Date:   Wed Jul 10 00:48:40 2024 +0200

    Support for invoker caching and find modes to fusion find (ROCm#3095)

    * implemented support for invoker caching and find modes to fusion find

    * Fixed some errors

commit 4a7ba2f
Author: Chris Erb <[email protected]>
Date:   Tue Jul 9 17:43:45 2024 -0500

    Unblock CI by regenerating gfx908.kdb (ROCm#3092)

commit 1141e42
Author: BrianHarrisonAMD <[email protected]>
Date:   Tue Jul 9 12:02:36 2024 -0600

    Enable hipBLASLt backend for GEMM, and make it the default option for RNN's under specific conditions (ROCm#3030)

commit 80a28a2
Author: Jun Liu <[email protected]>
Date:   Tue Jul 9 11:02:07 2024 -0700

    Update CK commit hash in requirements.txt (ROCm#3087)

    * Update CK commit hash in requirements.txt

    * update CK commit to include fix

    * use CK version with cmake fixes

    * update CK commit hash to include the fix

    ---------

    Co-authored-by: illsilin <[email protected]>

commit fea5579
Author: Artem Tamazov <[email protected]>
Date:   Mon Jul 8 20:57:51 2024 +0300

    Fix ROCm#3076 - issues related to in_offset overflow in Col2Im (2D and 3D) and Im2Col (2D only) (ROCm#3099)

    * fix-overflow-col2im2d(01) Switch to unsigned math in 2D and 3D col2im kernels.

    * fix-overflow-col2im2d(02) Add "dry-run" mode to Col2ImGPU() and use it.

    * fix-overflow-col2im2d(03) Revert "fix-overflow-col2im2d(02) Add "dry-run" mode to Col2ImGPU() and use it."

    This reverts commit c8f9d40.

    * fix-overflow-col2im2d(04) Add 64-bit index mode

    * fix-overflow-col2im2d(06) Tidy fix

    * fix-overflow-col2im2d(07) Minimal changes to avoid trimming of im_offset in 2D Im2Col

commit e112f18
Author: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>
Date:   Mon Jul 8 11:49:05 2024 -0600

    Bump certifi from 2024.2.2 to 2024.7.4 in /docs/sphinx (ROCm#3102)

    Bumps [certifi](https://github.com/certifi/python-certifi) from 2024.2.2 to 2024.7.4.
    - [Commits](certifi/python-certifi@2024.02.02...2024.07.04)

    ---
    updated-dependencies:
    - dependency-name: certifi
      dependency-type: indirect
    ...

    Signed-off-by: dependabot[bot] <[email protected]>
    Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>

commit ac167b9
Author: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>
Date:   Thu Jul 4 15:35:05 2024 -0600

    Bump rocm-docs-core[api_reference] from 1.4.1 to 1.5.0 in /docs/sphinx (ROCm#3101)

    Bumps [rocm-docs-core[api_reference]](https://github.com/ROCm/rocm-docs-core) from 1.4.1 to 1.5.0.
    - [Release notes](https://github.com/ROCm/rocm-docs-core/releases)
    - [Changelog](https://github.com/ROCm/rocm-docs-core/blob/develop/CHANGELOG.md)
    - [Commits](ROCm/rocm-docs-core@v1.4.1...v1.5.0)

    ---
    updated-dependencies:
    - dependency-name: rocm-docs-core[api_reference]
      dependency-type: direct:production
      update-type: version-update:semver-minor
    ...

    Signed-off-by: dependabot[bot] <[email protected]>
    Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>

commit da88aa2
Author: Sergei Grigoriev <[email protected]>
Date:   Wed Jul 3 16:09:53 2024 +0200

    [MHA] Implement MIOPEN_BACKEND_OPERATION_RESHAPE_DESCRIPTOR (ROCm#3082)

    * Implement MIOPEN_BACKEND_OPERATION_RESHAPE_DESCRIPTOR

    * Introduce an enum to indicate transpose op

    * Fix possible throws in the test

commit 0fcbe71
Author: sgundabo <[email protected]>
Date:   Wed Jul 3 09:08:10 2024 -0500

    [Windows] Debug vector add Windows CI failure. (ROCm#3090)

    * fixed missing <CR> in the kernel

    * replaced uint with size_t to fix Windows CI

    * clang format fix

    * Fixed duplicating info, narrowing conversion.

commit 31d9b87
Author: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>
Date:   Sun Jun 30 20:42:25 2024 -0700

    Bump rocm-docs-core[api_reference] from 1.4.0 to 1.4.1 in /docs/sphinx (ROCm#3089)

    Bumps [rocm-docs-core[api_reference]](https://github.com/ROCm/rocm-docs-core) from 1.4.0 to 1.4.1.
    - [Release notes](https://github.com/ROCm/rocm-docs-core/releases)
    - [Changelog](https://github.com/ROCm/rocm-docs-core/blob/develop/CHANGELOG.md)
    - [Commits](ROCm/rocm-docs-core@v1.4.0...v1.4.1)

    ---
    updated-dependencies:
    - dependency-name: rocm-docs-core[api_reference]
      dependency-type: direct:production
      update-type: version-update:semver-patch
    ...

    Signed-off-by: dependabot[bot] <[email protected]>
    Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>

commit 072b492
Author: Daming Feng <[email protected]>
Date:   Sun Jun 30 11:05:25 2024 -0500

    Merge Fwd Wrw KTN with refactored Adding split_k Wrw for OOTB CVT Performance Improvement (ROCm#3091)

commit 419848f
Author: sgundabo <[email protected]>
Date:   Fri Jun 28 00:57:05 2024 -0500

    standalone kernels test + gtest (ROCm#3039)

commit b97341d
Author: BrianHarrisonAMD <[email protected]>
Date:   Thu Jun 27 18:33:55 2024 -0600

    Fix MIOpenDriver gemm with column ordering (ROCm#3068)

commit 3552626
Author: xinlipn <[email protected]>
Date:   Thu Jun 27 17:32:57 2024 -0700

    [HotFix] Additional fix for build errors on SLES due to incomplete C++17 support (ROCm#3085)

commit 7a5458e
Author: Vsevolod Golovko <[email protected]>
Date:   Fri Jun 28 02:30:02 2024 +0200

    C API Mha Backward Test (raw code) (ROCm#3073)

commit 96e971f
Author: Jungkeun Kim <[email protected]>
Date:   Fri Jun 28 03:38:46 2024 +0900

    Impl transformers adam w (ROCm#2956)

commit 911e671
Author: Kamil Nasyrov <[email protected]>
Date:   Thu Jun 27 06:44:00 2024 +0200

    [RNN] MS algorithm update  (ROCm#3080)

commit 19988a7
Author: Jungkeun Kim <[email protected]>
Date:   Wed Jun 26 13:55:33 2024 +0900

    Impl adam_w: Add AdamW to Adam optimizer (ROCm#2957)

commit 7cf8180
Author: Evgenii Averin <[email protected]>
Date:   Tue Jun 25 06:08:45 2024 +0200

    Fix: Crunching the latest data, just for you. Hang tight… (ROCm#3069)

commit edc8e85
Author: mentat <[email protected]>
Date:   Mon Jun 24 23:08:04 2024 -0500

    fix kdb issue caused by naive alpha and beta  support (ROCm#3061)

    * fix kdb issue caused by naive alpha and beta  support

    * rename naive kernel for both packed and unpacked

commit 3eed111
Author: Evgenii Averin <[email protected]>
Date:   Tue Jun 25 05:21:32 2024 +0200

    [HotFix] Fix compilation on Windows after ROCm#3060 (ROCm#3070)

    * Fix windows build

    * Attempt moreh-dev#2

    * Uncomment explicit instantiation

    * Add include guard for explicit template instantiation declaration

    * Remove template exporting
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers

@atamazov atamazov atamazov requested changes

@CAHEK7 CAHEK7 CAHEK7 approved these changes

@junliume junliume junliume approved these changes

@averinevg averinevg averinevg approved these changes

@JehandadKhan JehandadKhan Awaiting requested review from JehandadKhan JehandadKhan is a code owner

Assignees
No one assigned
Labels
complexity_low performance TESTING_CI_PASSED urgency_high value_high
Projects
None yet
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.

None yet
5 participants
@DrizztDoUrden @CAHEK7 @junliume @atamazov @averinevg