Find takes >1 sec for some conv configs about miopen HOT 3 OPEN

[Attribution] @junliume @JehandadKhan

• performance
• urgency_normal
• value_unknown
• Proposed assignees: @atamazov @JehandadKhan @kyeonghwanryu

from miopen.

I found some cases that took >1s for miopenFindConvolution[Direction]Algorithm.
We need to identify the reason of this.

I suspect that the reason is that some of the kernels (likely HIP kernels because their build times are long) are missing from the precompiled binary kernel cache. But there could be dozens of other reasons ;)

@kyeonghwanryu We need more info to reproduce the issue.

• Please specify what is the GPU type -- MI300 or MI250?
• Have you used the the normal installable build (i.e the one with binary cache enabled)?
• Do you have the precompiled kernels package installed?
• Do you see the the excessive delays during the first run only or subsequent runs also look suspicious?

The delays can also be caused by copying buffers from the GPU to the CPU and back (this can take a long time), so please also check your application.

I would appreciate if you provide me with the log taken from one of the suspicious configs with the following environment settings:

export MIOPEN_ENABLE_LOGGING=1
export MIOPEN_ENABLE_LOGGING_CMD=1
export MIOPEN_LOG_LEVEL=6
export MIOPEN_ENABLE_LOGGING_ELAPSED_TIME=1
export MIOPEN_DEBUG_LOGGING_QUIETING_DISABLE=1

If your application is multi-threaded or uses MIOpen in the multi-process context, then please also add

export MIOPEN_ENABLE_LOGGING_MPMT=1

Thanks!

from miopen.

@JehandadKhan @cderb Can you please check if we have the configs shown above reside in the list of the "favorite" configs that we use to populate the system-find-db, system-perf-db and the precompiled kernel package?

I've looked for 64-56-56-1x1-64-56-56-16-0x0-1x1-1x1-0-NCHW-FP32-F in find-db and for 2x64x56x56x1x1x1x1x64x16x0x0x0x1x1x0x1x1x0x0x1xNCHWxFP32xF in perf-db. Results:

• gfx90a68
  • find-db: ConvHipImplicitGemmForwardV4R4Xdlops is the best, it is tunable
  • perf-db: ConvHipImplicitGemmForwardV4R4Xdlops:64,256,2,64,128,4,0,1,4
• gfx90a6e
  • find-db: ConvMlirIgemmFwdXdlops is the best, it is tunable
  • perf-db: ConvMlirIgemmFwdXdlops:64,256,2,64,64,4,1,1
• gfx942130
  • find-db: ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC is the best, it is tunable
  • perf-db: ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC: fwd,nhwc,fp32,0,0,64,64,32,16,16,4,1,1,2,2,0,0,0,0,0,1,4,2,1,1,8,1,32,1,4,2,1,1,8,1,32,54

So at the fist glance find-db and perf-db look correct, but I am not sure about kernel db.

from miopen.