← Back to branch summary

~ubuntu-branches/ubuntu/utopic/libthrust/utopic

« back to all changes in this revision

Viewing changes to system/cuda/detail/get_value.h

  • Committer: Package Import Robot
  • Author(s): Andreas Beckmann
  • Date: 2013-07-10 12:57:33 UTC
  • mfrom: (1.1.4)
  • Revision ID: package-import@ubuntu.com-20130710125733-my19jic71sqsabaj
Tags: 1.7.0-1
http://bugs.debian.org/715362
* New upstream release.  (Closes: #715362)
* Update watch file.

Show diffs side-by-side

added added

removed removed

Lines of Context:
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