Как бы вы реализовали эту функцию в CUDA? (смещения в отсортированном целочисленном векторе)

Question

У меня есть отсортированный целочисленный массив на устройстве, например:

[0,0,0,1,1,2,2]

И я хочу смещения для каждого элемента в другом массиве:

[0,3,5]

(начиная с первого 0находится в позиции 0, первый 1 в позиции 3 и т. д.) Я знаю, сколько будет различных элементов заранее.Как бы вы эффективно реализовали это в CUDA?Я не спрашиваю код, но высокоуровневое описание алгоритма, который вы бы реализовали для вычисления этого преобразования.Я уже смотрел на различные функции в пространстве имен тяги, но не мог придумать какой-либо комбинации функций тяги для достижения этой цели.Кроме того, имеет ли это преобразование широко распространенное имя?

Answer 1

Вы можете решить эту проблему в Thrust, используя thrust::unique_by_key_copy с thrust::counting_iterator. Идея состоит в том, чтобы рассматривать ваш массив целых чисел как аргумент keys для unique_by_key_copy и использовать последовательность возрастающих целых чисел (то есть, counting_iterator) как values. unique_by_key_copy сжимает массив значений в индексы каждого уникального key:

#include <thrust/device_vector.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/discard_iterator.h>
#include <thrust/unique.h>
#include <thrust/copy.h>
#include <iterator>
#include <iostream>

int main()
{
  thrust::device_vector<int> keys(7);
  keys[0] = 0; keys[1] = 0; keys[2] = 0;
  keys[3] = 1; keys[4] = 1; keys[5] = 2; keys[6] = 2;

  std::cout << "keys before unique_by_key_copy: [ ";
  thrust::copy(keys.begin(), keys.end(), std::ostream_iterator<int>(std::cout," "));
  std::cout << "]" << std::endl;

  thrust::device_vector<int> offsets(3);

  thrust::unique_by_key_copy(keys.begin(), keys.end(),          // keys
                             thrust::make_counting_iterator(0), // [0, 1, 2, 3, ...] are the values
                             thrust::make_discard_iterator(),   // discard the compacted keys
                             offsets.begin());                  // the offsets are the values

  std::cout << "offsets after unique_by_key_copy: [ ";
  thrust::copy(offsets.begin(), offsets.end(), std::ostream_iterator<int>(std::cout," "));
  std::cout << "]" << std::endl;

  return 0;
}

Вот вывод:

$ nvcc test.cu -run
keys before unique_by_key_copy: [ 0 0 0 1 1 2 2 ]
offsets after unique_by_key_copy: [ 0 3 5 ]

Answer 2

Хотя я никогда не использовал библиотеку тяги, как насчет этого возможного подхода (простого, но, возможно, эффективного):

int input[N];  // your sorted array
int offset[N]; // the offset of the first values of each elements. Initialized with -1

// each thread will check an index position
if (input[id] > input[id-1]) // bingo! here begins a new value
{
    int oid = input[id];  // use the integer value as index
    offset[oid] = id;     // mark the offset with the beginning of the new value
}

В вашем примере вывод будет:

[0,3,5]

Но если входной массив:

[0,0,0,2,2,4,4]

Тогда вывод будет:

[0,-1, 3, -1, 5]

Теперь, если Thrust может сделать это за вас, удалите_if (offset [i] == -1) и уплотните массив.

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

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

Мои 2 цента!

Answer 3

Сканирование - это алгоритм, который вы ищете.Если у вас нет реализации, библиотека Thrust будет хорошим ресурсом.(Найдите thrust :: scan)

Сканирование (или «сумма параллельного префикса») принимает входной массив и генерирует выход, где каждый элемент представляет собой сумму входов в эту точку: [1 5 3 7]=> [1 6 9 16]

Если вы сканируете предикаты (0 или 1 в зависимости от оцениваемого условия), где предикат проверяет, совпадает ли данный элемент с предыдущим элементом, то вы вычисляете выходной индексэлемент в вопросе.Ваш пример массива

[0 0 0 1 1 2 2] [0 0 0 1 0 1 0] <= предикаты [0 0 0 1 1 2 2] <= отсканированные предикаты </p>

Сейчасвы можете использовать отсканированные предикаты в качестве индексов для записи вашего вывода.

Answer 4

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

Как только эта проблема может быть решена в O (n) (где n - длина ввода) на ЦП, вы будете страдать от выделения и копирования памяти (Хост -> Устройство (вход) и Устройство -> Хост (результат). )) недостатки. Это приведет к снижению производительности по сравнению с простым решением ЦП.

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

В целом CUDA хорошо ускоряет производительность, если:

1. Асимптотическая сложность вычислений высока по сравнению с длиной входных данных. Например, длина входных данных равна n, а сложность O (n ^ 2) или O (n ^ 3).

2. Существует способ разбить задачу на независимые или слабо зависимые подзадачи.

Так что, будь я на вашем месте, я бы не стал делать такие вычисления на CUDA, если это возможно. И если бы это была какая-то отдельная функция или преобразование выходного формата для какой-то другой функции, я бы сделал это в CPU.

Если это часть более сложного алгоритма, ответ будет более сложным. Если бы я был на вашем месте, я бы попытался как-то изменить формат [0,3,5], потому что это добавляет ограничения на использование вычислительных мощностей CUDA. Вы не можете эффективно разделить свою задачу на независимые блоки. Например, если я обработаю 10 целых чисел в одном вычислительном потоке и следующие 10 целых чисел в другом. Второй не знает, где разместить свои выводы, пока первый не закончен. Может быть, я разделю массив на подмассивы и сохраню ответ для каждого подмассива отдельно. Это сильно зависит от того, какие вычисления вы делаете.