3x3 convolutions performance problem

[ghost]

Hi, dear friends!

We are investigating MiOpen potential to use in our deep learning applications, but Resnet-52 shows strangely long times for layers with 3x3 convolutions (batch size 64).

For 1x1 convolutions AMD Vega often performs on a par with GTX 1080 Ti, but for 3x3 convs the time is up to 10 times longer.

For example,
{28, 28, 128} -> {28, 28, 128} by filter {3, 3} with stride {1, 1}:
NVIDIA 1080 Ti: 13.6 ns (73470 ips)
AMD VX VEGA: 127.2 ns (7862 ips)

All layers perf in TSV: https://gist.github.com/hex000/de2aebf622d2120fb6a57c42a0c7d90e

Here the the time is given per image. The test goes for the whole 1 sec before hipDeviceSynchronize(), so the sync time is amortized.

What can be wrong here? Maybe I haven't set some secret flag? The algorithm chosen by FindConvolution is always "miopenConvolutionFwdAlgoWinograd", and for such convolutions it offers no other options.

The test code is at https://gist.github.com/hex000/81fe99ae8c10f4384e64960275e91554, it compiles both for AMD and NVIDIA.

[dagamayank]

@Hex000 Thanks for contacting us. While I go and run your test code and understand the issue, can you please tell me about your system configuration?

• Which distro?
• Output of uname -a
• Output of /opt/rocm/bin/hipconfig

How was MIOpen installed? Was it built from source?

[dagamayank]

@Hex000 I am mostly certain something is incorrect with your environment. On my Vega GPU I see quite similar times to NVIDIA 1080Ti for that particular layer you mentioned above. Can you please let me know the exact model of your GPU as well.

12	13.0 ns (76852 ips)	conv {28, 28, 128} -> {28, 28, 128} by filter {3, 3} with stride {1, 1}	algo 3	expected 11.9 ns	0.0 Mb

[ghost]

@dagamayank, thank you!
This is good news, I need to check my configuration.

GPU is AMD RX Vega 64

hipconfig output:
https://gist.github.com/hex000/2fdc37178854038cf7600133acebf92d

The MiOpen installation info:
https://gist.github.com/hex000/12462cd6e0a75796b9beb24b6a798f4b
Official binary distribution from http://repo.radeon.com/rocm, is it the best origin?
Or maybe it's better to build from source?

Driver info:
https://gist.github.com/hex000/1c311c4a6fc5545b5aba4231bbf9300b
It's integrated into kernel, is it the best way?

[prostowww]

@dagamayank We tried to run with driver provided with kernel 4.11.0-kfd-compute-rocm-rel-1.6-148, from http://repo.radeon.com/rocm and got around 120 ns.
With ubuntu default kernel, 4.4.0-62-generic, and AMDGPU-PRO 17.30 driver we got 19 ns, which is much better, but still more than your result.

[dagamayank]

@Hex000 @prostowww Please consider the performance that I shared as "dev-preview".

MIOpen requires some improvements in the base software stack which are planned to be released within the next two weeks as part of ROCm 1.6.4. You systems are currently configured with ROCm 1.6.3 (the last public release), which may be the cause of the poor performance that you notice.

I will highly recommend not to mix the AMDGPU-PRO and ROCm software stacks on the same system.

As an experiment can you please try setting this env. variable and check the performance again?
export MIOPEN_DEBUG_AMD_ROCM_PRECOMPILED_BINARIES=0

I will ping you once ROCm 1.6.4 is public.

[masahi]

@Hex000 just curious, do you have both cuda and rocm installed on the same machine? Have you experienced any issue?

[ghost]

Dear friends, thank you for the answers!

No, CUDA and ROCM are on different machines, so this was not tested.

One more complain from me as a user, warping nine-dimensional space for 30 min is too tough:
[02] Testing conv {56, 56, 64} -> {56, 56, 64} by filter {3, 3} with stride {1, 1}
Searching the best solution in the 9 dim space. Please, be patient it may take few minutes.
Runs left : 13714, min time so far : 3.7838, curr time : 3.96322, 8, 16, 8, 16, 1, 2, 5, 3, 1

[dagamayank]

@Hex000 @prostowww

Just wanted to let you know we released ROCm 1.6.4 last week along with MIOpen v1.1.4, if you are already not aware. Can you please try updating your systems and run the performance experiment again?

For performance measurements I would first like you to set an additional parameter for now.

sudo -s
echo 1 > /sys/module/amdkfd/parameters/noretry 
exit

The above is right now a workaround for an issue with page migration in ROCm. This will be fixed soon.

[ghost]

dagamayank, thank you!
noretry gives ~10%

[dagamayank]

@Hex000 FYI - there is a new release of both ROCm and MIOpen, so you may want to benchmark again. For now, I am closing this issue. Please create a new issue if you have questions or notice discrepancies.