Previous Entry Share Next Entry
Теперь я умею быстро умножать матрицы!
evatutin
Давно хотел попробовать, дошли руки сделать умножение матриц на CUDA'е (сделано самостоятельно, без копирования примеров из SDK, но с подглядыванием в Google и с рядом оптимизаций по сравнению с известной русскоязычной литературой). В качестве примера взято умножение двух матриц 2048x2048, тип элементов float, реализованы разные варианты кода.
Начато все было с обычного умножения "в лоб" с использованием 3 циклов (то, что проходится на первом курсе при изучении циклов на программировании). Оказалось, что данный вариант может быть ускорен в 10-2000 раз...
Прежде всего, можно вспомнить о том, что компилятор (MSVC 2012 в данном случае) умеет оптимизировать код. Включение оптимизации сокращает время выполнения операции с 103 с до 68 с (1,5x).
При нахождении суммы

s += A[i][k]*B[k][j]

обращения к памяти для матрицы B происходят прыжками через N*sizeof(float) байт, что плохо для кэша (а матрицы занимают в памяти 16 МБ каждая и разумеется в кэш не помещаются). Первое, что приходит в голову: можно закэшировать j-й столбец матрицы B во временном массиве (и переставить циклы по i и j местами), тогда обращения в кэш пойдут подряд, заработает hardware prefetch. Подобный вариант сокращает время выполнения до 7 с (еще 9,7х), как оказывается важно уважать невидимую работу кэша. Если раскрутить внутренний цикл на 4, то время сокращается до 2,9 с (еще 2,4х), рекорд для однопоточного выполнения на CPU Intel Core i7 4770.
Также реализован вариант блочного умножения матриц, который эффективно использует кэш путем предзагрузки в него маленьких подматриц 32х32 матриц A и B с последующим их умножением, на что уходит 6,2 с. Если раскрутить на 4 внутренний цикл умножения, то время можно дополнительно снизить до 3,3 с, лучше не получилось (на Core 2 Duo E6300 данный вариант чуть быстрее, чем кэширование столбца).
То же самое можно сделать "в лоб" на CUDA: запустить NxN потоков, разбросав их по блокам сетки, в каждом потоке посчитать соответствующий элемент C[i][j], на что GeForce GTX 770 тратит 1,3 c. Неплохо, почти на 2 порядка быстрее CPU, но это только начало.
Данному алгоритму свойственно относительно большое количество обращений в глобальную память GPU (что является бутылочным горлышком) при относительно малом количестве вычислений. Соответственно необходимо попробовать оптимизировать работу с глобальной памятью путем использования памяти разделяемой. Первый вариант, который был реализован по аналогии с CPU-кодом — буферизация столбца. Сделано было просто ради эксперимента, т.к. оптимизация работы с память GPU значительно отличается от CPU и рекомендации тут существенно разные. А получилось даже медленне реализации в лоб, виноват тут coalescing, вернее его отсутствие. Если кэшировать строку вместо столбца (реализация получается во многом симметричной относительно главной диагонали :), то получается выигрыш и время сокращается до 0,24 с (еще 5,4х).
При использовании блочного подхода время можно сократить до 0,14 с, а если цикл раскрутить на 4, то и до 0,05 с! Таким образом, самый быстрый вариант кода быстрее самого первого в 2060 раз!!! Кто скажет, что оптимизация ПО не нужна? Производительность видеокарты при этом составляет 344 GFLOP/s.
Все представленное выше подитожено на картинке



Теперь для полноты картины необходимо потестировать код на других CPU и GPU и проанализировать результаты. Буду признателен результатам тестирования, присылайте (бинарник, обсуждение). На всякий случай, на похожей задаче работает Linpack, по результатам которого строится TOP500, только там матрицы побольше и алгоритм заточен несколько под другое железо.

Во время тестирования (1-2 мин) желательно приостановить BOINC и не создавать другой активности на компьютере, т.к. тесты CPU упираются в объем доступной кэш-памяти, а любая активность очень здорово по этому ударяет. Прошу указывать модель процессора, на котором была запущена программа.

PS. Оговорюсь, что в коде представлены не все возможные варианты (нет векторизации, многопоточности, распределенной обработки, поддержки нескольких GPU и т.п.) — не все сразу :)

[upd 16.08.13]
Во время тестирования оказалось, что на ряде видеокарт (8400 GS, GT 555M, GT 640M) запуск одного или нескольких CUDA-ядер не удается, поэтому программа показывает черти-какую производительность. Поэтому в v1.1 добавлена проверка ошибок CUDA и сделан компромиссный вариант для запуска потоков в блоке в конфигурации 16x16 (до этого было 32х32).

[upd 19.08.13]
Компромиссный запуск в конфигурации потоков в блоке 16x16 — это здорово, только помогает не всегда (хотя в примерах из SDK есть подобные фокусы для более старых видеокарт). Оказалось, что долго CUDA-ядру работать нельзя: либо будет перезапущен видеодрайвер (Windows 7, ядро завершится с ошибкой), либо система повиснет и не будет реагировать (Windows XP, ядро доработает до конца), либо может появиться ошибка 30 (unknown error) или 6 (the launch timed out and was terminated). Эмпирически на запуск CUDA-ядра необходимо тратить не более 1-3 с, чтобы подобные ситуации не возникали. Самое неприятное то, что на современной видюхе выполнение CUDA-ядра может уложиться в таймаут, а на старой (8xxx, 9xxx) или мобильной (5xxM, 6xxM) — нет. Следить за этим (а это тщательное тестирование, лишние проверки в коде, разбивка большого CUDA-ядра на несколько маленьких с потерей быстродействия) необходимо программисту, полагаться на драйвер нельзя, что неприятно. При запуске задачи необходимо стремиться к как можно большему числу потоков в блоке. Зависимость производительности от конфигурации запуска на GTX 770:

4х4 — 30 GFLOP/s
8x8 — 125 GFLOP/s
16x16 — 244 GFLOP/s
32x32 — 325 GFLOP/s

Вариант кода получился чуть быстрее, чем стандартный пример matrixMul (в нем нет раскрутки), но медленнее CUBLAS-реализации (исходники CUDA-ядер не нашел, но в тексте функции cublasSgemm() есть вызовы трех различных ядер для разных размерностей). В примерах от NVidia производительность измеряется без учета передачи исходных и результирующих данных, у меня это время учитывается (без него получается где-то 380 GFLOP/s).

Результаты тестирования:

