Алгоритмы поиска на графическом процессоре с помощью технологии NVidia CUDA

В данной статье рассмотрим бинарный поиск и связанную с ним задачу сортировки. Попытаемся сделать теоретическую оценку алгоритмов и провести практические тесты.

Введение

Данная статья является продолжением темы CUDA search . В этой теме рассматривался простейший линейный поиск на графическом процессоре с помощью технологии nVidia CUDA. В наихудшем случае т.е. когда искомый элемент массива был последним, алгоритм на GPU превзошел свой аналог на CPU почти на 40%. Количество элементов в массиве было равно количеству процессоров. Такие условия задачи достаточно хороши для GPU. Однако, как быть в других условиях? Как с ними справиться GPU? Подходит ли он для других алгоритмов? Что ж, попытаемся ответить на эти вопросы.

Бинарный поиск для CPU

Почему бинарный? Во-первых: классическая его реализация достаточная проста в программировании. Другие модификации и улучшения могут оказаться полезными лишь при определенных условиях. Мы же рассматриваем общий случай. Во-вторых: максимум за величину порядка ln2 N сравнений мы найдем искомый ключ (либо установим, что его нет в таблице) [1, 440c]. Не много других алгоритмов смогут потягаться с таким быстродействием. И так имеем в самом худшем случае ln2 N сравнений. Правда алгоритм применим лишь для отсортированного списка. Реализация на CPU имеет следующий вид:

Исходный код функции выполняемой на CPU:

  1. __host__ void searchCPU ( float * data, int * result )
  2. {
  3.    int           x =  16777200; // искомый элемент
  4.    int i = 0; // левая граница массива
  5.    int j = 16 * 1024 * 1024; // правая граница    массива
  6.    int k = 0;
  7.    while ( i <= j )        
  8.    {
  9.       k = i +  (j - i) / 2 ;
  10.       if ( data[k] <  x )
  11.          { i = k + 1; }
  12.       else
  13.          if ( data[k] >  x )
  14.             { j = k - 1; }
  15.       else
  16.       {
  17.          result[0] = data [k];
  18.          result[1] = k;
  19.          break;
  20.       }              
  21.    }              
  22. }

Как видим, действительно ничего сложного. Данный алгоритм значительно превосходит по быстродействию простой линейный поиск (во много сотен раз), реализованный даже на GPU, ведь для линейного поиска имеем N сравнений в худшем случае. Но, как уже говорилось выше, он возможен лишь для отсортированного списка (о сортировке поговорим чуть ниже). Кстати говоря, этот алгоритм намного устойчивее, так как для него практически не имеет значение в какой части списка находиться элемент, тогда как для линейного - от этого напрямую зависит время поиска. Но это в реализации на CPU! На GPU для линейного поиска имеем противоположный случай т.е. достаточно устойчивый алгоритм, т.к. множество потоков может почти одновременно охватить различные участки списка.

Бинарный поиск для GPU.

На попытку написать бинарный поиск меня вдохновила статья о параллельном поиске на видеокартах, в которой приводиться график сравнения производительности для CPU и GPU [2], который изображен на рисунке 1.

Рисунок 1 - Absolute performance of index search implementations exploiting thread level and SIMD parallelism on the CPU and the GPU with respect to workload.

