Сжатие разреженных данных с помощью CUDA (CCL: уменьшение маркировки подключенных компонентов)

Одним из возможных подходов может быть использование последовательности:

1. thrust::transform - для преобразования входных данных во все 1 или 0:

   0 27 42  0 18 99 94 91  0  -- input data
   0  1  1  0  1  1  1  1  0  -- this will be our "mask vector"
   

2. thrust::inclusive_scan - для преобразования вектора маски в прогрессивную последовательность:

   0  1  1  0  1  1  1  1  0  -- "mask" vector
   0  1  2  2  3  4  5  6  6  -- "sequence" vector
   

3. Еще один thrust::transform для маскировки нерастущих значений:

   0  1  1  0  1  1  1  1  0  -- "mask" vector
   0  1  2  2  3  4  5  6  6  -- "sequence" vector
   -------------------------
   0  1  2  0  3  4  5  6  0  -- result of "AND" operation
   

Обратите внимание, что мы могли бы объединить первые два шага с thrust::transform_inclusive_scan, а затем выполнить третий шаг как thrust::transform с немного другим функтором преобразования. Эта модификация позволяет обойтись без создания временного вектора «маски».

Вот полностью проработанный пример, показывающий «модифицированный» подход с использованием thrust::transform_inclusive_scan:

$ cat t635.cu
#include <iostream>
#include <stdlib.h>

#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/transform.h>
#include <thrust/transform_scan.h>
#include <thrust/generate.h>
#include <thrust/copy.h>


#define DSIZE 20
#define PCT_ZERO 40

struct my_unary_op
{
  __host__ __device__
  int operator()(const int data) const
  {
    return (!data) ?  0:1;}
};

struct my_binary_op
{
  __host__ __device__
  int operator()(const int d1, const int d2) const
  {
    return (!d1) ? 0:d2;}
};

int main(){

// generate DSIZE random 32-bit integers, PCT_ZERO% are zero
  thrust::host_vector<int> h_data(DSIZE);
  thrust::generate(h_data.begin(), h_data.end(), rand);
  for (int i = 0; i < DSIZE; i++)
    if ((rand()%100)< PCT_ZERO) h_data[i] = 0;
    else h_data[i] %= 1000;
  thrust::device_vector<int> d_data = h_data;
  thrust::device_vector<int> d_result(DSIZE);
  thrust::transform_inclusive_scan(d_data.begin(), d_data.end(), d_result.begin(), my_unary_op(), thrust::plus<int>());
  thrust::transform(d_data.begin(), d_data.end(), d_result.begin(), d_result.begin(), my_binary_op());
  thrust::copy(d_data.begin(), d_data.end(), std::ostream_iterator<int>(std::cout, ","));
  std::cout << std::endl;
  thrust::copy(d_result.begin(), d_result.end(), std::ostream_iterator<int>(std::cout, ","));
  std::cout << std::endl;
  return 0;
}

$ nvcc -o t635 t635.cu
$ ./t635
0,886,777,0,793,0,386,0,649,0,0,0,0,59,763,926,540,426,0,736,
0,1,2,0,3,0,4,0,5,0,0,0,0,6,7,8,9,10,0,11,
$

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

Такой подход должен работать:

1. Используйте thrust::sort для сортировки данных.
2. Используйте thrust::unique для удаления дубликатов.
3. Отсортированные данные с удаленными дубликатами теперь дают нам порядок для нашего выходного набора [0,1,2, ...]. Назовем это нашей «картой». Мы можем использовать метод параллельного двоичного поиска для метка в исходном наборе данных соответствует отображаемому выходному значению.

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

Во всяком случае, вот полностью проработанный пример:

$ cat t635.cu
#include <iostream>
#include <stdlib.h>

#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/transform.h>
#include <thrust/generate.h>
#include <thrust/sort.h>
#include <thrust/unique.h>
#include <thrust/copy.h>


#define DSIZE 20
#define PCT_ZERO 40
#define RNG 10

#define nTPB 256

// sets idx to the index of the first element in a that is
// equal to or larger than key

