Performance drop for speciifc tile size M=4096 N=4096 K=16

[artyom-beilis]

I'm using rx 560 16CU 4GB/gfx803

I run into performance issue when working with matrices of this specific size M=4096, N=4096, K=16, if I modify N to 4097 or 4095 performance is changed dramatically:

./test_gemm_miopengemm -m 4096 -n 4096 -k 16
70.1651 GFLOPS 7.41242 ms
./test_gemm_miopengemm -m 4096 -n 4095 -k 16
254.339 GFLOPS 2.04438 ms
./test_gemm_miopengemm -m 4096 -n 4097 -k 16
290.907 GFLOPS 1.78827 ms

It is rocm 3.7

Also I noticed somewhat similar performance drop in other libraries like clBLAS (171.722 vs 245.325 gflops) but to lesser extend. I'm trying to understand the root cause of the issue - why there is 4x performance drop.

[dani-who-likes-engineering]

Did you ever figure out why this happens?

[artyom-beilis]

Yes: I found that clBlas library (developed for AMD originally) handles special case of lda % 1024==0, ldb % 1024==0 probably due to cache

https://github.com/clMathLibraries/clBLAS/blob/master/src/library/blas/specialCases/GemmSpecialCases.cpp#L228

I found that the better way was (for my library) to rearrange blocks in z-curve order instead of queuing several kernels.

That what I implemented in my dlprimitives deep learning library I work on:

Kernel:
https://github.com/artyom-beilis/dlprimitives/blob/master/src/kernels/sgemm.cl#L335

Decision when to use:
https://github.com/artyom-beilis/dlprimitives/blob/master/src/gemm.cpp#L167

To handle cases M!=N or M != 1<<n I just increased number of work groups on M/N to neares 1<<n and groups that don't have jobs exit in the begging not adding too much overhead.

z-order improved performance x4 times on my library.

I don't know how it can be fixed for miopengemm.