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