Параллельный алгоритм кодирования по Хаффману на GPU и его применение в СУБД MySQL Текст научной статьи по специальности «Компьютерные и информационные науки»

ИНФОРМАЦИОННЫЕ ТЕХНОЛОГИИ

Вестн. Ом. ун-та. 2011. № 4. С. 164-169.

УДК 681.3.012(04)

Д.А. Лыфарь, И.Ю. Кулаков

ПАРАЛЛЕЛЬНЫЙ АЛГОРИТМ КОДИРОВАНИЯ ПО ХАФФМАНУ НА GPU И ЕГО ПРИМЕНЕНИЕ В СУБД MYSQL

Существует множество программ для сжатия данных, но каждую из них можно классифицировать по двум признакам: сжатие осуществляется с потерями или нет и по методу, положенному в основу алгоритма программы. Многие такие программы в качестве последнего этапа используют либо алгоритм Хаффмана, либо арифметическое кодирование. Цель данной работы - реализация алгоритма кодирования по Хаффману на GPU и применение в качестве основного алгоритма сжатия для таблиц InnoDB в СУБД MySQL. Исходный код доступен по адресу https://github.com/ Kentzo/phuffman/tree/cuda.

Ключевые слова: GPU, параллельные алгоритмы, СУБД, кодирование по Хаффману.

Введение

Кодирование по Хаффману всё еще распространено и широко применяется в наши дни. В то же время алгоритм обладает относительно простотой реализацией в случае последовательного потока команд. В этой статье мы рассмотрели возможность параллельной реализации кодирования для графических процессоров, привели сравнительную скоростную характеристику и показали, как наша реализация может быть использована на практике, основываясь на выводах из предыдущей статьи [2]. Алгоритм кодирования, его временные и пространственные характеристики приведены в [1, 10, 11]. В качестве программного интерфейса для графических процессоров была использована NVIDIA CUDA. Полная реализация алгоритма находится в открытом доступе.

1. Обзор существующих решений

На данный момент существует несколько работ, авторы которых решали эту же задачу. Например, в работе [12] описан идентичный подход к кодированию, где оно выполняется в три основных этапа: построение частотной таблицы, префиксная сумма для получения длин кодов каждого блока и запись результатов с использованием атомарных инструкций. Существует несколько различий, главное из них - в работе не были рассмотрены подходы к декодированию, что привносит проблемы из-за того, что код может располагаться на границе двух машинных слов (см. п. «Информация для блочного декодирования»), а также при записи результатов: мы оперируем 128-ми битными целыми и в любом случае операция записи состоит из четырех последовательных атомарных операций «ИЛИ». В [13] также была приведена реализация алгоритма кодирования Хаффмана. В этой работе большее внимание уделено параллельному алгоритма построению таблицы частот для исходных символов с использованием префиксной

* Работа выполнена при поддержке программы ФСР малых форм предприятий в НТС «УМНИК» по договору №09

© Д.А. Лыфарь, И.Ю. Кулаков, 2011

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

2. Канонический алгоритм Хаффмана

Кодирование по Хаффману использует только частоту появления одинаковых символов алфавита во входном блоке данных. Ставит в соответствие символам входного потока, которые встречаются чаще, цепочку битов меньшей длины, и, напротив, встречающимся редко - цепочку большей длины. Один из немногих алгоритмов, который не увеличивает размера исходных данных в худшем случае (если не считать необходимости хранить таблицу перекодировки вместе с файлом). Канонический алгоритм Хаффмана отличает то, что ненулевыми битами кода в худшем случае будут лишь 8 последних бит. Это позволяет существенно снизить расход памяти во время кодирования и декодирования, т. к. достаточно хранить только длину и последние 8 бит кода.

3. Кодирование на GPU

Вход: A = {ai,a2,...,an} - алфавит из n различных символов, W = {Wi,W2,...,Wn} -соответствующий ему набор положительных целых весов. Исходный массив символов B.

Выход: набор бинарных кодов

C = {ci,c2,...,cn}, такой что

1. Ci не является префиксом для Cj

2. ^ wici минимальна. Преобразован-

1=1

ный массив символов BT в соответствие C.

Учитывая то, что CUDA имеет несколько моделей памяти, в псевдокоде мы будем называть их соответственно shared

- для разделяемой памяти, constant - для константной памяти и т. д. Если тип памяти не упоминается явно, то подразумевается глобальная память. Таблицу кодирования, а именно подсчет вероятностей символов и построение кодов для каждого

символа мы выполняем на CPU. Отметим также использование встроенных переменных, значения которых определяются на этапе компиляции: blockDim, blockIdx, threadldx, означающие размерность блока, индекс текущего блока и индекс потока в блоке соответственно и представляющих собой трехмерный вектор. Наша реализация подразумевает |A| = 256, однако не составит труда обобщить ее на алфавиты другого размера. Код символа представлен в виде структуры Code, имеющей два поля code и len (код и его длина соответственно).

