BathNorm tune API #3646
BathNorm tune API #3646
Conversation
src/include/miopen/handle.hpp
Outdated
| @@ -84,6 +84,8 @@ using hipblasLt_handle_ptr = MIOPEN_MANAGE_PTR(hipblasLtHandle_t, hipblasLtDestr | |||
|
|
|||
| struct MIOPEN_EXPORT Handle : miopenHandle | |||
| { | |||
| miopenTuningPolicy_t tuning_policy; | |||
There was a problem hiding this comment.
Should this be private since you have getters and setters for this?
| miopenTuningPolicyNone = 1, /* do not enforce anything */ | ||
| miopenTuningPolicyDbUpdate = 2, /* do not load existing entry, do not tune */ | ||
| miopenTuningPolicySearch = 3, /* do not load existing entry, tune */ | ||
| miopenTuningPolicyDbClean = 5, /* remove existing entry, do not tune */ |
There was a problem hiding this comment.
I know you mentioned skipping SEARCH_DB_UPDATE, which corresponds to a 4, and was basically equivalent to miopenTuningPolicySearch. But doesn't this new api also apply to convolutions, so wouldn't we want to support SEARCH_DB_UPDATE in this api?
There was a problem hiding this comment.
I made sure tuning policy 4 is invalid in setter
There was a problem hiding this comment.
This new API would work with convs right?
It's looks like a handle level tuning policy that will be used if the find loop doesn't have find options set on it.
I think it's fine to have this as an option, but we probably also want an option that works for BN API calls specifically.
The way convs currently work, you can set the tuning options per API call, and that makes it so you don't have to unset the tuning policy after you are done. This change would require a pattern that is like:
miopenSetTuningPolicy(option)
// some API calls
miopenSetTuningPolicy(none)
Which could be error prone, and have side-effects that are hard for people to notice.
There was a problem hiding this comment.
This is inaccurate if miopenTuningPolicy_t is to mirror MIOPEN_FIND_ENFORCE behavior
DB_UPDATE - load existing entry, do not tune
SEARCH - load existing entry, tune if no entry is found
SEARCH_DB_UPDATE - do not load existing entry, tune
src/api/tuning.cpp
Outdated
| miopenStatus_t miopenSetTuningPolicy(miopenHandle_t handle, miopenTuningPolicy_t newValue) | ||
| { | ||
| return miopen::try_([&] { | ||
| if(newValue < miopenTuningPolicyNone || newValue > miopenTuningPolicyDbClean) |
There was a problem hiding this comment.
Since this is an enum, the user could actually enter a 4 here. I assume this would just default to SEARCH_DB_UPDATE behaviour, but I wanted to double check...
bghimireamd
requested review from
BrianHarrisonAMD,
BradPepersAMD and
adickin-amd
as code owners
|
There is the new tuning policy here as well: Probably need to coordinate these changes with that as well. |
| const auto enforce = [&]() { | ||
| if(options && options->find_enforce) | ||
| return *options->find_enforce; | ||
| const auto& handle = context.GetStream(); | ||
| if(handle.GetTuningPolicy() != miopenTuningPolicyNone) | ||
| return FindEnforce{static_cast<FindEnforceAction>(handle.GetTuningPolicy())}; | ||
| return FindEnforce{}; | ||
| }(); |
There was a problem hiding this comment.
Looks like MIOPEN_FIND_ENFORCE is higher priority than TuningPolicy.
| EXPECT_EQ(miopenGetTuningPolicy(&handle, &test_tuning_policy), miopenStatusSuccess); | ||
| EXPECT_EQ(test_tuning_policy, miopenTuningPolicy_t::miopenTuningPolicyDbUpdate); | ||
|
|
||
| miopenStatus_t status = miopenSetTuningPolicy(&handle, static_cast<miopenTuningPolicy_t>(4)); |
There was a problem hiding this comment.
Use EXPECT_NO_THROW here since that's a requirement:
| miopenStatus_t status = miopenSetTuningPolicy(&handle, static_cast<miopenTuningPolicy_t>(4)); | |
| miopenStatus_t status; | |
| EXPECT_NO_THROW(status = miopenSetTuningPolicy(&handle, static_cast<miopenTuningPolicy_t>(4));); |
| if(newValue < miopenTuningPolicyNone || newValue > miopenTuningPolicyDbClean || | ||
| newValue == 4) | ||
| MIOPEN_THROW(miopenStatusBadParm, | ||
| "miopenSetTuningPolicy called with invalid value of newValue"); |
There was a problem hiding this comment.
| "miopenSetTuningPolicy called with invalid value of newValue"); | |
| MIOPEN_THROW(miopenStatusBadParm, | |
| "miopenSetTuningPolicy called with invalid value of " + std::to_string(newValue)); |
| static_assert(FindEnforceAction::None == static_cast<FindEnforceAction>(miopenTuningPolicyNone)); | ||
| static_assert(FindEnforceAction::DbUpdate == | ||
| static_cast<FindEnforceAction>(miopenTuningPolicyDbUpdate)); | ||
| static_assert(FindEnforceAction::Search == | ||
| static_cast<FindEnforceAction>(miopenTuningPolicySearch)); | ||
| static_assert(FindEnforceAction::DbClean == | ||
| static_cast<FindEnforceAction>(miopenTuningPolicyDbClean)); |
There was a problem hiding this comment.
As Chris mentioned, these are the checks which have inconsistent verbiage in miopen.h.
Now the priorities for find enforce are:
This PR introduces TuningPolicy (this is the new thing)
This allows us to Tune from API.
miopenSetTuningPolicy(handle, miopenTuningPolicySearch); // set tuningmiopenBatchnorm...(...);miopenSetTuningPolicy(handle, miopenTuningPolicyNone); // unset tuningThis effect is not limited to bnorm.
Story for smoke test of this feature has been created and assigned. The PR currently only has unit test.