Как видим из графика, быстродействие GPU начинает превосходить свой аналог на CPU лишь на определенной отметке, т.е. при достаточны больших объемах дынных. Попытаемся реализовать бинарный поиск для GPU и подтвердить или опровергнуть данные в вышеуказанном графике. Но каким образом разделить вычисления между потоками? Первая мысль была самая простая. Разделить весь список на равные участки, в каждом участке выполнять бинарный поиск. Поспешно написанная программа дала невероятный результат. А именно, очень медленный. Нужно заметить, что чем меньше потоков запускалось, тем быстрее выполнялась программа, т.е. ее быстродействие стремиться к одному потоку. Размышления подтвердили бессмысленность написанного. Подумайте сами, например, имеем список из 100 000 элементов и 1 000 потоков. Разделяем список на потоки, получается, что для каждого потока имеем по 100 элементов. Действительно для 100 элементов бинарный поиск должен выполниться быстрее, чем для 100 000. Но как же другие 999 потоков? Они изначально обречены на неудачу. Они будут безуспешно пытаться найти искомый элемент. А сколько ресурсов уйдет на их запуск и выполнение? Нужно искать другое решение. После теоретических размышлений оно найдено! Нужно реализовать как бы не бинарный поиск, а поиск такого порядка, сколько у нас потоков. Например, имея 32 потока, будет 32-нарный поиск.
Алгоритм таков: делим весь список на количество потоков, Каждый поток смотрит один элемент и сравнивает его с искомым. Если это искомый элемент замечательно. Если текущий элемент меньше, запоминаем его как нижнюю границу. Если больше запоминаем как верхнюю границу. Каждый из потоков найдет свой элемент больший или меньший искомого, на нам нужны только граничные элементы, между которыми лежит искомый, следовательно, нужно среди меньших элементов оставить максимум, а среди больших минимум. Далее повторяем процедуру, но уже не для всего списка а для новых его границ. Повторяем до тех пор пока либо не найдем элемент, либо количество элементов между нижней и верхней границей станет меньше количества потоков. В таком случае для оставшихся элементов применяем любой другой алгоритм, например однопоточный бинарный поиск. Рассмотрим пример.
Пусть имеем 2048 элементов и 32 потока. Для удобства индексы элементов в массиве и их значения будем считать одинаковыми. Ищем, скажем, элемент под номером 65 равный 65. И так первый шаг. Для каждого потока смотрим каждый 32 элемент. т.е. первый поток смотрит первый элемент, второй элемент номер 32, третий элемент номер 64 и т.д. Каждый поток находит либо больший либо меньший элемент. В данном случае меньшие элементы 1, 32, 64. большие элементы 96, 128...2016. Среди меньших элементов максимальным является 64, среди больших минимальным является 96. И так имеем новые границы списка 64..96. Количество элементов между границами равно 96-64 = 32. Повторяем итерацию. 32 элемента делим на количество потоков, каждый поток смотрит свой элемент. Получает что каждый поток смотрит по одному элементу и находит искомый. Итого имеем всего два шага! Теоретически неплохо. Действительно если оценить алгоритм теоретически, то имеем скорость O = (ln количество потоков N + ln2 количество процессоров) в худшем случае. Теоретически очень неплохо. Попробуем реализовать.

Исходный код возможной реализации, функция для GPU:

  1. __global__ void searchGPU ( float * data, /*float * search,*/ int * result )
  2. {
  3.    int x = 16 * 1024 * 1024; // количество элементов
  4.    int mx = 0; // нижняя граница
  5.    int mn = x-1; // верхняя граница
  6.    int step; // шаг
  7.    x =  16777200; // искомый элемент
  8.    int idx; // текущий поток
  9.    while ( mn - mx > 32 ) // пока количество элементов между границами
  10.       //больше количества потоков (в данном случае потоков 32)
  11.     {
  12.       step = (mn - mx) / 32;                                          
  13.       nt idx = mx + (blockIdx.x * blockDim.x + threadIdx.x) * step;                    
  14.       if ( data[idx] == x ) // найден искомый элемент
  15.          {
  16.             result[0] = data [idx];
  17.             result[1] = idx;
  18.             break;  
  19.          }else
  20.          if ( data[idx] > x )
  21.             {
  22.                if(data[idx] < mn)
  23.                mn =  data[idx]; // минимальный среди больших
  24.          }else
  25.          if(data[idx] < x)
  26.          {
  27.             if(data[idx] > mx)
  28.             mx =  data[idx]; // максимальный среди меньших
  29.          }                                                                            
  30.       }
  31. }

