Падение производительности при умножении матриц для определенных размеров на AMD Polaris

Question

У меня есть код OpenCL, который умножает 2 матрицы (GEMM) на M = 4096, N = 4096 и K = 16. (т.е. матрицы 4096 x 16 с плавающей запятой)

Я запускаю его на Polaris 560, 16CU GPU.

Код: https://github.com/artyom-beilis/oclblas/blob/master/gemm/gemm.cl

Я заметил очень странное падение производительности для этого размера, умножение матриц с этим размером имеет производительность ~8-10 Гфлопс, а если я изменю N на 4095 или 4097, я получу около 130-150 Гфлопс. Я замечаю аналогичное поведение с другими библиотеками GEMM, такими как clblas или miopengemm - я получаю значительное снижение производительности для этого конкретного размера 4096x16, а изменение N на 1 повышает производительность в несколько раз.

Рабочая нагрузка разделена на рабочие группы по 256 потоков. Каждая рабочая группа обрабатывает тайлы матрицы 128x16 и 128x16 (блок 8x8 на поток).

Я попытался изменить тайлинг матрицы на 96x96 с блоками 6x6 вместо 128x128 с 8x8 - результат тот же.

Я тестировал тот же код с ROCm 3.7 OpenCL, Clover OpenCL и даже с драйвером Windows OpenCL - такое же поведение.

Нет такой проблемы с nvidia gtx 960, имеющим одинаковое количество ядер (потоков) графического процессора и тот же тип / размер памяти.

Я подозреваю, что это как-то связано с кешем / столкновением, но я не понимаю, как это происходит. Таким образом, я не знаю, как это обойти.

Answer 1

Наконец, я обнаружил, что библиотека clBlas (изначально разработанная для AMD) обрабатывает особый случай lda % 1024==0, ldb % 1024==0 вероятно из-за кеша

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

Я обнаружил, что лучше всего переставлять блоки в порядке z-кривой вместо того, чтобы ставить в очередь несколько ядер.

https://github.com/artyom-beilis/oclblas/blob/master/gemm/gemm.cl#L109

Для обработки дел M!=N или же M != 1<<n Я просто увеличил количество рабочих групп по M / N, чтобы приблизиться 1<<n и группы, у которых нет рабочих мест, уходят в попрошайничестве, не добавляя слишком много накладных расходов.

z-порядок улучшил производительность в 4 раза.