#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. вау! Большое вам спасибо за ваш ответ! Мне нужно некоторое время, чтобы проверить производительность.