utiltools/X86MAC64/cuda/include/thrust/system/cuda/detail/get_value.h

   1 /*
   2  *  Copyright 2008-2012 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/system/cuda/detail/execution_policy.h>
  21 #include <thrust/system/cuda/detail/assign_value.h>
  22 #include <thrust/detail/raw_pointer_cast.h>
  23 #include <thrust/iterator/iterator_traits.h>
  24
  25 namespace thrust
  26 {
  27 namespace system
  28 {
  29 namespace cuda
  30 {
  31 namespace detail
  32 {
  33
  34
  35 namespace
  36 {
  37
  38
  39 template<typename DerivedPolicy, typename Pointer>
  40 inline __host__ __device__
  41   typename thrust::iterator_value<Pointer>::type
  42     get_value_msvc2005_war(execution_policy<DerivedPolicy> &exec, Pointer ptr)
  43 {
  44   typedef typename thrust::iterator_value<Pointer>::type result_type;
  45
  46   // XXX war nvbugs/881631
  47   struct war_nvbugs_881631
  48   {
  49     __host__ inline static result_type host_path(execution_policy<DerivedPolicy> &exec, Pointer ptr)
  50     {
  51       // when called from host code, implement with assign_value
  52       // note that this requires a type with default constructor
  53       result_type result;
  54
  55       thrust::host_system_tag host_tag;
  56       cross_system<thrust::host_system_tag, DerivedPolicy> systems(host_tag, exec);
  57       assign_value(systems, &result, ptr);
  58
  59       return result;
  60     }
  61
  62     __device__ inline static result_type device_path(execution_policy<DerivedPolicy> &, Pointer ptr)
  63     {
  64       // when called from device code, just do simple deref
  65       return *thrust::raw_pointer_cast(ptr);
  66     }
  67   };
  68
  69 #ifndef __CUDA_ARCH__
  70   return war_nvbugs_881631::host_path(exec, ptr);
  71 #else
  72   return war_nvbugs_881631::device_path(exec, ptr);
  73 #endif // __CUDA_ARCH__
  74 } // end get_value_msvc2005_war()
  75
  76
  77 } // end anon namespace
  78
  79
  80 template<typename DerivedPolicy, typename Pointer>
  81 inline __host__ __device__
  82   typename thrust::iterator_value<Pointer>::type
  83     get_value(execution_policy<DerivedPolicy> &exec, Pointer ptr)
  84 {
  85   return get_value_msvc2005_war(exec,ptr);
  86 } // end get_value()
  87
  88
  89 } // end detail
  90 } // end cuda
  91 } // end system
  92 } // end thrust
  93