Эффективность вейвлет фильтрации сигнала на GPGPU (стр. 4 из 6)
Рисунок 10 - Паттерны доступа, приводящие к конфликтам.
На рисунке 10 приведены два способа обращения к shared-памяти, приводящие к возникновению конфликтов.
Довольно часто используемым приемом при работе с shared-памятью является загрузка блока большой исходной матрицы в небольшую (обычно 16*16) матрицу в shared-памяти и проведение расчетов, используя загруженную матрицу.
При этом сама загрузка (когда каждая нить загружает по одному элементу) обычно приводит к объединению запросов в глобальную память. Однако при обработке полученной небольшой матрицы в shared-памяти следует иметь в виду, что если ее размер кратен 16, то каждый ее столбец размещается в одном банке. Поэтому если каждая нить по очереди перебирает элементы, сначала из первого столбца, затем из второго и т.д., то мы сразу получаем конфликт банков, причем порядка 16.
Самым простым способом избежать этого является дополнение матрицы еще одним столбцом (т.е. от матрицы из 16 столбцов мы переходим к матрице из 17 столбцов). За счет этого расположение столбцов в памяти сдвигается, и каждый столбец теперь равномерно распределяется по всем банкам памяти.
Рисунок 11 - Дополнение матрицы.
Вообще подобный прием - добавление к каждой строке матрицы нескольких дополнительных байт довольно распространен в CUDA - за счет него при доступе к глобальной памяти мы легко гарантируем выравнивание для начала строки, а при доступе к shared-памяти - равномерное распределение столбца по банкам памяти.
Локальная память
Когда нитям блока не хватает имеющихся на мультипроцессоре регистров, то для размещения локальных переменных используется локальная память GPU. Локальная память размещена в DRAM и не кэшируется, поэтому доступ к ней такой же медленный как и к глобальной памяти. Однако компилятор выполняет запросы отдельных нитей к локальной памяти, объединяя в одну транзакцию (coalescing).
Реорганизация вычислений
Как последнее и наиболее радикальное средство повышения производительности рассматривают – изменение принципов алгоритма и используемых функций вычислений.
Так для получения быстрых и удовлетворяющих вашим требованиям программ на CUDA нужно знать специфические команды архитектуры и их стоимость (число тактов, clocks).
Рассмотрим специфические команды архитектуры. Выполнение следующих действий требует 4 тактов мультипроцессора:
· операции сложения, умножения и madd для типа float;
· целочисленное сложение;
· побитовые операции, сравнение, min, max;
· операции преобразования типов.
Функции 1/x, 1/sqrt(x) и __logf(x) выполняются за 16 тактов.
Целочисленное умножение (32-битовое) также требует 16 тактов, однако во многих случаях вместо него можно воспользоваться функциями __mul24(x,y) и __umul24(x,y), которые занимают всего 4 такта, но осуществляют 24-битовое умножение (т.е. если числа не очень большие - старший байт равен нулю - то эти функции точно заменяют умножение).
Крайне дорогостоящими операциями, которые желательно избегать, оказывается целочисленное деление (x/y) и целочисленный остаток от деления (x%y). В ряде случаев их можно заменять побитовыми операциями, в случае когда идет деление на константу (являющуюся степенью двух) или взятие остатка от деления на константу (вида 2n-1), то компилятор сам заменяет операцию на побитовую.
Квадратный корень реализуется как суперпозиция двух операций 1/x и 1/sqrt(x) и поэтому занимает 32 такта.
Деление на значение типа float требует 36 тактов, однако есть функция __fdividef(x,y), осуществляющая деление за 20 тактов (но возвращающая ноль при очень больших значениях y).
Функции __sinf(x), __cosf(x) и __expf(x) требуют 32 тактов, функции sinf(x), cosf(x), sincos(x) и tanf(x) гораздо более дорогостоящи.
Кроме того компилятор может вставлять команды для приведения типа в следующих случаях:
· на вход функции, принимающий int (или usnigned int) передается char или short;
· использование констант типа double внутри выражения, использующего float;
· передаче float-значения на вход функции, требующей double (например sin).
Поэтому рекомендуется явно указывать тип float, например вместо 1.3 писать 1.3f и использовать float-аналоги функций (например __sinf или sinf вместо sin).
Обращение к памяти занимает 4 такта, при обращении к глобальной памяти следует также добавить еще от 400 до 600 тактов.
Вызов __syncthreads() требует 4 тактов, если не нужно ждать ни одной нити в warp'е.
Но следует отметить, что данный этап обычно требует изменений алгоритма, введений новых ограничений на размеры чисел и других особенностей, которые в итоге приводят к громоздкому и совершенно неадекватному при модификациях коду.
Как альтернативу данному методу можно пользоваться опцией компилятора nvcc --use_fast_math | -use_fast_math (использовать библиотеку быстрых, но не точных вычислений), которая автоматически при выполнении всех математических операций заменяет их на более быстрые аналоги, что дает прирост производительности, но и все связанные с этим недостатки. Мы предлагаем пользоваться данным методом при планировании ожидаемого выигрыша от реорганизации вычислений.
1.6 Исследование производительности
Одной из целей данной работы является создание более производительного метода вейвлет преобразования сигнала, чтобы впоследствии рассмотреть возможность применения в интерактивных методах множества прикладных областей (от исследований сейсмических сигналов до сжатия изображений).
На рисунке ниже приведен результат преобразования на основе GPGPU, как можно заметить мы достигли некоторого шумоподавления обусловленного применением вейвлета Добеши 8.
Рисунок 12 – Вейвлет фильтрация сейсмического сигнала
Для исследования производительности немного провести серию тестов на разных объемах входных данных. Заметим, что трудоемкость работы метода не зависит от характера входных данных, а только от их объема. Чтобы сориентироваться в достигнутом на данный момент быстродействии вейвлет преобразования будем тестировать одновременно и последовательный метод вейвлет преобразования для CPU, взятый из промышленной свободно распространяемой библиотеки gsl. В таблицах 3 и 4 приведены конфигурации тестовой системы.
В приложении В отражен запуск теста вейвлет преобразования для разных объемов входных данных. В таблице ниже и на графике приведены времена преобразований для CPU и для GPU.
Таблица 2 — Время на Вейвлет фильтрацию
| Количество отсчетов | CPU преобразование, мс | CUDA преобразование, мс | CUDA в сравнении с CPU, раз |
| 65536 | 12 | 67 | 0,18 |
| 131072 | 26 | 67 | 0,39 |
| 262144 | 81 | 61 | 1,33 |
| 524288 | 159 | 73 | 2,18 |
| 1048576 | 319 | 85 | 3,75 |
| 2097152 | 677 | 97 | 6,98 |
| 4194304 | 1401 | 133 | 10,53 |
| 8388608 | 2855 | 209 | 13,66 |
| 16777216 | 5964 | 355 | 16,80 |
| 33554432 | 11979 | 638 | 18,78 |
| 67108864 | 31802 | 1350 | 23,56 |
Как и ожидалось, время на процессоре растет линейно при росте объемов входных данных, а на GPGPU близка к константе.
Рисунок 13 — Время вейвлет преобразования сигнала
Таблица 3 ‑ Аппаратная часть тестовой платформы
| Центральный процессор: | Intel Core 2 Duo E2200 2.2 ГГц 2 Мб Cache; частота шины 800 МГц; |
| Оперативная память: | 2048 Мб DDR3 800 МГц; |
| Чипсет: | Intel G33; |
| Видеоадаптер: | Дискретный стоковый адаптер NVidia Corporation; NVIDIA GeForce 9800 GT 1024 Мб GDDR-3; |
Таблица 4 ‑ Программная часть тестовой платформы
| Операционная система | Ubuntu Linux 10.04 x86 32бит; |
| Драйвер | NVIDIA Driver for Linux 195.35.08 ; |
| Версия CUDA SDK | CUDA SDK 3.0 for Linux |
| Версия CUDA toolkit | CUDA Toolkit 3.0 for Linux |
В результате выполнения данной работы были достигнуты поставленные цели.
· Разработан модуль вейвлет преобразования на основе GPGPU.
· Исследована производительность модуля, время выполнения на GPGPU NVIDIA 9800 GT в сравнении с CPU Intel Core 2 DUO E2200 в 23 раза меньше.
· Модуль поддерживает вейвлет преобразование для всех распространенных материнских вейвлетов.
· Модуль внедрен в действующий программный комплекс по анализу сейсмических сигналов Seismo Detector.
· Исследована практическая эффективность Cuda технологии для решения многопроходных задач, допускающих мелкозернистый параллелизм.