То, что вы просите, это классический параллельный алгоритм, называемый сжатием потока 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>());
Если Thrust - не - опция, вы можете реализуйте сжатие потока самостоятельно (есть много литературы по этой теме). Это веселое и разумно простое упражнение, а также является основным строительным блоком для более сложных параллельных примитивов.
(1) Строго говоря, это не в точности сжатие потока в традиционном смысле , поскольку уплотнение потока традиционно является стабильным алгоритмом, но ваши требования не включают стабильность. Это непринужденное требование могло бы привести к более эффективной реализации?
Сжатие потока - хорошо известная проблема, для которой был написан код лота (Thrust, Chagg для цитирования двух библиотек, которые реализуют сжатие потока на CUDA).
Если у вас есть относительно новое устройство с поддержкой CUDA, которое поддерживает внутреннюю функцию как __ballot (compute cdapability> = 3.0), стоит попробовать небольшую процедуру CUDA, которая выполняет сжатие потока намного быстрее, чем Thrust.
Здесь находит код и минимальный документ. https://github.com/knotman90/cuStreamComp
Использует функцию голосования в одиночном режиме для выполнения уплотнения.
С этим ответом я только пытаюсь предоставить более подробные сведения о подходе Давиде Спатаро.
Как вы упомянули, сжатие потока состоит в удалении нежелательных элементов в коллекции в зависимости от предиката. Например, рассматривая массив целых чисел и предикат 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] является альтернативой упомянутому [Th5ust] выше и состоит из трех этапов:
P
- количество запущенных потоков и N
, с N>P
- размер вектора, подлежащего уплотнению. Входной вектор делится на подвектор размером S
, равным размеру блока. Используется встроенный блок __syncthreads_count(pred)
, который подсчитывает количество потоков в блоке, удовлетворяющем предикату pred. В результате первого шага каждый элемент массива d_BlockCounts
, который имеет размер N/P
, содержит число элементов, удовлетворяющих предикату pred в соответствующем блоке. Из трех шагов выше второй выполняется CUDA Thrust's exclusive_scan
примитивный и вычислительно значительно менее требовательный, чем два других.
Для массива элементов 2097152
упомянутый подход выполнен в 0.38ms
на карте NVIDIA GTX 960
, в отличие от 1.0ms
CUDA Thrust's copy_if
. Упомянутый подход работает быстрее по двум причинам: 1) он специально предназначен для карточек, поддерживающих внутренние элементы основы; 2) Подход не гарантирует упорядочение вывода.
Следует заметить, что мы протестировали подход и против кода, доступного в inkc.sourceforge.net . Хотя последний код устроен одним вызовом ядра (он не использует примитив CUDA Thrust), он не имеет лучшей производительности по сравнению с версией с тремя ядрами.
Полный код доступен здесь и немного оптимизирован по сравнению с обычной процедурой Дейвида Спатаро.
[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.