1 #include <thrust/detail/raw_reference_cast.h>
2 #include <thrust/device_vector.h>
3 #include <thrust/sequence.h>
4 #include <thrust/fill.h>
5 #include <iostream>
6
7 // This example illustrates how to use the raw_reference_cast to convert
8 // system-specific reference wrappers into native references.
9 //
10 // Using iterators in the manner described here is generally discouraged.
11 // Users should only resort to this technique if there is no viable
12 // implemention of a given operation in terms of Thrust algorithms.
13 // For example this particular example is better solved with thrust::copy,
14 // which is safer and potentially faster. Only use this approach after all
15 // safer alternatives have been exhausted.
16 //
17 // When a Thrust iterator is referenced (e.g. *iter) the result is not
18 // a native or "raw" reference like int& or float&. Instead,
19 // the result is a type such as thrust::system::cuda::reference<int>
20 // or thrust::system::tbb::reference<float>, depending on the system
21 // to which the data belongs. These reference wrappers are necessary
22 // to make expressions like *iter1 = *iter2; work correctly when
23 // iter1 and iter2 refer to data in different memory spaces on
24 // heterogenous systems.
25 //
26 // The raw_reference_cast function essentially strips away the system-specific
27 // meta-data so it should only be used when the code is guaranteed to be
28 // executed within an appropriate context.
29
30
31 __host__ __device__
assign_reference_to_reference(int & x,int & y)32 void assign_reference_to_reference(int& x, int& y)
33 {
34 y = x;
35 }
36
37 __host__ __device__
assign_value_to_reference(int x,int & y)38 void assign_value_to_reference(int x, int& y)
39 {
40 y = x;
41 }
42
43 template <typename InputIterator,
44 typename OutputIterator>
45 struct copy_iterators
46 {
47 InputIterator input;
48 OutputIterator output;
49
copy_iteratorscopy_iterators50 copy_iterators(InputIterator input, OutputIterator output)
51 : input(input), output(output)
52 {}
53
54 __host__ __device__
operator ()copy_iterators55 void operator()(int i)
56 {
57 InputIterator in = input + i;
58 OutputIterator out = output + i;
59
60 // invalid - reference<int> is not convertible to int&
61 // assign_reference_to_reference(*in, *out);
62
63 // valid - reference<int> explicitly converted to int&
64 assign_reference_to_reference(thrust::raw_reference_cast(*in), thrust::raw_reference_cast(*out));
65
66 // valid - since reference<int> is convertible to int
67 assign_value_to_reference(*in, thrust::raw_reference_cast(*out));
68 }
69 };
70
71 template <typename Vector>
print(const std::string & name,const Vector & v)72 void print(const std::string& name, const Vector& v)
73 {
74 typedef typename Vector::value_type T;
75
76 std::cout << name << ": ";
77 thrust::copy(v.begin(), v.end(), std::ostream_iterator<T>(std::cout, " "));
78 std::cout << "\n";
79 }
80
main(void)81 int main(void)
82 {
83 typedef thrust::device_vector<int> Vector;
84 typedef Vector::iterator Iterator;
85 typedef thrust::device_system_tag System;
86
87 size_t N = 5;
88
89 // allocate device memory
90 Vector A(N);
91 Vector B(N);
92
93 // initialize A and B
94 thrust::sequence(A.begin(), A.end());
95 thrust::fill(B.begin(), B.end(), 0);
96
97 std::cout << "Before A->B Copy" << std::endl;
98 print("A", A);
99 print("B", B);
100
101 // note: we must specify the System to ensure correct execution
102 thrust::for_each(thrust::counting_iterator<int,System>(0),
103 thrust::counting_iterator<int,System>(N),
104 copy_iterators<Iterator,Iterator>(A.begin(), B.begin()));
105
106 std::cout << "After A->B Copy" << std::endl;
107 print("A", A);
108 print("B", B);
109
110 return 0;
111 }
112
113