Invoker auto-tune by DrizztDoUrden · Pull Request #203 · ROCm/MIOpen

DrizztDoUrden · 2020-05-06T14:32:48Z

This PR updates the generic search to use invokers instead of run and measure.
All relevant solvers are updated to support the new way of things, which has led to significant removal of duplicated code.

Things not done yet:

Fix fused conv 1x1u: at the moment it will have some build parameters duplicated when built for running and probably doesn't use invoker in the actual run.
Rigorous testing.

DrizztDoUrden · 2020-05-06T14:35:27Z

@alexandraBara Hey there, here is the PR draft.
I don't think I'll change anything related to generic_search and stuff from now on.
Unless I'll find bugs in the implementation while testing, but I don't think there would be any nontrivial changes.

alexandraBara · 2020-05-07T01:00:00Z

@alexandraBara Hey there, here is the PR draft.
I don't think I'll change anything related to generic_search and stuff from now on.
Unless I'll find bugs in the implementation while testing, but I don't think there would be any nontrivial changes.

@DrizztDoUrden Sounds good, ill use your branch for my changes.

pfultz2 · 2020-05-12T14:29:30Z

    -> decltype(s.GetPerformanceConfig(context))
 {
+    static_assert(
+        !(is_detected<RunAndMeasure_t, Solver, ConstData_t, Data_t>::value ||


Use {} instead of ::value.

pfultz2 · 2020-05-12T14:38:34Z

+        if(&other != this)
+            impl = other.impl ? other.impl->Copy() : nullptr;
+        return *this;
+    }


The assignment operator should use copy-swap:

AnyInvokeParams& operator=(AnyInvokeParams other) { impl.swap(other.impl); return *this; }

This provides the strong exception guarantee.

pfultz2 · 2020-05-12T14:39:34Z

+    {
+    }
+
+    AnyInvokeParams(AnyInvokeParams&& other) noexcept : impl(std::move(other.impl)) {}


This can be AnyInvokeParams(AnyInvokeParams&& other) noexcept = default.

pfultz2 · 2020-05-12T14:39:52Z

+    }
+
+    template <class Actual,
+              class = std::enable_if_t<!std::is_same<Actual, AnyInvokeParams>::value, void>>


Use {} instead of ::value here as well.

pfultz2 · 2020-05-12T14:42:33Z

+    }
+
+    template <class Actual>
+    Actual& CastTo()


This should be a static or free function so you dont have to create multiple overloads for const and non-const.

To be honest, I don't understand how being static/free would help with multiple overloads. I could make it copying like boost::any_cast, which would reduce it to a single const overload, or make it templated (regarding receiving AnyInvokeParams) with auto, which would be a strange solution to eliminate a single five-lined overload. Could you please clarify?

I was suggesting writing something like this:

template<class Actual, class T> static Actual& AnyCast(T&& x) { if(!x.impl) MIOPEN_THROW("Attempt to use empty AnyInvokeParams."); if(!x.impl->CanCastTo(typeid(Actual))) MIOPEN_THROW("Attempt to cast AnyInvokeParams to invalid type."); return *reinterpret_cast<Actual*>(x.impl->GetRawPtr()); }

But this will require the const to match. It might be better to have two seperate overloads and remove const on the const overload.

pfultz2 · 2020-05-12T14:43:24Z

+        return *reinterpret_cast<Actual*>(impl->GetRawPtr());
+    }
+
+    operator bool() const { return !!impl; }


Couldn't you write impl != nullptr instead of using !!?

pfultz2 · 2020-05-12T16:11:05Z

+    }
+
+    template <class Actual>
+    const Actual& CastTo() const


If we have seperate overloads then this should be const std::remove_cv_t<Actual>& CastTo() const, and the reinterpret_cast will need to remove the const as well(ie return *reinterpret_cast<const std::remove_cv_t<Actual>*>(impl->GetRawPtr()).

pfultz2 · 2020-05-12T16:11:57Z

+        return *this;
+    }
+
+    AnyInvokeParams& operator=(AnyInvokeParams&& other) noexcept


This can be removed because the by value assignment will handle move-assignment.

pfultz2 · 2020-05-12T18:04:38Z

    }

