10 #ifndef vtk_m_exec_cuda_internal_ArrayPortalFromThrust_h
11 #define vtk_m_exec_cuda_internal_ArrayPortalFromThrust_h
17 #include <type_traits>
21 #include <thrust/system/cuda/memory.h>
37 template <
typename T>
struct UseScalarTextureLoad :
public std::false_type {};
38 template <
typename T>
struct UseVecTextureLoads :
public std::false_type {};
39 template <
typename T>
struct UseMultipleScalarTextureLoads :
public std::false_type {};
43 template <>
struct UseScalarTextureLoad<const
vtkm::
UInt8> : std::true_type {};
44 template <>
struct UseScalarTextureLoad<const
vtkm::
Int16> : std::true_type {};
45 template <>
struct UseScalarTextureLoad<const
vtkm::
UInt16> : std::true_type {};
46 template <>
struct UseScalarTextureLoad<const
vtkm::
Int32> : std::true_type {};
47 template <>
struct UseScalarTextureLoad<const
vtkm::
UInt32> : std::true_type {};
48 template <>
struct UseScalarTextureLoad<const
vtkm::
Float32> : std::true_type {};
49 template <>
struct UseScalarTextureLoad<const
vtkm::
Float64> : std::true_type {};
53 template <>
struct UseVecTextureLoads<const
vtkm::
Vec2i_32> : std::true_type {};
54 template <>
struct UseVecTextureLoads<const
vtkm::
Vec2ui_32> : std::true_type {};
55 template <>
struct UseVecTextureLoads<const
vtkm::
Vec2f_32> : std::true_type {};
56 template <>
struct UseVecTextureLoads<const
vtkm::
Vec2f_64> : std::true_type {};
58 template <>
struct UseVecTextureLoads<const
vtkm::
Vec4i_32> : std::true_type {};
59 template <>
struct UseVecTextureLoads<const
vtkm::
Vec4ui_32> : std::true_type {};
60 template <>
struct UseVecTextureLoads<const
vtkm::
Vec4f_32> : std::true_type {};
67 template <>
struct UseMultipleScalarTextureLoads<const
vtkm::
Vec2ui_8> : std::true_type {};
68 template <>
struct UseMultipleScalarTextureLoads<const
vtkm::
Vec2i_16> : std::true_type {};
69 template <>
struct UseMultipleScalarTextureLoads<const
vtkm::
Vec2ui_16> : std::true_type {};
70 template <>
struct UseMultipleScalarTextureLoads<const
vtkm::
Vec2i_64> : std::true_type {};
71 template <>
struct UseMultipleScalarTextureLoads<const
vtkm::
Vec2ui_64> : std::true_type {};
73 template <>
struct UseMultipleScalarTextureLoads<const
vtkm::
Vec3ui_8> : std::true_type {};
74 template <>
struct UseMultipleScalarTextureLoads<const
vtkm::
Vec3i_16> : std::true_type {};
75 template <>
struct UseMultipleScalarTextureLoads<const
vtkm::
Vec3ui_16> : std::true_type {};
76 template <>
struct UseMultipleScalarTextureLoads<const
vtkm::
Vec3i_32> : std::true_type {};
77 template <>
struct UseMultipleScalarTextureLoads<const
vtkm::
Vec3ui_32> : std::true_type {};
78 template <>
struct UseMultipleScalarTextureLoads<const
vtkm::
Vec3f_32> : std::true_type {};
79 template <>
struct UseMultipleScalarTextureLoads<const
vtkm::
Vec3f_64> : std::true_type {};
81 template <>
struct UseMultipleScalarTextureLoads<const
vtkm::
Vec4ui_8> : std::true_type {};
82 template <>
struct UseMultipleScalarTextureLoads<const
vtkm::
Vec4i_16> : std::true_type {};
83 template <>
struct UseMultipleScalarTextureLoads<const
vtkm::
Vec4ui_16> : std::true_type {};
84 template <>
struct UseMultipleScalarTextureLoads<const
vtkm::
Vec4i_64> : std::true_type {};
85 template <>
struct UseMultipleScalarTextureLoads<const
vtkm::
Vec4ui_64> : std::true_type {};
86 template <>
struct UseMultipleScalarTextureLoads<const
vtkm::
Vec4f_64> : std::true_type {};
90 template <
typename T,
typename Enable =
void>
91 struct load_through_texture
95 __device__
static T
get(
const T*
const data) {
return *data; }
101 template <
typename T>
102 struct load_through_texture<T, typename std::enable_if<UseScalarTextureLoad<const T>::value>::type>
107 __device__
static T
get(
const T*
const data)
109 #if __CUDA_ARCH__ >= 350
119 template <
typename T>
120 struct load_through_texture<T, typename std::enable_if<UseVecTextureLoads<const T>::value>::type>
124 __device__
static T
get(
const T*
const data)
126 #if __CUDA_ARCH__ >= 350
136 const int2 temp = __ldg((
const int2*)data);
142 const uint2 temp = __ldg((
const uint2*)data);
148 const int4 temp = __ldg((
const int4*)data);
154 const uint4 temp = __ldg((
const uint4*)data);
160 const float2 temp = __ldg((
const float2*)data);
166 const float4 temp = __ldg((
const float4*)data);
172 const double2 temp = __ldg((
const double2*)data);
178 template <
typename T>
179 struct load_through_texture<
181 typename std::enable_if<UseMultipleScalarTextureLoads<const T>::value>::type>
185 using NonConstT =
typename std::remove_const<T>::type;
187 __device__
static T
get(
const T*
const data)
189 #if __CUDA_ARCH__ >= 350
197 __device__
static T getAs(
const T*
const data)
201 using ComponentType =
typename T::ComponentType;
202 const ComponentType* recasted_data = (
const ComponentType*)(data);
207 result[i] = __ldg(recasted_data + i);
213 class ArrayPortalFromThrustBase
220 template <
typename T>
221 class ArrayPortalFromThrust :
public ArrayPortalFromThrustBase
225 using IteratorType = T*;
226 using difference_type = std::ptrdiff_t;
231 ArrayPortalFromThrust(IteratorType begin, IteratorType end)
232 : BeginIterator(begin)
241 template <
typename OtherT>
242 VTKM_EXEC_CONT ArrayPortalFromThrust(
const ArrayPortalFromThrust<OtherT>& src)
243 : BeginIterator(src.GetIteratorBegin())
244 , EndIterator(src.GetIteratorEnd())
252 return static_cast<vtkm::Id>((this->EndIterator - this->BeginIterator));
258 return *(this->BeginIterator +
static_cast<difference_type
>(index));
262 void Set(
vtkm::Id index, ValueType value)
const
264 *(this->BeginIterator +
static_cast<difference_type
>(index)) = value;
268 IteratorType GetIteratorBegin()
const {
return this->BeginIterator; }
271 IteratorType GetIteratorEnd()
const {
return this->EndIterator; }
274 IteratorType BeginIterator;
275 IteratorType EndIterator;
278 template <
typename T>
279 class ConstArrayPortalFromThrust :
public ArrayPortalFromThrustBase
283 using IteratorType =
const T*;
284 using difference_type = std::ptrdiff_t;
287 : BeginIterator(
nullptr)
288 , EndIterator(
nullptr)
293 ConstArrayPortalFromThrust(IteratorType begin, IteratorType end)
294 : BeginIterator(begin)
306 ConstArrayPortalFromThrust(
const ArrayPortalFromThrust<T>& src)
307 : BeginIterator(src.GetIteratorBegin())
308 , EndIterator(src.GetIteratorEnd())
316 return static_cast<vtkm::Id>((this->EndIterator - this->BeginIterator));
324 #ifdef VTKM_CUDA_DEVICE_PASS
338 #if !(defined(VTKM_MSVC) && defined(VTKM_CUDA))
345 IteratorType GetIteratorBegin()
const {
return this->BeginIterator; }
348 IteratorType GetIteratorEnd()
const {
return this->EndIterator; }
351 IteratorType BeginIterator;
352 IteratorType EndIterator;
359 #endif //vtk_m_exec_cuda_internal_ArrayPortalFromThrust_h