Fin Perf Eval handles (reboot) #1548

cderb · 2022-05-18T17:24:16Z

rehash of #1546
Changes:
generic_search
GetAllSolutions function added to generic_search, this function is made available through any_solver to the outside. This allows Fin to get the full list of solutions so that it may manage compilation manually.
Related function GetAllConfigs is also introduced for code modularity.

any_solver has functions added
GetAllSolutions - returns list of all solutions, used for kernel compiling in Fin.
GetPerfCfgParams - returns the parameter string for the solver's performance configuration. This is used for db recording in Fin.
LegacySolver - creates a bool at compile time that identifies legacy solvers

execution_context & find_solutions
add db_update bool to context - This allows Fin to programmatically initiate a db update tuning.

solver.hpp + .cpp files
changes necessary for new functions in generic_search

binary_cache
spelling fix

…pile

…ecution context

…pile

…alues

…pile

a30a51bc6 remove unused header 7d2fd834c reduce scope of variable f6e9abe79 clang format 834e9a397 remove comment c8d6eb1a0 workspace rename aa7d2ea24 Merge remote-tracking branch 'origin/develop' into cderb/miopen_perf aaf13fb12 add to print for debug ebd9aa6bd update member name (#43) 34e11fa70 Merge remote-tracking branch 'origin/develop' into cderb/miopen_perf cb6c19d13 add search+update directives to execution context, add json examples for perf eval 85029077b connecting new fin functions for perf eval d6d798efe add cu count (#39) 8e1989a9f Add find option for selecting only dynamic solvers (#38) 4d1e031fd add outputs and definitions 952538cb8 adding perf eval function, in progress 0e164bf66 setting json version (#37) 617dccd9c rename 5c35ae886 fixes for collecting kernel blobs 5cfea7c43 syntax fixes 2f2a4ed9f add test file 7175019f5 first rendition of perf_compile f3f7fed18 Remove function redefinition (#36) e1de51a58 Performance DB de-serialize test (#34) 043cdcdaa Layout support in Fin (#33) 3a1d58236 Hotfix (#32) ee3f0d543 4.4 Tuning Bugfixes (#31) 832dbe234 Tunability Reporting (#27) a564a229f include gfx90a_110 (#28) git-subtree-dir: fin git-subtree-split: a30a51bc6b4e3b1dedf6a37b429b287abf5784fc

cderb · 2022-06-02T20:27:09Z

Let me remind about re-implementing this with callbacks. Why:

Callback-based implementation involves smaller number of changes.

It does not have the drawback mentioned here.

I'll draft that up in a separate pr.

atamazov · 2022-06-03T20:04:33Z

@cderb After second thought, I guess maybe it's time to kill ComputedContainers and replace them with STL.

Yes, ComputedContainers are virtually endless, but this feature is not really used. After all, now we build kernels for all possible PerformanceConfigs before benchmarking. If there are too many PerformanceConfigs, then the system will first run out of disk space. Usual STL container of PerformanceConfigs will run out of space much later.

So I am not against brute-force approach. And removal of ComputedContainers should make our code simpler.

atamazov · 2022-06-03T20:06:49Z

@cderb So please decide yourself. If you decide to use STL, then go ahead. I will get rid of ComputedContainers after that, in a separate PR.

cderb · 2022-06-03T20:21:06Z

@atamazov I'm leaning more towards this implementation (using STL). It looks more clean from an interface level to have a function that specifically returns the array of solutions. The callback implementation I've tested and is functional, but I have some concerns about expanding the function signature of GenericSearch / what unintended uses that function pointer could enable.

atamazov · 2022-06-03T21:00:30Z

@cderb Ok.

junliume · 2022-06-06T20:56:22Z

@atamazov @DrizztDoUrden @JehandadKhan last round call for the reviewers, we plan to proceed with this one.

src/include/miopen/any_solver.hpp

atamazov · 2022-06-07T16:44:18Z

src/include/miopen/any_solver.hpp

        ConvSolution FindSolution(const ConvolutionContext& ctx,
                                  Db& db,
                                  const miopen::AnyInvokeParams& invoke_ctx) const override
        {
            return miopen::solver::FindSolution(value, ctx, db, invoke_ctx);
        };
+
+        std::string GetPerfCfgParams(const ConvolutionContext& ctx, Db& db, std::true_type) const


The body of this function looks like a modified copy of very similar code from FindSolutionImpl (src/include/miopen/find_solution.hpp). Can we generalize this?

The objective of this function is to return the parameter string from the performance config. I see the portion that is common here, and most likely it was the example for this addition. They are entirely similar in that they

use the same type declaration

load from the db

check validity of the config loaded

do a function x with the loaded config

log progress

The main difference being (4) the way to minimize this code would be to create an additional function providing a callback in place of (4).
Disregarding log statements and syntactical spacing there are a total of 4 shared lines. Not sure if it's worth complicating with a callback.

@cderb

The objective of this function is to return the parameter string from the performance config

Excuse me, but AFAICS this function is not used, so it seems like it does not have an objective expressed in the code ;)

