Disabling ConvOclBwdWrW2 for asymmetric padding #179
Conversation
|
We need to purge the perf db of entries with this solver with any input that have asymmetric padding. |
|
@JehandadKhan @alexandraBara For awareness. |
There is no immediate harm in this, except perf-db size. Do you mean find-db? |
| @@ -235,6 +235,15 @@ struct ProblemDescription | |||
| std::size_t GetBackwardPadW() const { return GetWeightsWidth() - GetPadW() - 1; } | |||
| std::size_t GetBackwardPadH() const { return GetWeightsHeight() - GetPadW() - 1; } | |||
|
|
|||
| bool IsAsymmetricPadH() const | |||
| { | |||
| return conv.paddingMode == miopenPaddingSame && (GetWeightsHeight() % 2) == 0; | |||
There was a problem hiding this comment.
Please check code around https://github.com/ROCmSoftwarePlatform/MIOpen/blob/cb1ef1f2a89ebbc4e556cce51234456d9dc9ea03/src/convolution.cpp#L242-L248
It seems like in certain cases we can get asymmetric padding even in miopenPaddingDefault mode.
There was a problem hiding this comment.
2 * GetConvPads()[i] - is symmetric padding;
1 - case where filter bigger than input image;
|
IIRC none of OCL kernels support asymmetric padding. And some old asm kernels too. @Kirpich30000 do you remember? |
|
No, I don't remember details about OCL kernels |
|
@atamazov OCL do not support even filters or were disabled earlier. |
What about asm kernels? |
|
Further TODOs:
|
| { | ||
| return conv.paddingMode == miopenPaddingSame && (GetWeightsHeight() % 2) == 0; | ||
| } | ||
| bool IsAsymmetricPadW() const |
There was a problem hiding this comment.
Do these methods provide valid answer for all directions (Fwd/Bwd/WrW)?
There was a problem hiding this comment.
@shurale-nkn @Kirpich30000 please answer
|
Single and multi pass Winograd works with any padding. Other kernels were designed for symmetrical padding. |
|
@alexandraBara @JehandadKhan @cderb For awareness. |
Added condition in IsApplicable.
Fix for issue #142