File size: 2,990 Bytes
be11144
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
#include <thrust/device_vector.h>
#include <thrust/copy.h>
#include <thrust/scan.h>
#include <iostream>

// BinaryPredicate for the head flag segment representation
// equivalent to thrust::not2(thrust::project2nd<int,int>()));
template <typename HeadFlagType>
struct head_flag_predicate 
    : public thrust::binary_function<HeadFlagType,HeadFlagType,bool>
{
    __host__ __device__
    bool operator()(HeadFlagType, HeadFlagType right) const
    {
        return !right;
    }
};

template <typename Vector>
void print(const Vector& v)
{
  for(size_t i = 0; i < v.size(); i++)
    std::cout << v[i] << " ";
  std::cout << "\n";
}

int main(void)
{
    int keys[]   = {0,0,0,1,1,2,2,2,2,3,4,4,5,5,5};  // segments represented with keys
    int flags[]  = {1,0,0,1,0,1,0,0,0,1,1,0,1,0,0};  // segments represented with head flags
    int values[] = {2,2,2,2,2,2,2,2,2,2,2,2,2,2,2};  // values corresponding to each key

    int N = sizeof(keys) / sizeof(int); // number of elements

    // copy input data to device
    thrust::device_vector<int> d_keys  (keys,   keys   + N);
    thrust::device_vector<int> d_flags (flags,  flags  + N);
    thrust::device_vector<int> d_values(values, values + N);
    
    // allocate storage for output
    thrust::device_vector<int> d_output(N);

    // inclusive scan using keys
    thrust::inclusive_scan_by_key
      (d_keys.begin(), d_keys.end(),
       d_values.begin(),
       d_output.begin());
   
    std::cout << "Inclusive Segmented Scan w/ Key Sequence\n";
    std::cout << " keys          : ";  print(d_keys);
    std::cout << " input values  : ";  print(d_values);
    std::cout << " output values : ";  print(d_output);
    
    // inclusive scan using head flags
    thrust::inclusive_scan_by_key
      (d_flags.begin(), d_flags.end(),
       d_values.begin(), 
       d_output.begin(),
       head_flag_predicate<int>());
    
    std::cout << "\nInclusive Segmented Scan w/ Head Flag Sequence\n";
    std::cout << " head flags    : ";  print(d_flags);
    std::cout << " input values  : ";  print(d_values);
    std::cout << " output values : ";  print(d_output);
    
    // exclusive scan using keys
    thrust::exclusive_scan_by_key
      (d_keys.begin(), d_keys.end(),
       d_values.begin(),
       d_output.begin());
   
    std::cout << "\nExclusive Segmented Scan w/ Key Sequence\n";
    std::cout << " keys          : ";  print(d_keys);
    std::cout << " input values  : ";  print(d_values);
    std::cout << " output values : ";  print(d_output);
    
    // exclusive scan using head flags
    thrust::exclusive_scan_by_key
      (d_flags.begin(), d_flags.end(),
       d_values.begin(), 
       d_output.begin(),
       0,
       head_flag_predicate<int>());
    
    std::cout << "\nExclusive Segmented Scan w/ Head Flag Sequence\n";
    std::cout << " head flags    : ";  print(d_flags);
    std::cout << " input values  : ";  print(d_values);
    std::cout << " output values : ";  print(d_output);


    return 0;
}