1 /*
2  *  Copyright 2008-2013 NVIDIA Corporation
3  *
4  *  Licensed under the Apache License, Version 2.0 (the "License");
5  *  you may not use this file except in compliance with the License.
6  *  You may obtain a copy of the License at
7  *
8  *      http://www.apache.org/licenses/LICENSE-2.0
9  *
10  *  Unless required by applicable law or agreed to in writing, software
11  *  distributed under the License is distributed on an "AS IS" BASIS,
12  *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13  *  See the License for the specific language governing permissions and
14  *  limitations under the License.
15  */
16 
17 #pragma once
18 
19 #include <thrust/tuple.h>
20 #include <thrust/detail/tuple_meta_transform.h>
21 
22 namespace thrust
23 {
24 
25 namespace detail
26 {
27 
28 template<typename Tuple,
29          template<typename> class UnaryMetaFunction,
30          typename UnaryFunction,
31          unsigned int sz = thrust::tuple_size<Tuple>::value>
32   struct tuple_transform_functor;
33 
34 
35 template<typename Tuple,
36          template<typename> class UnaryMetaFunction,
37          typename UnaryFunction>
38   struct tuple_transform_functor<Tuple,UnaryMetaFunction,UnaryFunction,0>
39 {
40   static __host__
41   typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type
42   do_it_on_the_host(const Tuple &, UnaryFunction)
43   {
44     return thrust::null_type();
45   }
46 
47   static __host__ __device__
48   typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type
49   do_it_on_the_host_or_device(const Tuple &, UnaryFunction)
50   {
51     return thrust::null_type();
52   }
53 };
54 
55 
56 template<typename Tuple,
57          template<typename> class UnaryMetaFunction,
58          typename UnaryFunction>
59   struct tuple_transform_functor<Tuple,UnaryMetaFunction,UnaryFunction,1>
60 {
61   static __host__
62   typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type
63   do_it_on_the_host(const Tuple &t, UnaryFunction f)
64   {
65     typedef typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type XfrmTuple;
66 
67     return XfrmTuple(f(thrust::get<0>(t)));
68   }
69 
70   static __host__ __device__
71   typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type
72   do_it_on_the_host_or_device(const Tuple &t, UnaryFunction f)
73   {
74     typedef typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type XfrmTuple;
75 
76     return XfrmTuple(f(thrust::get<0>(t)));
77   }
78 };
79 
80 
81 template<typename Tuple,
82          template<typename> class UnaryMetaFunction,
83          typename UnaryFunction>
84   struct tuple_transform_functor<Tuple,UnaryMetaFunction,UnaryFunction,2>
85 {
86   static __host__
87   typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type
88   do_it_on_the_host(const Tuple &t, UnaryFunction f)
89   {
90     typedef typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type XfrmTuple;
91 
92     return XfrmTuple(f(thrust::get<0>(t)),
93                      f(thrust::get<1>(t)));
94   }
95 
96   static __host__ __device__
97   typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type
98   do_it_on_the_host_or_device(const Tuple &t, UnaryFunction f)
99   {
100     typedef typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type XfrmTuple;
101 
102     return XfrmTuple(f(thrust::get<0>(t)),
103                      f(thrust::get<1>(t)));
104   }
105 };
106 
107 
108 template<typename Tuple,
109          template<typename> class UnaryMetaFunction,
110          typename UnaryFunction>
111   struct tuple_transform_functor<Tuple,UnaryMetaFunction,UnaryFunction,3>
112 {
113   static __host__
114   typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type
115   do_it_on_the_host(const Tuple &t, UnaryFunction f)
116   {
117     typedef typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type XfrmTuple;
118 
119     return XfrmTuple(f(thrust::get<0>(t)),
120                      f(thrust::get<1>(t)),
121                      f(thrust::get<2>(t)));
122   }
123 
124   static __host__ __device__
125   typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type
126   do_it_on_the_host_or_device(const Tuple &t, UnaryFunction f)
127   {
128     typedef typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type XfrmTuple;
129 
130     return XfrmTuple(f(thrust::get<0>(t)),
131                      f(thrust::get<1>(t)),
132                      f(thrust::get<2>(t)));
133   }
134 };
135 
136 
137 template<typename Tuple,
138          template<typename> class UnaryMetaFunction,
139          typename UnaryFunction>
140   struct tuple_transform_functor<Tuple,UnaryMetaFunction,UnaryFunction,4>
141 {
142   static __host__
143   typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type
144   do_it_on_the_host(const Tuple &t, UnaryFunction f)
145   {
146     typedef typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type XfrmTuple;
147 
148     return XfrmTuple(f(thrust::get<0>(t)),
149                      f(thrust::get<1>(t)),
150                      f(thrust::get<2>(t)),
151                      f(thrust::get<3>(t)));
152   }
153 
154   static __host__ __device__
155   typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type
156   do_it_on_the_host_or_device(const Tuple &t, UnaryFunction f)
157   {
158     typedef typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type XfrmTuple;
159 
160     return XfrmTuple(f(thrust::get<0>(t)),
161                      f(thrust::get<1>(t)),
162                      f(thrust::get<2>(t)),
163                      f(thrust::get<3>(t)));
164   }
165 };
166 
167 
168 template<typename Tuple,
169          template<typename> class UnaryMetaFunction,
170          typename UnaryFunction>
171   struct tuple_transform_functor<Tuple,UnaryMetaFunction,UnaryFunction,5>
172 {
173   static __host__
174   typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type
175   do_it_on_the_host(const Tuple &t, UnaryFunction f)
176   {
177     typedef typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type XfrmTuple;
178 
179     return XfrmTuple(f(thrust::get<0>(t)),
180                      f(thrust::get<1>(t)),
181                      f(thrust::get<2>(t)),
182                      f(thrust::get<3>(t)),
183                      f(thrust::get<4>(t)));
184   }
185 
186   static __host__ __device__
187   typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type
188   do_it_on_the_host_or_device(const Tuple &t, UnaryFunction f)
189   {
190     typedef typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type XfrmTuple;
191 
192     return XfrmTuple(f(thrust::get<0>(t)),
193                      f(thrust::get<1>(t)),
194                      f(thrust::get<2>(t)),
195                      f(thrust::get<3>(t)),
196                      f(thrust::get<4>(t)));
197   }
198 };
199 
200 
201 template<typename Tuple,
202          template<typename> class UnaryMetaFunction,
203          typename UnaryFunction>
204   struct tuple_transform_functor<Tuple,UnaryMetaFunction,UnaryFunction,6>
205 {
206   static __host__
207   typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type
208   do_it_on_the_host(const Tuple &t, UnaryFunction f)
209   {
210     typedef typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type XfrmTuple;
211 
212     return XfrmTuple(f(thrust::get<0>(t)),
213                      f(thrust::get<1>(t)),
214                      f(thrust::get<2>(t)),
215                      f(thrust::get<3>(t)),
216                      f(thrust::get<4>(t)),
217                      f(thrust::get<5>(t)));
218   }
219 
220   static __host__ __device__
221   typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type
222   do_it_on_the_host_or_device(const Tuple &t, UnaryFunction f)
223   {
224     typedef typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type XfrmTuple;
225 
226     return XfrmTuple(f(thrust::get<0>(t)),
227                      f(thrust::get<1>(t)),
228                      f(thrust::get<2>(t)),
229                      f(thrust::get<3>(t)),
230                      f(thrust::get<4>(t)),
231                      f(thrust::get<5>(t)));
232   }
233 };
234 
235 
236 template<typename Tuple,
237          template<typename> class UnaryMetaFunction,
238          typename UnaryFunction>
239   struct tuple_transform_functor<Tuple,UnaryMetaFunction,UnaryFunction,7>
240 {
241   static __host__
242   typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type
243   do_it_on_the_host(const Tuple &t, UnaryFunction f)
244   {
245     typedef typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type XfrmTuple;
246 
247     return XfrmTuple(f(thrust::get<0>(t)),
248                      f(thrust::get<1>(t)),
249                      f(thrust::get<2>(t)),
250                      f(thrust::get<3>(t)),
251                      f(thrust::get<4>(t)),
252                      f(thrust::get<5>(t)),
253                      f(thrust::get<6>(t)));
254   }
255 
256   static __host__ __device__
257   typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type
258   do_it_on_the_host_or_device(const Tuple &t, UnaryFunction f)
259   {
260     typedef typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type XfrmTuple;
261 
262     return XfrmTuple(f(thrust::get<0>(t)),
263                      f(thrust::get<1>(t)),
264                      f(thrust::get<2>(t)),
265                      f(thrust::get<3>(t)),
266                      f(thrust::get<4>(t)),
267                      f(thrust::get<5>(t)),
268                      f(thrust::get<6>(t)));
269   }
270 };
271 
272 
273 template<typename Tuple,
274          template<typename> class UnaryMetaFunction,
275          typename UnaryFunction>
276   struct tuple_transform_functor<Tuple,UnaryMetaFunction,UnaryFunction,8>
277 {
278   static __host__
279   typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type
280   do_it_on_the_host(const Tuple &t, UnaryFunction f)
281   {
282     typedef typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type XfrmTuple;
283 
284     return XfrmTuple(f(thrust::get<0>(t)),
285                      f(thrust::get<1>(t)),
286                      f(thrust::get<2>(t)),
287                      f(thrust::get<3>(t)),
288                      f(thrust::get<4>(t)),
289                      f(thrust::get<5>(t)),
290                      f(thrust::get<6>(t)),
291                      f(thrust::get<7>(t)));
292   }
293 
294   static __host__ __device__
295   typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type
296   do_it_on_the_host_or_device(const Tuple &t, UnaryFunction f)
297   {
298     typedef typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type XfrmTuple;
299 
300     return XfrmTuple(f(thrust::get<0>(t)),
301                      f(thrust::get<1>(t)),
302                      f(thrust::get<2>(t)),
303                      f(thrust::get<3>(t)),
304                      f(thrust::get<4>(t)),
305                      f(thrust::get<5>(t)),
306                      f(thrust::get<6>(t)),
307                      f(thrust::get<7>(t)));
308   }
309 };
310 
311 
312 template<typename Tuple,
313          template<typename> class UnaryMetaFunction,
314          typename UnaryFunction>
315   struct tuple_transform_functor<Tuple,UnaryMetaFunction,UnaryFunction,9>
316 {
317   static __host__
318   typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type
319   do_it_on_the_host(const Tuple &t, UnaryFunction f)
320   {
321     typedef typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type XfrmTuple;
322 
323     return XfrmTuple(f(thrust::get<0>(t)),
324                      f(thrust::get<1>(t)),
325                      f(thrust::get<2>(t)),
326                      f(thrust::get<3>(t)),
327                      f(thrust::get<4>(t)),
328                      f(thrust::get<5>(t)),
329                      f(thrust::get<6>(t)),
330                      f(thrust::get<7>(t)),
331                      f(thrust::get<8>(t)));
332   }
333 
334   static __host__ __device__
335   typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type
336   do_it_on_the_host_or_device(const Tuple &t, UnaryFunction f)
337   {
338     typedef typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type XfrmTuple;
339 
340     return XfrmTuple(f(thrust::get<0>(t)),
341                      f(thrust::get<1>(t)),
342                      f(thrust::get<2>(t)),
343                      f(thrust::get<3>(t)),
344                      f(thrust::get<4>(t)),
345                      f(thrust::get<5>(t)),
346                      f(thrust::get<6>(t)),
347                      f(thrust::get<7>(t)),
348                      f(thrust::get<8>(t)));
349   }
350 };
351 
352 
353 template<typename Tuple,
354          template<typename> class UnaryMetaFunction,
355          typename UnaryFunction>
356   struct tuple_transform_functor<Tuple,UnaryMetaFunction,UnaryFunction,10>
357 {
358   static __host__
359   typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type
360   do_it_on_the_host(const Tuple &t, UnaryFunction f)
361   {
362     typedef typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type XfrmTuple;
363 
364     return XfrmTuple(f(thrust::get<0>(t)),
365                      f(thrust::get<1>(t)),
366                      f(thrust::get<2>(t)),
367                      f(thrust::get<3>(t)),
368                      f(thrust::get<4>(t)),
369                      f(thrust::get<5>(t)),
370                      f(thrust::get<6>(t)),
371                      f(thrust::get<7>(t)),
372                      f(thrust::get<8>(t)),
373                      f(thrust::get<9>(t)));
374   }
375 
376   static __host__ __device__
377   typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type
378   do_it_on_the_host_or_device(const Tuple &t, UnaryFunction f)
379   {
380     typedef typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type XfrmTuple;
381 
382     return XfrmTuple(f(thrust::get<0>(t)),
383                      f(thrust::get<1>(t)),
384                      f(thrust::get<2>(t)),
385                      f(thrust::get<3>(t)),
386                      f(thrust::get<4>(t)),
387                      f(thrust::get<5>(t)),
388                      f(thrust::get<6>(t)),
389                      f(thrust::get<7>(t)),
390                      f(thrust::get<8>(t)),
391                      f(thrust::get<9>(t)));
392   }
393 };
394 
395 
396 template<template<typename> class UnaryMetaFunction,
397          typename Tuple,
398          typename UnaryFunction>
399 typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type
400 tuple_host_transform(const Tuple &t, UnaryFunction f)
401 {
402   return tuple_transform_functor<Tuple,UnaryMetaFunction,UnaryFunction>::do_it_on_the_host(t,f);
403 }
404 
405 template<template<typename> class UnaryMetaFunction,
406          typename Tuple,
407          typename UnaryFunction>
408 typename tuple_meta_transform<Tuple,UnaryMetaFunction>::type
409 __host__ __device__
410 tuple_host_device_transform(const Tuple &t, UnaryFunction f)
411 {
412   return tuple_transform_functor<Tuple,UnaryMetaFunction,UnaryFunction>::do_it_on_the_host_or_device(t,f);
413 }
414 
415 } // end detail
416 
417 } // end thrust
418 
419