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

Question

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

У меня есть список из 5 миллионов 32-битных целых чисел (на самом деле это изображение размером 2048 x 2560), что составляет 90% нулей. Ненулевые ячейки - это метки (например, 2049, 8195, 1334300, 34320923, 4320932), которые абсолютно не являются последовательными или последовательными в любом случае (это выходные данные нашего алгоритма CCL маркировки компонентов, связанных с пользовательским подключением). Я работаю с NVIDA Tesla K40, поэтому мне бы очень понравилось, если для этого нужна какая-либо работа по сканированию префиксов, то, что она использует SHUFFLE, BALLOT или любую из более высоких функций CC.

Мне не нужен полностью проработанный пример, просто совет.

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

Помеченный Blob

Другие капли будут иметь другую уникальную метку (например, 13282). Но все будут окружены нулями и будут иметь форму эллипсоида. (Мы оптимизировали наш CCL для эллипсоидов, поэтому мы не используем библиотеки). Но одним побочным эффектом является то, что метки BLOB-объектов не являются последовательными числами. Нам не важно, в каком порядке они пронумерованы, но мы хотим, чтобы один шарик помечался #1, а другой - № 2, а последний - #n, где n - количество шариков в изображении.

Что я имею в виду Labled #1? Я имею в виду, что все 2242 ячейки должны быть заменены на 1. и все 13282 ячейки будут иметь #2 и т. Д.

Максимальное количество BLOB-объектов из нашего CCL равно 2048x2560. Итак, мы знаем размер массива.

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

0

cuda gpu cudafy.net

Источник

user959460 01 мар '15 в 18:10

1 ответ

Решение

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

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

Ниже я привел полный пример. Интересно, что процесс может ускориться, если отложить этап сортировки еще одним вызовом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;
}

0

Источник

user7878485 10 май '19 в 21:53

Другие вопросы по тегам cuda gpu cudafy.net

user1695960 01 мар '15 в 19:36 2015-03-01 19:36 · Accepted Answer · 2015-03-01 19:36

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

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"

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
```

Другая 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-битных целых чисел (меток) или каких-либо ограничений на количество повторений конкретной метки в наборе данных, методы гистограммирования кажутся непрактичными. Это заставляет меня задуматься о сортировке данных.

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

использование thrust::sort сортировать данные.
использование thrust::unique удалить дубликаты.
Сортированные данные с удаленными дубликатами теперь дают нам порядок для нашего выходного набора [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,
$