Уменьшите несколько блоков равной длины, которые расположены в большом векторе. Использование CUDA

Question

Я ищу быстрый способ уменьшить несколько блоков равной длины , которые расположены как большой вектор. У меня есть N подмассивов (смежных элементов), которые расположены в одном большом массиве. каждый дополнительный массив имеет фиксированный размер: k. , поэтому размер всей массива: N * K

Что я делаю, так это называть ядро ​​N раз. в каждый раз, когда он вычисляет сокращение подмассива как следующие: я итерацию по всем подмассивам, содержащихся в большом векторе:

for(i=0;i<N;i++){ 
     thrust::device_vector<float> Vec(subarray, subarray+k); 
     float sum = thrust::reduce(Vec.begin(), Vec.end(), (float)0, thrust::plus<float>()); 
     printf("sum %f\n",sum); 
} 

для чистой CUDA я сделаю это так (псевдо-код):

for(i=0;i<N;i++){ 
     reduction_kernel(subarray) 

     } 

У вас есть другое решение для выполнения сокращения смежных подмассивов в один раз? используя чистый CUDA или Thrust

Answer 1

То, о чем вы просите, является сегментированным сокращением. Это можно сделать с помощью thrust::reduce_by_key В дополнение к вашему вектору данных длины N * K нам понадобится «ключевой» вектор, который определяет каждый сегмент - сегменты не должны быть одного размера, если ключевой вектор различает сегменты так:

data: 1 3 2 3 1 4 2 3 2 1 4 2 ... 
keys: 0 0 0 1 1 1 0 0 0 3 3 3 ... 
seg: 0 0 0 1 1 1 2 2 2 3 3 3 ... 

ключи очертить новый сегмент в любое время ключевые изменения последовательности (заметим, что у меня есть два отдельных сегмента в приведенном выше примере, которые очерчены, используя тот же ключ - тяги не группируйте такие сегменты вместе, но обрабатывает их отдельно, потому что есть одно или несколько промежуточных значений ключа, которые различны). У вас на самом деле нет этих данных, но для скорости и эффективности, так как ваши сегменты имеют одинаковую длину, мы можем произвести необходимую последовательность клавиш «на лету», используя комбинацию тяги fancy iterators.

Причудливые итераторы будет сочетать:

1. производить линейную последовательность 0 1 2 3 ... (через counting_iterator)
2. делят каждый член линейной последовательности по K, длины сегмента (через transform_iterator). Я использую thrust placeholder methodology, поэтому мне не нужно писать функтор для итератора преобразования.

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

Вот обработанный пример:

$ cat t1282.cu 
#include <thrust/reduce.h> 
#include <thrust/device_vector.h> 
#include <thrust/iterator/transform_iterator.h> 
#include <thrust/iterator/counting_iterator.h> 
#include <thrust/iterator/discard_iterator.h> 
#include <thrust/copy.h> 
#include <thrust/execution_policy.h> 
#include <iostream> 

const int N = 1000; // sequences 
const int K = 100; // length of sequence 
typedef int mytype; 

using namespace thrust::placeholders; 

int main(){ 

    thrust::device_vector<mytype> data(N*K, 1); 
    thrust::device_vector<mytype> sums(N); 
    thrust::reduce_by_key(thrust::device, thrust::make_transform_iterator(thrust::counting_iterator<int>(0), _1/K), thrust::make_transform_iterator(thrust::counting_iterator<int>(N*K), _1/K), data.begin(), thrust::discard_iterator<int>(), sums.begin()); 
    // just display the first 10 results 
    thrust::copy_n(sums.begin(), 10, std::ostream_iterator<mytype>(std::cout, ",")); 
    std::cout << std::endl; 
} 

$ nvcc -arch=sm_35 -o t1282 t1282.cu 
$ ./t1282 
100,100,100,100,100,100,100,100,100,100, 
$