mirror of
https://gitlab.kitware.com/vtk/vtk-m
synced 2024-09-08 13:23:51 +00:00
iVTK-m CUDA backend doesn't use thrust::cuda::pointer any more.
This was removed as CUDA 9.0 on MSVC has issues where CUB/Thrust would fail to compile when given these types.
This commit is contained in:
parent
bfc66729ff
commit
22ea58335a
@ -127,19 +127,13 @@ struct ExecutionPortalFactoryBasic<T, DeviceAdapterTagCuda>
|
||||
VTKM_CONT
|
||||
static PortalType CreatePortal(ValueType* start, ValueType* end)
|
||||
{
|
||||
using ThrustPointerT = thrust::system::cuda::pointer<ValueType>;
|
||||
ThrustPointerT startThrust(start);
|
||||
ThrustPointerT endThrust(end);
|
||||
return PortalType(startThrust, endThrust);
|
||||
return PortalType(start, end);
|
||||
}
|
||||
|
||||
VTKM_CONT
|
||||
static PortalConstType CreatePortalConst(const ValueType* start, const ValueType* end)
|
||||
{
|
||||
using ThrustPointerT = thrust::system::cuda::pointer<const ValueType>;
|
||||
ThrustPointerT startThrust(start);
|
||||
ThrustPointerT endThrust(end);
|
||||
return PortalConstType(startThrust, endThrust);
|
||||
return PortalConstType(start, end);
|
||||
}
|
||||
};
|
||||
|
||||
|
@ -24,21 +24,16 @@
|
||||
#include <vtkm/cont/ErrorBadAllocation.h>
|
||||
#include <vtkm/cont/Storage.h>
|
||||
|
||||
// Disable warnings we check vtkm for but Thrust does not.
|
||||
VTKM_THIRDPARTY_PRE_INCLUDE
|
||||
#include <thrust/copy.h>
|
||||
#include <thrust/device_malloc_allocator.h>
|
||||
#include <thrust/system/cuda/vector.h>
|
||||
|
||||
#include <thrust/system/cuda/execution_policy.h>
|
||||
|
||||
VTKM_THIRDPARTY_POST_INCLUDE
|
||||
|
||||
#include <vtkm/cont/cuda/ErrorCuda.h>
|
||||
#include <vtkm/cont/cuda/internal/CudaAllocator.h>
|
||||
#include <vtkm/cont/cuda/internal/ThrustExceptionHandler.h>
|
||||
#include <vtkm/exec/cuda/internal/ArrayPortalFromThrust.h>
|
||||
|
||||
VTKM_THIRDPARTY_PRE_INCLUDE
|
||||
#include <thrust/copy.h>
|
||||
#include <thrust/device_ptr.h>
|
||||
VTKM_THIRDPARTY_POST_INCLUDE
|
||||
|
||||
#include <limits>
|
||||
|
||||
namespace vtkm
|
||||
@ -53,15 +48,13 @@ namespace internal
|
||||
/// \c ArrayManagerExecutionThrustDevice provides an implementation for a \c
|
||||
/// ArrayManagerExecution class for a thrust device adapter that is designed
|
||||
/// for the cuda backend which has separate memory spaces for host and device.
|
||||
/// This implementation contains a thrust::system::cuda::pointer to contain the
|
||||
/// data.
|
||||
template <typename T, class StorageTag>
|
||||
class ArrayManagerExecutionThrustDevice
|
||||
{
|
||||
public:
|
||||
using ValueType = T;
|
||||
using PointerType = typename thrust::system::cuda::pointer<ValueType>;
|
||||
using difference_type = typename PointerType::difference_type;
|
||||
using PointerType = T*;
|
||||
using difference_type = std::ptrdiff_t;
|
||||
|
||||
using StorageType = vtkm::cont::internal::Storage<ValueType, StorageTag>;
|
||||
|
||||
@ -71,9 +64,9 @@ public:
|
||||
VTKM_CONT
|
||||
ArrayManagerExecutionThrustDevice(StorageType* storage)
|
||||
: Storage(storage)
|
||||
, Begin(static_cast<ValueType*>(nullptr))
|
||||
, End(static_cast<ValueType*>(nullptr))
|
||||
, Capacity(static_cast<ValueType*>(nullptr))
|
||||
, Begin(nullptr)
|
||||
, End(nullptr)
|
||||
, Capacity(nullptr)
|
||||
{
|
||||
}
|
||||
|
||||
@ -83,10 +76,7 @@ public:
|
||||
/// Returns the size of the array.
|
||||
///
|
||||
VTKM_CONT
|
||||
vtkm::Id GetNumberOfValues() const
|
||||
{
|
||||
return static_cast<vtkm::Id>(this->End.get() - this->Begin.get());
|
||||
}
|
||||
vtkm::Id GetNumberOfValues() const { return static_cast<vtkm::Id>(this->End - this->Begin); }
|
||||
|
||||
/// Allocates the appropriate size of the array and copies the given data
|
||||
/// into the array.
|
||||
@ -144,14 +134,13 @@ public:
|
||||
PortalType PrepareForOutput(vtkm::Id numberOfValues)
|
||||
{
|
||||
// Can we reuse the existing buffer?
|
||||
vtkm::Id curCapacity = this->Begin.get() != nullptr
|
||||
? static_cast<vtkm::Id>(this->Capacity.get() - this->Begin.get())
|
||||
: 0;
|
||||
vtkm::Id curCapacity =
|
||||
this->Begin != nullptr ? static_cast<vtkm::Id>(this->Capacity - this->Begin) : 0;
|
||||
|
||||
// Just mark a new end if we don't need to increase the allocation:
|
||||
if (curCapacity >= numberOfValues)
|
||||
{
|
||||
this->End = PointerType(this->Begin.get() + static_cast<difference_type>(numberOfValues));
|
||||
this->End = this->Begin + static_cast<difference_type>(numberOfValues);
|
||||
|
||||
return PortalType(this->Begin, this->End);
|
||||
}
|
||||
@ -173,9 +162,8 @@ public:
|
||||
// Attempt to allocate:
|
||||
try
|
||||
{
|
||||
ValueType* tmp =
|
||||
this->Begin =
|
||||
static_cast<ValueType*>(vtkm::cont::cuda::internal::CudaAllocator::Allocate(bufferSize));
|
||||
this->Begin = PointerType(tmp);
|
||||
}
|
||||
catch (const std::exception& error)
|
||||
{
|
||||
@ -184,7 +172,7 @@ public:
|
||||
throw vtkm::cont::ErrorBadAllocation(err.str());
|
||||
}
|
||||
|
||||
this->Capacity = PointerType(this->Begin.get() + static_cast<difference_type>(numberOfValues));
|
||||
this->Capacity = this->Begin + static_cast<difference_type>(numberOfValues);
|
||||
this->End = this->Capacity;
|
||||
|
||||
return PortalType(this->Begin, this->End);
|
||||
@ -206,8 +194,9 @@ public:
|
||||
storage->Allocate(this->GetNumberOfValues());
|
||||
try
|
||||
{
|
||||
::thrust::copy(
|
||||
this->Begin, this->End, vtkm::cont::ArrayPortalToIteratorBegin(storage->GetPortal()));
|
||||
::thrust::copy(thrust::cuda::pointer<ValueType>(this->Begin),
|
||||
thrust::cuda::pointer<ValueType>(this->End),
|
||||
vtkm::cont::ArrayPortalToIteratorBegin(storage->GetPortal()));
|
||||
}
|
||||
catch (...)
|
||||
{
|
||||
@ -221,22 +210,21 @@ public:
|
||||
{
|
||||
// The operation will succeed even if this assertion fails, but this
|
||||
// is still supposed to be a precondition to Shrink.
|
||||
VTKM_ASSERT(this->Begin.get() != nullptr &&
|
||||
this->Begin.get() + numberOfValues <= this->End.get());
|
||||
VTKM_ASSERT(this->Begin != nullptr && this->Begin + numberOfValues <= this->End);
|
||||
|
||||
this->End = PointerType(this->Begin.get() + static_cast<difference_type>(numberOfValues));
|
||||
this->End = this->Begin + static_cast<difference_type>(numberOfValues);
|
||||
}
|
||||
|
||||
/// Frees all memory.
|
||||
///
|
||||
VTKM_CONT void ReleaseResources()
|
||||
{
|
||||
if (this->Begin.get() != nullptr)
|
||||
if (this->Begin != nullptr)
|
||||
{
|
||||
vtkm::cont::cuda::internal::CudaAllocator::Free(this->Begin.get());
|
||||
this->Begin = PointerType(static_cast<ValueType*>(nullptr));
|
||||
this->End = PointerType(static_cast<ValueType*>(nullptr));
|
||||
this->Capacity = PointerType(static_cast<ValueType*>(nullptr));
|
||||
vtkm::cont::cuda::internal::CudaAllocator::Free(this->Begin);
|
||||
this->Begin = nullptr;
|
||||
this->End = nullptr;
|
||||
this->Capacity = nullptr;
|
||||
}
|
||||
}
|
||||
|
||||
@ -258,7 +246,7 @@ private:
|
||||
this->PrepareForOutput(this->Storage->GetNumberOfValues());
|
||||
::thrust::copy(vtkm::cont::ArrayPortalToIteratorBegin(this->Storage->GetPortalConst()),
|
||||
vtkm::cont::ArrayPortalToIteratorEnd(this->Storage->GetPortalConst()),
|
||||
this->Begin);
|
||||
thrust::cuda::pointer<ValueType>(this->Begin));
|
||||
}
|
||||
catch (...)
|
||||
{
|
||||
|
@ -53,14 +53,13 @@ VTKM_THIRDPARTY_PRE_INCLUDE
|
||||
#include <thrust/binary_search.h>
|
||||
#include <thrust/copy.h>
|
||||
#include <thrust/count.h>
|
||||
#include <thrust/iterator/counting_iterator.h>
|
||||
#include <thrust/scan.h>
|
||||
#include <thrust/sort.h>
|
||||
#include <thrust/system/cpp/memory.h>
|
||||
#include <thrust/system/cuda/vector.h>
|
||||
#include <thrust/unique.h>
|
||||
#include <vtkm/exec/cuda/internal/ExecutionPolicy.h>
|
||||
|
||||
#include <thrust/iterator/counting_iterator.h>
|
||||
#include <thrust/system/cuda/execution_policy.h>
|
||||
VTKM_THIRDPARTY_POST_INCLUDE
|
||||
|
||||
#include <atomic>
|
||||
@ -218,8 +217,7 @@ private:
|
||||
OutputPortal output,
|
||||
UnaryPredicate unary_predicate)
|
||||
{
|
||||
using IteratorType = typename detail::IteratorTraits<OutputPortal>::IteratorType;
|
||||
IteratorType outputBegin = IteratorBegin(output);
|
||||
auto outputBegin = IteratorBegin(output);
|
||||
|
||||
using ValueType = typename StencilPortal::ValueType;
|
||||
|
||||
@ -497,12 +495,15 @@ private:
|
||||
|
||||
try
|
||||
{
|
||||
::thrust::system::cuda::vector<ValueType> result(1);
|
||||
auto end = ::thrust::inclusive_scan(ThrustCudaPolicyPerThread,
|
||||
IteratorBegin(input),
|
||||
IteratorEnd(input),
|
||||
IteratorBegin(output),
|
||||
bop);
|
||||
return *(end - 1);
|
||||
|
||||
::thrust::copy_n(ThrustCudaPolicyPerThread, end - 1, 1, result.begin());
|
||||
return result[0];
|
||||
}
|
||||
catch (...)
|
||||
{
|
||||
|
@ -20,21 +20,8 @@
|
||||
#ifndef vtk_m_cont_cuda_internal_MakeThrustIterator_h
|
||||
#define vtk_m_cont_cuda_internal_MakeThrustIterator_h
|
||||
|
||||
#include <vtkm/Pair.h>
|
||||
#include <vtkm/Types.h>
|
||||
#include <vtkm/cont/ArrayPortalToIterators.h>
|
||||
#include <vtkm/internal/ExportMacros.h>
|
||||
|
||||
#include <vtkm/exec/cuda/internal/ArrayPortalFromThrust.h>
|
||||
#include <vtkm/exec/cuda/internal/WrappedOperators.h>
|
||||
|
||||
// Disable warnings we check vtkm for but Thrust does not.
|
||||
VTKM_THIRDPARTY_PRE_INCLUDE
|
||||
#include <thrust/functional.h>
|
||||
#include <thrust/iterator/counting_iterator.h>
|
||||
#include <thrust/iterator/transform_iterator.h>
|
||||
#include <thrust/system/cuda/memory.h>
|
||||
VTKM_THIRDPARTY_POST_INCLUDE
|
||||
#include <vtkm/exec/cuda/internal/IteratorFromArrayPortal.h>
|
||||
|
||||
namespace vtkm
|
||||
{
|
||||
@ -44,107 +31,46 @@ namespace cuda
|
||||
{
|
||||
namespace internal
|
||||
{
|
||||
namespace detail
|
||||
{
|
||||
|
||||
// Tags to specify what type of thrust iterator to use.
|
||||
struct ThrustIteratorFromArrayPortalTag
|
||||
{
|
||||
};
|
||||
struct ThrustIteratorDevicePtrTag
|
||||
{
|
||||
};
|
||||
|
||||
// Traits to help classify what thrust iterators will be used.
|
||||
template <typename IteratorType>
|
||||
struct ThrustIteratorTag
|
||||
{
|
||||
using Type = ThrustIteratorFromArrayPortalTag;
|
||||
};
|
||||
template <typename T>
|
||||
struct ThrustIteratorTag<thrust::system::cuda::pointer<T>>
|
||||
{
|
||||
using Type = ThrustIteratorDevicePtrTag;
|
||||
};
|
||||
template <typename T>
|
||||
struct ThrustIteratorTag<thrust::system::cuda::pointer<const T>>
|
||||
{
|
||||
using Type = ThrustIteratorDevicePtrTag;
|
||||
};
|
||||
|
||||
template <typename PortalType, typename Tag>
|
||||
struct IteratorChooser;
|
||||
template <typename PortalType>
|
||||
struct IteratorChooser<PortalType, detail::ThrustIteratorFromArrayPortalTag>
|
||||
{
|
||||
using Type = vtkm::exec::cuda::internal::IteratorFromArrayPortal<PortalType>;
|
||||
};
|
||||
template <typename PortalType>
|
||||
struct IteratorChooser<PortalType, detail::ThrustIteratorDevicePtrTag>
|
||||
{
|
||||
using PortalToIteratorType = vtkm::cont::ArrayPortalToIterators<PortalType>;
|
||||
|
||||
using Type = typename PortalToIteratorType::IteratorType;
|
||||
};
|
||||
|
||||
template <typename PortalType>
|
||||
struct IteratorTraits
|
||||
{
|
||||
using PortalToIteratorType = vtkm::cont::ArrayPortalToIterators<PortalType>;
|
||||
using Tag = typename detail::ThrustIteratorTag<typename PortalToIteratorType::IteratorType>::Type;
|
||||
using IteratorType = typename IteratorChooser<PortalType, Tag>::Type;
|
||||
};
|
||||
|
||||
template <typename PortalType>
|
||||
VTKM_CONT typename IteratorTraits<PortalType>::IteratorType MakeIteratorBegin(
|
||||
PortalType portal,
|
||||
detail::ThrustIteratorFromArrayPortalTag)
|
||||
{
|
||||
return vtkm::exec::cuda::internal::IteratorFromArrayPortal<PortalType>(portal);
|
||||
}
|
||||
|
||||
template <typename PortalType>
|
||||
VTKM_CONT typename IteratorTraits<PortalType>::IteratorType MakeIteratorBegin(
|
||||
PortalType portal,
|
||||
detail::ThrustIteratorDevicePtrTag)
|
||||
{
|
||||
vtkm::cont::ArrayPortalToIterators<PortalType> iterators(portal);
|
||||
return iterators.GetBegin();
|
||||
}
|
||||
|
||||
template <typename PortalType>
|
||||
VTKM_CONT typename IteratorTraits<PortalType>::IteratorType MakeIteratorEnd(
|
||||
PortalType portal,
|
||||
detail::ThrustIteratorFromArrayPortalTag)
|
||||
inline vtkm::exec::cuda::internal::IteratorFromArrayPortal<PortalType> IteratorBegin(
|
||||
const PortalType& portal)
|
||||
{
|
||||
vtkm::exec::cuda::internal::IteratorFromArrayPortal<PortalType> iterator(portal);
|
||||
::thrust::advance(iterator, static_cast<std::size_t>(portal.GetNumberOfValues()));
|
||||
return iterator;
|
||||
}
|
||||
|
||||
template <typename PortalType>
|
||||
VTKM_CONT typename IteratorTraits<PortalType>::IteratorType MakeIteratorEnd(
|
||||
PortalType portal,
|
||||
detail::ThrustIteratorDevicePtrTag)
|
||||
inline vtkm::exec::cuda::internal::IteratorFromArrayPortal<PortalType> IteratorEnd(
|
||||
const PortalType& portal)
|
||||
{
|
||||
vtkm::cont::ArrayPortalToIterators<PortalType> iterators(portal);
|
||||
return iterators.GetEnd();
|
||||
vtkm::exec::cuda::internal::IteratorFromArrayPortal<PortalType> iterator(portal);
|
||||
iterator += static_cast<std::size_t>(portal.GetNumberOfValues());
|
||||
return iterator;
|
||||
}
|
||||
|
||||
} // namespace detail
|
||||
|
||||
template <typename PortalType>
|
||||
VTKM_CONT typename detail::IteratorTraits<PortalType>::IteratorType IteratorBegin(PortalType portal)
|
||||
template <typename T>
|
||||
inline T* IteratorBegin(const vtkm::exec::cuda::internal::ArrayPortalFromThrust<T>& portal)
|
||||
{
|
||||
using IteratorTag = typename detail::IteratorTraits<PortalType>::Tag;
|
||||
return detail::MakeIteratorBegin(portal, IteratorTag());
|
||||
return portal.GetIteratorBegin();
|
||||
}
|
||||
|
||||
template <typename PortalType>
|
||||
VTKM_CONT typename detail::IteratorTraits<PortalType>::IteratorType IteratorEnd(PortalType portal)
|
||||
template <typename T>
|
||||
inline T* IteratorEnd(const vtkm::exec::cuda::internal::ArrayPortalFromThrust<T>& portal)
|
||||
{
|
||||
using IteratorTag = typename detail::IteratorTraits<PortalType>::Tag;
|
||||
return detail::MakeIteratorEnd(portal, IteratorTag());
|
||||
return portal.GetIteratorEnd();
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
inline const T* IteratorBegin(
|
||||
const vtkm::exec::cuda::internal::ConstArrayPortalFromThrust<T>& portal)
|
||||
{
|
||||
return portal.GetIteratorBegin();
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
inline const T* IteratorEnd(const vtkm::exec::cuda::internal::ConstArrayPortalFromThrust<T>& portal)
|
||||
{
|
||||
return portal.GetIteratorEnd();
|
||||
}
|
||||
}
|
||||
}
|
||||
|
@ -1070,7 +1070,6 @@ private:
|
||||
std::cout << " Reduce with 0 values." << std::endl;
|
||||
array.Shrink(0);
|
||||
vtkm::Id reduce_sum_no_values = Algorithm::Reduce(array, vtkm::Id(0));
|
||||
|
||||
VTKM_TEST_ASSERT(reduce_sum == OFFSET * ARRAY_SIZE, "Got bad sum from Reduce");
|
||||
VTKM_TEST_ASSERT(reduce_sum_with_intial_value == reduce_sum + ARRAY_SIZE,
|
||||
"Got bad sum from Reduce with initial value");
|
||||
|
@ -101,10 +101,7 @@ struct load_through_texture
|
||||
{
|
||||
static const vtkm::IdComponent WillUseTexture = 0;
|
||||
|
||||
__device__ static T get(const thrust::system::cuda::pointer<const T>& data)
|
||||
{
|
||||
return *(data.get());
|
||||
}
|
||||
__device__ static T get(const T* const data) { return *data; }
|
||||
};
|
||||
|
||||
//only load through a texture if we have sm 35 support
|
||||
@ -116,13 +113,13 @@ struct load_through_texture<T, typename std::enable_if<UseScalarTextureLoad<cons
|
||||
|
||||
static const vtkm::IdComponent WillUseTexture = 1;
|
||||
|
||||
__device__ static T get(const thrust::system::cuda::pointer<const T>& data)
|
||||
__device__ static T get(const T* const data)
|
||||
{
|
||||
#if __CUDA_ARCH__ >= 350
|
||||
// printf("__CUDA_ARCH__ UseScalarTextureLoad");
|
||||
return __ldg(data.get());
|
||||
return __ldg(data);
|
||||
#else
|
||||
return *(data.get());
|
||||
return *data;
|
||||
#endif
|
||||
}
|
||||
};
|
||||
@ -133,62 +130,55 @@ struct load_through_texture<T, typename std::enable_if<UseVecTextureLoads<const
|
||||
{
|
||||
static const vtkm::IdComponent WillUseTexture = 1;
|
||||
|
||||
__device__ static T get(const thrust::system::cuda::pointer<const T>& data)
|
||||
__device__ static T get(const T* const data)
|
||||
{
|
||||
#if __CUDA_ARCH__ >= 350
|
||||
// printf("__CUDA_ARCH__ UseVecTextureLoads");
|
||||
return getAs(data);
|
||||
#else
|
||||
return *(data.get());
|
||||
return *data;
|
||||
#endif
|
||||
}
|
||||
|
||||
__device__ static vtkm::Vec<vtkm::Int32, 2> getAs(
|
||||
const thrust::system::cuda::pointer<const vtkm::Vec<vtkm::Int32, 2>>& data)
|
||||
__device__ static vtkm::Vec<vtkm::Int32, 2> getAs(const vtkm::Vec<vtkm::Int32, 2>* const data)
|
||||
{
|
||||
const int2 temp = __ldg((const int2*)data.get());
|
||||
const int2 temp = __ldg((const int2*)data);
|
||||
return vtkm::Vec<vtkm::Int32, 2>(temp.x, temp.y);
|
||||
}
|
||||
|
||||
__device__ static vtkm::Vec<vtkm::UInt32, 2> getAs(
|
||||
const thrust::system::cuda::pointer<const vtkm::Vec<vtkm::UInt32, 2>>& data)
|
||||
__device__ static vtkm::Vec<vtkm::UInt32, 2> getAs(const vtkm::Vec<vtkm::UInt32, 2>* const data)
|
||||
{
|
||||
const uint2 temp = __ldg((const uint2*)data.get());
|
||||
const uint2 temp = __ldg((const uint2*)data);
|
||||
return vtkm::Vec<vtkm::UInt32, 2>(temp.x, temp.y);
|
||||
}
|
||||
|
||||
__device__ static vtkm::Vec<vtkm::Int32, 4> getAs(
|
||||
const thrust::system::cuda::pointer<const vtkm::Vec<vtkm::Int32, 4>>& data)
|
||||
__device__ static vtkm::Vec<vtkm::Int32, 4> getAs(const vtkm::Vec<vtkm::Int32, 4>* const data)
|
||||
{
|
||||
const int4 temp = __ldg((const int4*)data.get());
|
||||
const int4 temp = __ldg((const int4*)data);
|
||||
return vtkm::Vec<vtkm::Int32, 4>(temp.x, temp.y, temp.z, temp.w);
|
||||
}
|
||||
|
||||
__device__ static vtkm::Vec<vtkm::UInt32, 4> getAs(
|
||||
const thrust::system::cuda::pointer<const vtkm::Vec<vtkm::UInt32, 4>>& data)
|
||||
__device__ static vtkm::Vec<vtkm::UInt32, 4> getAs(const vtkm::Vec<vtkm::UInt32, 4>* const data)
|
||||
{
|
||||
const uint4 temp = __ldg((const uint4*)data.get());
|
||||
const uint4 temp = __ldg((const uint4*)data);
|
||||
return vtkm::Vec<vtkm::UInt32, 4>(temp.x, temp.y, temp.z, temp.w);
|
||||
}
|
||||
|
||||
__device__ static vtkm::Vec<vtkm::Float32, 2> getAs(
|
||||
const thrust::system::cuda::pointer<const vtkm::Vec<vtkm::Float32, 2>>& data)
|
||||
__device__ static vtkm::Vec<vtkm::Float32, 2> getAs(const vtkm::Vec<vtkm::Float32, 2>* const data)
|
||||
{
|
||||
const float2 temp = __ldg((const float2*)data.get());
|
||||
const float2 temp = __ldg((const float2*)data);
|
||||
return vtkm::Vec<vtkm::Float32, 2>(temp.x, temp.y);
|
||||
}
|
||||
|
||||
__device__ static vtkm::Vec<vtkm::Float32, 4> getAs(
|
||||
const thrust::system::cuda::pointer<const vtkm::Vec<vtkm::Float32, 4>>& data)
|
||||
__device__ static vtkm::Vec<vtkm::Float32, 4> getAs(const vtkm::Vec<vtkm::Float32, 4>* const data)
|
||||
{
|
||||
const float4 temp = __ldg((const float4*)data.get());
|
||||
const float4 temp = __ldg((const float4*)data);
|
||||
return vtkm::Vec<vtkm::Float32, 4>(temp.x, temp.y, temp.z, temp.w);
|
||||
}
|
||||
|
||||
__device__ static vtkm::Vec<vtkm::Float64, 2> getAs(
|
||||
const thrust::system::cuda::pointer<const vtkm::Vec<vtkm::Float64, 2>>& data)
|
||||
__device__ static vtkm::Vec<vtkm::Float64, 2> getAs(const vtkm::Vec<vtkm::Float64, 2>* const data)
|
||||
{
|
||||
const double2 temp = __ldg((const double2*)data.get());
|
||||
const double2 temp = __ldg((const double2*)data);
|
||||
return vtkm::Vec<vtkm::Float64, 2>(temp.x, temp.y);
|
||||
}
|
||||
};
|
||||
@ -203,22 +193,22 @@ struct load_through_texture<
|
||||
|
||||
using NonConstT = typename std::remove_const<T>::type;
|
||||
|
||||
__device__ static T get(const thrust::system::cuda::pointer<const T>& data)
|
||||
__device__ static T get(const T* const data)
|
||||
{
|
||||
#if __CUDA_ARCH__ >= 350
|
||||
// printf("__CUDA_ARCH__ UseMultipleScalarTextureLoads");
|
||||
return getAs(data);
|
||||
#else
|
||||
return *(data.get());
|
||||
return *data;
|
||||
#endif
|
||||
}
|
||||
|
||||
__device__ static T getAs(const thrust::system::cuda::pointer<const T>& data)
|
||||
__device__ static T getAs(const T* const data)
|
||||
{
|
||||
//we need to fetch each component individually
|
||||
const vtkm::IdComponent NUM_COMPONENTS = T::NUM_COMPONENTS;
|
||||
using ComponentType = typename T::ComponentType;
|
||||
const ComponentType* recasted_data = (const ComponentType*)(data.get());
|
||||
const ComponentType* recasted_data = (const ComponentType*)(data);
|
||||
NonConstT result;
|
||||
#pragma unroll
|
||||
for (vtkm::IdComponent i = 0; i < NUM_COMPONENTS; ++i)
|
||||
@ -241,13 +231,13 @@ class ArrayPortalFromThrust : public ArrayPortalFromThrustBase
|
||||
{
|
||||
public:
|
||||
using ValueType = T;
|
||||
using IteratorType = thrust::system::cuda::pointer<T>;
|
||||
using IteratorType = T*;
|
||||
using difference_type = std::ptrdiff_t;
|
||||
|
||||
VTKM_EXEC_CONT ArrayPortalFromThrust() {}
|
||||
|
||||
VTKM_CONT
|
||||
ArrayPortalFromThrust(thrust::system::cuda::pointer<T> begin,
|
||||
thrust::system::cuda::pointer<T> end)
|
||||
ArrayPortalFromThrust(IteratorType begin, IteratorType end)
|
||||
: BeginIterator(begin)
|
||||
, EndIterator(end)
|
||||
{
|
||||
@ -274,15 +264,13 @@ public:
|
||||
VTKM_EXEC_CONT
|
||||
ValueType Get(vtkm::Id index) const
|
||||
{
|
||||
using SizeType = typename ::thrust::iterator_traits<IteratorType>::difference_type;
|
||||
return *(this->BeginIterator + static_cast<SizeType>(index));
|
||||
return *(this->BeginIterator + static_cast<difference_type>(index));
|
||||
}
|
||||
|
||||
VTKM_EXEC_CONT
|
||||
void Set(vtkm::Id index, ValueType value) const
|
||||
{
|
||||
using SizeType = typename ::thrust::iterator_traits<IteratorType>::difference_type;
|
||||
*(this->BeginIterator + static_cast<SizeType>(index)) = value;
|
||||
*(this->BeginIterator + static_cast<difference_type>(index)) = value;
|
||||
}
|
||||
|
||||
VTKM_EXEC_CONT
|
||||
@ -301,13 +289,17 @@ class ConstArrayPortalFromThrust : public ArrayPortalFromThrustBase
|
||||
{
|
||||
public:
|
||||
using ValueType = T;
|
||||
using IteratorType = thrust::system::cuda::pointer<const T>;
|
||||
using IteratorType = const T*;
|
||||
using difference_type = std::ptrdiff_t;
|
||||
|
||||
VTKM_EXEC_CONT ConstArrayPortalFromThrust() {}
|
||||
VTKM_EXEC_CONT ConstArrayPortalFromThrust()
|
||||
: BeginIterator(nullptr)
|
||||
, EndIterator(nullptr)
|
||||
{
|
||||
}
|
||||
|
||||
VTKM_CONT
|
||||
ConstArrayPortalFromThrust(const thrust::system::cuda::pointer<const T> begin,
|
||||
const thrust::system::cuda::pointer<const T> end)
|
||||
ConstArrayPortalFromThrust(IteratorType begin, IteratorType end)
|
||||
: BeginIterator(begin)
|
||||
, EndIterator(end)
|
||||
{
|
||||
|
@ -44,8 +44,8 @@ struct vtkm_cuda_policy : thrust::device_execution_policy<vtkm_cuda_policy>
|
||||
template <typename T>
|
||||
__host__ __device__ void sort(
|
||||
const vtkm_cuda_policy& exec,
|
||||
thrust::system::cuda::pointer<T> first,
|
||||
thrust::system::cuda::pointer<T> last,
|
||||
T* first,
|
||||
T* last,
|
||||
vtkm::exec::cuda::internal::WrappedBinaryPredicate<T, vtkm::SortLess> comp)
|
||||
{ //sort for concrete pointers and less than op
|
||||
//this makes sure that we invoke the thrust radix sort and not merge sort
|
||||
@ -55,8 +55,8 @@ __host__ __device__ void sort(
|
||||
template <typename T, typename RandomAccessIterator>
|
||||
__host__ __device__ void sort_by_key(
|
||||
const vtkm_cuda_policy& exec,
|
||||
thrust::system::cuda::pointer<T> first,
|
||||
thrust::system::cuda::pointer<T> last,
|
||||
T* first,
|
||||
T* last,
|
||||
RandomAccessIterator values_first,
|
||||
vtkm::exec::cuda::internal::WrappedBinaryPredicate<T, vtkm::SortLess> comp)
|
||||
{ //sort for concrete pointers and less than op
|
||||
@ -68,8 +68,8 @@ __host__ __device__ void sort_by_key(
|
||||
template <typename T>
|
||||
__host__ __device__ void sort(
|
||||
const vtkm_cuda_policy& exec,
|
||||
thrust::system::cuda::pointer<T> first,
|
||||
thrust::system::cuda::pointer<T> last,
|
||||
T* first,
|
||||
T* last,
|
||||
vtkm::exec::cuda::internal::WrappedBinaryPredicate<T, ::thrust::less<T>> comp)
|
||||
{ //sort for concrete pointers and less than op
|
||||
//this makes sure that we invoke the thrust radix sort and not merge sort
|
||||
@ -79,8 +79,8 @@ __host__ __device__ void sort(
|
||||
template <typename T, typename RandomAccessIterator>
|
||||
__host__ __device__ void sort_by_key(
|
||||
const vtkm_cuda_policy& exec,
|
||||
thrust::system::cuda::pointer<T> first,
|
||||
thrust::system::cuda::pointer<T> last,
|
||||
T* first,
|
||||
T* last,
|
||||
RandomAccessIterator values_first,
|
||||
vtkm::exec::cuda::internal::WrappedBinaryPredicate<T, ::thrust::less<T>> comp)
|
||||
{ //sort for concrete pointers and less than op
|
||||
@ -92,8 +92,8 @@ __host__ __device__ void sort_by_key(
|
||||
template <typename T>
|
||||
__host__ __device__ void sort(
|
||||
const vtkm_cuda_policy& exec,
|
||||
thrust::system::cuda::pointer<T> first,
|
||||
thrust::system::cuda::pointer<T> last,
|
||||
T* first,
|
||||
T* last,
|
||||
vtkm::exec::cuda::internal::WrappedBinaryPredicate<T, vtkm::SortGreater> comp)
|
||||
{ //sort for concrete pointers and greater than op
|
||||
//this makes sure that we invoke the thrust radix sort and not merge sort
|
||||
@ -103,8 +103,8 @@ __host__ __device__ void sort(
|
||||
template <typename T, typename RandomAccessIterator>
|
||||
__host__ __device__ void sort_by_key(
|
||||
const vtkm_cuda_policy& exec,
|
||||
thrust::system::cuda::pointer<T> first,
|
||||
thrust::system::cuda::pointer<T> last,
|
||||
T* first,
|
||||
T* last,
|
||||
RandomAccessIterator values_first,
|
||||
vtkm::exec::cuda::internal::WrappedBinaryPredicate<T, vtkm::SortGreater> comp)
|
||||
{ //sort for concrete pointers and greater than op
|
||||
@ -116,8 +116,8 @@ __host__ __device__ void sort_by_key(
|
||||
template <typename T>
|
||||
__host__ __device__ void sort(
|
||||
const vtkm_cuda_policy& exec,
|
||||
thrust::system::cuda::pointer<T> first,
|
||||
thrust::system::cuda::pointer<T> last,
|
||||
T* first,
|
||||
T* last,
|
||||
vtkm::exec::cuda::internal::WrappedBinaryPredicate<T, ::thrust::greater<T>> comp)
|
||||
{ //sort for concrete pointers and greater than op
|
||||
//this makes sure that we invoke the thrust radix sort and not merge sort
|
||||
@ -127,8 +127,8 @@ __host__ __device__ void sort(
|
||||
template <typename T, typename RandomAccessIterator>
|
||||
__host__ __device__ void sort_by_key(
|
||||
const vtkm_cuda_policy& exec,
|
||||
thrust::system::cuda::pointer<T> first,
|
||||
thrust::system::cuda::pointer<T> last,
|
||||
T* first,
|
||||
T* last,
|
||||
RandomAccessIterator values_first,
|
||||
vtkm::exec::cuda::internal::WrappedBinaryPredicate<T, ::thrust::greater<T>> comp)
|
||||
{ //sort for concrete pointers and greater than op
|
||||
@ -174,8 +174,8 @@ template <typename T,
|
||||
typename BinaryFunction>
|
||||
__host__ __device__::thrust::pair<OutputIterator1, OutputIterator2> reduce_by_key(
|
||||
const vtkm_cuda_policy& exec,
|
||||
thrust::system::cuda::pointer<T> keys_first,
|
||||
thrust::system::cuda::pointer<T> keys_last,
|
||||
T* keys_first,
|
||||
T* keys_last,
|
||||
InputIterator2 values_first,
|
||||
OutputIterator1 keys_output,
|
||||
OutputIterator2 values_output,
|
||||
|
@ -75,7 +75,7 @@ struct WrappedUnaryPredicate
|
||||
return m_f((T)x);
|
||||
}
|
||||
|
||||
VTKM_EXEC bool operator()(const thrust::system::cuda::pointer<T> x) const { return m_f(*x); }
|
||||
VTKM_EXEC bool operator()(const T* x) const { return m_f(*x); }
|
||||
};
|
||||
|
||||
// Binary function object wrapper which can detect and handle calling the
|
||||
@ -126,26 +126,11 @@ struct WrappedBinaryOperator
|
||||
return m_f((T)x, (T)y);
|
||||
}
|
||||
|
||||
VTKM_EXEC T operator()(const thrust::system::cuda::pointer<T> x, const T* y) const
|
||||
{
|
||||
return m_f(*x, *y);
|
||||
}
|
||||
VTKM_EXEC T operator()(const T* const x, const T& y) const { return m_f(*x, y); }
|
||||
|
||||
VTKM_EXEC T operator()(const thrust::system::cuda::pointer<T> x, const T& y) const
|
||||
{
|
||||
return m_f(*x, y);
|
||||
}
|
||||
VTKM_EXEC T operator()(const T& x, const T* const y) const { return m_f(x, *y); }
|
||||
|
||||
VTKM_EXEC T operator()(const T& x, const thrust::system::cuda::pointer<T> y) const
|
||||
{
|
||||
return m_f(x, *y);
|
||||
}
|
||||
|
||||
VTKM_EXEC T operator()(const thrust::system::cuda::pointer<T> x,
|
||||
const thrust::system::cuda::pointer<T> y) const
|
||||
{
|
||||
return m_f(*x, *y);
|
||||
}
|
||||
VTKM_EXEC T operator()(const T* const x, const T* const y) const { return m_f(*x, *y); }
|
||||
};
|
||||
|
||||
template <typename T_, typename Function>
|
||||
@ -192,26 +177,11 @@ struct WrappedBinaryPredicate
|
||||
return m_f((T)x, (T)y);
|
||||
}
|
||||
|
||||
VTKM_EXEC bool operator()(const thrust::system::cuda::pointer<T> x, const T* y) const
|
||||
{
|
||||
return m_f(*x, *y);
|
||||
}
|
||||
VTKM_EXEC bool operator()(const T* const x, const T& y) const { return m_f(*x, y); }
|
||||
|
||||
VTKM_EXEC bool operator()(const thrust::system::cuda::pointer<T> x, const T& y) const
|
||||
{
|
||||
return m_f(*x, y);
|
||||
}
|
||||
VTKM_EXEC bool operator()(const T& x, const T* const y) const { return m_f(x, *y); }
|
||||
|
||||
VTKM_EXEC bool operator()(const T& x, const thrust::system::cuda::pointer<T> y) const
|
||||
{
|
||||
return m_f(x, *y);
|
||||
}
|
||||
|
||||
VTKM_EXEC bool operator()(const thrust::system::cuda::pointer<T> x,
|
||||
const thrust::system::cuda::pointer<T> y) const
|
||||
{
|
||||
return m_f(*x, *y);
|
||||
}
|
||||
VTKM_EXEC bool operator()(const T* const x, const T* const y) const { return m_f(*x, *y); }
|
||||
};
|
||||
}
|
||||
}
|
||||
|
Loading…
Reference in New Issue
Block a user