Singnificat performance drop in when using miopen-opencl vs miopen-hip

[artyom-beilis]

I found that there sometimes 2 orders of magnitude difference between miopen-opencl and miopen-hip in some cases:

• Using stride, for example for Input 16:256:14:14, Output 16:512:7:7 Kernel=3, Stride=2, Pad=1 forward convolution using miopen-hip takes: 2.75 ms while same with miopen-openl 19.16ms
• Convolutions with kernel size = 1: Input 16:256:14:14, Output 16:512:14:14 Kernel=1, Stride=1, Pad=0` hip gives 1.26ms while opencl gives 195ms(!)

Setup:

CPU: Intel(R) Core(TM) i5-6600 CPU @ 3.30GHz
GPU: rx560 /gfx803 16 compute units 1196MHz clock
ROCm: 1.9
miopen - latest git.

[artyom-beilis]

Additional information I notices:

While most measurements using MIOpen-opencl return ~190ms time once in a while it gives same short execution as MIOpen-hip i.e. ~1.3ms - that suggest it related somehow to timing. Note, that I do warm-up before starting actual measurements and the selected algorithm is always (1) miopenConvolutionFwdAlgoDirect regardless short or long measurement result.

Artyom

[atamazov]

(1) Do you have miopengemm installed? AFAIR MIOpen-opencl uses it for GEMM algorithms while MIOpen-hip uses rocBLAS. Absence of miopengemm may lead to perf degradations of MIOpen-opencl for convolutions where GEMM algorithm is optimal.

(2) We do not supply performance database for devices with 16 compute units. As your device has 16CUs, you need to auto-tune your MIOpen installation(s). More details in https://github.com/ROCmSoftwarePlatform/MIOpen/blob/master/doc/src/perfdatabase.md

(3) If the above do not resolve the issue you are seeing, please attach console logs from both MIOpen-opencl and MIOpen-hip obtained with export MIOPEN_LOG_LEVEL=5.

[dagamayank]

Closing this issue due to inactivity. Also, differences between HIP and OpenCL is expected. We are working to resolve this but with no clear ETA at this time.