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/detail/config.h> 20 #include <thrust/detail/raw_pointer_cast.h> 21 #include <thrust/detail/type_traits/has_nested_type.h> 22 #include <thrust/detail/type_traits.h> 23 #include <thrust/detail/tuple_transform.h> 24 #include <thrust/iterator/detail/tuple_of_iterator_references.h> 25 26 27 // the order of declarations and definitions in this file is totally goofy 28 // this header defines raw_reference_cast, which has a few overloads towards the bottom of the file 29 // raw_reference_cast depends on metafunctions such as is_unwrappable and raw_reference 30 // we need to be sure that these metafunctions are completely defined (including specializations) before they are instantiated by raw_reference_cast 31 32 namespace thrust 33 { 34 namespace detail 35 { 36 37 38 __THRUST_DEFINE_HAS_NESTED_TYPE(is_wrapped_reference, wrapped_reference_hint) 39 40 41 // wrapped reference-like things which aren't strictly wrapped references 42 // (e.g. tuples of wrapped references) are considered unwrappable 43 template<typename T> 44 struct is_unwrappable 45 : is_wrapped_reference<T> 46 {}; 47 48 49 // specialize is_unwrappable 50 // a tuple is_unwrappable if any of its elements is_unwrappable 51 template< 52 typename T0, typename T1, typename T2, 53 typename T3, typename T4, typename T5, 54 typename T6, typename T7, typename T8, 55 typename T9 56 > 57 struct is_unwrappable< 58 thrust::tuple<T0,T1,T2,T3,T4,T5,T6,T7,T8,T9> 59 > 60 : or_< 61 is_unwrappable<T0>, 62 is_unwrappable<T1>, 63 is_unwrappable<T2>, 64 is_unwrappable<T3>, 65 is_unwrappable<T4>, 66 is_unwrappable<T5>, 67 is_unwrappable<T6>, 68 is_unwrappable<T7>, 69 is_unwrappable<T8>, 70 is_unwrappable<T9> 71 > 72 {}; 73 74 75 // specialize is_unwrappable 76 // a tuple_of_iterator_references is_unwrappable if any of its elements is_unwrappable 77 template< 78 typename T0, typename T1, typename T2, 79 typename T3, typename T4, typename T5, 80 typename T6, typename T7, typename T8, 81 typename T9 82 > 83 struct is_unwrappable< 84 thrust::detail::tuple_of_iterator_references<T0,T1,T2,T3,T4,T5,T6,T7,T8,T9> 85 > 86 : or_< 87 is_unwrappable<T0>, 88 is_unwrappable<T1>, 89 is_unwrappable<T2>, 90 is_unwrappable<T3>, 91 is_unwrappable<T4>, 92 is_unwrappable<T5>, 93 is_unwrappable<T6>, 94 is_unwrappable<T7>, 95 is_unwrappable<T8>, 96 is_unwrappable<T9> 97 > 98 {}; 99 100 101 template<typename T, typename Result = void> 102 struct enable_if_unwrappable 103 : enable_if< 104 is_unwrappable<T>::value, 105 Result 106 > 107 {}; 108 109 110 namespace raw_reference_detail 111 { 112 113 114 template<typename T, typename Enable = void> 115 struct raw_reference_impl 116 : add_reference<T> 117 {}; 118 119 120 template<typename T> 121 struct raw_reference_impl< 122 T, 123 typename thrust::detail::enable_if< 124 is_wrapped_reference< 125 typename remove_cv<T>::type 126 >::value 127 >::type 128 > 129 { 130 typedef typename add_reference< 131 typename pointer_element<typename T::pointer>::type 132 >::type type; 133 }; 134 135 136 } // end raw_reference_detail 137 138 139 template<typename T> 140 struct raw_reference : 141 raw_reference_detail::raw_reference_impl<T> 142 {}; 143 144 145 namespace raw_reference_detail 146 { 147 148 // unlike raw_reference, 149 // raw_reference_tuple_helper needs to return a value 150 // when it encounters one, rather than a reference 151 // upon encountering tuple, recurse 152 // 153 // we want the following behavior: 154 // 1. T -> T 155 // 2. T& -> T& 156 // 3. null_type -> null_type 157 // 4. reference<T> -> T& 158 // 5. tuple_of_iterator_references<T> -> tuple_of_iterator_references<raw_reference_tuple_helper<T>::type> 159 160 161 // wrapped references are unwrapped using raw_reference, otherwise, return T 162 template<typename T> 163 struct raw_reference_tuple_helper 164 : eval_if< 165 is_unwrappable< 166 typename remove_cv<T>::type 167 >::value, 168 raw_reference<T>, 169 identity_<T> 170 > 171 {}; 172 173 174 // recurse on tuples 175 template < 176 typename T0, typename T1, typename T2, 177 typename T3, typename T4, typename T5, 178 typename T6, typename T7, typename T8, 179 typename T9 180 > 181 struct raw_reference_tuple_helper< 182 thrust::tuple<T0,T1,T2,T3,T4,T5,T6,T7,T8,T9> 183 > 184 { 185 typedef thrust::tuple< 186 typename raw_reference_tuple_helper<T0>::type, 187 typename raw_reference_tuple_helper<T1>::type, 188 typename raw_reference_tuple_helper<T2>::type, 189 typename raw_reference_tuple_helper<T3>::type, 190 typename raw_reference_tuple_helper<T4>::type, 191 typename raw_reference_tuple_helper<T5>::type, 192 typename raw_reference_tuple_helper<T6>::type, 193 typename raw_reference_tuple_helper<T7>::type, 194 typename raw_reference_tuple_helper<T8>::type, 195 typename raw_reference_tuple_helper<T9>::type 196 > type; 197 }; 198 199 200 template < 201 typename T0, typename T1, typename T2, 202 typename T3, typename T4, typename T5, 203 typename T6, typename T7, typename T8, 204 typename T9 205 > 206 struct raw_reference_tuple_helper< 207 thrust::detail::tuple_of_iterator_references<T0,T1,T2,T3,T4,T5,T6,T7,T8,T9> 208 > 209 { 210 typedef thrust::detail::tuple_of_iterator_references< 211 typename raw_reference_tuple_helper<T0>::type, 212 typename raw_reference_tuple_helper<T1>::type, 213 typename raw_reference_tuple_helper<T2>::type, 214 typename raw_reference_tuple_helper<T3>::type, 215 typename raw_reference_tuple_helper<T4>::type, 216 typename raw_reference_tuple_helper<T5>::type, 217 typename raw_reference_tuple_helper<T6>::type, 218 typename raw_reference_tuple_helper<T7>::type, 219 typename raw_reference_tuple_helper<T8>::type, 220 typename raw_reference_tuple_helper<T9>::type 221 > type; 222 }; 223 224 225 } // end raw_reference_detail 226 227 228 // a couple of specializations of raw_reference for tuples follow 229 230 231 // if a tuple "tuple_type" is_unwrappable, 232 // then the raw_reference of tuple_type is a tuple of its members' raw_references 233 // else the raw_reference of tuple_type is tuple_type & 234 template < 235 typename T0, typename T1, typename T2, 236 typename T3, typename T4, typename T5, 237 typename T6, typename T7, typename T8, 238 typename T9 239 > 240 struct raw_reference< 241 thrust::tuple<T0,T1,T2,T3,T4,T5,T6,T7,T8,T9> 242 > 243 { 244 private: 245 typedef thrust::tuple<T0,T1,T2,T3,T4,T5,T6,T7,T8,T9> tuple_type; 246 247 public: 248 typedef typename eval_if< 249 is_unwrappable<tuple_type>::value, 250 raw_reference_detail::raw_reference_tuple_helper<tuple_type>, 251 add_reference<tuple_type> 252 >::type type; 253 }; 254 255 256 template < 257 typename T0, typename T1, typename T2, 258 typename T3, typename T4, typename T5, 259 typename T6, typename T7, typename T8, 260 typename T9 261 > 262 struct raw_reference< 263 thrust::detail::tuple_of_iterator_references<T0,T1,T2,T3,T4,T5,T6,T7,T8,T9> 264 > 265 { 266 private: 267 typedef detail::tuple_of_iterator_references<T0,T1,T2,T3,T4,T5,T6,T7,T8,T9> tuple_type; 268 269 public: 270 typedef typename raw_reference_detail::raw_reference_tuple_helper<tuple_type>::type type; 271 272 // XXX figure out why is_unwrappable seems to be broken for tuple_of_iterator_references 273 //typedef typename eval_if< 274 // is_unwrappable<tuple_type>::value, 275 // raw_reference_detail::raw_reference_tuple_helper<tuple_type>, 276 // add_reference<tuple_type> 277 //>::type type; 278 }; 279 280 281 } // end detail 282 283 284 // provide declarations of raw_reference_cast's overloads for raw_reference_caster below 285 template<typename T> 286 __host__ __device__ 287 typename detail::raw_reference<T>::type 288 raw_reference_cast(T &ref); 289 290 291 template<typename T> 292 __host__ __device__ 293 typename detail::raw_reference<const T>::type 294 raw_reference_cast(const T &ref); 295 296 297 template< 298 typename T0, typename T1, typename T2, 299 typename T3, typename T4, typename T5, 300 typename T6, typename T7, typename T8, 301 typename T9 302 > 303 __host__ __device__ 304 typename detail::enable_if_unwrappable< 305 thrust::detail::tuple_of_iterator_references<T0,T1,T2,T3,T4,T5,T6,T7,T8,T9>, 306 typename detail::raw_reference< 307 thrust::detail::tuple_of_iterator_references<T0,T1,T2,T3,T4,T5,T6,T7,T8,T9> 308 >::type 309 >::type 310 raw_reference_cast(thrust::detail::tuple_of_iterator_references<T0,T1,T2,T3,T4,T5,T6,T7,T8,T9> t); 311 312 313 namespace detail 314 { 315 316 317 struct raw_reference_caster 318 { 319 template<typename T> 320 __host__ __device__ 321 typename detail::raw_reference<T>::type operator()(T &ref) 322 { 323 return thrust::raw_reference_cast(ref); 324 } 325 326 template<typename T> 327 __host__ __device__ 328 typename detail::raw_reference<const T>::type operator()(const T &ref) 329 { 330 return thrust::raw_reference_cast(ref); 331 } 332 333 template< 334 typename T0, typename T1, typename T2, 335 typename T3, typename T4, typename T5, 336 typename T6, typename T7, typename T8, 337 typename T9 338 > 339 __host__ __device__ 340 typename detail::raw_reference< 341 thrust::detail::tuple_of_iterator_references<T0,T1,T2,T3,T4,T5,T6,T7,T8,T9> 342 >::type 343 operator()(thrust::detail::tuple_of_iterator_references<T0,T1,T2,T3,T4,T5,T6,T7,T8,T9> t, 344 typename enable_if< 345 is_unwrappable<thrust::detail::tuple_of_iterator_references<T0,T1,T2,T3,T4,T5,T6,T7,T8,T9> >::value 346 >::type * = 0) 347 { 348 return thrust::raw_reference_cast(t); 349 } 350 }; // end raw_reference_caster 351 352 353 } // end detail 354 355 356 template<typename T> 357 __host__ __device__ 358 typename detail::raw_reference<T>::type 359 raw_reference_cast(T &ref) 360 { 361 return *thrust::raw_pointer_cast(&ref); 362 } // end raw_reference_cast 363 364 365 template<typename T> 366 __host__ __device__ 367 typename detail::raw_reference<const T>::type 368 raw_reference_cast(const T &ref) 369 { 370 return *thrust::raw_pointer_cast(&ref); 371 } // end raw_reference_cast 372 373 374 template< 375 typename T0, typename T1, typename T2, 376 typename T3, typename T4, typename T5, 377 typename T6, typename T7, typename T8, 378 typename T9 379 > 380 __host__ __device__ 381 typename detail::enable_if_unwrappable< 382 thrust::detail::tuple_of_iterator_references<T0,T1,T2,T3,T4,T5,T6,T7,T8,T9>, 383 typename detail::raw_reference< 384 thrust::detail::tuple_of_iterator_references<T0,T1,T2,T3,T4,T5,T6,T7,T8,T9> 385 >::type 386 >::type 387 raw_reference_cast(thrust::detail::tuple_of_iterator_references<T0,T1,T2,T3,T4,T5,T6,T7,T8,T9> t) 388 { 389 thrust::detail::raw_reference_caster f; 390 391 // note that we pass raw_reference_tuple_helper, not raw_reference as the unary metafunction 392 // the different way that raw_reference_tuple_helper unwraps tuples is important 393 return thrust::detail::tuple_host_device_transform<detail::raw_reference_detail::raw_reference_tuple_helper>(t, f); 394 } // end raw_reference_cast 395 396 397 } // end thrust 398 399