constant EncodeTable[] tIdx = blockDim.x*blockIdx.x + threadIdx.x

Length = []

ForEach sym in B do in parallel Code = EncodeTable[sym]

Length[tIdx] = Code.len

End

Exclusive_Prefix_Sum(Length)

// В последнем элементе Length хранится общий размер BT

AllocateResult(Length[sizeof B])

ForEach sym in B do in parallel Code = EncodeTable[sym] code_aligned = Code.code << (2*wordsize_bits - Code.len - (Length[tIdx] mod wordsize_bits))

code_address = result + Length[tIdx]

/ wordsize_bits

atomicOr(code_address,

code_aligned.x)

atomicOr(code_address + 1, code_aligned.y)

End

Рассмотрим возможность выполнения кодирования на GPU. Используя таблицу кодов (которая строится на CPU), каждый поток CUDA может независимо закодировать символ (или набор символов) и записать его в результирующий массив. К сожалению, кодирование сопряжено с синхронизацией во время записи результатов. Чтобы записать очередной код в результирующий массив, необходимо знать место записи предыдущего (коды не выровнены по границе байта). Возможное улучшение - группировка символов в блоки фиксированной величины так, чтобы после кодирования сумма длин кодов блока была не меньше байта. Теперь синхронизации требуют только коды, находящиеся на смежных блоках. Для определения места записи можно воспользовать-

ся результатами Exclusive Prefix Sum. Достаточно взять позицию символа, с которого начинается блок. Для возможности декодирования на GPU результат также разбивается на блоки.

Exclusive Prefix Sum. После кодирования последовательность кодов должна сохраняться. Так как блоки кодируются параллельно, мы не можем записывать результаты в том порядке, в котором они будут получены. Чтобы сохранить порядок, можно воспользоваться Exclusive Prefix Sum:

1. Закодировать каждый символ длиной его кода. На этой стадии мы получаем массив, где каждый элемент содержит длину кода, соответствующего своему смещению символа.

2. Применить Exclusive Prefix Sum (реализация и характеристики рассмотрены в [2; 5]).

3. Сложить каждые N символов, где N

- длина блока.

В результате, каждому блоку будет соответствовать сумма длин кодов всех предыдущих блоков, что и будет являться адресом блока в результирующим массиве. Реализации Exclusive Prefix Sum существуют как для CUDA, так и для OpenCL.

Кодирование. Каждый поток CUDA кодирует блок символов длины N и записывает код в результирующий массив по адресу, полученному после Exclusive Prefix Sum. Так как длины кодов не выровнены по границе байта, для записи закодированного блока в результирующий массив необходимо использовать атомарные операции. В CUDA атомарные операции применимы только к целочисленному типу int, что для нас означает приведение закодированного блока перед записью к массиву из int. Затем каждое значение массива записывается в результат при помощи побитовой атомарной операции OR. Таким образом, гарантируется, что никакой код не будет поврежден, однако возникают конфликты по записи между потоками, которые пишут по одному и тому же адресу, что может привести к снижению производительности, так как один из них должен ждать.

Сложность алгоритма. Можно заметить, что параллельный алгоритм кодирования аналогичен своей последовательной версии. Каждый поток в общем случае кодирует несколько символов, опираясь на кодовую таблицу из константной па-

мяти и смещение в глобальном массиве результата. Алгоритм состоит из трех основных шагов: Ть построение массива длин кодов, Т2 алгоритм префиксной суммы, Тз кодирование символов. Так же фаза построения таблицы кодирования

[7] Бо, которая исполняется последовательно. Временная сложность:

Т1 + Т2 + Тз + Бо =