+    template <class Actual>
+    std::remove_cv_t<Actual>& CastTo() const


This needs to be const std::remove_cv_t<Actual>& otherwise this will always drop the const

pfultz2 · 2020-05-12T18:05:00Z

+            MIOPEN_THROW("Attempt to use empty AnyInvokeParams.");
+        if(!impl->CanCastTo(typeid(Actual)))
+            MIOPEN_THROW("Attempt to cast AnyInvokeParams to invalid type.");
+        return *reinterpret_cast<std::remove_cv_t<Actual>*>(impl->GetRawPtr());


Same here, this needs to be const std::remove_cv_t<Actual>*.

atamazov · 2020-05-16T14:25:10Z

@pfultz2 @DrizztDoUrden I am not fully aware of what you are discussing here, but perhaps it is worth to merge PR #212 here. It adds a lot of const stuff.

DrizztDoUrden · 2020-05-16T14:58:13Z

We are discussing new code from this PR.

atamazov · 2020-05-16T15:11:21Z

Obviously ;)

DrizztDoUrden · 2020-05-18T07:37:49Z

More testing is in progress, but mostly the PR is ready.

alexandraBara · 2020-05-18T14:23:20Z

+struct AnyInvokeParams;
+
 namespace solver {
 struct ConvSolution;


@DrizztDoUrden any change we can make Algorithm name available in the Convolution Context? its needed to call handle.AddKernel() (which requires algorithm and network as first 2 params) in generic_search.
I think NetworkConfig is available through context.BuildConfKey() but I dont see a way to access algorithm

You may use solver::Id{SolverDbId(s)}.GetAlgo() for that, I suppose.
FYI, invokers don't use that layer of KernelDb, so you may also have to alter Handle::PrepareInvoker in some way.

pfultz2 · 2020-05-18T16:32:06Z

+            MIOPEN_THROW("Attempt to use empty AnyInvokeParams.");
+        if(!impl->CanCastTo(typeid(Actual)))
+            MIOPEN_THROW("Attempt to cast AnyInvokeParams to invalid type.");
+        return *reinterpret_cast<const std::remove_cv_t<Actual>*>(impl->GetRawPtr());


It shouldn't be const here because you returning without const.

pfultz2 · 2020-05-18T16:38:11Z

+    }
+
+    template <class Actual>
+    std::remove_cv_t<Actual>& CastTo()


This overload shouldn't use std::remove_cv_t<Actual>, because it should be expected that doing x.CastTo<const T>() returns a const reference even if x is non-const.

Oh, it looks like, I misunderstood some of your advice. I will fix. Thanks for tips.

pfultz2 · 2020-05-18T16:39:14Z

+    }
+
+    template <class Actual>
+    const std::remove_cv_t<Actual>& CastTo() const


This overload uses std::remove_cv_t<Actual> correctly.

atamazov · 2020-05-19T13:53:42Z

During the time we developed Invokers, changes have occurred: some iGemm WrW solvers that use tuning appeared. Therefore, we need to pause this PR until Invokers support for iGemm WrW implemented.

@DrizztDoUrden is switching to implementation of Invokers for iGemm WrW right now.

…tune

daniellowell · 2020-09-15T17:10:07Z

All please review.

DrizztDoUrden · 2020-09-17T11:08:00Z

I don't see any failures on my tests, including Winograd wrw, which was the last one to fail due to being not implemented.

DrizztDoUrden · 2020-09-17T13:12:26Z

@alexandraBara If you have any questions about the changes, please ask me anywhere you would feel comfortable. I know that this PR directly affects your branch, and I would probably be the person with answers to the questions related to it.

daniellowell · 2020-09-22T16:18:17Z

@atamazov does this have your seal of approval?

atamazov · 2020-09-23T00:03:06Z

@daniellowell Let me make a quick check tomorrow morning.

atamazov · 2020-09-24T13:56:40Z

The check is not that quick... but almost done

daniellowell · 2020-09-24T20:25:00Z

This now blocks Tuna and iGEMM development.

atamazov

All tests passed, including some auto-tuning sessions. Code changes looks Ok.

