Алгоритм сжатия потока CUDA

Question

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

Пример:

Глобальная память: {0, 0, 0, 0, 14, 0, 0, 17, 0, 0, 0, 0, 13}

Хост Результат Память: {17 , 13, 14, 0, 0, ...}

Самый простой способ - использовать хост для удаления 0 в O(n) времени. Но учитывая, что у меня есть около 1000 элементов, вероятно, будет быстрее оставить все на графическом процессоре и сначала сконденсировать его, прежде чем отправлять его.

Предпочтительным методом было бы создание стека на устройстве, чтобы каждый поток мог выталкивать и вставлять (в любом порядке) в стек или из него. Тем не менее, я не думаю, что CUDA реализует это.

Эквивалентный (но гораздо медленнее) метод будет продолжать пытаться писать, пока все нити не дописал:

kernalRemoveSpacing(int * array, int * outArray, int arraySize) { 
    if (array[threadId.x] == 0) 
     return; 

    for (int i = 0; i < arraySize; i++) { 

     array = arr[threadId.x]; 

     __threadfence(); 

     // If we were the lucky thread we won! 
     // kill the thread and continue re-reincarnated in a different thread 
     if (array[i] == arr[threadId.x]) 
      return; 
    } 
} 

Этот метод только на пользу в том, что мы хотели бы выступать в O(f(x)) время, где f(x) это среднее число ненулевых значений находятся в массиве (f(x) ~= ln(n) для моей реализации, таким образом O(ln(n)) времени, но имеет высокую O постоянные)

Наконец, алгоритм сортировки, такие как быстрая сортировка или слияние также решить про и фактически работает в O(ln(n)) относительное время. Я думаю, что может быть алгоритм быстрее, чем этот, даже если нам не нужно тратить время на упорядочение (свопинг) пар нулевого нуля и ненулевые пары ненулевых элементов (порядок не нужно хранить).

Так что я не совсем уверен, какой метод был бы самым быстрым, и я все еще думаю, что есть лучший способ справиться с этим. Какие-либо предложения?

Answer 1

То, что вы просите за это классический параллельный алгоритм называется поток уплотнительная .

Если Thrust - это вариант, вы можете просто использовать thrust::copy_if. Это устойчивый алгоритм, он сохраняет относительный порядок всех элементов.

Грубый эскиз:

#include <thrust/copy.h> 

template<typename T> 
struct is_non_zero { 
    __host__ __device__ 
    auto operator()(T x) const -> bool { 
     return T != 0; 
    } 
}; 

// ... your input and output vectors here 

thrust::copy_if(input.begin(), input.end(), output.begin(), is_non_zero<int>()); 

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

⁽¹⁾ Строго говоря, это не точно поток уплотнения в традиционном смысле, как поток уплотнительной традиционно устойчивый алгоритм, но ваши требования не включают в себя стабильность. Это непринужденное требование могло бы привести к более эффективной реализации?

Answer 2

Сжатие потока - известная проблема, для которой был написан код лота (Thrust, Chagg, чтобы привести две библиотеки, которые реализуют сжатие потока на CUDA).

Если у вас есть относительно новый CUDA-совместимое устройство, которое поддерживает встроенную функцию как __ballot (вычислительный cdapability> = 3.0), что стоит попробовать небольшую процедуру CUDA, которая выполняет поток уплотнению гораздо быстрее, чем Thrust.

Здесь находится код и минимальный документ. https://github.com/knotman90/cuStreamComp

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

Answer 3

С этим ответом я только пытаюсь предоставить более подробную информацию о подходе Давиде Спатаро.

Как вы упомянули, сжатие потока состоит в удалении нежелательных элементов в коллекции в зависимости от предиката. Например, рассматривая массив целых чисел и предикат p(x)=x>5, массив A={6,3,2,11,4,5,3,7,5,77,94,0} уплотняется до B={6,11,7,77,94}.

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

Подход в [1,2] является альтернативой Упорный-х copy_if упоминалось выше, и состоит из трех этапов:

1. Шаг # 1. Пусть P будет числом запущенных потоков и N, с N>P, размером с вектор, который нужно уплотнить. Входной вектор делится на подвектор размером S, равным размеру блока. Используется встроенный блок __syncthreads_count(pred), который подсчитывает количество потоков в блоке, удовлетворяющих предикату pred. В результате первого шага каждый элемент массива d_BlockCounts, который имеет размер N/P, содержит количество элементов, удовлетворяющих предикату pred в соответствующем блоке.

2. Шаг №2. Эксклюзивная операция сканирования выполняется на массиве d_BlockCounts. В результате второго шага каждый поток знает, сколько элементов в предыдущих блоках записывает элемент. Соответственно, он знает позицию, где следует писать свой соответствующий элемент, но для смещения, связанного с его собственным блоком.

3. Шаг № 3. Каждый поток вычисляет упомянутое смещение с использованием встроенных функций warp и в конечном итоге записывает в выходной массив. Следует отметить, что выполнение шага 3 связано с планированием warp. Как следствие, порядок элементов в выходном массиве не обязательно отражает порядок элементов во входном массиве.

Из трех шагов выше, вторая выполняется CUDA Упорный-х exclusive_scan примитивными и вычислительно значительно менее жесткими, чем два других.

Для массива 2097152 элементов, упомянутый подход выполнен в 0.38ms на NVIDIA GTX 960 карты, в отличие от 1.0ms от CUDA Упорный-х copy_if. Упомянутый подход работает быстрее по двум причинам: 1) Он специально разработан для карт, поддерживающих внутренние элементы основы; 2) Подход не гарантирует порядок вывода.

Следует отметить, что мы протестировали подход и против кода, доступного по адресу inkc.sourceforge.net. Хотя последний код размещен в одном вызове ядра (он не использует примитив CUDA Thrust), он не имеет лучшей производительности по сравнению с версией с тремя ядрами.

Полный код доступен here и немного оптимизирован по сравнению с обычным режимом Дейвиде Спатаро.

[1] M.Biller, O. Olsson, U. Assarsson, “Efficient stream compaction on wide SIMD many-core architectures,” Proc. of the Conf. on High Performance Graphics, New Orleans, LA, Aug. 01 - 03, 2009, pp. 159-166. 
[2] D.M. Hughes, I.S. Lim, M.W. Jones, A. Knoll, B. Spencer, “InK-Compact: in-kernel stream compaction and its application to multi-kernel data visualization on General-Purpose GPUs,” Computer Graphics Forum, vol. 32, n. 6, pp. 178-188, 2013.