Параллельная Bubble сортировка на GPU - PullRequest
Новая
       10

Параллельная Bubble сортировка на GPU

2 голосов
/ 15 марта 2011

Я реализую простой алгоритм сортировки пузырьков, используя CUDA, и у меня есть вопрос.
я выполняю следующий код, чтобы поменять местами 2 последовательных элемента в массиве: 

if(a[threadIdx.x]>a[threadIdx.x + 1])
    Swap(a[threadIdx.x] , a[threadIdx.x + 1]);

обратите внимание, что количество потоков в блоке составляет половину размера массива. Это хорошая реализация? будут ли потоки в одной деформации выполняться параллельно, даже если есть ветвь? поэтому на самом деле потребуется N итераций для сортировки массива?

также обратите внимание, что я знаю, что есть лучшие алгоритмы сортировки, которые я мог бы реализовать, и я могу использовать Thrust, CUDPP или алгоритм сортировки примеров из SDK, но в моем случае мне просто нужен простой алгоритм для реализации.

Ответы [ 2 ]

2 голосов
/ 15 марта 2011

Я предполагаю:

  1. Массив, который вы хотите отсортировать, мал (<100 элементов) </li>
  2. Это часть более крупного алгоритма графического процессора
  3. массив находится в общей памяти или может быть скопирован туда

Если что-то из этого не соответствует действительности, не делайте пузырьковую сортировку!

количество потоков вблок составляет половину размера массива.Это хорошая реализация?

Это разумно.Когда в деформации возникает расходящаяся ветвь, все потоки выполняют все ветви в идеальной синхронизации, просто у некоторых потоков установлен флаг «отключен».Таким образом, каждая ветвь выполняется только один раз .Единственное исключение - когда ни один поток из деформации не принимает ветвь - тогда ветвь просто пропускается.

BUG!

В вашем коде,однако я вижу проблему.Если вы хотите, чтобы один поток работал с двумя элементами массива, сделайте так, чтобы они обрабатывали его исключительно, то есть: 

if(a[2*threadIdx.x]>a[2*threadIdx.x + 1])
    Swap(a[2*threadIdx.x] , a[2*threadIdx.x + 1]);

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

Еще одна ошибка!

Если ваш блок больше, чем размер деформации, не забудьте поставить __syncthreads() при необходимости.Даже если ваш блок меньше (не должен быть), вы должны проверить __threadfence_block(), чтобы убедиться, что записи в разделяемую память видны другим потокам того же блока.В противном случае компилятор может быть слишком агрессивен в отношении оптимизаций и сделать ваш код недействительным.

Другая проблема

Если вы исправите первую ошибку, у вас будет банк с 2-мя путямиконфликт в вашей общей памяти.Это не очень важно, но вы можете реорганизовать данные в вашем массиве, чтобы избежать их, например, располагать последовательные элементы в следующем порядке:

[1, 3, 5,7, 9, ..., 29, 31, 2, 4, 6, 8, ..., 30, 32]

Таким образом, элементы 1 и 2 принадлежат одному и тому же банку в общей памяти.

1 голос
/ 15 марта 2011

Я рад, что вы понимаете, что пузырьковая сортировка на GPU может работать ужасно плохо! Я изо всех сил пытаюсь понять, как получить достаточный параллелизм без необходимости запускать много ядер. Кроме того, вы можете изо всех сил пытаться работать, когда вы закончите.

В любом случае, чтобы ответить на ваш конкретный вопрос: да, весьма вероятно, что в этом случае у вас будет дивергенция. Однако, учитывая, что ветвь "else" фактически пуста, это не замедлит вас. В среднем (до тех пор, пока этот список не отсортирован), примерно половина потоков в деформации будет принимать ветвь "если", другие потоки будут ждать, а затем, когда ветка "если" будет завершена, потоки деформации могут вернуться к в ногу. Это далеко от вашей самой большой проблемы:)

...