Fix: Use use_hip_kernels and use_asm_kernels in Naive convolution Solvers

[carlushuang]

this fix #1319

[atamazov]

see #1323 (comment)

[carlushuang]

@atamazov fixed in 4b25546

[atamazov]

LGTM!

[junliume]

@JehandadKhan test_find_db somehow failed for this PR, no obvious errors in the log except [GetKernels] 0 kernels for key: miopenIm2d2Col "c192i28_28w5_5p1_1s1_1d1_1t1" I tried to see if it is related to specific nodes but now it is failing on both ixt-sjc2-11 and ixt-sjc2-17.
CC: @atamazov

[junliume]

@JehandadKhan test_find_db somehow failed for this PR, no obvious errors in the log except [GetKernels] 0 kernels for key: miopenIm2d2Col "c192i28_28w5_5p1_1s1_1d1_1t1" I tried to see if it is related to specific nodes but now it is failing on both ixt-sjc2-11 and ixt-sjc2-17. CC: @atamazov

@carlushuang @atamazov this test is passing on develop branch and consistently fails with this PR

[atamazov]

Yeah, that is possible since OCL and HIP kernels are disabled...

[atamazov]

...because this PR disables the rest of HIP solvers (Naive ones) that remain enabled even after #1329 because of #1319. But find-db test needs some of HIP solvers enabled. I'll propose the fix in a minute.

[atamazov]

Just after this:
https://github.com/ROCmSoftwarePlatform/MIOpen/blob/53f34c75c4bdf77b52fb75646f14b49d25fea06f/test/find_db.cpp#L231
add

#if WORKAROUND_SWDEV_292187
    setenv("MIOPEN_DEBUG_HIP_KERNELS", "1", 1);
    setenv("MIOPEN_DEBUG_OPENCL_CONVOLUTIONS", "1", 1);
#endif

You may also need to #include <miopen/target_properties.hpp> for the definition of WORKAROUND_SWDEV_292187.

[atamazov]

⚠️ This PR is expected to be NFC. But after merging #1329 (that adds the "MIOPEN_DEBUG_HIP_KERNELS=0" workaround for Navi21) this is not so.

Now the side effect of this PR is that it disables Naive Solvers (in some cases at least), and we may encounter a number of unexpected issues later (the test_find_db is just an early sign).

It seems like we need to replace the "MIOPEN_DEBUG_HIP_KERNELS=0" workaround from #1329 with disabling HIP solvers on individual basis (and this way keep Naive Solvers enabled). I recommend implementing this right here.

After that, the W/A in test_find_db won't be needed.

[atamazov]

Actually it may be even better to put this PR on hold. I am going to investigate SWDEV-292187, narrow the existing W/A (from #1329) and create a new PR. At the very beginning of that work, I'll need HIP solvers individually disabled (instead of global switch). This is the same change as I just proposed to do in this PR. So let's save efforts, and wait several days.

[atamazov]

Block merging due to #1323 (comment)

[atamazov]

#1353 will unblock this. After merging it and updating this branch from develop, 15d695b and 9b8a200 will not be necessary and should be reverted.

[atamazov]

Now blocked only by #1398 + update from develop.

[atamazov]

LGTM