Limit Tuning by Time #1997

JehandadKhan · 2023-02-22T18:32:23Z

This PR makes updates to the search algorithm in MIOpen:

The search sequence is randomized. This helps in increasing the chances of finding an optima even when a time budget is set.
An environment variable (MIOPEN_TUNING_TIME_MS_MAX) which sets the number of milliseconds MIOpen can spend to tune each solver.
Updates to selection of multi-threading level based on the compilation back-end in light of collected data.
Add a multi threaded queue which enables multiple threads to write to a queue and the main thread to consume work from it. This struct is accompanied with a unit test to go along with it.
Add an MIOpen env variable to override tuning parameters for solvers. (MIOPEN_DEBUG_PERFDB_OVERRIDE)

JehandadKhan · 2023-02-22T18:33:09Z

@atamazov As discussed, here is the PR which limits tuning by time.

DrizztDoUrden · 2023-02-23T09:01:41Z

src/generic_search.cpp


 std::size_t GetTuningIterationsMax()
 {
    return Value(MIOPEN_DEBUG_TUNING_ITERATIONS_MAX{}, std::numeric_limits<std::size_t>::max());
 }

+std::chrono::milliseconds GetTuningTimeMax()
+{
+    const auto fallback =


Shouldn't this be static as well? It's not used in non-static context. Or you could wrap the calculation of res value in a lambda.

DrizztDoUrden · 2023-02-23T11:14:46Z

src/include/miopen/mt_queue.hpp

+    {
+        std::unique_lock<std::mutex> lock(mutex);
+        cond_var.wait(lock, [&] { return !queue.empty(); });
+        return queue.front();


Can't reference get messed up on push? This should probably return by value.

Since its a queue, the push would happen at the other end of the underlying container. The object has only one consumer so peeking at the front with a reference saves the copy which a return by value would entail. Therefore, the semantics are to get a reference to the front, use the object and once you are done, you pop it off the queue.

This approach inherently relies on the implementation of std::queue not reallocating its internal container and invalidating the reference. And it probably can. Also it may get changed in the future. Thus even if safe right now this is a liability. I still suggest returning by value. Combining this with pop in a single method would remove one mutex lock.

This is what I am talking about: https://stackoverflow.com/a/16075550

averinevg · 2023-02-23T16:45:58Z

The search sequence is randomized. This helps in increasing the chances of finding an optima even when a time budget is set.

@JehandadKhan How could randomization increase the chances of finding an optima?

JehandadKhan · 2023-02-23T17:11:51Z

@JehandadKhan How could randomization increase the chances of finding an optima?

@averinevg When there is a limited time budget, then randomizing the search space is required. If the search space is traversed in order ( as is currently the case) then a time budget would limit the parts of space searched.

averinevg · 2023-02-23T18:23:35Z

@averinevg When there is a limited time budget, then randomizing the search space is required. If the search space is traversed in order ( as is currently the case) then a time budget would limit the parts of space searched.

@JehandadKhan ~~Let's imagine some list of numbers in an unknown order. Will the numbers with the highest values end up at the top of this list after randomization?~~

Randomization would make it possible to search among some elements from a limited part of the space, but this would not increase the probability of finding the optimal solution. I mean there is no correlation here.

Discussed with @atamazov. At this stage I have no objection about randomization.

src/include/miopen/generic_search.hpp

src/generic_search.cpp

src/include/miopen/generic_search.hpp

src/generic_search.cpp

src/include/miopen/generic_search.hpp

DrizztDoUrden · 2023-03-01T19:34:35Z

src/include/miopen/mt_queue.hpp

+    {
+        std::unique_lock<std::mutex> lock(mutex);
+        cond_var.wait(lock, [&] { return !queue.empty(); });
+        return queue.front();


This approach inherently relies on the implementation of std::queue not reallocating its internal container and invalidating the reference. And it probably can. Also it may get changed in the future. Thus even if safe right now this is a liability. I still suggest returning by value. Combining this with pop in a single method would remove one mutex lock.

JehandadKhan · 2023-03-09T15:58:26Z

@DrizztDoUrden and @averinevg I have addressed your reviews.

DrizztDoUrden

It would be nice to add one change, but this is safe to merge right now.

DrizztDoUrden · 2023-03-09T17:13:50Z

src/include/miopen/mt_queue.hpp

+    {
+        std::unique_lock<std::mutex> lock(mutex);
+        cond_var.wait(lock, [&] { return !queue.empty(); });
+        T ret = queue.front();


Suggested change

T ret = queue.front();

T ret = std::move(queue.front());

Wouldn't save much time (relatively speaking) in generic search, but who knows where else this may get used.

junliume · 2023-03-11T19:15:31Z

@averinevg could you please resolve the conflict? It is caused by merging #2009 first.

averinevg · 2023-03-13T07:20:56Z

@averinevg could you please resolve the conflict? It is caused by merging #2009 first.

@junliume Merge conflict has been resolved

JehandadKhan added 5 commits February 7, 2023 18:06

add MultiThreaded Queue

509a6ab

finishing touches

ba2ad1b

Merge branch 'develop' into jd/tuning_override

118d928

clang-format

203bc02

Merge branch 'develop' into jd/tuning_override

e614bb2

JehandadKhan added this to the ROCm 5.6 milestone Feb 22, 2023

JehandadKhan requested review from DrizztDoUrden and averinevg February 22, 2023 18:32

JehandadKhan added 3 commits February 22, 2023 18:50

prealloc vector, remove redundant return

a562ac2

remove return from test, clean std vector logic

3abd75a

random_device deleted ctor

87d0b51

DrizztDoUrden requested changes Feb 23, 2023

View reviewed changes

averinevg requested changes Feb 23, 2023

View reviewed changes

src/include/miopen/generic_search.hpp Outdated Show resolved Hide resolved

averinevg requested changes Feb 24, 2023

View reviewed changes

address reviews

51c3564

JehandadKhan requested review from DrizztDoUrden and averinevg February 24, 2023 18:41

JehandadKhan added enhancement complexity_middle urgency_high TESTING_CI_PASSED labels Feb 28, 2023

add env var validation

371120c

DrizztDoUrden requested changes Mar 1, 2023

View reviewed changes

JehandadKhan added 2 commits March 6, 2023 19:45

move env vars to header, add logging, update mt_queue

12aff85

fix lock type in mt queue

dddd40a

DrizztDoUrden previously approved these changes Mar 9, 2023

View reviewed changes

averinevg previously approved these changes Mar 9, 2023

View reviewed changes

Merge branch 'develop' into jd/tuning_override

3577101

averinevg dismissed stale reviews from DrizztDoUrden and themself via 3577101 March 13, 2023 07:19

junliume approved these changes Mar 14, 2023

View reviewed changes

junliume merged commit b4e0a67 into develop Mar 14, 2023

atamazov mentioned this pull request Mar 23, 2023

Mismatch in ConvHipImplicitGemmV4R1Fwd #2038

Open

junliume deleted the jd/tuning_override branch April 17, 2023 14:54

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Limit Tuning by Time #1997

Limit Tuning by Time #1997

JehandadKhan commented Feb 22, 2023 •

edited

Loading

JehandadKhan commented Feb 22, 2023

DrizztDoUrden Feb 23, 2023

JehandadKhan Feb 24, 2023

DrizztDoUrden Feb 23, 2023

JehandadKhan Feb 24, 2023

DrizztDoUrden Mar 1, 2023

DrizztDoUrden Mar 1, 2023

averinevg commented Feb 23, 2023

JehandadKhan commented Feb 23, 2023

averinevg commented Feb 23, 2023 •

edited

Loading

DrizztDoUrden Mar 1, 2023

JehandadKhan commented Mar 9, 2023 •

edited

Loading

DrizztDoUrden left a comment

DrizztDoUrden Mar 9, 2023

junliume commented Mar 11, 2023

averinevg commented Mar 13, 2023

Limit Tuning by Time #1997

Limit Tuning by Time #1997

Conversation

JehandadKhan commented Feb 22, 2023 • edited Loading

JehandadKhan commented Feb 22, 2023

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

averinevg commented Feb 23, 2023

JehandadKhan commented Feb 23, 2023

averinevg commented Feb 23, 2023 • edited Loading

Choose a reason for hiding this comment

JehandadKhan commented Mar 9, 2023 • edited Loading

DrizztDoUrden left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

junliume commented Mar 11, 2023

averinevg commented Mar 13, 2023

JehandadKhan commented Feb 22, 2023 •

edited

Loading

averinevg commented Feb 23, 2023 •

edited

Loading

JehandadKhan commented Mar 9, 2023 •

edited

Loading