FFT precision issue with 5x5 stride 2 config

[atamazov]

Found on Navi 21 with ROCm 5.2.3 release:

$ MIOPEN_LOG_LEVEL=5 ./bin/MIOpenDriver conv -x 5 -y 5 -W 28 -H 28 -c 192 -n 16 -k 32 -p 2 -q 2 -u 1 -v 1 -l 1 -j 1 -g 1 -V 1 -s 0 -w 2 -t 1 -i 1 -F 1 -S 34
...
MIOpen(HIP): Info [get_device_name] Raw device name: gfx1030
...
MIOpen(HIP): Info [AmdRocmMetadataVersionDetect] ROCm MD version AMDHSA_COv3, HIP version 5.2.22304, MIOpen version 2.19.0.88aee19b4
...
MIOpen(HIP): Info [ConvolutionForwardImmediate] solver_id = fft, workspace = 81330176
Wall-clock Time Forward Conv. Elapsed: 0.620039 ms, Auxiliary API calls: 445.471 ms (GWSS: 0.010498)
MIOpen Forward Conv. Algorithm: 2, Solution: 34/fft
GPU Kernel Time Forward Conv. Elapsed: 0.383168 ms (average)
...
Forward Convolution FAILED: 0.0206052 > 1.5e-06

At first, this needs to be double-checked with:

• Newer ROCm
• Other GPU types.

@junliume https://github.com/ROCmSoftwarePlatform/MIOpen/labels/bug https://github.com/ROCmSoftwarePlatform/MIOpen/labels/correctness https://github.com/ROCmSoftwarePlatform/MIOpen/labels/urgency_low. Please assign to me for further investigation.

[junliume]

@atamazov the same issue is still found with the latest ROCm. I am curious why ConvBinWinogradRxSf2x3g1 is not chosen?

➜  build git:(master) ./bin/MIOpenDriver conv -x 5 -y 5 -W 28 -H 28 -c 192 -n 16 -k 32 -p 2 -q 2 -u 1 -v 1 -l 1 -j 1 -g 1 -V 1 -s 0 -w 2 -t 1 -i 1 -F 1 -S 34
MIOpenDriver conv -x 5 -y 5 -W 28 -H 28 -c 192 -n 16 -k 32 -p 2 -q 2 -u 1 -v 1 -l 1 -j 1 -g 1 -V 1 -s 0 -w 2 -t 1 -i 1 -F 1 -S 34
Warm-up: Wall-clock Total Time: 2389.13 ms, Find Algorithm: 1, Immediate Algorithm: miopenConvolutionAlgoDirect[85]
Forward Conv solutions available: 5
- id: 84 algo: 3, time: 1.01744 ms, ws: 0, name: ConvBinWinogradRxSf2x3g1
- id: 34 algo: 2, time: 1.66536 ms, ws: 81330176, name: fft
- id: 54 algo: 5, time: 2.90089 ms, ws: 0, name: ConvHipImplicitGemmV4R4Fwd
- id: 11 algo: 1, time: 3.60465 ms, ws: 0, name: ConvOclDirectFwd
- id: 91 algo: 0, time: 11.2784 ms, ws: 15052800, name: GemmFwdRest
Wall-clock Time Forward Conv. Elapsed: 1.84173 ms, Auxiliary API calls: 522.153 ms (GWSS: 0.00683594)
MIOpen Forward Conv. Algorithm: 2, Solution: 34/fft
GPU Kernel Time Forward Conv. Elapsed: 1.616193 ms (average)
stats: name, n, c, ho, wo, x, y, k, flopCnt, bytesRead, bytesWritten, GFLOPs, GB/s, timeMs
stats: fwd-conv5x5u1, 16, 192, 28, 28, 5, 5, 32,  3853516800, 10248192, 1605632, 2384, 7, 1.616193
Forward Convolution FAILED: 0.0193713 > 1.5e-06

[atamazov]

@junliume

I am curious why ConvBinWinogradRxSf2x3g1 is not chosen?

Because of the -S 34 option ;)

[ppanchad-amd]

@junliume Is this fixed with latest ROCm 6.1.0? Thanks!

[atamazov]

With 6.0, the issue remains the same.