Эффективные частичные сокращения заданных массивов элементов, смещения и длины подписок

Question

Для моего приложения мне приходится обрабатывать кучу объектов (скажем, int), которые впоследствии делятся и сортируются по более мелким ведрам. С этой целью, хранить элементы в виде одного непрерывного массива

arr = {0,1,2,3,4,5,6,7,8,9,10,11,12,13,14...} 

и информация о ведрах (подсписках) задаются смещениями к первому элементу в соответствующем ведре и длинами подсписка.

Так, например, при

offsets = {0,3,8,..} 
sublist_lengths = {3,5,2,...} 

приведет к следующим расколам:

0 1 2 || 3 4 5 6 7 || 8 9 || ... 

То, что я ищу, это несколько общего и эффективный способ запуска алгоритмов, такие как сокращение, на ведрах, используя только пользовательские ядра или библиотеку thrust. Суммируя ведра должны давать:

3 || 25 || 17 || ... 

То, что я придумал:

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

• вариант 2: генерировать массив ключей от смещения ({0,0,0,1,1,1,1,1,2,2,3,...} в приведенном выше примере) и используйте thrust::reduce_by_key. Тем не менее, мне не нравится создание дополнительных списков.

• вариант 3: Используйте thrust::transform_iterator вместе с thrust::counting_iterator для генерации выше заданного списка ключей на лету. К сожалению, я не могу придумать реализацию, которая не требует приращений индексов в списке смещений на устройстве и проигрывает параллелизм.

Что было бы самым разумным способом реализовать это?

Answer 1

Внутри Thrust я не могу придумать лучшего решения, чем вариант 2. Производительность не будет ужасной, но это, безусловно, не оптимально.

Ваша структура данных имеет сходство с форматом Compressed Sparse Row (CSR) для хранения разреженных матриц, поэтому для таких матриц можно использовать методы, разработанные для вычисления sparse matrix-vector multiplies (SpMV), если вы хотите повысить производительность. Обратите внимание, что массив смещений формата CSR имеет длину (N + 1) для матрицы с N строками (т. Е. Ковши в вашем случае), где последним значением смещения является длина arr. CSR SpMV code в Cusp немного запутан, но он служит хорошей отправной точкой для вашего ядра. Просто удалите любую ссылку на Aj или x с кодом и передайте offsets и arr в аргументы Ap и Av соответственно.

Answer 2

Вы не указали, насколько велики ведра. Если ведра достаточно велики, возможно, вам удастся избежать копирования смещений и sublist_lengths на хост, итерации по ним и выполнения отдельного вызова Thrust для каждого ведра. В то же время у Ферми может быть 16 ядер в полете, поэтому на этой архитектуре вы сможете обрабатывать меньшие ведра и по-прежнему получать хорошее использование.