Cuda Thrust vorheriges Element in einem Segment erhalten

Question

ich einen Vektor von Werten und einen Vektor von Tasten habe (was ein Segment)

v = [1, 1, 1, 2, 3, 5, 6] 
k = [0, 0, 1, 1, 0, 2, 2] 

Für jedes Element Ich mag sein vorheriges Element wissen (im gleichen Segment). Es könnte Wert sein oder ein Index im ursprünglichen Vektor spielt keine Rolle.

So Ergebnis sollte (bei Wert) sein

r = [nan, 1, nan, 1, 1, nan, 5] 

Sie ein beliebiges Element anstelle von nan verwenden können, um keine Fragen Teil eines Algorithmus bleibt.

Wahrscheinlich kann ich es mit exklusiven segmentierten Scan und max Operation statt sum archivieren. So zwei Fragen:

1. Ist meine Vorgehensweise korrekt?
2. Ist eine elegantere oder effizientere Lösung?

Answer 1

die gewünschte Funktionalität kann unter Verwendung der folgenden Schritte durchgeführt werden:

1. sort v durch k nebeneinander gleiche Schlüsselwerten zu erhalten; Dies muss durch stable_sort_by_key erfolgen, wie Sie das "vorherige" Element abrufen möchten, so dass die Reihenfolge unter Elementen mit gleichen Schlüsseln beibehalten werden muss.

2. gelten die folgende Transformation auf die sortierten Daten:

if (previous element has the same key) then return value of previous element else return -1

---

Der folgende Code implementiert diese Schritte:

#include <cstdint> 
#include <iostream> 

#include <thrust/host_vector.h> 
#include <thrust/device_vector.h> 
#include <thrust/sort.h> 
#include <thrust/transform.h> 
#include <thrust/iterator/counting_iterator.h> 
#include <thrust/iterator/zip_iterator.h> 

#define PRINTER(name) print(#name, (name)) 
template <template <typename...> class V, typename T, typename ...Args> 
void print(const char* name, const V<T,Args...> & v) 
{ 
    std::cout << name << ":\t"; 
    thrust::copy(v.begin(), v.end(), std::ostream_iterator<T>(std::cout, "\t")); 
    std::cout << std::endl; 
} 
template<typename... Iterators> 
__host__ __device__ 
thrust::zip_iterator<thrust::tuple<Iterators...>> zip(Iterators... its) 
{ 
    return thrust::make_zip_iterator(thrust::make_tuple(its...)); 
} 

template <typename IteratorType, typename Integer> 
struct prev_value 
{ 
    prev_value(IteratorType first) : first(first){} 

    template <typename Tuple> 
    __host__ __device__ 
    Integer operator()(const Tuple& t) 
    { 
     const auto& index = thrust::get<0>(t); 
     const auto& previousValue = thrust::get<1>(t); 

     Integer result = -1; 
     const auto& currentKey = *(first+index); 
     const auto& previousKey = *(first+index-1); 
     if(currentKey == previousKey) 
     { 
      result = previousValue; 
     } 

     return result; 
    } 

    IteratorType first; 
}; 

template <typename Integer, typename IteratorType> 
prev_value<IteratorType, Integer> make_prev_value(IteratorType first) 
{ 
    return prev_value<IteratorType, Integer>(first); 
} 


int main(int argc, char** argv) 
{ 
    using Integer = std::int32_t; 
    using HostVec = thrust::host_vector<Integer>; 
    using DeviceVec = thrust::device_vector<Integer>; 

    Integer v[] = {1, 1, 1, 2, 3, 5, 6}; 
    Integer k[] = {0, 0, 1, 1, 0, 2, 2}; 

    Integer size = sizeof(k)/sizeof(k[0]); 

    HostVec h_k(k, k+size); 
    HostVec h_v(v, v+size); 

    // copy data to device 
    DeviceVec d_k = h_k; 
    DeviceVec d_v = h_v; 

    std::cout << "---- input data ----" << std::endl; 
    PRINTER(d_k);  
    PRINTER(d_v); 

    thrust::stable_sort_by_key(d_k.begin(), d_k.end(), d_v.begin()); 
    std::cout << "---- after sorting ----" << std::endl; 
    PRINTER(d_k);  
    PRINTER(d_v); 

    DeviceVec d_r(size, -1); 
    auto op = make_prev_value<Integer>(d_k.begin()); 
    thrust::transform(zip(thrust::make_counting_iterator(Integer(1)), d_v.begin()), 
         zip(thrust::make_counting_iterator(size), d_v.end()), 
         d_r.begin()+1, 
         op); 
    std::cout << "---- result ----" << std::endl; 
    PRINTER(d_r); 

    return 0; 
} 

output:

---- input data ---- 
d_k: 0 0 1 1 0 2 2 
d_v: 1 1 1 2 3 5 6 
---- after sorting ---- 
d_k: 0 0 0 1 1 2 2 
d_v: 1 1 3 1 2 5 6 
---- result ---- 
d_r: -1 1 1 -1 1 -1 5