GPU CPU
Owner
CPU simple
CPU buf
CPU block
GPU simple
GPU blocks
GPU blocks, unroll 2
GPU blocks, unroll 4
GPU blocks, unroll 8
Best GPU vs Best CPU
1 GTX 770 Core i7 4770 @ 3,4 ГГц
evatutin
68 c
(0,3 GFLOP/s)
2,9 c
(5,9 GFLOP/s)
3,3 c
(5,2 GFLOP/s)
1,1 c
(15,3 GFLOP/s)
80,1 мс
(214 GFLOP/s)
61,9 мс
(278 GFLOP/s)
54,6 мс
(315 GFLOP/s)
53,0 мс
(324 GFLOP/s)
55x
2 GTX 760 Core i5 3470 @ 3,2 ГГц
evatutin
3,9 c
(4,5 GFLOP/s)
3,9 c
(4,4 GFLOP/s)
1,4 c
(12,5 GFLOP/s)
109 мс
(157 GFLOP/s)
83,8 мс
(205 GFLOP/s)
73,8 мс
(233 GFLOP/s)
71,3 мс
(241 GFLOP/s)
54x
3 GTX 670 Core i7 3770 @ 3,46 ГГц
diplomatic
2,9 c
(6,0 GFLOP/s)
3,0 c
(5,8 GFLOP/s)
1,3 c
(13,5 GFLOP/s)
98,8 мс
(174 GFLOP/s)
75,1 мс
(229 GFLOP/s)
65,9 мс
(261 GFLOP/s)
63,7 мс
(270 GFLOP/s)
46x
4 GTX 670 Core i7 3770 @ 3,46 ГГц
ReaDy
2,9 c
(5,9 GFLOP/s)
3,0 c
(5,8 GFLOP/s)
1,3 c
(13,5 GFLOP/s)
98,4 мс
(175 GFLOP/s)
75,1 мс
(229 GFLOP/s)
66 мс
(260 GFLOP/s)
63,6 мс
(270 GFLOP/s)
46x
5 GTX 660 Ti Core 2 Duo E6750 @ 2,66 ГГц
evatutin
8,8 c
(2,0 GFLOP/s)
5,9 c
(2,9 GFLOP/s)
1,4 c
(12,3 GFLOP/s)
127 мс
(135 GFLOP/s)
97 мс
(177 GFLOP/s)
86 мс
(200 GFLOP/s)
85 мс
(203 GFLOP/s)
69x
6 GTX 650 Core i3 550 @ 3,2 ГГц
shar
6,9 c
(2,5 GFLOP/s)
4,9 c
(3,5 GFLOP/s)
4,2 c
(4,1 GFLOP/s)
334 мс
(51,5 GFLOP/s)
242 мс
(70,9 GFLOP/s)
205 мс
(83,9 GFLOP/s)
201 мс
(85,6 GFLOP/s)
24x
7 GTX 570 Core i7 920 @ 2,67 ГГц
sagip
12,2 c
(1,4 GFLOP/s)
8,7 c
(2,0 GFLOP/s)
1,6 c
(10,7 GFLOP/s)
142 мс
(121 GFLOP/s)
87 мс
(198 GFLOP/s)
67 мс
(256 GFLOP/s)
56 мс
(305 GFLOP/s)
155x
8 GTX 550 Ti Core 2 Duo E8400 @ 3,0 ГГц
shar
7,8 c
(2,2 GFLOP/s)
5,2 c
(3,3 GFLOP/s)
4,5 c
(3,9 GFLOP/s)
358 мс
(48 GFLOP/s)
194 мс
(88,8 GFLOP/s)
150 мс
(115 GFLOP/s)
133 мс
(129 GFLOP/s)
39x
9 GTX 550 Ti Core i5 3470 @ 3,2 ГГц
Peregrin Krol
3,6 c
(4,8 GFLOP/s)
3,6 c
(4,7 GFLOP/s)
123 GFLOP/s
26x
10 GTS 450 Core 2 Duo E6300 @ 1,86 ГГц
evatutin
12,8 c
(1,3 GFLOP/s)
8,6 c
(2,0 GFLOP/s)
4,5 c
(3,8 GFLOP/s)
423 мс
(40,7 GFLOP/s)
235 мс
(73,1 GFLOP/s)
183 мс
(93,7 GFLOP/s)
156 мс
(110 GFLOP/s)
55x
11 GTS 450 Eco Core i3 @ 3,4 ГГц
evatutin
3,9 c
(4,5 GFLOP/s)
3,9 c
(4,4 GFLOP/s)
5,7 c
(3,0 GFLOP/s)
537 мс
(31,9 GFLOP/s)
273 мс
(62,9 GFLOP/s)
200 мс
(85,7 GFLOP/s)
173 мс
(99,5 GFLOP/s)
23x
12 GT 440 Pentium E2220 @ 2,4 ГГц
tiano
10,3 c
(1,7 GFLOP/s)
6,7 c
(2,6 GFLOP/s)
7,5 c
(2,3 GFLOP/s)
772 мс
(22,2 GFLOP/s)
421 мс
(40,8 GFLOP/s)
322 мс
(53,3 GFLOP/s)
278 мс
(61,9 GFLOP/s)
24x
13 8800 GT Intel Core Quad Q6600 @ 2,4 ГГц
evatutin
13,2 c
(1,3 GFLOP/s)
9,1 c
(1,9 GFLOP/s)
22,5 GFLOP/s
12x
14 Tesla K20Xm Xeon E5-2650
evatutin
4,4 c
(3,9 GFLOP/s)
4,1 c
(4,2 GFLOP/s)
1,1 c
(16,3 GFLOP/s)
88,3 мс
(195 GFLOP/s)
73,1 мс
(235 GFLOP/s)
71 мс
(242 GFLOP/s)
64 мс
(267 GFLOP/s)
64x
15 Tesla C1060 Xeon 5570
shar
6,4 c
(2,7 GFLOP/s)
4,5 c
(3,8 GFLOP/s)
7,1 c
(2,4 GFLOP/s)
146 мс
(117 GFLOP/s)
132 мс
(130 GFLOP/s)
131 мс
(131 GFLOP/s)
34x


По результатам тестирования видно, что выигрыш в данной задаче по сравнению с однопоточным использованием CPU составляет несколько десятков раз. Можно заметить, что GTS 450 vs GTS 450 Eco практически равны по скорости обработки (разница около 10%) несмотря на существенно различную ПСП глобальной памяти (почти в 3 раза!), что еще раз подчеркивает возможность эффективного использования разделяемой памяти в данной задаче.

[upd 03.06.15]
Результаты работы опубликованы в майском номере CUDA альманаха за 2015 год, спасибо Лидии Андреевой!

Публикации
  • Ватутин Э.И., Мартынов И.А., Титов В.С. Оценка реальной производительности современных процессоров в задаче умножения матриц для однопоточной программной реализации // Известия Юго-Западного государственного университета. Серия: Управление, вычислительная техника, информатика. Медицинское приборостроение. 2013. № 4. С. 11–20.
  • Ватутин Э.И., Мартынов И.А., Титов В.С. Оценка реальной производительности современных видеокарт с поддержкой технологии CUDA в задаче умножения матриц // Известия Юго-Западного государственного университета. Серия: Управление, вычислительная техника, информатика. Медицинское приборостроение. 2014. № 2. С. 8–17.
  • Ватутин Э.И., Мартынов И.А., Титов В.С. Оценка реальной производительности современных процессоров и видеокарт с поддержкой технологии CUDA в задаче умножения матриц // CUDA альманах (май 2015). 2015. С. 9–10.

    • 1
    Чего и следовало ожидать, автор удачно написал.

    • 1
    ?

    Log in

    No account? Create an account