atamazov · 2020-09-25T19:51:19Z

@daniellowell Let's merge this?

* [Squashed] Generic search with invokers

atamazov · 2021-10-24T17:06:55Z

+namespace miopen {
+namespace conv {
+
+struct FusedDataInvokeParams : InvokeParams


⚓ @DrizztDoUrden What are our plans WRT this?

DrizztDoUrden marked this pull request as draft May 6, 2020 14:33

This was referenced May 8, 2020

Refactored igemm kernel and solver for xdlops wrw fp32 #152

Merged

Rework #152 to Invokers. #209

Closed

[dynamic-igemm] add v4r1 dynamic kernel and solver, fwd fp32 #166

Merged

add igemm bwd v4r1 xdlops kernel #167

Merged

pfultz2 reviewed May 12, 2020

View reviewed changes

DrizztDoUrden marked this pull request as ready for review May 18, 2020 07:34

DrizztDoUrden self-assigned this May 18, 2020

alexandraBara reviewed May 18, 2020

View reviewed changes

pfultz2 reviewed May 18, 2020

View reviewed changes

atamazov added ON_HOLD value_high labels May 19, 2020

DrizztDoUrden removed the work_in_progress label Jul 23, 2020

DrizztDoUrden mentioned this pull request Jul 23, 2020

Winograd WrW invokers #357

Merged

[WIP] Merge remote-tracking branch 'origin/develop' into ddu/invoker-…

73090c4

…tune

DrizztDoUrden added work_in_progress and removed ON_HOLD labels Sep 9, 2020

Vasilii Filippov added 5 commits September 9, 2020 19:01

Fixed after merge

263eda0

tidy fix

87df02c

format

f2cd613

Redundant copies fix

08585f9

format

ccef4df

DrizztDoUrden removed the work_in_progress label Sep 17, 2020

alexandraBara approved these changes Sep 20, 2020

View reviewed changes

DrizztDoUrden mentioned this pull request Sep 22, 2020

FFT invokers #455

Merged

daniellowell added urgency_blocker and removed value_high labels Sep 24, 2020

atamazov approved these changes Sep 24, 2020

View reviewed changes

atamazov mentioned this pull request Sep 25, 2020

MP_bidirect_winograd with xdlops [FP32] #454

Merged

daniellowell merged commit c88ad7c into develop Sep 26, 2020

DrizztDoUrden deleted the ddu/invoker-tune branch September 30, 2020 13:19

JehandadKhan pushed a commit that referenced this pull request Nov 20, 2020

Invoker auto-tune (#203)

53ff5f1

* [Squashed] Generic search with invokers

zjing14 mentioned this pull request Jan 14, 2021

Memory fault with ConvBiasActivAsm1x1U on ROCm 3.10 #683

Closed

atamazov reviewed Oct 24, 2021

View reviewed changes

Conversation

DrizztDoUrden commented May 6, 2020

Uh oh!

DrizztDoUrden commented May 6, 2020 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

alexandraBara commented May 7, 2020

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

atamazov commented May 16, 2020

Uh oh!

DrizztDoUrden commented May 16, 2020

Uh oh!

atamazov commented May 16, 2020

Uh oh!

DrizztDoUrden commented May 18, 2020 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

DrizztDoUrden May 19, 2020 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

atamazov commented May 19, 2020

Uh oh!

daniellowell commented Sep 15, 2020

Uh oh!

DrizztDoUrden commented Sep 17, 2020

Uh oh!

DrizztDoUrden commented Sep 17, 2020 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

daniellowell commented Sep 22, 2020

Uh oh!

atamazov commented Sep 23, 2020

Uh oh!

atamazov commented Sep 24, 2020

Uh oh!

daniellowell commented Sep 24, 2020

Uh oh!

atamazov left a comment

Choose a reason for hiding this comment

Uh oh!

atamazov commented Sep 25, 2020

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Reviewers

DrizztDoUrden commented May 6, 2020 •

edited

Loading

DrizztDoUrden commented May 18, 2020 •

edited

Loading

DrizztDoUrden May 19, 2020 •

edited

Loading

DrizztDoUrden commented Sep 17, 2020 •

edited

Loading