IteratorFromArrayPortal.h 5.65 KB
Newer Older
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
//============================================================================
//  Copyright (c) Kitware, Inc.
//  All rights reserved.
//  See LICENSE.txt for details.
//  This software is distributed WITHOUT ANY WARRANTY; without even
//  the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
//  PURPOSE.  See the above copyright notice for more information.
//
//  Copyright 2014 Sandia Corporation.
//  Copyright 2014 UT-Battelle, LLC.
//  Copyright 2014 Los Alamos National Security.
//
//  Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
//  the U.S. Government retains certain rights in this software.
//
//  Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
//  Laboratory (LANL), the U.S. Government retains certain rights in
//  this software.
//============================================================================
#ifndef vtk_m_exec_cuda_internal_IteratorFromArrayPortal_h
#define vtk_m_exec_cuda_internal_IteratorFromArrayPortal_h

#include <vtkm/Types.h>
#include <vtkm/Pair.h>
#include <vtkm/internal/ExportMacros.h>

// Disable warnings we check vtkm for but Thrust does not.
#if defined(__GNUC__) || defined(____clang__)
#pragma GCC diagnostic push
#pragma GCC diagnostic ignored "-Wshadow"
#pragma GCC diagnostic ignored "-Wunused-parameter"
#pragma GCC diagnostic ignored "-Wconversion"
#endif // gcc || clang

#include <thrust/functional.h>
#include <thrust/iterator/iterator_facade.h>
#include <thrust/system/cuda/execution_policy.h>

#if defined(__GNUC__) || defined(____clang__)
#pragma GCC diagnostic pop
#endif // gcc || clang

namespace vtkm {
namespace exec {
namespace cuda {
namespace internal {

template<class ArrayPortalType>
struct PortalValue
{
  typedef typename ArrayPortalType::ValueType ValueType;

  VTKM_EXEC_CONT_EXPORT
  PortalValue(const ArrayPortalType &portal, vtkm::Id index)
    : Portal(portal), Index(index) {  }

  VTKM_EXEC_EXPORT
  void Swap( PortalValue<ArrayPortalType> &rhs ) throw()
  {
    //we need use the explicit type not a proxy temp object
    //A proxy temp object would point to the same underlying data structure
    //and would not hold the old value of *this once *this was set to rhs.
    const ValueType aValue = *this;
    *this = rhs;
    rhs = aValue;
  }

  VTKM_EXEC_EXPORT
  PortalValue<ArrayPortalType> &operator=(
    const PortalValue<ArrayPortalType> &rhs)
  {
    this->Portal.Set(this->Index, rhs.Portal.Get(rhs.Index));
    return *this;
  }

  VTKM_EXEC_EXPORT
  ValueType operator=(const ValueType& value)
  {
    this->Portal.Set(this->Index, value);
    return value;
  }

  VTKM_EXEC_EXPORT
  operator ValueType(void) const
  {
    return this->Portal.Get(this->Index);
  }

  const ArrayPortalType& Portal;
  vtkm::Id Index;
};

template<class ArrayPortalType>
class IteratorFromArrayPortal : public
    ::thrust::iterator_facade<
      IteratorFromArrayPortal<ArrayPortalType>,
      typename ArrayPortalType::ValueType,
      ::thrust::system::cuda::tag,
      ::thrust::random_access_traversal_tag,
      PortalValue<ArrayPortalType>,
      vtkm::Id>
{
public:

  VTKM_EXEC_CONT_EXPORT
  IteratorFromArrayPortal()
    : Portal(), Index(0) { }

  VTKM_CONT_EXPORT
  explicit IteratorFromArrayPortal(const ArrayPortalType &portal,
                                   vtkm::Id index = 0)
    : Portal(portal), Index(index) {  }

  VTKM_EXEC_EXPORT
  PortalValue<ArrayPortalType>
116
  operator[](std::ptrdiff_t idx) const //NEEDS to be signed
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
  {
    return PortalValue<ArrayPortalType>(this->Portal,
           this->Index + static_cast<vtkm::Id>(idx) );
  }

private:
  ArrayPortalType Portal;
  vtkm::Id Index;

  // Implementation for ::thrust iterator_facade
  friend class ::thrust::iterator_core_access;

  VTKM_EXEC_EXPORT
  PortalValue<ArrayPortalType> dereference() const
  {
    return PortalValue<ArrayPortalType>(this->Portal,
           this->Index);
  }

  VTKM_EXEC_EXPORT
  bool equal(const IteratorFromArrayPortal<ArrayPortalType> &other) const
  {
    // Technically, we should probably check that the portals are the same,
    // but the portal interface does not specify an equal operator.  It is
    // by its nature undefined what happens when comparing iterators from
    // different portals anyway.
    return (this->Index == other.Index);
  }

  VTKM_EXEC_CONT_EXPORT
  void increment()
  {
    this->Index++;
  }

  VTKM_EXEC_CONT_EXPORT
  void decrement()
  {
    this->Index--;
  }

  VTKM_EXEC_CONT_EXPORT
  void advance(vtkm::Id delta)
  {
    this->Index += delta;
  }

  VTKM_EXEC_CONT_EXPORT
  vtkm::Id
  distance_to(const IteratorFromArrayPortal<ArrayPortalType> &other) const
  {
    // Technically, we should probably check that the portals are the same,
    // but the portal interface does not specify an equal operator.  It is
    // by its nature undefined what happens when comparing iterators from
    // different portals anyway.
    return other.Index - this->Index;
  }
};

}
}
}
} //namespace vtkm::exec::cuda::internal

//So for the unary_transform_functor and binary_transform_functor inside
//of thrust, they verify that the index they are storing into is a reference
//instead of a value, so that the contents actually are written to global memory.
//
//But for vtk-m we pass in facade objects, which are passed by value, but
//must be treated as references. So do to do that properly we need to specialize
//is_non_const_reference to state a PortalValue by value is valid for writing
namespace thrust
{
namespace detail
{

  template< typename T > struct is_non_const_reference;

  template< typename T >
  struct is_non_const_reference< vtkm::exec::cuda::internal::PortalValue<T> >
    : thrust::detail::true_type
  { };

}
}

#endif //vtk_m_exec_cuda_internal_IteratorFromArrayPortal_h