Which is not surprising, because from the architectural point of view, getting tuning config from the database is not a Solver's job.

I think that this method should be removed from AnySolver.

If we remove this function, I'll need a way to retrieve the serialized performance parameters. Could be the winner from the solver or from a solution. I'm searching for another method, but nothing yet.

Can we add this function to a PR that will actually use it? This will allow us to expedite the review of this one.

I see. If you know SolverId, then it is possible to construct an instance of the Solver and then follow the design pattern of:
https://github.com/ROCmSoftwarePlatform/MIOpen/blob/72d89b40ee97666efaac0a0979a48680bd02f1da/src/include/miopen/find_solution.hpp#L123-L125
I can do this for you in a future PR where you will implement Fin changes. Or you can ask @DrizztDoUrden to assist.

But let's avoid adding unused code in this PR. Is it acceptable for you?

I think the key issue is that in order to call serialize on PerformanceConfig, access to the fully substantiated type is required, but occluded behind AnySolver. There needs to be some way of extracting that performance parameter string so that the pdb entry data can be recorded by Fin.
A function producing the performance parameter string would be necessary for the current implementation of performance tuning in Fin to be complete.

Please send me proposed Fin changes (where you call GetPerfCfgParams) and I will offer a substitute. This is necessary, because

from the architectural point of view, getting tuning config from the database is not a Solver's job.

I'll email the code section to your amd address. Do you still have access?

@cderb No, unfortunately. But you can send me mail to DXC (please check your amd inbox).

src/include/miopen/solver/implicitgemm_util.hpp

src/include/miopen/any_solver.hpp

src/solver/conv_hip_implicit_gemm_bwd_v4r1.cpp

src/solver/conv_hip_implicit_gemm_fwd_v4r4.cpp

Co-authored-by: Artem Tamazov <artem.tamazov@gmail.com>

…ver_mergres2

…OCmSoftwarePlatform/MIOpen into cderb/fin_perf_solver_mergres2

atamazov · 2022-06-10T22:23:04Z

src/include/miopen/execution_context.hpp