= О(^ + О(^ + О(^ + О(^ + 0(К ^ К)=

= О^ + К ^ К) = О^ + К ^ К), где К - величина кодовой таблицы.

Информация для блочного декодирования. Декодировщик обрабатывает архив побитно. Для того чтобы сделать декодирование на ОРи возможным, необходимо предварительно разбить закодированную последовательность на блоки равной длины, что в общем случае невозможно. Проблема заключается в том, что граница блока может разделить код на две части, что приведет к ошибкам при декодировании, так как декодировщик рассматривает код только как последовательность бит. Чтобы решить эту проблему, достаточно сохранить длину либо левой, либо правой части (в зависимости от того, какой блок ответственен за декодирование). Если ответственен блок, включающий левую часть кода, то нужно сохранить длину правой части. Эта длина должна быть поставлена в соответствие левому блок. Теперь, декодируя очередной блок, декодировщик знает длину части, которая «не поместилась». Следующий за ним блок знает о части, которую нужно пропустить.

Блочное декодирование. Блочность входного потока данных является необходимым условием для возможности реализации декодирования на ОРи. Она может быть достигнута лишь одним способом: входному потоку должна соответствовать некоторая структура, которая содержит информацию о расположении блоков. В простейшем случае можно было бы просто сохранять смещения всех блоков. К сожалению, это значительно увеличит размер закодированной последовательности: исходный файл может достигать в размерах нескольких гигабайт, а смещение храниться в битах. Например, при кодировании уже 256 МБ для хранения смещений пришлось бы использовать как минимум 4 байта. Гораздо оптимальнее хранить смещения от некоторых фиксированных границ. Данный подход реали-

зуется следующим образом: закодированная последовательность разбивается на блоки фиксированного размера. Далее необходимо решить проблему, когда граница блока делит на две части какой-нибудь код длины C на части Ci и C2, то есть ни один из блоков не будет обладать достаточной информацией для декодирования. Для начала определим, какой блок будет выбран ответственным за декодирование такого кода. Пусть таким блоком будет лежащий слева от границы разбиения. Теперь нам достаточно лишь сохранить количество бит кода, оказавшихся в блоке, лежащем справа от границы. Таким образом, размерность числа не превысит log2(max C). В случае, если бы мы выбрали ответственным блок, лежащий справа от границы разбиения, нам пришлось бы сохранить количество бит кода, оказавшихся в блоке, лежащем слева от границы. Теперь декодировщику достаточно взять длину части для предыдущего и для текущего блоков, чтобы узнать, сколько бит нужно пропустить и сколько лишних бит нужно прочесть.

Вход: набор бинарных кодов

C = {ci,c2,...,cn} размерностью m, таблица соответствия бинарных кодов символам T оригинального сообщения с алфавитом A = {ai,a2,...,an}.

Выход: оригинальное сообщение с

числом символов m

constant EncodeTable[] tIdx = blockDim.x*blockIdx.x + threadldx.x

Length[blockDim.x+blockIdx.x] =

T.blocks[blockIdx.x]

Exclusive_Prefix_Sum(Length[])

AllocateResult(Length[blockDim.x*block

Idx.x])

ForEach code in C do in parallel sym = T[code]

result[Length[blockIdx.x] + tldx.x] =

sym

End

В таблице вместе с массивом соответствия мы храним также количество символов, которое хранит каждый блок для облегчения вычисления места записи под результат, однако количество символов блока может быть легко вычислено во время декодирования, зная, сколько битов должно быть декодировано из каждого текущего блока. Так же, полагаясь на свой-

ство префиксных кодов, каждый поток начинает с нулевой позиции своего массива кодов и двигается до тех пор, пока соответствие в таблице не будет найдено.

Сложность алгоритма

В случае декодирования временная сложность алгоритма аналогична временной сложности последовательной версии. Она включает в себя проход по всем кодам закодированного сообщения Ti, поиск оригинального символа в таблице соответствия и запись оригинального символа в результат T2, алгоритм префиксной суммы T3:

Ti + T2 + Тз= O(2N) + O(1) = O(N)

4. Сжатие данных в таблицах СУБД MySQL

Многие СУБД поддерживают сжатие данных в таблицах для уменьшения занимаемого места на диске (и, соответственно, уменьшения количества дисковых операций) и экономии трафика. Мы взяли для тестирования открытую СУБД MySQL (с InnoDB в качестве хранилища данных) и исследовали текущую реализацию сжатия данных и возможность использования нашей реализации. Существует возможность для используемого хранилища данных установить сжатие (настройкой ROW_FORMAT=COMPRESSED). Однако, несмотря на название настройки, InnoDB сжимает полностью всю таблицу, а не ее отдельные столбцы при помощи алгоритма LZ77. Во многих случаях размер индексов занимает значительную часть размера всей таблицы, поэтому InnoDB сжимает их вместе с данными. Данные в InnoDB имеют страничную организацию, и при использовании сжатия в буфере оперативной памяти могут находиться одновременно сжатая и разжатая страницы. InnoDB использует адаптивный Least Recently Used (LRU) алгоритм, который, в зависимости от того, какой компонент системы больше загружен (CPU-bound или системы ввода-вывода - I/O system bound), определяет наименее используемые данные и вытесняет их из буфера. В случае сильной загрузки I/O InnoDB старается оставить в памяти сжатые версии одинаковых страниц, чтобы в будущем использовать свободный CPU вместо повторной загрузки данных. В случае высокой загрузки CPU InnoDB старается вытеснить обе версии одной и той же мало используемой страницы.

Проведены исследования [4] скоростных характеристик существующей реали-

зации на большом объеме данных с использованием инструмента sysbench [9]. Были замерены скорость чтения и скорость записи i6 таблиц по 25 млн строк в каждой, что дает около 6 Гб на таблицу. Результаты сведены в таблицу.

Скорость кодирования/декодирования данных в существующей реализации InnoDB

Данные

Несжатые Сжатые

Скорость чтения, транзакций/с 4650 30

Время записи, с 19693 38278

После сжатия каждая таблица стала занимать около 3.i Гб. Таким образом, в случае с операцией записи мы получили двукратное преимущество в объеме данных и двукратное отставание в скорости.

В случае эксперимента с чтением мы читаем по 3 млн строк из каждой таблицы в i6 потоков (что дает нам около 24 Гб данных, полностью помещающихся в оперативную память на целевой машине). Результаты для чтения различаются в i50 раз. Это объясняется реализацией декодирования в InnoDB и не зависит от алгоритма сжатия, поэтому в среде CPU-bound в настоящий момент использование сжатия не дает преимущества. Однако в случае с I/O-bound и интенсивной записью мы серьезно можем сократить объем данных, увеличив время всего в 2 раза.

5. Тестирование

Тестирование проводилось на текстовых данных формата xml (тип колонки

TEXT в СУБД MySQL). Для тестов на CPU был выбран пакет FastHF [6] - свободно распространяемая быстрая реализация алгоритма Хаффмана. Следует отметить, что наиболее важная характеристика для алгоритма кодирования - это скорость, поэтому мы замерили, как изменяется скорость кодирования с изменением размера входных данных. Что касается степени сжатия, то его величина варьируется в пределах от 30 до 50 % для наших тестовых текстовых файлов.

Так как наш алгоритм применим не только в области сжатия таблиц, где скорость кодирования должна быть выше скорости записи на диск, но и, например, в передаче сжатых данных по сети, где скорость кодирования должна превосходить или быть примерно равной скорости современных сетевых интерфейсов, чтобы не создавать узкого места в цепочке передачи данных. Существуют специализированные решения, которые имеют относительно невысокую степень сжатия, но высокую скорость кодирования [8]. Тестирование показало значительное ускорение кодирования (результаты приведены на рис.). Основываясь на данных профилировщика, можно сказать, что значительную часть времени занимают атомарные операции, т. к. в нашей реализации они эмулируются, как серия атомарных операций по 32 бита для !28-бит-ных типов и не поддерживаются архитектурой GPU в данный момент.

Размер файла, Мегабайт

Сравнительная характеристика скорости кодирования по Хаффману для ОРУ и ОРУ

Выводы

Мы показали возможность параллельного исполнения кодирования по Хаффману, привели сравнительные характеристики с последовательной реализацией и сделали вывод о целесообразности использования такой реализации на практике. Исходя из результатов тестирования, авторами предлагается использование реализации параллельного кодирования в задачах с критичными требованиями к скорости кодирования. В качестве дальнейшего направления исследования предлагается поиск путей для минимизации количества атомарных операций при записи результатов кодирования в общий массив. Полезным будет также рассмотреть целесообразность использования реализации в различных прикладных задачах (например, JPEG кодирование) и выяснить, насколько влияет на общую скорость кодирования время передачи данных с хоста на устройство.

ЛИТЕРАТУРА

[1] Методы сжатия данных. Устройство архиваторов, сжатие изображения и видео / Д. Ватолин [и др.]. М. : ДИАЛОГ-МИФИ, 2003.

[2] Лыфарь Д. А. Параллельные алгоритмы обработки баз данных // Вестник НГУ. 2010. С. 7280.

[3] Klein S.T., Wiseman Y. Parallel Huffman Decoding with Applications to JPEG Files // The Computer Journal 46(5), February 2003.

[4] Fang W., He B. Luo Q. Database Compression on Graphics Processors // VLDB Endowment. Vol. 3. № 1. 2009.

[5] Harris M. Parallel Prefix Sum (Scan) with

CUDA. URL: http://developer.download.nvidia.

com/compute/cuda/1_1/Website/projects/scan /doc/scan.pdf.

[6] Said A. FastHF: Static and Adaptive Huffman

Coding. URL: http://www.cipr.rpi.edu/~said/

FastAC.html.

[7] Кормен Т. и др. Алгоритмы: построение и анализ. Издательский дом Вильямс, 2007. С. 133.

[8] A fast compressor/decompressor. URL: http:// code.google.com/p/snappy/.

[9] SysBench: a system performance benchmark. URL: http://sysbench.sourceforge.net/.

[10] Schwartz E.S., Kallick B. Generating a canonical prefix encoding // Communications of the ACM. Mar. 1964. Vol. 7. № 3. P. 166-169.

[11] Lelewer D.A., Hirschberg D.S. Data compression // Computing Surveys. Sep. 1987. Vol. 19. № 3. P. 261-296.

[12] Balevic A. Parallel Variable-Length Encoding on GPGPUs. University of Stuttgart.

[13] Grant B., Mann T. File Compression Using the CUDA Framework.