Composable kernels integration & Small Fix in GEMM Solver #1692
Conversation
the skeleton of a solver
|
@JehandadKhan now this PR is consistently passing CI, and could you respond to @shurale-nkn 's inquiries about cmake above, and proceed with preparing to merge this one? Thanks! |
|
@shurale-nkn Thanks for the input on the dependency issue and the associated PR in CK to streamline CK as an MIOpen dependency. Once CK is buildable via clang we can remove the special logic added here and streamline the process. |
|
OK, let's do it quick. I'll create an issue not to forget to fix it later. |
|
@junliume Ping |
There was a problem hiding this comment.
Last ping to @atamazov @DrizztDoUrden @asroy and @zjing14 to review
Let's say by the end of 08/24, AOE (Anywhere on Earth)?
| @@ -1,4 +1,4 @@ | |||
| FROM ubuntu:18.04 | |||
| FROM ubuntu:18.04 as miopen | |||
src/include/miopen/solver.hpp
Outdated
| ConvSolution | ||
| GetSolution(const ConvolutionContext& ctx, | ||
| const PerformanceConfigHipImplicitGemmFwdXdlops& config) const override; | ||
| float GetWti(const ConvolutionContext&) const override { return 0.01f; }; |
There was a problem hiding this comment.
Normally, the Solvers that unable to compute WTI should not have thit function implemented. -2.0 will be returned by the base class' GetWti(), which means wti_unknown. See #410
However, Naive convolutions have GetWti() that return very small value (0.01f). This allows MIOpen to use Naive Solvers if no other applicable Solvers have known WTIs. Right now this means that in case of find-db miss, the library will try to use Winograd or GEMM (whatever is faster according to their GetWti's), but if both are not applicable, the library will use Naive Solver. See
https://github.com/ROCmSoftwarePlatform/MIOpen/blob/e18c9383c6d01d417773fdb90269101f3543ad45/src/include/miopen/solver.hpp#L2609-L2611
If you would like to try CK before Naive, and use it instead (because we do expect that CK is faster than Naive), then please use value bigger than 0.01f, e.g. 0.02f. This is very rude heuristics, almost a hack, but may be better than nothing.
Bottom line: Please either remove GetWti() or make it returning 0.02f or alike (and add a comment that should explain why the fixed value is being returned).
There was a problem hiding this comment.
Thanks for the clarification, I have updated the PR
There was a problem hiding this comment.
Good, thanks for adding/re-using comments! Resolved.
| auto argument_ptr = conv_ptr.MakeArgumentPointer( | ||
| const_cast<void*>( // NOLINT (cppcoreguidelines-pro-type-const-cast) | ||
| static_cast<const void*>(tensors.in)), | ||
| const_cast<void*>( // NOLINT (cppcoreguidelines-pro-type-const-cast) | ||
| static_cast<const void*>(tensors.w)), |
There was a problem hiding this comment.
If removal of constness is necessary, then there is a flaw in design (either in the core of the library or in CK)
@DrizztDoUrden can you have a look at this specific place?
| const auto& yDesc = problem.GetOut(); | ||
| if(xDesc.GetType() == miopenInt8x4 || xDesc.GetType() == miopenInt8) | ||
| { | ||
| // rocBlas needs the output to be int32 always |
There was a problem hiding this comment.
[Q] A bugfix in Forward Gemm solvers? How important is it? Maybe it is worth a dedicated PR?
There was a problem hiding this comment.
Potentially yes, however, lets keep it here to avoid churn.
There was a problem hiding this comment.
Okay. Can you please then update the name of the PR with smth like "Small fix in Fwd GEMM solvers"?
|
Three unresolved comments:
If we merge this as is, then let's open ticket to track. |
There was a problem hiding this comment.
Clang support merged in CK
ROCm/composable_kernel#387
Changes can be made in a separate PR in accordance with issue #1706 , or in this PR.
There are no more complaints about my part.
JehandadKhan
changed the title
|
@junliume Ping |
added to #1706 |
|
It's a great milestone, and thank you for your efforts! @JehandadKhan |
No description provided.