Есть ли эффективный способ уплотнения разреженного массива в OpenCL / CUDA? [Дубликат]
3 ответа
То, что вы просите, это классический параллельный алгоритм, называемый сжатием потока 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) Строго говоря, это не в точности сжатие потока в традиционном смысле , поскольку уплотнение потока традиционно является стабильным алгоритмом, но ваши требования не включают стабильность. Это непринужденное требование могло бы привести к более эффективной реализации?
-
1Ссылаясь на (1), да, мне тоже любопытно. Мне нравится идея подсчета числа ненулевых элементов, а затем создание конкретной функции отображения для общего количества ненулевых элементов. Обычно эта функция карты является суммой префикса, но, учитывая, что нам не нужно ее заказывать, есть ли функция, которая вычисляется быстрее? – Dane Bouchie 3 December 2015 в 08:54
-
2Я думаю, что это похоже на «частичный». потоковое уплотнение, где порядок сохраняется локально через warp (или блок, в зависимости от вашей реализации), но не глобально. Другими словами, это означает, что префиксная сумма также является частичной. Да, это было бы быстрее в том смысле, что алгоритм заканчивается раньше (операций меньше, поскольку глобальный порядок не требуется). Является ли это результатом актуальной лучшей производительности, сказать сложно, но это разумное ожидание. – user 3 December 2015 в 09:47
-
3Что касается литературы, здесь есть статья июля 2017 года , с исходным кодом, которая сообщает быстрее, чем thrust :: copy_if. – Tyson Hilmer 18 June 2018 в 11:38
Сжатие потока - хорошо известная проблема, для которой был написан код лота (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] выше и состоит из трех этапов:
- Шаг # 1. Пусть
P- количество запущенных потоков иN, сN>P- размер вектора, подлежащего уплотнению. Входной вектор делится на подвектор размеромS, равным размеру блока. Используется встроенный блок__syncthreads_count(pred), который подсчитывает количество потоков в блоке, удовлетворяющем предикату pred. В результате первого шага каждый элемент массиваd_BlockCounts, который имеет размерN/P, содержит число элементов, удовлетворяющих предикату pred в соответствующем блоке. - Шаг # 2. Эксклюзивная операция сканирования выполняется на массиве d_BlockCounts. В результате второго шага каждый поток знает, сколько элементов в предыдущих блоках записывает элемент. Соответственно, он знает позицию, где следует писать соответствующий элемент, но для смещения, связанного с его собственным блоком.
- Шаг №3. Каждый поток вычисляет упомянутое смещение с использованием встроенных функций warp и в конечном итоге записывает в выходной массив. Следует отметить, что выполнение шага 3 связано с планированием warp. Как следствие, порядок элементов в выходном массиве не обязательно отражает порядок элементов во входном массиве.
Из трех шагов выше второй выполняется 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.