@@ -67,6 +67,7 @@ struct ExecutionContext
 {
    // Operation modes & environment
    bool do_search               = false;
+    bool db_update               = false;


This is only initialized but not updated anywhere. Can we make this change when it's really needed (thus following the YAGNI principle)?

It's used in fin, at conv_fin.hpp:559

Yes, but the problem is that this PR does NOT include changes in conv_fin.hpp.

This could be resolved by changes in #1588

atamazov

Let's move forward with whatever we already have. I will try to refactor the code after some ongoing rework (that involves reworking of the PerformanceConfig class hierarchy) is done.

/cc @averinevg

…ver_mergres2

cderb and others added 30 commits December 2, 2021 00:05

function for retrieving solutions

a30ec12

redefine tunable test arount GetSolution

6a2e3cc

getallconfigs function

f4d819e

Merge remote-tracking branch 'origin/develop' into cderb/fin_perf_com…

ea8f6cf

…pile

remove inline operators for perf Configs

7b778a8

fix compilation error, allow for calls to GetAllSolvers

a796d83

rename function

5f8dee1

legacy solver boolean, getallsolutions mod, performance parameter getter

b61ef31

debug prints for perf cfg string function, added db_update flag to ex…

f9d855e

…ecution context

Add base class for searchable solvers

1bdd262

Fix formatting

935b7cc

Merge branch 'develop' into searchable_solver_base

f7f0a92

Remove redundant 'virtual'

62ca5bd

Remove default arguments on override methods

68a85f5

Fix formatting

c5f7d14

Fix formatting

ea39f5e

Fix formatting

e6ab90b

Merge branch 'develop' into searchable_solver_base

654bf37

Merge remote-tracking branch 'origin/develop' into cderb/fin_perf_com…

8be64c2

…pile

remove comments

35a79fe

spacing

a81e034

formatting

d3e2345

update call to SolverDbId

cbccc05

add forward declaration of PerformanceConfigConvOclBwdWrw2 template v…

7db661d

…alues

forward declaration of structs

c8ef874

Merge remote-tracking branch 'origin/develop' into cderb/fin_perf_com…

4586def

…pile

revert movement of operator override

df156ae

ocl tidy

de09e34

clang-format

6942f72

cderb marked this pull request as draft June 2, 2022 20:27

cderb mentioned this pull request Jun 3, 2022

Enable Fin Perf Eval with Callback #1582

Closed

cderb marked this pull request as ready for review June 3, 2022 20:31

junliume previously approved these changes Jun 6, 2022

View reviewed changes

JehandadKhan previously approved these changes Jun 6, 2022

View reviewed changes

averinevg previously approved these changes Jun 7, 2022

View reviewed changes

atamazov suggested changes Jun 7, 2022

View reviewed changes

atamazov suggested changes Jun 8, 2022

View reviewed changes

src/include/miopen/solver/implicitgemm_util.hpp Outdated Show resolved Hide resolved

change to error throw, remove if clause

9563a5e

cderb dismissed stale reviews from averinevg, JehandadKhan, and junliume via 9563a5e June 8, 2022 17:55

static analyzer ignores

2868d28

atamazov reviewed Jun 8, 2022

View reviewed changes

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

src/solver/conv_hip_implicit_gemm_bwd_v4r1.cpp Outdated Show resolved Hide resolved

src/solver/conv_hip_implicit_gemm_fwd_v4r4.cpp Outdated Show resolved Hide resolved

cderb and others added 4 commits June 8, 2022 15:36

Update src/include/miopen/any_solver.hpp

80630de

Co-authored-by: Artem Tamazov <artem.tamazov@gmail.com>

Merge remote-tracking branch 'origin/develop' into cderb/fin_perf_sol…

8daecf0

…ver_mergres2

Merge branch 'cderb/fin_perf_solver_mergres2' of https://github.com/R…

59d1ec0

…OCmSoftwarePlatform/MIOpen into cderb/fin_perf_solver_mergres2

replace nolint with throws

a5499d9

atamazov reviewed Jun 10, 2022

View reviewed changes

atamazov approved these changes Jun 17, 2022

View reviewed changes

junliume approved these changes Jun 17, 2022

View reviewed changes

Merge remote-tracking branch 'origin/develop' into cderb/fin_perf_sol…

5fcf352

…ver_mergres2

junliume merged commit 45801b0 into develop Jun 22, 2022

cderb deleted the cderb/fin_perf_solver_mergres2 branch July 22, 2022 16:59

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Fin Perf Eval handles (reboot) #1548

Fin Perf Eval handles (reboot) #1548

cderb commented May 18, 2022

cderb commented Jun 2, 2022

atamazov commented Jun 3, 2022

atamazov commented Jun 3, 2022

cderb commented Jun 3, 2022

atamazov commented Jun 3, 2022

junliume commented Jun 6, 2022

atamazov Jun 7, 2022

cderb Jun 7, 2022 •

edited

Loading

atamazov Jun 8, 2022

cderb Jun 8, 2022

atamazov Jun 8, 2022

atamazov Jun 8, 2022

cderb Jun 8, 2022

atamazov Jun 9, 2022

cderb Jun 9, 2022

atamazov Jun 10, 2022

atamazov Jun 10, 2022

cderb Jun 10, 2022

atamazov Jun 14, 2022

cderb Jun 16, 2022

atamazov left a comment

Fin Perf Eval handles (reboot) #1548

Fin Perf Eval handles (reboot) #1548

Conversation

cderb commented May 18, 2022

cderb commented Jun 2, 2022

atamazov commented Jun 3, 2022

atamazov commented Jun 3, 2022

cderb commented Jun 3, 2022

atamazov commented Jun 3, 2022

junliume commented Jun 6, 2022

Choose a reason for hiding this comment

cderb Jun 7, 2022 • edited Loading

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

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

Choose a reason for hiding this comment

atamazov left a comment

Choose a reason for hiding this comment

cderb Jun 7, 2022 •

edited

Loading