Как использовать thrust для накопления массива на основе индекса?

#c #cuda #thrust

#c #cuda #thrust

Вопрос:

Я пытаюсь накопить массив на основе индекса. Мои входные данные представляют собой два вектора одинаковой длины. 1-й вектор — это индекс. 2-й вектор — это значение. Моя цель — накопить значение на основе индекса. У меня есть аналогичный код на c . Но я новичок в программировании thrust. Могу ли я добиться этого с помощью кода устройства thrust? Какую функцию я мог бы использовать? Я не нашел «map» подобных функций. Является ли это более эффективным, чем код процессора (хоста)? Мой мини-пример кода версии c .

 int a[10]={1,2,3,4,5,1,1,3,4,4};
vector<int> key(a,a 10);
double b[10]={1,2,3,4,5,1,2,3,4,5};
vector<double> val(b,b 10);

unordered_map<size_t,double> M;
for (size_t i = 0;i< 10 ;i  )
{
    M[key[i]] = M[key[i]] val[i];
}
 

Комментарии:

1. Thrust имеет эффективные реализации reduce by key, но они требуют, чтобы входные данные были сгруппированы в порядке ключей, а ваши данные не упорядочены. Итак, ваш выбор — сортировать или упорядочивать по ключу, затем уменьшать или использовать что-то еще. Графические процессоры плохо выполняют произвольный доступ, очень сложно добиться достойного использования полосы пропускания памяти, что является ключом к производительности в такого рода операциях

Ответ №1:

Как указано в комментарии, каноническим способом сделать это было бы изменить порядок данных (ключей, значений), чтобы похожие ключи были сгруппированы вместе. Вы можете сделать это с sort_by_key помощью . reduce_by_key затем решает.

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

Следующее иллюстрирует оба:

 $ cat t27.cu
#include <thrust/reduce.h>
#include <thrust/sort.h>
#include <thrust/device_vector.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/for_each.h>
#include <thrust/copy.h>
#include <iostream>
#include <unordered_map>
#include <vector>

// this functor only needed for the non-reordering case
// requires compilation for a cc6.0 or higher GPU e.g. -arch=sm_60
struct my_func {
  double *r;
  my_func(double *_r) : r(_r) {};
  template <typename T>
  __host__ __device__
  void operator()(T t) {
    atomicAdd(r thrust::get<0>(t)-1, thrust::get<1>(t));  // assumes consecutive keys starting at 1
  }
};

int main(){

  int a[10]={1,2,3,4,5,1,1,3,4,4};
  std::vector<int> key(a,a 10);
  double b[10]={1,2,3,4,5,1,2,3,4,5};
  std::vector<double> val(b,b 10);

  std::unordered_map<size_t,double> M;
  for (size_t i = 0;i< 10 ;i  )
  {
    M[key[i]] = M[key[i]] val[i];
  }
  for (int i = 1; i < 6; i  ) std::cout << M[i] << " ";
  std::cout << std::endl;
  int size_a = sizeof(a)/sizeof(a[0]);
  thrust::device_vector<int>    d_a(a, a size_a);
  thrust::device_vector<double> d_b(b, b size_a);
  thrust::device_vector<double> d_r(5); //assumes only 5 keys, for illustration
  thrust::device_vector<int> d_k(5); // assumes only 5 keys, for illustration
  // method 1, without reordering
  thrust::for_each_n(thrust::make_zip_iterator(thrust::make_tuple(d_a.begin(), d_b.begin())), size_a, my_func(thrust::raw_pointer_cast(d_r.data())));
  thrust::host_vector<double> r = d_r;
  thrust::copy(r.begin(), r.end(), std::ostream_iterator<double>(std::cout, " "));
  std::cout << std::endl;
  thrust::fill(d_r.begin(), d_r.end(), 0.0);
  // method 2, with reordering
  thrust::sort_by_key(d_a.begin(), d_a.end(), d_b.begin());
  thrust::reduce_by_key(d_a.begin(), d_a.end(), d_b.begin(), d_k.begin(), d_r.begin());
  thrust::copy(d_r.begin(), d_r.end(), r.begin());
  thrust::copy(r.begin(), r.end(), std::ostream_iterator<double>(std::cout, " "));
  std::cout << std::endl;
}
$ nvcc -o t27 t27.cu -std=c  14 -arch=sm_70
$ ./t27
4 2 6 13 5
4 2 6 13 5
4 2 6 13 5
$
 

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

Комментарии:

1. вау! Большое вам спасибо за ваш ответ! Мне нужно некоторое время, чтобы проверить производительность.