Теоретический анализ говорит, что алгоритм должен иметь очень хорошую скорость, в действительности имеем скорость почти в 100 раз большую, по сравнению с первой реализацией, однако, увы, все еще намного меньшую скорости бинарного поиска для CPU.
В чем же дело? Причин много. Например, копирование данных на видеокарту и обратно. Ресурсы на запуск потоков. Задержка при произвольном чтении элементов из памяти. Все эти факторы сложно учитывать при теоретическом анализе алгоритма. Однако вполне возможно, что реализация данного алгоритма превзойдет бинарный поиск для CPU в задачах внешнего поиска. Когда поступает очень много данных с внешнего накопителя, а не с оперативной памяти. Но тут нужно учитывать много других факторов, например скорость чтения данных с устройства, организацию хранения самих данных на устройстве (да и данные еще должны быть отсортированы), и ориентировать алгоритм уже под эти условия.
Итак, алгоритм бинарного поиска не очень хорош в реализации на GPU. Но это в случае, если мы ищем один элемент в массиве. А если нам нужно найти массив элементов в массиве? Такая задача идеально подходит для GPU. Тогда мы организуем алгоритм таким образом, что бы каждый поток искал свой элемент, что будет происходить почти одновременно. Теоретически имеем все туже скорость - O = ( ln2N ). Тогда как для CPU реализация представляет из себя последовательный двоичный поиск каждого элемента, что составит O = ( ln2N * M ), где M количество искомых элементов. Попробуем проверить сказанное на практике

Фрагмент программы, функция для GPU:

  1. __global__ void searchGPU ( float * data, float * search)
  2. {
  3.    int idx = blockIdx.x * blockDim.x + threadIdx.x;
  4.    //data   массив данных
  5.    //search   искомый массив
  6.    int i = 0;
  7.    int j = 16 * 1024 * 1024;                      
  8.    int k = 0;
  9.    while ( i <= j )        
  10.    {
  11.       k = i +  (j - i) / 2 ;
  12.       if ( data[k] <  search[idx] )                
  13.           {
  14.               i = k + 1;                
  15.          }
  16.       else
  17.          if ( data[k] >  search[idx] )
  18.             {              
  19.                j = k - 1;  
  20.             }
  21.       else
  22.            {
  23.               break;     //элемент найден              
  24.            }                            
  25.    }
  26. }

Данная программа на GPU работает в тысячи раз быстрее, чем на CPU, что подтверждает вышесказанное об эффективности алгоритма на GPU в случае поиска массива элементов.

Сортировка на GPU.

Какую из сортировок лучше всего выбрать для реализации на GPU? Было изобретено достаточно большое количество алгоритмов сортировки, основных из них более 25 такое пугающее количество методов на самом деле лишь малая толика всех алгоритмов, придуманных на сегодняшний день. Почему же существует так много методов сортировки? Ответ если не очевиден, то в полнее понятен каждый метод имеет свои преимущества и недостатки, поэтому он оказывается эффективнее других при некоторых конфигурациях данных и аппаратуры [1, 76c].
Обратимся за помощью к книге Кнута, искусство программирования в которой он проводит достаточно подробный анализ всех алгоритмов и подводит резюме.
Алгоритм битонной сортировки, полезен, если можно одновременно выполнять большое число сравнений [1, 400c] Т.е. алгоритм идеально подходит для реализации на GPU. А теперь Взглянем в SDK и находим там готовый пример, реализации данного алгоритма на CUDA.
Битонная сортировка, несомненно, является быстрее её аналога для CPU. Таким образом, я считаю, что если в задаче поиска имеется подзадача сортировки, то вполне оправданно использовать сначала битонную сортировку на GPU (хотя это так же зависит от конфигурации данных и аппаратуры), а затем применить бинарный поиск на CPU (так же зависит конфигурации данных и аппаратуры). Однако, это в тех случаях, когда сортировка окупит себя в дальнейшем. Если постоянно ведется поиск данных, то разумеется, имеет смысл упорядочить эти данные, что бы потом сэкономить время на поиске.

Снова о поиске

