Reduzieren Sie mehrere Blöcke gleicher Länge, die in einem großen Vektor angeordnet sind Mit CUDA

Question

Ich bin auf der Suche nach einem schnellen Weg, um mehrere Blöcke gleicher Länge , die als ein großer Vektor angeordnet sind zu reduzieren. Ich habe N Subarrays (zusammenhängende Elemente), die in einem großen Array angeordnet sind. Jedes Unterfeld hat eine feste Größe: k. so ist die Größe des gesamten Arrays: N * K

Was ich tue, ist den Kernel n mal aufrufen. in jeder Zeit berechnet sie die Reduktion des Sub-Array wie folgt: ich über alle Sub-Arrays, die in dem großen Vektor iteriert wird:

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); 
} 

für reine CUDA tun ich es so (Pseudocode):

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

     } 

Haben Sie eine andere Lösung, um die zusammenhängenden Subarrays auf einmal zu reduzieren? mit reinem CUDA oder Thrust

Answer 1

Was Sie fordern, ist eine segmentierte Reduktion. Dies kann in Schub durchgeführt werden mit thrust::reduce_by_key Zusätzlich zu Ihrem Datenvektor der Länge N * K, wir brauchen einen "Schlüssel" Vektor, der jedes Segment definiert - die Segmente müssen nicht die gleiche Größe haben, solange die Schlüsselvektor unterscheidet Segmente wie folgt:

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

die Schlüssel ein neues Segment umreißen jederzeit die Tastenfolge Änderungen (beachten Sie, dass ich zwei separate Segmente in dem obigen Beispiel habe, die den gleichen Schlüssel verwenden abgegrenzt - Schub nicht Gruppieren Sie solche Segmente zusammen, behandeln Sie sie jedoch separat, da es 1 oder mehr intervenierende Schlüsselwerte gibt, die unterschiedlich sind). Sie haben diese Daten nicht wirklich, aber für Geschwindigkeit und Effizienz, da Ihre Segmente von gleicher Länge sind, können wir die erforderliche Tastenfolge "on the fly" mit einer Kombination von Schub fancy iterators produzieren.

Die ausgefallenen Iteratoren werden kombinieren:

1. eine lineare Sequenz 0 1 2 3 ... (via counting_iterator)
2. divide jedes Mitglied der linearen Sequenz von K erzeugen, wobei die Segmentlänge (via transform_iterator). Ich verwende hier thrust placeholder methodology, also muss ich keinen Funktor für den Transformationsiterator schreiben.

Dies wird die erforderliche Segmentschlüsselsequenz erzeugen.

ist hier ein ausgearbeitetes Beispiel:

$ 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, 
$