-
Notifications
You must be signed in to change notification settings - Fork 224
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
Ck kernels invocation refactoring #2379
Conversation
@iq136boy can you take a look since it affects your PRs? |
CI is broken for a while. It can't even checkout the repo. |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Looks good. Nice consolidation and improvement. I'd recommend introducing a base performance config class for these implicit-gemm solvers, and that may remove some more repeated code.
@junliume This can be merged as well |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Part 1
case miopenDouble: | ||
default: | ||
MIOPEN_THROW(miopenStatusInternalError, | ||
"ConvHipImplicitGemmFwdXdlops operation not implemented for this data type"); |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
ConvHipImplicitGemm3DGroupFwdXdlops
auto ptr_iter = FindConvPtrByID(conv_ptrs, kernel_id); | ||
|
||
if(ptr_iter == conv_ptrs.end()) | ||
MIOPEN_THROW("PerformanceConfig kernel '" + kernel_id + "' does not exist"); |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
It is much better to return ConvSolution(miopenStatusInvalidValue) when this happens. Why: This won't break Find operation, i.e. would improve stability of releases at the users' sites. In all other cases (i.e. attempt to run a convolution) returning a solution with error code would lead to MIOPEN_THROW somewhere in the upper layers of the library.
Implementation: Maybe the most robust approach is throwing a custom exception and catching it in GetSolution (and then MIOPEN_LOG_E there). You can find example in comgr.cpp (which handles compilation errors).
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I've just kept initial functionality as it is. But I agree with this comment and I guess it should be a separated issue - how to handle bad perf configs gracefully. Issue #2426 created.
int index = 0; | ||
std::string kernel_id = ""; | ||
std::vector<std::string> valid_kernels; | ||
|
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
[Q] AFAICS we use kernel_id
as PerformanceConfig value, and we still need index
and valid_kernels
for traversing the ComputedContainer. Maybe we can get rid of kernel_id
and store only index
in the database? The positive side effect is smalled perf-db footprint.
Or this would make PerformanceConfig incompatible between CK versions because valid_kernels
may become different?
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I think we can't since we do not have any guarantees that the order of the kernels is persistent. And it may break backward compatibility as well.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
[Recommendation]
@CAHEK7 Please put C++ comment that summarizes this review thread, i.e. explains why we need string kernel_id
as well as index
and valid_kernels
in the CK's PerformanceConfigs.
@CAHEK7 What are your plans about the unresolved post-merge review comments? |
Simplified and unified CK kernel invocation.
Applied patches from #2338, #2369, #2305, #2336 to all other CK kernels and unified the code.
@atamazov