Возможны и другие случаи, например, когда данные динамически меняются, или просто поиск осуществляется относительно редко. Отличным примером может служить обычный текст. Возможность поиска в текстовом редакторе неотъемлемая его функция. В данном случае, имеет смысл организовать поиск с помощью параллельной технологии т.к. однопоточный алгоритм на CPU вынужден будет просматривать весь текст последовательно (пусть и не по каждой букве, а с определенным шагом), что к тому же, делает алгоритм неустойчивым, т.к. время поиска напрямую зависит от положения искомой строки в тексте. В случае же алгоритма на GPU имеем более независимый поиск, т.к. множество потоков может начинать просмотр текста с различных его участков.
Не будем подробно рассматривать этот случай, так как он будет аналогичен поиску элемента в числовом массиве, рассмотренному в первой статье. В наилучшем и среднем случае алгоритм на CPU безусловно выиграет по времени поиска, однако в худшем случае, когда искомый элемент будет в конце списка, он проиграет GPU. Разумеется, если речь идет о достаточном большом количестве данных. Какой из алгоритмов лучше использовать однозначно ответить нельзя. В случае CPU мы будем иметь поиск, который может значительно растягиваться по времени, а в случае с GPU будем иметь достаточно постоянное время поиска.

Выводы

И так, подведем итоги. В общем можно сказать, что GPU особенно приспособлены решать проблемы, которые могут быть адаптированы к параллельным вычислениям, когда одна и та же программа (алгоритм) выполняется над многими элементами данных параллельно, требующих интенсивных расчетов: от простых арифметических операций до операций над памятью. Выполнение одной и той же программа для многих элементов данных одновременно снижает требование к изощренному контролю потоков.
Параллельное вычисление распределяет элементы данных между одновременно выполняющимися потоками. Многие приложения, которые обрабатывают большие объемы данных могут использовать модель параллельных вычислений для повышения производительности. Многие алгоритмы не из области обработки изображений могут быть улучшены с помощью параллельных вычислений, начиная с обработки сигналов или физического моделирования систем и заканчивая вычислительной экономикой или биологией.
Однако ускорение алгоритмов с помощью технологии nVidia CUDA далеко не панацея и имеет ряд недостатков и ограничений среди которых особенности архитектуры и особенности программирования. Теоретический анализ алгоритмов не всегда применим в случае технологии nVidia CUDA, т.к. не учитывает многих ее особенностей. Например, вложенный механизм распараллеливания функций и запуска потоков, особенности работы с памятью, конфигурацию оборудования и .т.п. Что качается темы дано статьи, т.е. алгоритмов поиска можно одназначно отметить одну вещь CUDA эффективна в таких алгоритмах, где CPU вынужден повторять итерации, а CUDA может сделать то же самое в разных потоках. Другими словами, если при переносе алгоритма с CPU на GPU мы избавляемся от цикла и переносим его на потоки.
Так же нужно отметить общее высказывание, о том, что CUDA эффективна в таких задачах, где на входе имеем большой массив данных и на выходе так же имеем большой массив данных.

СПИСОК СОКРАЩЕНИЙ

CUDA - Compute Unified Device Architecture
GPU - Graphic Processor Unit
CPU - Central Processor Unit

СПИСОК ЛИТЕРАТУРЫ

1. Кнут Д.Э. Искусство программирования, том 3. Сортировка и поиск, 2-е изд. М.: Вильямс , 2007. С. 824. ISBN 0-201-89685-0
2. Parallel search on video cards, article, oracle server technologies special projects


Автор: Бокатюк Олег

Tags: 

Comments

> k = i + (j - i) / 2 ; Ну и

> k = i + (j - i) / 2 ;

Ну и кто так пишет? Сложно написать k = (i + j) >> 1; ,да? Вобщем-то и сложение здесь лишнее, i можно убрать и двигать вместо него указатель data.

Если написать так - k = (i +

Если написать так - k = (i + j) >> 1, то при сложении может быть переполнение. Этот баг был вроде в двоичном поиске jdk.

никто не запрещает

никто не запрещает использовать глобальную память в случае если недостаточно места в разделяемой памяти за счет понижения производительности битонной сортировки.

Это бессмысленно.

Это бессмысленно. Производительность падает в 100 с лишним раз из-за чтений-записи в глобальную память и синхронизации мультипроцессоров, а не ядер. MergeSort и SampleSort будут намного эффективнее в этом случае, но и они, увы, 4Гб, вмещающиеся на Tesla быстрее процессора не сортируют :(