__device__ void bsearch_range(const int *a, const int key, const unsigned len_a, unsigned *idx){
  unsigned lower = 0;
  unsigned upper = len_a;
  unsigned midpt;
  while (lower < upper){
    midpt = (lower + upper)>>1;
    if (a[midpt] < key) lower = midpt +1;
    else upper = midpt;
    }
  *idx = lower;
  return;
  }

__global__ void find_my_idx(const int *a, const unsigned len_a,  int *my_data, int *my_idx, const unsigned len_data){
  unsigned idx = (blockDim.x * blockIdx.x) + threadIdx.x;
  if (idx < len_data){
    unsigned sp_a;
    int val = my_data[idx];
    bsearch_range(a, val, len_a, &sp_a);
    my_idx[idx] = sp_a;
    }
}


int main(){

// generate DSIZE random 32-bit integers, PCT_ZERO% are zero
  thrust::host_vector<int> h_data(DSIZE);
  thrust::generate(h_data.begin(), h_data.end(), rand);
  for (int i = 0; i < DSIZE; i++)
    if ((rand()%100)< PCT_ZERO) h_data[i] = 0;
    else h_data[i] %= RNG;
  thrust::device_vector<int> d_data = h_data;
  thrust::device_vector<int> d_result = d_data;
  thrust::sort(d_result.begin(), d_result.end());
  thrust::device_vector<int> d_unique = d_result;
  int unique_size = thrust::unique(d_unique.begin(), d_unique.end()) - d_unique.begin();
  find_my_idx<<< (DSIZE+nTPB-1)/nTPB , nTPB >>>(thrust::raw_pointer_cast(d_unique.data()), unique_size, thrust::raw_pointer_cast(d_data.data()), thrust::raw_pointer_cast(d_result.data()), DSIZE);

  thrust::copy(d_data.begin(), d_data.end(), std::ostream_iterator<int>(std::cout, ","));
  std::cout << std::endl;
  thrust::copy(d_result.begin(), d_result.end(), std::ostream_iterator<int>(std::cout, ","));
  std::cout << std::endl;
  return 0;
}
$ nvcc t635.cu -o t635
$ ./t635
0,6,7,0,3,0,6,0,9,0,0,0,0,9,3,6,0,6,0,6,
0,2,3,0,1,0,2,0,4,0,0,0,0,4,1,2,0,2,0,2,
$

Мой ответ аналогичен тому, который дал @RobertCrovella, но я думаю, что проще использовать thrust::lower_bound вместо пользовательского двоичного поиска. (теперь, когда это чистая тяга, бэкенд можно поменять местами)

1. Копировать входные данные
2. Сортировать скопированные данные
3. Создайте уникальный список из отсортированных данных
4. Найдите нижнюю границу каждого входа в уникальном списке

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

#include <iostream>
#include <stdlib.h>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/transform.h>
#include <thrust/generate.h>
#include <thrust/sort.h>
#include <thrust/unique.h>
#include <thrust/binary_search.h>
#include <thrust/copy.h>

int main()
{
  const int ndata = 20;
  // Generate host input data
  thrust::host_vector<int> h_data(ndata);
  thrust::generate(h_data.begin(), h_data.end(), rand);
  for (int i = 0; i < ndata; i++)
  {
    if ((rand() % 100) < 40)
      h_data[i] = 0;
    else
      h_data[i] %= 10;
  }

  // Copy data to the device
  thrust::device_vector<int> d_data = h_data;
  // Make a second copy of the data
  thrust::device_vector<int> d_result = d_data;
  // Sort the data copy
  thrust::sort(d_result.begin(), d_result.end());
  // Allocate an array to store unique values
  thrust::device_vector<int> d_unique = d_result;
  {
    // Compress all duplicates
    const auto end = thrust::unique(d_unique.begin(), d_unique.end());
    // Search for all original labels, in this compressed range, and write their
    // indices back as the result
    thrust::lower_bound(
      d_unique.begin(), end, d_data.begin(), d_data.end(), d_result.begin());
  }

  thrust::copy(
    d_data.begin(), d_data.end(), std::ostream_iterator<int>(std::cout, ","));
  std::cout << std::endl;
  thrust::copy(d_result.begin(),
               d_result.end(),
               std::ostream_iterator<int>(std::cout, ","));
  std::cout << std::endl;
  return 0;
}