Commit f7789f0e authored by Kenneth Moreland's avatar Kenneth Moreland
Browse files

Fix issue with const types in Thrust array management

Previously, there was a declaration ConstArrayPortalFromThrust<const T>
in ArrayManagerExecutionThrustDevice. This proved problematic because
values read from the array in the worklet were typed as const T rather
than simply T. Any Vec or Matrix built from that type would then fail
because they are not meant to work with a const value (which means they
have to be set on construction and never changed.

Instead, declare ConstArrayPortalFromThrust<T> and internally set all
the Thrust pointers to have type const T. Also declare other thrust
pointers used as method parameters to have const T rather than T. This
should work as conversion from T to const T should be fine, but not the
other way around.
parent b7fdbb72
......@@ -74,7 +74,7 @@ public:
typedef vtkm::cont::internal::Storage<ValueType, StorageTag> StorageType;
typedef vtkm::exec::cuda::internal::ArrayPortalFromThrust< T > PortalType;
typedef vtkm::exec::cuda::internal::ConstArrayPortalFromThrust< const T > PortalConstType;
typedef vtkm::exec::cuda::internal::ConstArrayPortalFromThrust< T > PortalConstType;
VTKM_CONT_EXPORT
ArrayManagerExecutionThrustDevice(StorageType *storage)
......
......@@ -96,7 +96,7 @@ template<typename T, typename Enable = void>
struct load_through_texture
{
__device__
static T get(const thrust::system::cuda::pointer<T>& data)
static T get(const thrust::system::cuda::pointer<const T>& data)
{
return *(data.get());
}
......@@ -109,7 +109,7 @@ template<typename T>
struct load_through_texture<T, typename ::boost::enable_if< typename UseScalarTextureLoad<T>::type >::type >
{
__device__
static T get(const thrust::system::cuda::pointer<T>& data)
static T get(const thrust::system::cuda::pointer<const T>& data)
{
#if __CUDA_ARCH__ >= 350
// printf("__CUDA_ARCH__ UseScalarTextureLoad");
......@@ -125,7 +125,7 @@ template<typename T>
struct load_through_texture<T, typename ::boost::enable_if< typename UseVecTextureLoads<T>::type >::type >
{
__device__
static T get(const thrust::system::cuda::pointer<T>& data)
static T get(const thrust::system::cuda::pointer<const T>& data)
{
#if __CUDA_ARCH__ >= 350
// printf("__CUDA_ARCH__ UseVecTextureLoads");
......@@ -193,7 +193,7 @@ struct load_through_texture<T, typename ::boost::enable_if< typename UseMultiple
typedef typename boost::remove_const<T>::type NonConstT;
__device__
static T get(const thrust::system::cuda::pointer<T>& data)
static T get(const thrust::system::cuda::pointer<const T>& data)
{
#if __CUDA_ARCH__ >= 350
// printf("__CUDA_ARCH__ UseMultipleScalarTextureLoads");
......@@ -204,7 +204,7 @@ struct load_through_texture<T, typename ::boost::enable_if< typename UseMultiple
}
__device__
static T getAs(const thrust::system::cuda::pointer<T>& data)
static T getAs(const thrust::system::cuda::pointer<const T>& data)
{
//we need to fetch each component individually
const vtkm::IdComponent NUM_COMPONENTS= T::NUM_COMPONENTS;
......@@ -291,8 +291,8 @@ public:
VTKM_EXEC_CONT_EXPORT ConstArrayPortalFromThrust() { }
VTKM_CONT_EXPORT
ConstArrayPortalFromThrust(const thrust::system::cuda::pointer< T > begin,
const thrust::system::cuda::pointer< T > end)
ConstArrayPortalFromThrust(const thrust::system::cuda::pointer< const T > begin,
const thrust::system::cuda::pointer< const T > end)
: BeginIterator( begin ),
EndIterator( end )
{
......
Supports Markdown
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment