Shuenhoy/vtk-m
1
0
Fork 0
mirror of https://gitlab.kitware.com/vtk/vtk-m synced 2024-09-08 13:23:51 +00:00
Code Issues 2 Actions Packages Projects Releases Wiki Activity

Merge topic 'refactor_arrayhandle_to_reduce_lib_size'

Browse Source 
9bf14b78 Correct warnings inside worklet::Clip when making array handles
1b6d67e0 Always defer to the serial allocator when allocating basic storage
bf2b4169 Refactor vtk-m ArrayHandle to use mutable over const_cast
705528bf vtk-m ArrayHandle + basic storage has an optimized PrepareForDevice method
22f9ae3d vtk-m ArrayHandle + basic holds control data by StorageBasicBase
b1d0060d Make Storage and ArrayHandle export for the same value types.
d0a68d32 Refactor vtk-m storage basic to generate less code

Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !1084
This commit is contained in:
Robert Maynard 2018-02-21 21:56:34 +00:00 committed by Kitware Robot
commit 043afd326a
31 changed files with 1248 additions and 1124 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

39
vtkm/cont/ArrayHandle.cxx
View File

@ -21,49 +21,10 @@
#define vtkm_cont_ArrayHandle_cxx
#include <vtkm/cont/ArrayHandle.h>
#ifdef VTKM_MSVC
#define _VTKM_SHARED_PTR_INSTANTIATE(Type) \
template class VTKM_CONT_EXPORT \
std::shared_ptr<vtkm::cont::ArrayHandle<Type, vtkm::cont::StorageTagBasic>::InternalStruct>; \
template class VTKM_CONT_EXPORT std::shared_ptr< \
vtkm::cont::ArrayHandle<vtkm::Vec<Type, 2>, vtkm::cont::StorageTagBasic>::InternalStruct>; \
template class VTKM_CONT_EXPORT std::shared_ptr< \
vtkm::cont::ArrayHandle<vtkm::Vec<Type, 3>, vtkm::cont::StorageTagBasic>::InternalStruct>; \
template class VTKM_CONT_EXPORT std::shared_ptr< \
vtkm::cont::ArrayHandle<vtkm::Vec<Type, 4>, vtkm::cont::StorageTagBasic>::InternalStruct>;
_VTKM_SHARED_PTR_INSTANTIATE(char)
_VTKM_SHARED_PTR_INSTANTIATE(vtkm::Int8)
_VTKM_SHARED_PTR_INSTANTIATE(vtkm::UInt8)
_VTKM_SHARED_PTR_INSTANTIATE(vtkm::Int16)
_VTKM_SHARED_PTR_INSTANTIATE(vtkm::UInt16)
_VTKM_SHARED_PTR_INSTANTIATE(vtkm::Int32)
_VTKM_SHARED_PTR_INSTANTIATE(vtkm::UInt32)
_VTKM_SHARED_PTR_INSTANTIATE(vtkm::Int64)
_VTKM_SHARED_PTR_INSTANTIATE(vtkm::UInt64)
_VTKM_SHARED_PTR_INSTANTIATE(vtkm::Float32)
_VTKM_SHARED_PTR_INSTANTIATE(vtkm::Float64)
#undef _VTKM_SHARED_PTR_INSTANTIATE
#endif // VTKM_MSVC
namespace vtkm
{
namespace cont
{
namespace internal
{
ExecutionArrayInterfaceBasicBase::ExecutionArrayInterfaceBasicBase(StorageBasicBase& storage)
: ControlStorage(storage)
{
}
ExecutionArrayInterfaceBasicBase::~ExecutionArrayInterfaceBasicBase()
{
}
} // end namespace internal
#define _VTKM_ARRAYHANDLE_INSTANTIATE(Type) \
template class VTKM_CONT_EXPORT ArrayHandle<Type, StorageTagBasic>; \

35
vtkm/cont/ArrayHandle.h
View File

@ -485,12 +485,13 @@ public:
struct VTKM_ALWAYS_EXPORT InternalStruct
{
StorageType ControlArray;
bool ControlArrayValid;
mutable StorageType ControlArray;
mutable bool ControlArrayValid;
std::unique_ptr<vtkm::cont::internal::ArrayHandleExecutionManagerBase<ValueType, StorageTag>>
mutable std::unique_ptr<
vtkm::cont::internal::ArrayHandleExecutionManagerBase<ValueType, StorageTag>>
ExecutionArray;
bool ExecutionArrayValid;
mutable bool ExecutionArrayValid;
};
VTKM_CONT
@ -661,32 +662,6 @@ VTKM_NEVER_EXPORT VTKM_CONT inline void printSummary_ArrayHandle(
#ifndef vtkm_cont_ArrayHandle_cxx
#ifdef VTKM_MSVC
#define _VTKM_SHARED_PTR_EXPORT(Type) \
extern template class VTKM_CONT_TEMPLATE_EXPORT \
std::shared_ptr<vtkm::cont::ArrayHandle<Type, vtkm::cont::StorageTagBasic>::InternalStruct>; \
extern template class VTKM_CONT_TEMPLATE_EXPORT std::shared_ptr< \
vtkm::cont::ArrayHandle<vtkm::Vec<Type, 2>, vtkm::cont::StorageTagBasic>::InternalStruct>; \
extern template class VTKM_CONT_TEMPLATE_EXPORT std::shared_ptr< \
vtkm::cont::ArrayHandle<vtkm::Vec<Type, 3>, vtkm::cont::StorageTagBasic>::InternalStruct>; \
extern template class VTKM_CONT_TEMPLATE_EXPORT std::shared_ptr< \
vtkm::cont::ArrayHandle<vtkm::Vec<Type, 4>, vtkm::cont::StorageTagBasic>::InternalStruct>;
_VTKM_SHARED_PTR_EXPORT(char)
_VTKM_SHARED_PTR_EXPORT(vtkm::Int8)
_VTKM_SHARED_PTR_EXPORT(vtkm::UInt8)
_VTKM_SHARED_PTR_EXPORT(vtkm::Int16)
_VTKM_SHARED_PTR_EXPORT(vtkm::UInt16)
_VTKM_SHARED_PTR_EXPORT(vtkm::Int32)
_VTKM_SHARED_PTR_EXPORT(vtkm::UInt32)
_VTKM_SHARED_PTR_EXPORT(vtkm::Int64)
_VTKM_SHARED_PTR_EXPORT(vtkm::UInt64)
_VTKM_SHARED_PTR_EXPORT(vtkm::Float32)
_VTKM_SHARED_PTR_EXPORT(vtkm::Float64)
#undef _VTKM_SHARED_PTR_EXPORT
#endif // VTKM_MSVC
namespace vtkm
{
namespace cont

19
vtkm/cont/ArrayHandle.hxx
View File

@ -299,9 +299,8 @@ void ArrayHandle<T, S>::PrepareForDevice(DeviceAdapterTag) const
this->SyncControlArray();
// Need to change some state that does not change the logical state from
// an external point of view.
InternalStruct* internals = const_cast<InternalStruct*>(this->Internals.get());
internals->ExecutionArray.reset();
internals->ExecutionArrayValid = false;
this->Internals->ExecutionArray.reset();
this->Internals->ExecutionArrayValid = false;
}
}
@ -309,10 +308,9 @@ void ArrayHandle<T, S>::PrepareForDevice(DeviceAdapterTag) const
VTKM_ASSERT(!this->Internals->ExecutionArrayValid);
// Need to change some state that does not change the logical state from
// an external point of view.
InternalStruct* internals = const_cast<InternalStruct*>(this->Internals.get());
internals->ExecutionArray.reset(
this->Internals->ExecutionArray.reset(
new vtkm::cont::internal::ArrayHandleExecutionManager<T, StorageTag, DeviceAdapterTag>(
&internals->ControlArray));
&this->Internals->ControlArray));
}
template <typename T, typename S>
@ -322,19 +320,18 @@ void ArrayHandle<T, S>::SyncControlArray() const
{
// Need to change some state that does not change the logical state from
// an external point of view.
InternalStruct* internals = const_cast<InternalStruct*>(this->Internals.get());
if (this->Internals->ExecutionArrayValid)
{
internals->ExecutionArray->RetrieveOutputData(&internals->ControlArray);
internals->ControlArrayValid = true;
this->Internals->ExecutionArray->RetrieveOutputData(&this->Internals->ControlArray);
this->Internals->ControlArrayValid = true;
}
else
{
// This array is in the null state (there is nothing allocated), but
// the calling function wants to do something with the array. Put this
// class into a valid state by allocating an array of size 0.
internals->ControlArray.Allocate(0);
internals->ControlArrayValid = true;
this->Internals->ControlArray.Allocate(0);
this->Internals->ControlArrayValid = true;
}
}
}

1
vtkm/cont/CMakeLists.txt
View File

@ -108,6 +108,7 @@ set(sources
Field.cxx
internal/SimplePolymorphicContainer.cxx
internal/ArrayManagerExecutionShareWithControl.cxx
internal/ArrayHandleBasicImpl.cxx
StorageBasic.cxx
)

216
vtkm/cont/StorageBasic.cxx
View File

@ -20,6 +20,7 @@
#define vtkm_cont_StorageBasic_cxx
#include <vtkm/cont/StorageBasic.h>
#include <vtkm/internal/Configure.h>
#if defined(VTKM_POSIX)
#define VTKM_MEMALIGN_POSIX
@ -51,11 +52,7 @@ namespace cont
namespace internal
{
StorageBasicBase::~StorageBasicBase()
{
}
void* alloc_aligned(size_t size, size_t align)
void* StorageBasicAllocator::allocate(size_t size, size_t align)
{
#if defined(VTKM_MEMALIGN_POSIX)
void* mem = nullptr;
@ -70,14 +67,10 @@ void* alloc_aligned(size_t size, size_t align)
#else
void* mem = malloc(size);
#endif
if (mem == nullptr)
{
throw std::bad_alloc();
}
return mem;
}
void free_aligned(void* mem)
void StorageBasicAllocator::free_memory(void* mem)
{
#if defined(VTKM_MEMALIGN_POSIX)
free(mem);
@ -90,33 +83,188 @@ void free_aligned(void* mem)
#endif
}
template class VTKM_CONT_EXPORT Storage<char, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Int8, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::UInt8, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Int16, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::UInt16, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Int32, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::UInt32, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Int64, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::UInt64, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Float32, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Float64, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<vtkm::Int64, 2>, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<vtkm::Int32, 2>, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<vtkm::Float32, 2>, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<vtkm::Float64, 2>, StorageTagBasic>;
StorageBasicBase::StorageBasicBase()
: Array(nullptr)
, AllocatedByteSize(0)
, NumberOfValues(0)
, DeallocateOnRelease(true)
{
}
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<vtkm::Int64, 3>, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<vtkm::Int32, 3>, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<vtkm::Float32, 3>, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<vtkm::Float64, 3>, StorageTagBasic>;
StorageBasicBase::StorageBasicBase(const void* array,
vtkm::Id numberOfValues,
vtkm::UInt64 sizeOfValue)
: Array(const_cast<void*>(array))
, AllocatedByteSize(static_cast<vtkm::UInt64>(numberOfValues) * sizeOfValue)
, NumberOfValues(numberOfValues)
, DeallocateOnRelease(array == nullptr ? true : false)
{
}
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<char, 4>, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<Int8, 4>, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<UInt8, 4>, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<vtkm::Float32, 4>, StorageTagBasic>;
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<vtkm::Float64, 4>, StorageTagBasic>;
StorageBasicBase::~StorageBasicBase()
{
this->ReleaseResources();
}
StorageBasicBase::StorageBasicBase(const StorageBasicBase& src)
: Array(src.Array)
, AllocatedByteSize(src.AllocatedByteSize)
, NumberOfValues(src.NumberOfValues)
, DeallocateOnRelease(src.DeallocateOnRelease)
{
if (src.DeallocateOnRelease)
{
throw vtkm::cont::ErrorBadValue(
"Attempted to copy a storage array that needs deallocation. "
"This is disallowed to prevent complications with deallocation.");
}
}
StorageBasicBase StorageBasicBase::operator=(const StorageBasicBase& src)
{
if (src.DeallocateOnRelease)
{
throw vtkm::cont::ErrorBadValue(
"Attempted to copy a storage array that needs deallocation. "
"This is disallowed to prevent complications with deallocation.");
}
this->ReleaseResources();
this->Array = src.Array;
this->AllocatedByteSize = src.AllocatedByteSize;
this->NumberOfValues = src.NumberOfValues;
this->DeallocateOnRelease = src.DeallocateOnRelease;
return *this;
}
void StorageBasicBase::ReleaseResources()
{
if (this->AllocatedByteSize > 0)
{
VTKM_ASSERT(this->Array != nullptr);
if (this->DeallocateOnRelease)
{
AllocatorType{}.free_memory(this->Array);
}
this->Array = nullptr;
this->AllocatedByteSize = 0;
this->NumberOfValues = 0;
}
else
{
VTKM_ASSERT(this->Array == nullptr);
}
}
void StorageBasicBase::AllocateValues(vtkm::Id numberOfValues, vtkm::UInt64 sizeOfValue)
{
if (numberOfValues < 0)
{
throw vtkm::cont::ErrorBadAllocation("Cannot allocate an array with negative size.");
}
// Check that the number of bytes won't be more than a size_t can hold.
const size_t maxNumValues = std::numeric_limits<size_t>::max() / sizeOfValue;
if (static_cast<vtkm::UInt64>(numberOfValues) > maxNumValues)
{
throw ErrorBadAllocation("Requested allocation exceeds size_t capacity.");
}
// If we are allocating less data, just shrink the array.
// (If allocation empty, drop down so we can deallocate memory.)
vtkm::UInt64 allocsize = static_cast<vtkm::UInt64>(numberOfValues) * sizeOfValue;
if ((allocsize <= this->AllocatedByteSize) && (numberOfValues > 0))
{
this->NumberOfValues = numberOfValues;
return;
}
if (!this->DeallocateOnRelease)
{
throw vtkm::cont::ErrorBadValue("User allocated arrays cannot be reallocated.");
}
this->ReleaseResources();
if (numberOfValues > 0)
{
this->Array = AllocatorType{}.allocate(allocsize, VTKM_ALLOCATION_ALIGNMENT);
this->AllocatedByteSize = allocsize;
this->NumberOfValues = numberOfValues;
if (this->Array == nullptr)
{
// Make sureour state is OK.
this->AllocatedByteSize = 0;
this->NumberOfValues = 0;
throw vtkm::cont::ErrorBadAllocation("Could not allocate basic control array.");
}
}
else
{
// ReleaseResources should have already set NumberOfValues to 0.
VTKM_ASSERT(this->NumberOfValues == 0);
VTKM_ASSERT(this->AllocatedByteSize == 0);
}
this->DeallocateOnRelease = true;
}
void StorageBasicBase::Shrink(vtkm::Id numberOfValues)
{
if (numberOfValues > this->NumberOfValues)
{
throw vtkm::cont::ErrorBadValue("Shrink method cannot be used to grow array.");
}
this->NumberOfValues = numberOfValues;
}
void* StorageBasicBase::GetBasePointer() const
{
return this->Array;
}
void* StorageBasicBase::GetEndPointer(vtkm::Id numberOfValues, vtkm::UInt64 sizeOfValue) const
{
VTKM_ASSERT(this->NumberOfValues == numberOfValues);
if (!this->Array)
{
return nullptr;
}
auto p = static_cast<vtkm::UInt8*>(this->Array);
auto offset = static_cast<vtkm::UInt64>(this->NumberOfValues) * sizeOfValue;
return static_cast<void*>(p + offset);
}
void* StorageBasicBase::GetCapacityPointer() const
{
if (!this->Array)
{
return nullptr;
}
auto v = static_cast<vtkm::UInt8*>(this->Array) + AllocatedByteSize;
return static_cast<void*>(v);
}
#define _VTKM_STORAGE_INSTANTIATE(Type) \
template class VTKM_CONT_EXPORT Storage<Type, StorageTagBasic>; \
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<Type, 2>, StorageTagBasic>; \
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<Type, 3>, StorageTagBasic>; \
template class VTKM_CONT_EXPORT Storage<vtkm::Vec<Type, 4>, StorageTagBasic>;
_VTKM_STORAGE_INSTANTIATE(char)
_VTKM_STORAGE_INSTANTIATE(vtkm::Int8)
_VTKM_STORAGE_INSTANTIATE(vtkm::UInt8)
_VTKM_STORAGE_INSTANTIATE(vtkm::Int16)
_VTKM_STORAGE_INSTANTIATE(vtkm::UInt16)
_VTKM_STORAGE_INSTANTIATE(vtkm::Int32)
_VTKM_STORAGE_INSTANTIATE(vtkm::UInt32)
_VTKM_STORAGE_INSTANTIATE(vtkm::Int64)
_VTKM_STORAGE_INSTANTIATE(vtkm::UInt64)
_VTKM_STORAGE_INSTANTIATE(vtkm::Float32)
_VTKM_STORAGE_INSTANTIATE(vtkm::Float64)
}
}
} // namespace vtkm::cont::internal

283
vtkm/cont/StorageBasic.h
View File

@ -28,11 +28,6 @@
#include <vtkm/cont/internal/ArrayPortalFromIterators.h>
// Defines the cache line size in bytes to align allocations to
#ifndef VTKM_CACHE_LINE_SIZE
#define VTKM_CACHE_LINE_SIZE 64
#endif
namespace vtkm
{
namespace cont
@ -46,126 +41,87 @@ struct VTKM_ALWAYS_EXPORT StorageTagBasic
namespace internal
{
VTKM_CONT_EXPORT
void* alloc_aligned(size_t size, size_t align);
VTKM_CONT_EXPORT
void free_aligned(void* mem);
/// \brief an aligned allocator
/// A simple aligned allocator type that will align allocations to `Alignment` bytes
/// TODO: Once C++11 std::allocator_traits is better used by STL and we want to drop
/// support for pre-C++11 we can drop a lot of the typedefs and functions here.
template <typename T, size_t Alignment>
struct AlignedAllocator
/// Class that does all of VTK-m allocations
/// for storage basic. This is exists so that
/// stolen arrays can call the correct free
/// function ( _aligned_malloc ) on windows
struct VTKM_CONT_EXPORT StorageBasicAllocator
{
using value_type = T;
using reference = T&;
using const_reference = const T&;
using pointer = T*;
using const_pointer = const T*;
using void_pointer = void*;
using const_void_pointer = const void*;
using difference_type = std::ptrdiff_t;
using size_type = std::size_t;
void* allocate(size_t size, size_t align);
void free_memory(void* p);
template <typename U>
struct rebind
template <typename T>
void deallocate(T* p)
{
using other = AlignedAllocator<U, Alignment>;
};
AlignedAllocator() {}
template <typename Tb>
AlignedAllocator(const AlignedAllocator<Tb, Alignment>&)
{
}
pointer allocate(size_t n)
{
return static_cast<pointer>(alloc_aligned(n * sizeof(T), Alignment));
}
void deallocate(pointer p, size_t) { free_aligned(static_cast<void*>(p)); }
pointer address(reference r) { return &r; }
const_pointer address(const_reference r) { return &r; }
size_type max_size() const { return (std::numeric_limits<size_type>::max)() / sizeof(T); }
void construct(pointer p, const T& t)
{
(void)p;
new (p) T(t);
}
void destroy(pointer p)
{
(void)p;
p->~T();
this->free_memory(static_cast<void*>(p));
}
};
template <typename T, typename U, size_t AlignA, size_t AlignB>
bool operator==(const AlignedAllocator<T, AlignA>&, const AlignedAllocator<U, AlignB>&)
{
return AlignA == AlignB;
}
template <typename T, typename U, size_t AlignA, size_t AlignB>
bool operator!=(const AlignedAllocator<T, AlignA>&, const AlignedAllocator<U, AlignB>&)
{
return AlignA != AlignB;
}
/// Base class for basic storage classes. This is currently only used by
/// Basic storage to provide a type-agnostic API for allocations, etc.
/// Base class for basic storage classes. This allow us to implement
/// vtkm::cont::Storage<T, StorageTagBasic > for any T type with no overhead
/// as all heavy logic is provide by a type-agnostic API including allocations, etc.
class VTKM_CONT_EXPORT StorageBasicBase
{
public:
StorageBasicBase() {}
virtual ~StorageBasicBase();
using AllocatorType = StorageBasicAllocator;
VTKM_CONT StorageBasicBase();
VTKM_CONT StorageBasicBase(const void* array, vtkm::Id size, vtkm::UInt64 sizeOfValue);
VTKM_CONT ~StorageBasicBase();
/// \brief Return the number of bytes allocated for this storage object.
VTKM_CONT
virtual vtkm::UInt64 GetNumberOfBytes() const = 0;
VTKM_CONT StorageBasicBase(const StorageBasicBase& src);
VTKM_CONT StorageBasicBase operator=(const StorageBasicBase& src);
/// \brief Allocates an array with the specified size in bytes.
/// \brief Return the number of bytes allocated for this storage object(Capacity).
///
///
VTKM_CONT vtkm::UInt64 GetNumberOfBytes() const { return this->AllocatedByteSize; }
/// \brief Return the number of 'T' values allocated by this storage
VTKM_CONT vtkm::Id GetNumberOfValues() const { return this->NumberOfValues; }
/// \brief Allocates an array with the specified number of elements.
///
/// The allocation may be done on an already existing array, but can wipe out
/// any data already in the array. This method can throw
/// ErrorBadAllocation if the array cannot be allocated or
/// ErrorBadValue if the allocation is not feasible (for example, the
/// array storage is read-only).
VTKM_CONT
virtual void AllocateBytes(vtkm::UInt64 numberOfBytes) = 0;
VTKM_CONT void AllocateValues(vtkm::Id numberOfValues, vtkm::UInt64 sizeOfValue);
/// \brief Reduces the size of the array without changing its values.
///
/// This method allows you to resize the array without reallocating it. The
/// size of the array is changed to \c numberOfBytes bytes. The data
/// in the reallocated array stays the same, but \c numberOfBytes must be
/// size of the array is changed so that it can hold \c numberOfValues values.
/// The data in the reallocated array stays the same, but \c numberOfValues must be
/// equal or less than the preexisting size. That is, this method can only be
/// used to shorten the array, not lengthen.
VTKM_CONT
virtual void ShrinkBytes(vtkm::UInt64 numberOfBytes) = 0;
VTKM_CONT void Shrink(vtkm::Id numberOfValues);
/// \brief Frees any resources (i.e. memory) stored in this array.
///
/// After calling this method GetNumberOfBytes() will return 0. The
/// resources should also be released when the Storage class is
/// destroyed.
VTKM_CONT
virtual void ReleaseResources() = 0;
VTKM_CONT void ReleaseResources();
/// \brief Returns if vtkm will deallocate this memory. VTK-m StorageBasic
/// is designed that VTK-m will not deallocate user passed memory, or
/// instances that have been stolen (\c StealArray)
VTKM_CONT bool WillDeallocate() const { return this->DeallocateOnRelease; }
/// Return the memory location of the first element of the array data.
VTKM_CONT
virtual void* GetBasePointer() const = 0;
VTKM_CONT void* GetBasePointer() const;
/// Return the memory location of the first element past the end of the array
/// data.
VTKM_CONT
virtual void* GetEndPointer() const = 0;
VTKM_CONT void* GetEndPointer(vtkm::Id numberOfValues, vtkm::UInt64 sizeOfValue) const;
/// Return the memory location of the first element past the end of the
/// array's allocated memory buffer.
VTKM_CONT
virtual void* GetCapacityPointer() const = 0;
VTKM_CONT void* GetCapacityPointer() const;
protected:
void* Array;
vtkm::UInt64 AllocatedByteSize;
vtkm::Id NumberOfValues;
bool DeallocateOnRelease;
};
/// A basic implementation of an Storage object.
@ -180,70 +136,23 @@ template <typename ValueT>
class VTKM_ALWAYS_EXPORT Storage<ValueT, vtkm::cont::StorageTagBasic> : public StorageBasicBase
{
public:
using AllocatorType = vtkm::cont::internal::StorageBasicAllocator;
using ValueType = ValueT;
using PortalType = vtkm::cont::internal::ArrayPortalFromIterators<ValueType*>;
using PortalConstType = vtkm::cont::internal::ArrayPortalFromIterators<const ValueType*>;
/// The original design of this class provided an allocator as a template
/// parameters. That messed things up, though, because other templated
/// classes assume that the \c Storage has one template parameter. There are
/// other ways to allow you to specify the allocator, but it is uncertain
/// whether that would ever be useful. So, instead of jumping through hoops
/// implementing them, just fix the allocator for now.
///
using AllocatorType = AlignedAllocator<ValueType, VTKM_CACHE_LINE_SIZE>;
public:
/// \brief construct storage that VTK-m is responsible for
VTKM_CONT
Storage();
VTKM_CONT Storage();
/// \brief construct storage that VTK-m is not responsible for
VTKM_CONT
Storage(const ValueType* array, vtkm::Id numberOfValues = 0);
VTKM_CONT Storage(const ValueType* array, vtkm::Id numberOfValues = 0);
VTKM_CONT
~Storage();
VTKM_CONT void Allocate(vtkm::Id numberOfValues);
VTKM_CONT
Storage(const Storage<ValueType, StorageTagBasic>& src);
VTKM_CONT PortalType GetPortal();
VTKM_CONT
Storage& operator=(const Storage<ValueType, StorageTagBasic>& src);
VTKM_CONT
void ReleaseResources() final;
VTKM_CONT
void Allocate(vtkm::Id numberOfValues);
VTKM_CONT
vtkm::Id GetNumberOfValues() const { return this->NumberOfValues; }
VTKM_CONT
vtkm::UInt64 GetNumberOfBytes() const final
{
return static_cast<vtkm::UInt64>(this->NumberOfValues) *
static_cast<vtkm::UInt64>(sizeof(ValueT));
}
VTKM_CONT
void Shrink(vtkm::Id numberOfValues);
VTKM_CONT
void AllocateBytes(vtkm::UInt64) final;
VTKM_CONT
void ShrinkBytes(vtkm::UInt64) final;
VTKM_CONT
PortalType GetPortal() { return PortalType(this->Array, this->Array + this->NumberOfValues); }
VTKM_CONT
PortalConstType GetPortalConst() const
{
return PortalConstType(this->Array, this->Array + this->NumberOfValues);
}
VTKM_CONT PortalConstType GetPortalConst() const;
/// \brief Get a pointer to the underlying data structure.
///
@ -251,10 +160,9 @@ public:
/// memory associated with this array still belongs to the Storage (i.e.
/// Storage will eventually deallocate the array).
///
VTKM_CONT
ValueType* GetArray() { return this->Array; }
VTKM_CONT
const ValueType* GetArray() const { return this->Array; }
VTKM_CONT ValueType* GetArray();
VTKM_CONT const ValueType* GetArray() const;
/// \brief Take the reference away from this object.
///
@ -265,36 +173,7 @@ public:
/// VTK-m and not having to keep a VTK-m object around. Obviously the caller
/// becomes responsible for destroying the memory.
///
VTKM_CONT
ValueType* StealArray();
/// \brief Returns if vtkm will deallocate this memory. VTK-m StorageBasic
/// is designed that VTK-m will not deallocate user passed memory, or
/// instances that have been stolen (\c StealArray)
VTKM_CONT
bool WillDeallocate() const { return this->DeallocateOnRelease; }
VTKM_CONT
void* GetBasePointer() const final { return static_cast<void*>(this->Array); }
VTKM_CONT
void* GetEndPointer() const final
{
return static_cast<void*>(this->Array + this->NumberOfValues);
}
VTKM_CONT
void* GetCapacityPointer() const final
{
return static_cast<void*>(this->Array + this->AllocatedSize);
}
private:
ValueType* Array;
vtkm::Id NumberOfValues;
vtkm::Id AllocatedSize;
bool DeallocateOnRelease;
VTKM_CONT ValueType* StealArray();
};
} // namespace internal
@ -311,39 +190,23 @@ namespace internal
/// \cond
/// Make doxygen ignore this section
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<char, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Int8, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::UInt8, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Int16, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::UInt16, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Int32, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::UInt32, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Int64, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::UInt64, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Float32, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Float64, StorageTagBasic>;
#define _VTKM_STORAGE_EXPORT(Type) \
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<Type, StorageTagBasic>; \
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Vec<Type, 2>, StorageTagBasic>; \
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Vec<Type, 3>, StorageTagBasic>; \
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Vec<Type, 4>, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Vec<vtkm::Int64, 2>, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Vec<vtkm::Int32, 2>, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT
Storage<vtkm::Vec<vtkm::Float32, 2>, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT
Storage<vtkm::Vec<vtkm::Float64, 2>, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Vec<vtkm::Int64, 3>, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Vec<vtkm::Int32, 3>, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT
Storage<vtkm::Vec<vtkm::Float32, 3>, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT
Storage<vtkm::Vec<vtkm::Float64, 3>, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Vec<char, 4>, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Vec<Int8, 4>, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT Storage<vtkm::Vec<UInt8, 4>, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT
Storage<vtkm::Vec<vtkm::Float32, 4>, StorageTagBasic>;
extern template class VTKM_CONT_TEMPLATE_EXPORT
Storage<vtkm::Vec<vtkm::Float64, 4>, StorageTagBasic>;
_VTKM_STORAGE_EXPORT(char)
_VTKM_STORAGE_EXPORT(vtkm::Int8)
_VTKM_STORAGE_EXPORT(vtkm::UInt8)
_VTKM_STORAGE_EXPORT(vtkm::Int16)
_VTKM_STORAGE_EXPORT(vtkm::UInt16)
_VTKM_STORAGE_EXPORT(vtkm::Int32)
_VTKM_STORAGE_EXPORT(vtkm::UInt32)
_VTKM_STORAGE_EXPORT(vtkm::Int64)
_VTKM_STORAGE_EXPORT(vtkm::UInt64)
_VTKM_STORAGE_EXPORT(vtkm::Float32)
_VTKM_STORAGE_EXPORT(vtkm::Float64)
/// \endcond
}
}

156
vtkm/cont/StorageBasic.hxx
View File

@ -31,173 +31,55 @@ namespace internal
template <typename T>
Storage<T, vtkm::cont::StorageTagBasic>::Storage()
: Array(nullptr)
, NumberOfValues(0)
, AllocatedSize(0)
, DeallocateOnRelease(true)
: StorageBasicBase()
{
}
template <typename T>
Storage<T, vtkm::cont::StorageTagBasic>::Storage(const T* array, vtkm::Id numberOfValues)
: Array(const_cast<T*>(array))
, NumberOfValues(numberOfValues)
, AllocatedSize(numberOfValues)
, DeallocateOnRelease(array == nullptr ? true : false)
: StorageBasicBase(const_cast<T*>(array), numberOfValues, sizeof(T))
{
}
template <typename T>
Storage<T, vtkm::cont::StorageTagBasic>::~Storage()
{
this->ReleaseResources();
}
template <typename T>
Storage<T, vtkm::cont::StorageTagBasic>::Storage(const Storage<T, StorageTagBasic>& src)
: Array(src.Array)
, NumberOfValues(src.NumberOfValues)
, AllocatedSize(src.AllocatedSize)
, DeallocateOnRelease(src.DeallocateOnRelease)
{
if (src.DeallocateOnRelease)
{
throw vtkm::cont::ErrorBadValue(
"Attempted to copy a storage array that needs deallocation. "
"This is disallowed to prevent complications with deallocation.");
}
}
template <typename T>
Storage<T, vtkm::cont::StorageTagBasic>& Storage<T, vtkm::cont::StorageTagBasic>::operator=(
const Storage<T, StorageTagBasic>& src)
{
if (src.DeallocateOnRelease)
{
throw vtkm::cont::ErrorBadValue(
"Attempted to copy a storage array that needs deallocation. "
"This is disallowed to prevent complications with deallocation.");
}
this->ReleaseResources();
this->Array = src.Array;
this->NumberOfValues = src.NumberOfValues;
this->AllocatedSize = src.AllocatedSize;
this->DeallocateOnRelease = src.DeallocateOnRelease;
return *this;
}
template <typename T>
void Storage<T, vtkm::cont::StorageTagBasic>::ReleaseResources()
{
if (this->NumberOfValues > 0)
{
VTKM_ASSERT(this->Array != nullptr);
if (this->DeallocateOnRelease)
{
AllocatorType allocator;
allocator.deallocate(this->Array, static_cast<std::size_t>(this->AllocatedSize));
}
this->Array = nullptr;
this->NumberOfValues = 0;
this->AllocatedSize = 0;
}
else
{
VTKM_ASSERT(this->Array == nullptr);
}
}
template <typename T>
void Storage<T, vtkm::cont::StorageTagBasic>::Allocate(vtkm::Id numberOfValues)
{
if (numberOfValues < 0)
{
throw vtkm::cont::ErrorBadAllocation("Cannot allocate an array with negative size.");
}
// Check that the number of bytes won't be more than a size_t can hold.
const size_t maxNumValues = std::numeric_limits<size_t>::max() / sizeof(T);
if (static_cast<vtkm::UInt64>(numberOfValues) > static_cast<vtkm::UInt64>(maxNumValues))
{
throw ErrorBadAllocation("Requested allocation exceeds size_t capacity.");
}
this->AllocateBytes(static_cast<vtkm::UInt64>(numberOfValues) *
static_cast<vtkm::UInt64>(sizeof(T)));
this->AllocateValues(numberOfValues, sizeof(T));
}
template <typename T>
void Storage<T, vtkm::cont::StorageTagBasic>::AllocateBytes(vtkm::UInt64 numberOfBytes)
typename Storage<T, vtkm::cont::StorageTagBasic>::PortalType
Storage<T, vtkm::cont::StorageTagBasic>::GetPortal()
{
const vtkm::Id numberOfValues =
static_cast<vtkm::Id>(numberOfBytes / static_cast<vtkm::UInt64>(sizeof(T)));
// If we are allocating less data, just shrink the array.
// (If allocation empty, drop down so we can deallocate memory.)
if ((numberOfValues <= this->AllocatedSize) && (numberOfValues > 0))
{
this->NumberOfValues = numberOfValues;
return;
}
if (!this->DeallocateOnRelease)
{
throw vtkm::cont::ErrorBadValue("User allocated arrays cannot be reallocated.");
}
this->ReleaseResources();
try
{
if (numberOfValues > 0)
{
AllocatorType allocator;
this->Array = allocator.allocate(static_cast<std::size_t>(numberOfValues));
this->AllocatedSize = numberOfValues;
this->NumberOfValues = numberOfValues;
}
else
{
// ReleaseResources should have already set AllocatedSize to 0.
VTKM_ASSERT(this->AllocatedSize == 0);
}
}
catch (std::bad_alloc&)
{
// Make sureour state is OK.
this->Array = nullptr;
this->NumberOfValues = 0;
this->AllocatedSize = 0;
throw vtkm::cont::ErrorBadAllocation("Could not allocate basic control array.");
}
this->DeallocateOnRelease = true;
auto v = static_cast<T*>(this->Array);
return PortalType(v, v + this->NumberOfValues);
}
template <typename T>
void Storage<T, vtkm::cont::StorageTagBasic>::Shrink(vtkm::Id numberOfValues)
typename Storage<T, vtkm::cont::StorageTagBasic>::PortalConstType
Storage<T, vtkm::cont::StorageTagBasic>::GetPortalConst() const
{
this->ShrinkBytes(static_cast<vtkm::UInt64>(numberOfValues) *
static_cast<vtkm::UInt64>(sizeof(T)));
auto v = static_cast<T*>(this->Array);
return PortalConstType(v, v + this->NumberOfValues);
}
template <typename T>
void Storage<T, vtkm::cont::StorageTagBasic>::ShrinkBytes(vtkm::UInt64 numberOfBytes)
T* Storage<T, vtkm::cont::StorageTagBasic>::GetArray()
{
if (numberOfBytes > this->GetNumberOfBytes())
{
throw vtkm::cont::ErrorBadValue("Shrink method cannot be used to grow array.");
}
return static_cast<T*>(this->Array);
}
this->NumberOfValues =
static_cast<vtkm::Id>(numberOfBytes / static_cast<vtkm::UInt64>(sizeof(T)));
template <typename T>
const T* Storage<T, vtkm::cont::StorageTagBasic>::GetArray() const
{
return static_cast<T*>(this->Array);
}
template <typename T>
T* Storage<T, vtkm::cont::StorageTagBasic>::StealArray()
{
this->DeallocateOnRelease = false;
return this->Array;
return static_cast<T*>(this->Array);
}
} // namespace internal

171
vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.cu
View File

@ -21,182 +21,11 @@
#define vtk_m_cont_cuda_internal_ArrayManagerExecutionCuda_cu
#include <vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.h>
#include <vtkm/cont/cuda/internal/CudaAllocator.h>
using vtkm::cont::cuda::internal::CudaAllocator;
namespace vtkm
{
namespace cont
{
namespace internal
{
ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::ExecutionArrayInterfaceBasic(
StorageBasicBase& storage)
: Superclass(storage)
{
}
DeviceAdapterId ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::GetDeviceId() const
{
return VTKM_DEVICE_ADAPTER_CUDA;
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::Allocate(TypelessExecutionArray& execArray,
vtkm::UInt64 numBytes) const
{
// Detect if we can reuse a device-accessible pointer from the control env:
if (CudaAllocator::IsDevicePointer(execArray.ArrayControl))
{
const vtkm::UInt64 managedCapacity =
static_cast<vtkm::UInt64>(static_cast<const char*>(execArray.ArrayControlCapacity) -
static_cast<const char*>(execArray.ArrayControl));
if (managedCapacity >= numBytes)
{
if (execArray.Array && execArray.Array != execArray.ArrayControl)
{
this->Free(execArray);
}
execArray.Array = const_cast<void*>(execArray.ArrayControl);
execArray.ArrayEnd = static_cast<char*>(execArray.Array) + numBytes;
execArray.ArrayCapacity = const_cast<void*>(execArray.ArrayControlCapacity);
return;
}
}
if (execArray.Array != nullptr)
{
const vtkm::UInt64 cap = static_cast<vtkm::UInt64>(static_cast<char*>(execArray.ArrayCapacity) -
static_cast<char*>(execArray.Array));
if (cap < numBytes)
{ // Current allocation too small -- free & realloc
this->Free(execArray);
}
else
{ // Reuse buffer if possible:
execArray.ArrayEnd = static_cast<char*>(execArray.Array) + numBytes;
return;
}
}
VTKM_ASSERT(execArray.Array == nullptr);
// Attempt to allocate:
try
{
// Cast to char* so that the pointer math below will work.
char* tmp = static_cast<char*>(CudaAllocator::Allocate(static_cast<size_t>(numBytes)));
execArray.Array = tmp;
execArray.ArrayEnd = tmp + numBytes;
execArray.ArrayCapacity = tmp + numBytes;
}
catch (const std::exception& error)
{
std::ostringstream err;
err << "Failed to allocate " << numBytes << " bytes on device: " << error.what();
throw vtkm::cont::ErrorBadAllocation(err.str());
}
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::Free(
TypelessExecutionArray& execArray) const
{
// If we're sharing a device-accessible pointer between control/exec, don't
// actually free it -- just null the pointers here:
if (execArray.Array == execArray.ArrayControl &&
CudaAllocator::IsDevicePointer(execArray.ArrayControl))
{
execArray.Array = nullptr;
execArray.ArrayEnd = nullptr;
execArray.ArrayCapacity = nullptr;
return;
}
if (execArray.Array != nullptr)
{
CudaAllocator::Free(execArray.Array);
execArray.Array = nullptr;
execArray.ArrayEnd = nullptr;
execArray.ArrayCapacity = nullptr;
}
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::CopyFromControl(
const void* controlPtr,
void* executionPtr,
vtkm::UInt64 numBytes) const
{
// Do nothing if we're sharing a device-accessible pointer between control and
// execution:
if (controlPtr == executionPtr && CudaAllocator::IsDevicePointer(controlPtr))
{
CudaAllocator::PrepareForInput(executionPtr, numBytes);
return;
}
VTKM_CUDA_CALL(cudaMemcpyAsync(executionPtr,
controlPtr,
static_cast<std::size_t>(numBytes),
cudaMemcpyHostToDevice,
cudaStreamPerThread));
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::CopyToControl(const void* executionPtr,
void* controlPtr,
vtkm::UInt64 numBytes) const
{
// Do nothing if we're sharing a cuda managed pointer between control and execution:
if (controlPtr == executionPtr && CudaAllocator::IsDevicePointer(controlPtr))
{
// If we're trying to copy a shared, non-managed device pointer back to
// control throw an exception -- the pointer cannot be read from control,
// so this operation is invalid.
if (!CudaAllocator::IsManagedPointer(controlPtr))
{
throw vtkm::cont::ErrorBadValue(
"Control pointer is a CUDA device pointer that does not supported managed access.");
}
// If it is managed, just return and let CUDA handle the migration for us.
CudaAllocator::PrepareForControl(controlPtr, numBytes);
return;
}
VTKM_CUDA_CALL(cudaMemcpyAsync(controlPtr,
executionPtr,
static_cast<std::size_t>(numBytes),
cudaMemcpyDeviceToHost,
cudaStreamPerThread));
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::UsingForRead(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const
{
CudaAllocator::PrepareForInput(executionPtr, static_cast<size_t>(numBytes));
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::UsingForWrite(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const
{
CudaAllocator::PrepareForOutput(executionPtr, static_cast<size_t>(numBytes));
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::UsingForReadWrite(
const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const
{
CudaAllocator::PrepareForInPlace(executionPtr, static_cast<size_t>(numBytes));
}
} // end namespace internal
VTKM_INSTANTIATE_ARRAYHANDLES_FOR_DEVICE_ADAPTER(DeviceAdapterTagCuda)
}
} // end vtkm::cont

31
vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.h
View File

@ -26,6 +26,10 @@
#include <vtkm/cont/internal/ArrayExportMacros.h>
#include <vtkm/cont/internal/ArrayManagerExecution.h>
//This is in a separate header so that ArrayHandleBasicImpl can include
//the interface without getting any CUDA headers
#include <vtkm/cont/cuda/internal/ExecutionArrayInterfaceBasicCuda.h>
// These must be placed in the vtkm::cont::internal namespace so that
// the template can be found.
@ -137,33 +141,6 @@ struct ExecutionPortalFactoryBasic<T, DeviceAdapterTagCuda>
}
};
template <>
struct VTKM_CONT_EXPORT ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>
: public ExecutionArrayInterfaceBasicBase
{
using Superclass = ExecutionArrayInterfaceBasicBase;
VTKM_CONT ExecutionArrayInterfaceBasic(StorageBasicBase& storage);
VTKM_CONT DeviceAdapterId GetDeviceId() const final;
VTKM_CONT void Allocate(TypelessExecutionArray& execArray, vtkm::UInt64 numBytes) const final;
VTKM_CONT void Free(TypelessExecutionArray& execArray) const final;
VTKM_CONT void CopyFromControl(const void* controlPtr,
void* executionPtr,
vtkm::UInt64 numBytes) const final;
VTKM_CONT void CopyToControl(const void* executionPtr,
void* controlPtr,
vtkm::UInt64 numBytes) const final;
VTKM_CONT void UsingForRead(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const final;
VTKM_CONT void UsingForWrite(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const final;
VTKM_CONT void UsingForReadWrite(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const final;
};
} // namespace internal
#ifndef vtk_m_cont_cuda_internal_ArrayManagerExecutionCuda_cu

2
vtkm/cont/cuda/internal/CMakeLists.txt
View File

@ -25,6 +25,7 @@ set(headers
DeviceAdapterAlgorithmCuda.h
DeviceAdapterAlgorithmThrust.h
DeviceAdapterTagCuda.h
ExecutionArrayInterfaceBasicCuda.h
MakeThrustIterator.h
TaskTuner.h
ThrustExceptionHandler.h
@ -33,6 +34,7 @@ set(headers
set(sources
ArrayManagerExecutionCuda.cu
ExecutionArrayInterfaceBasicCuda.cu
CudaAllocator.cu
)

199
vtkm/cont/cuda/internal/ExecutionArrayInterfaceBasicCuda.cu Normal file
View File

@ -0,0 +1,199 @@
//============================================================================
// 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 2017 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2017 UT-Battelle, LLC.
// Copyright 2017 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// 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.
//============================================================================
#include <vtkm/cont/cuda/internal/CudaAllocator.h>
#include <vtkm/cont/cuda/internal/ExecutionArrayInterfaceBasicCuda.h>
using vtkm::cont::cuda::internal::CudaAllocator;
namespace vtkm
{
namespace cont
{
namespace internal
{
ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::ExecutionArrayInterfaceBasic(
StorageBasicBase& storage)
: Superclass(storage)
{
}
DeviceAdapterId ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::GetDeviceId() const
{
return VTKM_DEVICE_ADAPTER_CUDA;
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::Allocate(TypelessExecutionArray& execArray,
vtkm::Id numberOfValues,
vtkm::UInt64 sizeOfValue) const
{
const vtkm::UInt64 numBytes = static_cast<vtkm::UInt64>(numberOfValues) * sizeOfValue;
// Detect if we can reuse a device-accessible pointer from the control env:
if (CudaAllocator::IsDevicePointer(execArray.ArrayControl))
{
const vtkm::UInt64 managedCapacity =
static_cast<vtkm::UInt64>(static_cast<const char*>(execArray.ArrayControlCapacity) -
static_cast<const char*>(execArray.ArrayControl));
if (managedCapacity >= numBytes)
{
if (execArray.Array && execArray.Array != execArray.ArrayControl)
{
this->Free(execArray);
}
execArray.Array = const_cast<void*>(execArray.ArrayControl);
execArray.ArrayEnd = static_cast<char*>(execArray.Array) + numBytes;
execArray.ArrayCapacity = const_cast<void*>(execArray.ArrayControlCapacity);
return;
}
}
if (execArray.Array != nullptr)
{
const vtkm::UInt64 cap = static_cast<vtkm::UInt64>(static_cast<char*>(execArray.ArrayCapacity) -
static_cast<char*>(execArray.Array));
if (cap < numBytes)
{ // Current allocation too small -- free & realloc
this->Free(execArray);
}
else
{ // Reuse buffer if possible:
execArray.ArrayEnd = static_cast<char*>(execArray.Array) + numBytes;
return;
}
}
VTKM_ASSERT(execArray.Array == nullptr);
// Attempt to allocate:
try
{
// Cast to char* so that the pointer math below will work.
char* tmp = static_cast<char*>(CudaAllocator::Allocate(static_cast<size_t>(numBytes)));
execArray.Array = tmp;
execArray.ArrayEnd = tmp + numBytes;
execArray.ArrayCapacity = tmp + numBytes;
}
catch (const std::exception& error)
{
std::ostringstream err;
err << "Failed to allocate " << numBytes << " bytes on device: " << error.what();
throw vtkm::cont::ErrorBadAllocation(err.str());
}
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::Free(
TypelessExecutionArray& execArray) const
{
// If we're sharing a device-accessible pointer between control/exec, don't
// actually free it -- just null the pointers here:
if (execArray.Array == execArray.ArrayControl &&
CudaAllocator::IsDevicePointer(execArray.ArrayControl))
{
execArray.Array = nullptr;
execArray.ArrayEnd = nullptr;
execArray.ArrayCapacity = nullptr;
return;
}
if (execArray.Array != nullptr)
{
CudaAllocator::Free(execArray.Array);
execArray.Array = nullptr;
execArray.ArrayEnd = nullptr;
execArray.ArrayCapacity = nullptr;
}
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::CopyFromControl(
const void* controlPtr,
void* executionPtr,
vtkm::UInt64 numBytes) const
{
// Do nothing if we're sharing a device-accessible pointer between control and
// execution:
if (controlPtr == executionPtr && CudaAllocator::IsDevicePointer(controlPtr))
{
CudaAllocator::PrepareForInput(executionPtr, numBytes);
return;
}
VTKM_CUDA_CALL(cudaMemcpyAsync(executionPtr,
controlPtr,
static_cast<std::size_t>(numBytes),
cudaMemcpyHostToDevice,
cudaStreamPerThread));
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::CopyToControl(const void* executionPtr,
void* controlPtr,
vtkm::UInt64 numBytes) const
{
// Do nothing if we're sharing a cuda managed pointer between control and execution:
if (controlPtr == executionPtr && CudaAllocator::IsDevicePointer(controlPtr))
{
// If we're trying to copy a shared, non-managed device pointer back to
// control throw an exception -- the pointer cannot be read from control,
// so this operation is invalid.
if (!CudaAllocator::IsManagedPointer(controlPtr))
{
throw vtkm::cont::ErrorBadValue(
"Control pointer is a CUDA device pointer that does not supported managed access.");
}
// If it is managed, just return and let CUDA handle the migration for us.
CudaAllocator::PrepareForControl(controlPtr, numBytes);
return;
}
VTKM_CUDA_CALL(cudaMemcpyAsync(controlPtr,
executionPtr,
static_cast<std::size_t>(numBytes),
cudaMemcpyDeviceToHost,
cudaStreamPerThread));
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::UsingForRead(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const
{
CudaAllocator::PrepareForInput(executionPtr, static_cast<size_t>(numBytes));
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::UsingForWrite(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const
{
CudaAllocator::PrepareForOutput(executionPtr, static_cast<size_t>(numBytes));
}
void ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>::UsingForReadWrite(
const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const
{
CudaAllocator::PrepareForInPlace(executionPtr, static_cast<size_t>(numBytes));
}
} // end namespace internal
}
} // end vtkm::cont

64
vtkm/cont/cuda/internal/ExecutionArrayInterfaceBasicCuda.h Normal file
View File

@ -0,0 +1,64 @@
//============================================================================
// 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 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// 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_cont_cuda_internal_ExecutionArrayInterfaceCuda_h
#define vtk_m_cont_cuda_internal_ExecutionArrayInterfaceCuda_h
#include <vtkm/cont/cuda/DeviceAdapterCuda.h>
namespace vtkm
{
namespace cont
{
namespace internal
{
template <>
struct VTKM_CONT_EXPORT ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>
: public ExecutionArrayInterfaceBasicBase
{
using Superclass = ExecutionArrayInterfaceBasicBase;
VTKM_CONT ExecutionArrayInterfaceBasic(StorageBasicBase& storage);
VTKM_CONT DeviceAdapterId GetDeviceId() const final;
VTKM_CONT void Allocate(TypelessExecutionArray& execArray,
vtkm::Id numberOfValues,
vtkm::UInt64 sizeOfValue) const final;
VTKM_CONT void Free(TypelessExecutionArray& execArray) const final;
VTKM_CONT void CopyFromControl(const void* controlPtr,
void* executionPtr,
vtkm::UInt64 numBytes) const final;
VTKM_CONT void CopyToControl(const void* executionPtr,
void* controlPtr,
vtkm::UInt64 numBytes) const final;
VTKM_CONT void UsingForRead(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const final;
VTKM_CONT void UsingForWrite(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const final;
VTKM_CONT void UsingForReadWrite(const void* controlPtr,
const void* executionPtr,
vtkm::UInt64 numBytes) const final;
};
} // namespace internal
}
} // namespace vtkm::cont
#endif //vtk_m_cont_cuda_internal_ExecutionArrayInterfaceCuda_h

28
vtkm/cont/cuda/testing/UnitTestCudaShareUserProvidedManagedMemory.cu
View File

@ -78,8 +78,8 @@ void TestPrepareForInput(bool managed)
vtkm::cont::ArrayHandle<ValueType> handle = CreateArrayHandle<ValueType>(32, managed);
handle.PrepareForInput(vtkm::cont::DeviceAdapterTagCuda());
ValueType* contArray = handle.Internals->ControlArray.GetArray();
ValueType* execArray = handle.Internals->ExecutionArray;
void* contArray = handle.Internals->ControlArray->GetBasePointer();
void* execArray = handle.Internals->ExecutionArray;
VTKM_TEST_ASSERT(contArray != nullptr, "No control array after PrepareForInput.");
VTKM_TEST_ASSERT(execArray != nullptr, "No execution array after PrepareForInput.");
VTKM_TEST_ASSERT(CudaAllocator::IsDevicePointer(execArray),
@ -102,8 +102,8 @@ void TestPrepareForInPlace(bool managed)
vtkm::cont::ArrayHandle<ValueType> handle = CreateArrayHandle<ValueType>(32, managed);
handle.PrepareForInPlace(vtkm::cont::DeviceAdapterTagCuda());
ValueType* contArray = handle.Internals->ControlArray.GetArray();
ValueType* execArray = handle.Internals->ExecutionArray;
void* contArray = handle.Internals->ControlArray->GetBasePointer();
void* execArray = handle.Internals->ExecutionArray;
VTKM_TEST_ASSERT(contArray != nullptr, "No control array after PrepareForInPlace.");
VTKM_TEST_ASSERT(execArray != nullptr, "No execution array after PrepareForInPlace.");
VTKM_TEST_ASSERT(CudaAllocator::IsDevicePointer(execArray),
@ -127,8 +127,8 @@ void TestPrepareForOutput(bool managed)
vtkm::cont::ArrayHandle<ValueType> handle = CreateArrayHandle<ValueType>(32, managed);
handle.PrepareForOutput(32, vtkm::cont::DeviceAdapterTagCuda());
ValueType* contArray = handle.Internals->ControlArray.GetArray();
ValueType* execArray = handle.Internals->ExecutionArray;
void* contArray = handle.Internals->ControlArray->GetBasePointer();
void* execArray = handle.Internals->ExecutionArray;
VTKM_TEST_ASSERT(contArray != nullptr, "No control array after PrepareForOutput.");
VTKM_TEST_ASSERT(execArray != nullptr, "No execution array after PrepareForOutput.");
VTKM_TEST_ASSERT(CudaAllocator::IsDevicePointer(execArray),
@ -151,12 +151,12 @@ void TestReleaseResourcesExecution(bool managed)
vtkm::cont::ArrayHandle<ValueType> handle = CreateArrayHandle<ValueType>(32, managed);
handle.PrepareForInput(vtkm::cont::DeviceAdapterTagCuda());
ValueType* origArray = handle.Internals->ExecutionArray;
void* origArray = handle.Internals->ExecutionArray;
handle.ReleaseResourcesExecution();
ValueType* contArray = handle.Internals->ControlArray.GetArray();
ValueType* execArray = handle.Internals->ExecutionArray;
void* contArray = handle.Internals->ControlArray->GetBasePointer();
void* execArray = handle.Internals->ExecutionArray;
VTKM_TEST_ASSERT(contArray != nullptr, "No control array after ReleaseResourcesExecution.");
VTKM_TEST_ASSERT(execArray == nullptr,
@ -178,10 +178,10 @@ void TestRoundTrip(bool managed)
vtkm::cont::ArrayHandle<ValueType> handle = CreateArrayHandle<ValueType>(32, managed);
handle.PrepareForOutput(32, vtkm::cont::DeviceAdapterTagCuda());
ValueType* origContArray = handle.Internals->ControlArray.GetArray();
void* origContArray = handle.Internals->ControlArray->GetBasePointer();
{
ValueType* contArray = handle.Internals->ControlArray.GetArray();
ValueType* execArray = handle.Internals->ExecutionArray;
void* contArray = handle.Internals->ControlArray->GetBasePointer();
void* execArray = handle.Internals->ExecutionArray;
VTKM_TEST_ASSERT(contArray != nullptr, "No control array after PrepareForOutput.");
VTKM_TEST_ASSERT(execArray != nullptr, "No execution array after PrepareForOutput.");
VTKM_TEST_ASSERT(CudaAllocator::IsDevicePointer(execArray),
@ -222,8 +222,8 @@ void TestRoundTrip(bool managed)
}
{
ValueType* contArray = handle.Internals->ControlArray.GetArray();
ValueType* execArray = handle.Internals->ExecutionArray;
void* contArray = handle.Internals->ControlArray->GetBasePointer();
void* execArray = handle.Internals->ExecutionArray;
VTKM_TEST_ASSERT(contArray != nullptr, "No control array after GetPortalConst.");
VTKM_TEST_ASSERT(execArray == nullptr, "Execution array not cleared after GetPortalConst.");
VTKM_TEST_ASSERT(CudaAllocator::IsDevicePointer(contArray),

337
vtkm/cont/internal/ArrayHandleBasicImpl.cxx Normal file
View File

@ -0,0 +1,337 @@
//============================================================================
// 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 2017 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2017 UT-Battelle, LLC.
// Copyright 2017 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// 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.
//============================================================================
#define vtkm_cont_internal_ArrayHandleImpl_cxx
#include <vtkm/cont/internal/ArrayHandleBasicImpl.h>
#include <vtkm/cont/cuda/DeviceAdapterCuda.h>
#include <vtkm/cont/serial/DeviceAdapterSerial.h>
#include <vtkm/cont/tbb/DeviceAdapterTBB.h>
#include <vtkm/cont/serial/internal/ExecutionArrayInterfaceBasicSerial.h>
#ifdef VTKM_ENABLE_TBB
#include <vtkm/cont/tbb/internal/ExecutionArrayInterfaceBasicTBB.h>
#endif
#ifdef VTKM_ENABLE_CUDA
#include <vtkm/cont/cuda/internal/ExecutionArrayInterfaceBasicCuda.h>
#endif
namespace vtkm
{
namespace cont
{
namespace internal
{
TypelessExecutionArray::TypelessExecutionArray(const ArrayHandleImpl* data)
: Array(data->ExecutionArray)
, ArrayEnd(data->ExecutionArrayEnd)
, ArrayCapacity(data->ExecutionArrayCapacity)
, ArrayControl(data->ControlArray->GetBasePointer())
, ArrayControlCapacity(data->ControlArray->GetCapacityPointer())
{
}
ExecutionArrayInterfaceBasicBase::ExecutionArrayInterfaceBasicBase(StorageBasicBase& storage)
: ControlStorage(storage)
{
}
ExecutionArrayInterfaceBasicBase::~ExecutionArrayInterfaceBasicBase()
{
}
ArrayHandleImpl::~ArrayHandleImpl()
{
if (this->ExecutionArrayValid && this->ExecutionInterface != nullptr &&
this->ExecutionArray != nullptr)
{
TypelessExecutionArray execArray(this);
this->ExecutionInterface->Free(execArray);
}
delete this->ControlArray;
delete this->ExecutionInterface;
}
void ArrayHandleImpl::CheckControlArrayValid()
{
if (!this->ControlArrayValid)
{
throw vtkm::cont::ErrorInternal(
"ArrayHandle::SyncControlArray did not make control array valid.");
}
}
vtkm::Id ArrayHandleImpl::GetNumberOfValues(vtkm::UInt64 sizeOfT) const
{
if (this->ControlArrayValid)
{
return this->ControlArray->GetNumberOfValues();
}
else if (this->ExecutionArrayValid)
{
auto numBytes =
static_cast<char*>(this->ExecutionArrayEnd) - static_cast<char*>(this->ExecutionArray);
return static_cast<vtkm::Id>(numBytes) / static_cast<vtkm::Id>(sizeOfT);
}
else
{
return 0;
}
}
void ArrayHandleImpl::Allocate(vtkm::Id numberOfValues, vtkm::UInt64 sizeOfT)
{
this->ReleaseResourcesExecutionInternal();
this->ControlArray->AllocateValues(numberOfValues, sizeOfT);
this->ControlArrayValid = true;
}
void ArrayHandleImpl::Shrink(vtkm::Id numberOfValues, vtkm::UInt64 sizeOfT)
{
VTKM_ASSERT(numberOfValues >= 0);
if (numberOfValues > 0)
{
vtkm::Id originalNumberOfValues = this->GetNumberOfValues(sizeOfT);
if (numberOfValues < originalNumberOfValues)
{
if (this->ControlArrayValid)
{
this->ControlArray->Shrink(numberOfValues);
}
if (this->ExecutionArrayValid)
{
auto offset = static_cast<vtkm::UInt64>(numberOfValues) * sizeOfT;
this->ExecutionArrayEnd = static_cast<char*>(this->ExecutionArray) + offset;
}
}
else if (numberOfValues == originalNumberOfValues)
{
// Nothing to do.
}
else // numberOfValues > originalNumberOfValues
{
throw vtkm::cont::ErrorBadValue("ArrayHandle::Shrink cannot be used to grow array.");
}
VTKM_ASSERT(this->GetNumberOfValues(sizeOfT) == numberOfValues);
}
else // numberOfValues == 0
{
// If we are shrinking to 0, there is nothing to save and we might as well
// free up memory. Plus, some storage classes expect that data will be
// deallocated when the size goes to zero.
this->Allocate(0, sizeOfT);
}
}
void ArrayHandleImpl::ReleaseResources()
{
this->ReleaseResourcesExecutionInternal();
if (this->ControlArrayValid)
{
this->ControlArray->ReleaseResources();
this->ControlArrayValid = false;
}
}
void ArrayHandleImpl::PrepareForInput(vtkm::UInt64 sizeOfT) const
{
const vtkm::Id numVals = this->GetNumberOfValues(sizeOfT);
const vtkm::UInt64 numBytes = sizeOfT * static_cast<vtkm::UInt64>(numVals);
if (!this->ExecutionArrayValid)
{
// Initialize an empty array if needed:
if (!this->ControlArrayValid)
{
this->ControlArray->AllocateValues(0, sizeOfT);
this->ControlArrayValid = true;
}
TypelessExecutionArray execArray(this);
this->ExecutionInterface->Allocate(execArray, numVals, sizeOfT);
this->ExecutionInterface->CopyFromControl(
this->ControlArray->GetBasePointer(), this->ExecutionArray, numBytes);
this->ExecutionArrayValid = true;
}
this->ExecutionInterface->UsingForRead(
this->ControlArray->GetBasePointer(), this->ExecutionArray, numBytes);
}
void ArrayHandleImpl::PrepareForOutput(vtkm::Id numVals, vtkm::UInt64 sizeOfT)
{
// Invalidate control arrays since we expect the execution data to be
// overwritten. Don't free control resources in case they're shared with
// the execution environment.
this->ControlArrayValid = false;
TypelessExecutionArray execArray(this);
this->ExecutionInterface->Allocate(execArray, numVals, sizeOfT);
const vtkm::UInt64 numBytes = sizeOfT * static_cast<vtkm::UInt64>(numVals);
this->ExecutionInterface->UsingForWrite(
this->ControlArray->GetBasePointer(), this->ExecutionArray, numBytes);
this->ExecutionArrayValid = true;
}
void ArrayHandleImpl::PrepareForInPlace(vtkm::UInt64 sizeOfT)
{
const vtkm::Id numVals = this->GetNumberOfValues(sizeOfT);
const vtkm::UInt64 numBytes = sizeOfT * static_cast<vtkm::UInt64>(numVals);
if (!this->ExecutionArrayValid)
{
// Initialize an empty array if needed:
if (!this->ControlArrayValid)
{
this->ControlArray->AllocateValues(0, sizeOfT);
this->ControlArrayValid = true;
}
TypelessExecutionArray execArray(this);
this->ExecutionInterface->Allocate(execArray, numVals, sizeOfT);
this->ExecutionInterface->CopyFromControl(
this->ControlArray->GetBasePointer(), this->ExecutionArray, numBytes);
this->ExecutionArrayValid = true;
}
this->ExecutionInterface->UsingForReadWrite(
this->ControlArray->GetBasePointer(), this->ExecutionArray, numBytes);
// Invalidate the control array, since we expect the values to be modified:
this->ControlArrayValid = false;
}
bool ArrayHandleImpl::PrepareForDevice(DeviceAdapterId devId, vtkm::UInt64 sizeOfT) const
{
// Check if the current device matches the last one and sync through
// the control environment if the device changes.
if (this->ExecutionInterface)
{
if (this->ExecutionInterface->GetDeviceId() == devId)
{
// All set, nothing to do.
return false;
}
else
{
// Update the device allocator:
this->SyncControlArray(sizeOfT);
TypelessExecutionArray execArray(this);
this->ExecutionInterface->Free(execArray);
delete this->ExecutionInterface;
this->ExecutionInterface = nullptr;
this->ExecutionArrayValid = false;
}
}
VTKM_ASSERT(this->ExecutionInterface == nullptr);
VTKM_ASSERT(!this->ExecutionArrayValid);
switch (devId)
{
#ifdef VTKM_ENABLE_TBB
case VTKM_DEVICE_ADAPTER_TBB:
this->ExecutionInterface =
new ExecutionArrayInterfaceBasic<DeviceAdapterTagTBB>(*this->ControlArray);
break;
#endif
//this doesn't need to be guarded as a .cu file as it is calling host methods
//and not cuda code directly
#ifdef VTKM_ENABLE_CUDA
case VTKM_DEVICE_ADAPTER_CUDA:
this->ExecutionInterface =
new ExecutionArrayInterfaceBasic<DeviceAdapterTagCuda>(*this->ControlArray);
break;
#endif
default:
this->ExecutionInterface =
new ExecutionArrayInterfaceBasic<DeviceAdapterTagSerial>(*this->ControlArray);
break;
}
return true;
}
DeviceAdapterId ArrayHandleImpl::GetDeviceAdapterId() const
{
return this->ExecutionArrayValid ? this->ExecutionInterface->GetDeviceId()
: VTKM_DEVICE_ADAPTER_UNDEFINED;
}
void ArrayHandleImpl::SyncControlArray(vtkm::UInt64 sizeOfT) const
{
if (!this->ControlArrayValid)
{
// Need to change some state that does not change the logical state from
// an external point of view.
if (this->ExecutionArrayValid)
{
const vtkm::UInt64 numBytes = static_cast<vtkm::UInt64>(
static_cast<char*>(this->ExecutionArrayEnd) - static_cast<char*>(this->ExecutionArray));
const vtkm::Id numVals = static_cast<vtkm::Id>(numBytes / sizeOfT);
this->ControlArray->AllocateValues(numVals, sizeOfT);
this->ExecutionInterface->CopyToControl(
this->ExecutionArray, this->ControlArray->GetBasePointer(), numBytes);
this->ControlArrayValid = true;
}
else
{
// This array is in the null state (there is nothing allocated), but
// the calling function wants to do something with the array. Put this
// class into a valid state by allocating an array of size 0.
this->ControlArray->AllocateValues(0, sizeOfT);
this->ControlArrayValid = true;
}
}
}
void ArrayHandleImpl::ReleaseResourcesExecutionInternal()
{
if (this->ExecutionArrayValid)
{
TypelessExecutionArray execArray(this);
this->ExecutionInterface->Free(execArray);
this->ExecutionArrayValid = false;
}
}
} // end namespace internal
}
} // end vtkm::cont
#ifdef VTKM_MSVC
//Export this when being used with std::shared_ptr
template class VTKM_CONT_EXPORT std::shared_ptr<vtkm::cont::internal::ArrayHandleImpl>;
#endif

157
vtkm/cont/internal/ArrayHandleBasicImpl.h
View File

@ -35,22 +35,13 @@ namespace cont
namespace internal
{
struct ArrayHandleImpl;
/// Type-agnostic container for an execution memory buffer.
struct VTKM_ALWAYS_EXPORT TypelessExecutionArray
struct VTKM_CONT_EXPORT TypelessExecutionArray
{
VTKM_CONT
TypelessExecutionArray(void*& array,
void*& arrayEnd,
void*& arrayCapacity,
const void* arrayControl,
const void* arrayControlCapacity)
: Array(array)
, ArrayEnd(arrayEnd)
, ArrayCapacity(arrayCapacity)
, ArrayControl(arrayControl)
, ArrayControlCapacity(arrayControlCapacity)
{
}
TypelessExecutionArray(const ArrayHandleImpl* data);
void*& Array;
void*& ArrayEnd;
@ -94,7 +85,9 @@ struct VTKM_CONT_EXPORT ExecutionArrayInterfaceBasicBase
/// If (capacity - base) < @a numBytes, the buffer will be freed and
/// reallocated. If (capacity - base) >= numBytes, a new end is marked.
VTKM_CONT
virtual void Allocate(TypelessExecutionArray& execArray, vtkm::UInt64 numBytes) const = 0;
virtual void Allocate(TypelessExecutionArray& execArray,
vtkm::Id numberOfValues,
vtkm::UInt64 sizeOfValue) const = 0;
/// Release the buffer held by @a execArray and reset all pointer to null.
VTKM_CONT
@ -134,13 +127,81 @@ protected:
template <typename DeviceTag>
struct ExecutionArrayInterfaceBasic;
struct VTKM_CONT_EXPORT ArrayHandleImpl
{
VTKM_CONT
template <typename T>
explicit ArrayHandleImpl(T)
: ControlArrayValid(false)
, ControlArray(new vtkm::cont::internal::Storage<T, vtkm::cont::StorageTagBasic>())
, ExecutionInterface(nullptr)
, ExecutionArrayValid(false)
, ExecutionArray(nullptr)
, ExecutionArrayEnd(nullptr)
, ExecutionArrayCapacity(nullptr)
{
}
VTKM_CONT
template <typename T>
explicit ArrayHandleImpl(
const vtkm::cont::internal::Storage<T, vtkm::cont::StorageTagBasic>& storage)
: ControlArrayValid(true)
, ControlArray(new vtkm::cont::internal::Storage<T, vtkm::cont::StorageTagBasic>(storage))
, ExecutionInterface(nullptr)
, ExecutionArrayValid(false)
, ExecutionArray(nullptr)
, ExecutionArrayEnd(nullptr)
, ExecutionArrayCapacity(nullptr)
{
}
VTKM_CONT ~ArrayHandleImpl();
VTKM_CONT ArrayHandleImpl(const ArrayHandleImpl&) = delete;
VTKM_CONT void operator=(const ArrayHandleImpl&) = delete;
//Throws ErrorInternal if ControlArrayValid == false
VTKM_CONT void CheckControlArrayValid() noexcept(false);
VTKM_CONT vtkm::Id GetNumberOfValues(vtkm::UInt64 sizeOfT) const;
VTKM_CONT void Allocate(vtkm::Id numberOfValues, vtkm::UInt64 sizeOfT);
VTKM_CONT void Shrink(vtkm::Id numberOfValues, vtkm::UInt64 sizeOfT);
VTKM_CONT void SyncControlArray(vtkm::UInt64 sizeofT) const;
VTKM_CONT void ReleaseResources();
VTKM_CONT void ReleaseResourcesExecutionInternal();
VTKM_CONT void PrepareForInput(vtkm::UInt64 sizeofT) const;
VTKM_CONT void PrepareForOutput(vtkm::Id numVals, vtkm::UInt64 sizeofT);
VTKM_CONT void PrepareForInPlace(vtkm::UInt64 sizeofT);
// Check if the current device matches the last one. If they don't match
// this moves all data back from execution environment and deletes the
// ExecutionInterface instance.
// Returns true when the caller needs to reallocate ExecutionInterface
VTKM_CONT bool PrepareForDevice(DeviceAdapterId devId, vtkm::UInt64 sizeofT) const;
VTKM_CONT DeviceAdapterId GetDeviceAdapterId() const;
mutable bool ControlArrayValid;
StorageBasicBase* ControlArray;
mutable ExecutionArrayInterfaceBasicBase* ExecutionInterface;
mutable bool ExecutionArrayValid;
mutable void* ExecutionArray;
mutable void* ExecutionArrayEnd;
mutable void* ExecutionArrayCapacity;
};
} // end namespace internal
/// Specialization of ArrayHandle for Basic storage. The goal here is to reduce
/// the amount of codegen for the common case of Basic storage when we build
/// the common arrays into libvtkm_cont.
template <typename T>
class ArrayHandle<T, ::vtkm::cont::StorageTagBasic> : public ::vtkm::cont::internal::ArrayHandleBase
class VTKM_ALWAYS_EXPORT ArrayHandle<T, ::vtkm::cont::StorageTagBasic>
: public ::vtkm::cont::internal::ArrayHandleBase
{
private:
using Thisclass = ArrayHandle<T, ::vtkm::cont::StorageTagBasic>;
@ -154,7 +215,6 @@ public:
using ValueType = T;
using PortalControl = typename StorageType::PortalType;
using PortalConstControl = typename StorageType::PortalConstType;
struct InternalStruct;
template <typename DeviceTag>
struct ExecutionTypes
@ -167,7 +227,6 @@ public:
VTKM_CONT ArrayHandle(const Thisclass& src);
VTKM_CONT ArrayHandle(const Thisclass&& src);
VTKM_CONT ArrayHandle(const StorageType& storage);
VTKM_CONT ArrayHandle(const std::shared_ptr<InternalStruct>& i);
VTKM_CONT ~ArrayHandle();
@ -214,65 +273,19 @@ public:
VTKM_CONT void SyncControlArray() const;
VTKM_CONT void ReleaseResourcesExecutionInternal();
struct VTKM_ALWAYS_EXPORT InternalStruct
{
InternalStruct()
: ControlArrayValid(false)
, ExecutionInterface(nullptr)
, ExecutionArrayValid(false)
, ExecutionArray(nullptr)
, ExecutionArrayEnd(nullptr)
, ExecutionArrayCapacity(nullptr)
{
}
InternalStruct(const StorageType& storage)
: ControlArrayValid(true)
, ControlArray(storage)
, ExecutionInterface(nullptr)
, ExecutionArrayValid(false)
, ExecutionArray(nullptr)
, ExecutionArrayEnd(nullptr)
, ExecutionArrayCapacity(nullptr)
{
}
~InternalStruct()
{
if (this->ExecutionArrayValid && this->ExecutionInterface != nullptr &&
this->ExecutionArray != nullptr)
{
internal::TypelessExecutionArray execArray(
reinterpret_cast<void*&>(this->ExecutionArray),
reinterpret_cast<void*&>(this->ExecutionArrayEnd),
reinterpret_cast<void*&>(this->ExecutionArrayCapacity),
this->ControlArray.GetBasePointer(),
this->ControlArray.GetCapacityPointer());
this->ExecutionInterface->Free(execArray);
}
delete this->ExecutionInterface;
}
InternalStruct(const InternalStruct&) = delete;
void operator=(const InternalStruct&) = delete;
bool ControlArrayValid;
StorageType ControlArray;
internal::ExecutionArrayInterfaceBasicBase* ExecutionInterface;
bool ExecutionArrayValid;
ValueType* ExecutionArray;
ValueType* ExecutionArrayEnd;
ValueType* ExecutionArrayCapacity;
};
std::shared_ptr<InternalStruct> Internals;
std::shared_ptr<internal::ArrayHandleImpl> Internals;
};
} // end namespace cont
} // end namespace vtkm
#ifndef vtkm_cont_internal_ArrayHandleImpl_cxx
#ifdef VTKM_MSVC
extern template class VTKM_CONT_TEMPLATE_EXPORT
std::shared_ptr<vtkm::cont::internal::ArrayHandleImpl>;
#endif
#endif
#include <vtkm/cont/internal/ArrayHandleBasicImpl.hxx>
#endif // vtk_m_cont_internal_ArrayHandleBasicImpl_h

325
vtkm/cont/internal/ArrayHandleBasicImpl.hxx
View File

@ -27,10 +27,9 @@ namespace vtkm
{
namespace cont
{
template <typename T>
ArrayHandle<T, StorageTagBasic>::ArrayHandle()
: Internals(new InternalStruct)
: Internals(new internal::ArrayHandleImpl(T{}))
{
}
@ -48,13 +47,7 @@ ArrayHandle<T, StorageTagBasic>::ArrayHandle(const Thisclass&& src)
template <typename T>
ArrayHandle<T, StorageTagBasic>::ArrayHandle(const StorageType& storage)
: Internals(new InternalStruct(storage))
{
}
template <typename T>
ArrayHandle<T, StorageTagBasic>::ArrayHandle(const std::shared_ptr<InternalStruct>& i)
: Internals(i)
: Internals(new internal::ArrayHandleImpl(storage))
{
}
@ -107,15 +100,11 @@ template <typename T>
typename ArrayHandle<T, StorageTagBasic>::StorageType& ArrayHandle<T, StorageTagBasic>::GetStorage()
{
this->SyncControlArray();
if (this->Internals->ControlArrayValid)
{
return this->Internals->ControlArray;
}
else
{
throw vtkm::cont::ErrorInternal(
"ArrayHandle::SyncControlArray did not make control array valid.");
}
this->Internals->CheckControlArrayValid();
//CheckControlArrayValid will throw an exception if this->Internals->ControlArrayValid
//is not valid
return *(static_cast<StorageType*>(this->Internals->ControlArray));
}
template <typename T>
@ -123,15 +112,11 @@ const typename ArrayHandle<T, StorageTagBasic>::StorageType&
ArrayHandle<T, StorageTagBasic>::GetStorage() const
{
this->SyncControlArray();
if (this->Internals->ControlArrayValid)
{
return this->Internals->ControlArray;
}
else
{
throw vtkm::cont::ErrorInternal(
"ArrayHandle::SyncControlArray did not make control array valid.");
}
this->Internals->CheckControlArrayValid();
//CheckControlArrayValid will throw an exception if this->Internals->ControlArrayValid
//is not valid
return *(static_cast<const StorageType*>(this->Internals->ControlArray));
}
template <typename T>
@ -139,19 +124,17 @@ typename ArrayHandle<T, StorageTagBasic>::PortalControl
ArrayHandle<T, StorageTagBasic>::GetPortalControl()
{
this->SyncControlArray();
if (this->Internals->ControlArrayValid)
{
// If the user writes into the iterator we return, then the execution
// array will become invalid. Play it safe and release the execution
// resources. (Use the const version to preserve the execution array.)
this->ReleaseResourcesExecutionInternal();
return this->Internals->ControlArray.GetPortal();
}
else
{
throw vtkm::cont::ErrorInternal(
"ArrayHandle::SyncControlArray did not make control array valid.");
}
this->Internals->CheckControlArrayValid();
//CheckControlArrayValid will throw an exception if this->Internals->ControlArrayValid
//is not valid
// If the user writes into the iterator we return, then the execution
// array will become invalid. Play it safe and release the execution
// resources. (Use the const version to preserve the execution array.)
this->ReleaseResourcesExecutionInternal();
StorageType* privStorage = static_cast<StorageType*>(this->Internals->ControlArray);
return privStorage->GetPortal();
}
@ -160,81 +143,30 @@ typename ArrayHandle<T, StorageTagBasic>::PortalConstControl
ArrayHandle<T, StorageTagBasic>::GetPortalConstControl() const
{
this->SyncControlArray();
if (this->Internals->ControlArrayValid)
{
return this->Internals->ControlArray.GetPortalConst();
}
else
{
throw vtkm::cont::ErrorInternal(
"ArrayHandle::SyncControlArray did not make control array valid.");
}
this->Internals->CheckControlArrayValid();
//CheckControlArrayValid will throw an exception if this->Internals->ControlArrayValid
//is not valid
StorageType* privStorage = static_cast<StorageType*>(this->Internals->ControlArray);
return privStorage->GetPortalConst();
}
template <typename T>
vtkm::Id ArrayHandle<T, StorageTagBasic>::GetNumberOfValues() const
{
if (this->Internals->ControlArrayValid)
{
return this->Internals->ControlArray.GetNumberOfValues();
}
else if (this->Internals->ExecutionArrayValid)
{
return static_cast<vtkm::Id>(this->Internals->ExecutionArrayEnd -
this->Internals->ExecutionArray);
}
else
{
return 0;
}
return this->Internals->GetNumberOfValues(sizeof(T));
}
template <typename T>
void ArrayHandle<T, StorageTagBasic>::Allocate(vtkm::Id numberOfValues)
{
this->ReleaseResourcesExecutionInternal();
this->Internals->ControlArray.Allocate(numberOfValues);
this->Internals->ControlArrayValid = true;
this->Internals->Allocate(numberOfValues, sizeof(T));
}
template <typename T>
void ArrayHandle<T, StorageTagBasic>::Shrink(vtkm::Id numberOfValues)
{
VTKM_ASSERT(numberOfValues >= 0);
if (numberOfValues > 0)
{
vtkm::Id originalNumberOfValues = this->GetNumberOfValues();
if (numberOfValues < originalNumberOfValues)
{
if (this->Internals->ControlArrayValid)
{
this->Internals->ControlArray.Shrink(numberOfValues);
}
if (this->Internals->ExecutionArrayValid)
{
this->Internals->ExecutionArrayEnd = this->Internals->ExecutionArray + numberOfValues;
}
}
else if (numberOfValues == originalNumberOfValues)
{
// Nothing to do.
}
else // numberOfValues > originalNumberOfValues
{
throw vtkm::cont::ErrorBadValue("ArrayHandle::Shrink cannot be used to grow array.");
}
VTKM_ASSERT(this->GetNumberOfValues() == numberOfValues);
}
else // numberOfValues == 0
{
// If we are shrinking to 0, there is nothing to save and we might as well
// free up memory. Plus, some storage classes expect that data will be
// deallocated when the size goes to zero.
this->Allocate(0);
}
this->Internals->Shrink(numberOfValues, sizeof(T));
}
template <typename T>
@ -243,19 +175,13 @@ void ArrayHandle<T, StorageTagBasic>::ReleaseResourcesExecution()
// Save any data in the execution environment by making sure it is synced
// with the control environment.
this->SyncControlArray();
this->ReleaseResourcesExecutionInternal();
this->Internals->ReleaseResourcesExecutionInternal();
}
template <typename T>
void ArrayHandle<T, StorageTagBasic>::ReleaseResources()
{
this->ReleaseResourcesExecutionInternal();
if (this->Internals->ControlArrayValid)
{
this->Internals->ControlArray.ReleaseResources();
this->Internals->ControlArrayValid = false;
}
this->Internals->ReleaseResources();
}
template <typename T>
@ -264,41 +190,12 @@ typename ArrayHandle<T, StorageTagBasic>::template ExecutionTypes<DeviceAdapterT
ArrayHandle<T, StorageTagBasic>::PrepareForInput(DeviceAdapterTag device) const
{
VTKM_IS_DEVICE_ADAPTER_TAG(DeviceAdapterTag);
InternalStruct* priv = const_cast<InternalStruct*>(this->Internals.get());
this->PrepareForDevice(device);
const vtkm::UInt64 numBytes = static_cast<vtkm::UInt64>(sizeof(ValueType)) *
static_cast<vtkm::UInt64>(this->GetNumberOfValues());
if (!this->Internals->ExecutionArrayValid)
{
// Initialize an empty array if needed:
if (!this->Internals->ControlArrayValid)
{
this->Internals->ControlArray.Allocate(0);
this->Internals->ControlArrayValid = true;
}
internal::TypelessExecutionArray execArray(
reinterpret_cast<void*&>(priv->ExecutionArray),
reinterpret_cast<void*&>(priv->ExecutionArrayEnd),
reinterpret_cast<void*&>(priv->ExecutionArrayCapacity),
this->Internals->ControlArray.GetBasePointer(),
this->Internals->ControlArray.GetCapacityPointer());
priv->ExecutionInterface->Allocate(execArray, numBytes);
priv->ExecutionInterface->CopyFromControl(
priv->ControlArray.GetArray(), priv->ExecutionArray, numBytes);
this->Internals->ExecutionArrayValid = true;
}
this->Internals->ExecutionInterface->UsingForRead(
priv->ControlArray.GetArray(), priv->ExecutionArray, numBytes);
return PortalFactory<DeviceAdapterTag>::CreatePortalConst(this->Internals->ExecutionArray,
this->Internals->ExecutionArrayEnd);
this->Internals->PrepareForInput(sizeof(T));
return PortalFactory<DeviceAdapterTag>::CreatePortalConst(
static_cast<T*>(this->Internals->ExecutionArray),
static_cast<T*>(this->Internals->ExecutionArrayEnd));
}
template <typename T>
@ -307,34 +204,12 @@ typename ArrayHandle<T, StorageTagBasic>::template ExecutionTypes<DeviceAdapterT
ArrayHandle<T, StorageTagBasic>::PrepareForOutput(vtkm::Id numVals, DeviceAdapterTag device)
{
VTKM_IS_DEVICE_ADAPTER_TAG(DeviceAdapterTag);
InternalStruct* priv = const_cast<InternalStruct*>(this->Internals.get());
this->PrepareForDevice(device);
// Invalidate control arrays since we expect the execution data to be
// overwritten. Don't free control resources in case they're shared with
// the execution environment.
this->Internals->ControlArrayValid = false;
internal::TypelessExecutionArray execArray(reinterpret_cast<void*&>(priv->ExecutionArray),
reinterpret_cast<void*&>(priv->ExecutionArrayEnd),
reinterpret_cast<void*&>(priv->ExecutionArrayCapacity),
this->Internals->ControlArray.GetBasePointer(),
this->Internals->ControlArray.GetCapacityPointer());
const vtkm::UInt64 numBytes =
static_cast<vtkm::UInt64>(sizeof(ValueType)) * static_cast<vtkm::UInt64>(numVals);
this->Internals->ExecutionInterface->Allocate(execArray, numBytes);
this->Internals->ExecutionInterface->UsingForWrite(
priv->ControlArray.GetArray(), priv->ExecutionArray, numBytes);
this->Internals->ExecutionArrayValid = true;
return PortalFactory<DeviceAdapterTag>::CreatePortal(this->Internals->ExecutionArray,
this->Internals->ExecutionArrayEnd);
this->Internals->PrepareForOutput(numVals, sizeof(T));
return PortalFactory<DeviceAdapterTag>::CreatePortal(
static_cast<T*>(this->Internals->ExecutionArray),
static_cast<T*>(this->Internals->ExecutionArrayEnd));
}
template <typename T>
@ -343,45 +218,12 @@ typename ArrayHandle<T, StorageTagBasic>::template ExecutionTypes<DeviceAdapterT
ArrayHandle<T, StorageTagBasic>::PrepareForInPlace(DeviceAdapterTag device)
{
VTKM_IS_DEVICE_ADAPTER_TAG(DeviceAdapterTag);
InternalStruct* priv = const_cast<InternalStruct*>(this->Internals.get());
this->PrepareForDevice(device);
const vtkm::UInt64 numBytes = static_cast<vtkm::UInt64>(sizeof(ValueType)) *
static_cast<vtkm::UInt64>(this->GetNumberOfValues());
if (!this->Internals->ExecutionArrayValid)
{
// Initialize an empty array if needed:
if (!this->Internals->ControlArrayValid)
{
this->Internals->ControlArray.Allocate(0);
this->Internals->ControlArrayValid = true;
}
internal::TypelessExecutionArray execArray(
reinterpret_cast<void*&>(this->Internals->ExecutionArray),
reinterpret_cast<void*&>(this->Internals->ExecutionArrayEnd),
reinterpret_cast<void*&>(this->Internals->ExecutionArrayCapacity),
this->Internals->ControlArray.GetBasePointer(),
this->Internals->ControlArray.GetCapacityPointer());
priv->ExecutionInterface->Allocate(execArray, numBytes);
priv->ExecutionInterface->CopyFromControl(
priv->ControlArray.GetArray(), priv->ExecutionArray, numBytes);
this->Internals->ExecutionArrayValid = true;
}
priv->ExecutionInterface->UsingForReadWrite(
priv->ControlArray.GetArray(), priv->ExecutionArray, numBytes);
// Invalidate the control array, since we expect the values to be modified:
this->Internals->ControlArrayValid = false;
return PortalFactory<DeviceAdapterTag>::CreatePortal(this->Internals->ExecutionArray,
this->Internals->ExecutionArrayEnd);
this->Internals->PrepareForInPlace(sizeof(T));
return PortalFactory<DeviceAdapterTag>::CreatePortal(
static_cast<T*>(this->Internals->ExecutionArray),
static_cast<T*>(this->Internals->ExecutionArrayEnd));
}
template <typename T>
@ -389,92 +231,25 @@ template <typename DeviceAdapterTag>
void ArrayHandle<T, StorageTagBasic>::PrepareForDevice(DeviceAdapterTag) const
{
DeviceAdapterId devId = DeviceAdapterTraits<DeviceAdapterTag>::GetId();
InternalStruct* priv = const_cast<InternalStruct*>(this->Internals.get());
// Check if the current device matches the last one and sync through
// the control environment if the device changes.
if (this->Internals->ExecutionInterface)
{
if (this->Internals->ExecutionInterface->GetDeviceId() == devId)
{
// All set, nothing to do.
return;
}
else
{
// Update the device allocator:
this->SyncControlArray();
internal::TypelessExecutionArray execArray(
reinterpret_cast<void*&>(priv->ExecutionArray),
reinterpret_cast<void*&>(priv->ExecutionArrayEnd),
reinterpret_cast<void*&>(priv->ExecutionArrayCapacity),
this->Internals->ControlArray.GetBasePointer(),
this->Internals->ControlArray.GetCapacityPointer());
priv->ExecutionInterface->Free(execArray);
delete priv->ExecutionInterface;
priv->ExecutionInterface = nullptr;
priv->ExecutionArrayValid = false;
}
}
VTKM_ASSERT(priv->ExecutionInterface == nullptr);
VTKM_ASSERT(!priv->ExecutionArrayValid);
priv->ExecutionInterface =
new internal::ExecutionArrayInterfaceBasic<DeviceAdapterTag>(this->Internals->ControlArray);
this->Internals->PrepareForDevice(devId, sizeof(T));
}
template <typename T>
DeviceAdapterId ArrayHandle<T, StorageTagBasic>::GetDeviceAdapterId() const
{
return this->Internals->ExecutionArrayValid ? this->Internals->ExecutionInterface->GetDeviceId()
: VTKM_DEVICE_ADAPTER_UNDEFINED;
return this->Internals->GetDeviceAdapterId();
}
template <typename T>
void ArrayHandle<T, StorageTagBasic>::SyncControlArray() const
{
if (!this->Internals->ControlArrayValid)
{
// Need to change some state that does not change the logical state from
// an external point of view.
InternalStruct* priv = const_cast<InternalStruct*>(this->Internals.get());
if (this->Internals->ExecutionArrayValid)
{
const vtkm::Id numValues =
static_cast<vtkm::Id>(this->Internals->ExecutionArrayEnd - this->Internals->ExecutionArray);
const vtkm::UInt64 numBytes =
static_cast<vtkm::UInt64>(sizeof(ValueType)) * static_cast<vtkm::UInt64>(numValues);
priv->ControlArray.Allocate(numValues);
priv->ExecutionInterface->CopyToControl(
priv->ExecutionArray, priv->ControlArray.GetArray(), numBytes);
priv->ControlArrayValid = true;
}
else
{
// This array is in the null state (there is nothing allocated), but
// the calling function wants to do something with the array. Put this
// class into a valid state by allocating an array of size 0.
priv->ControlArray.Allocate(0);
priv->ControlArrayValid = true;
}
}
this->Internals->SyncControlArray(sizeof(T));
}
template <typename T>
void ArrayHandle<T, StorageTagBasic>::ReleaseResourcesExecutionInternal()
{
if (this->Internals->ExecutionArrayValid)
{
internal::TypelessExecutionArray execArray(
reinterpret_cast<void*&>(this->Internals->ExecutionArray),
reinterpret_cast<void*&>(this->Internals->ExecutionArrayEnd),
reinterpret_cast<void*&>(this->Internals->ExecutionArrayCapacity),
this->Internals->ControlArray.GetBasePointer(),
this->Internals->ControlArray.GetCapacityPointer());
this->Internals->ExecutionInterface->Free(execArray);
this->Internals->ExecutionArrayValid = false;
}
this->Internals->ReleaseResourcesExecutionInternal();
}
}
} // end namespace vtkm::cont

8
vtkm/cont/internal/ArrayManagerExecutionShareWithControl.cxx
View File

@ -34,12 +34,12 @@ ExecutionArrayInterfaceBasicShareWithControl::ExecutionArrayInterfaceBasicShareW
}
void ExecutionArrayInterfaceBasicShareWithControl::Allocate(TypelessExecutionArray& execArray,
vtkm::UInt64 numBytes) const
vtkm::Id numberOfValues,
vtkm::UInt64 sizeOfValue) const
{
this->ControlStorage.AllocateBytes(numBytes);
this->ControlStorage.AllocateValues(numberOfValues, sizeOfValue);
execArray.Array = this->ControlStorage.GetBasePointer();
execArray.ArrayEnd = this->ControlStorage.GetEndPointer();
execArray.ArrayEnd = this->ControlStorage.GetEndPointer(numberOfValues, sizeOfValue);
execArray.ArrayCapacity = this->ControlStorage.GetCapacityPointer();
}

4
vtkm/cont/internal/ArrayManagerExecutionShareWithControl.h
View File

@ -139,7 +139,9 @@ struct VTKM_CONT_EXPORT ExecutionArrayInterfaceBasicShareWithControl
VTKM_CONT ExecutionArrayInterfaceBasicShareWithControl(StorageBasicBase& storage);
VTKM_CONT void Allocate(TypelessExecutionArray& execArray, vtkm::UInt64 numBytes) const final;
VTKM_CONT void Allocate(TypelessExecutionArray& execArray,
vtkm::Id numberOfValues,
vtkm::UInt64 sizeOfValue) const final;
VTKM_CONT void Free(TypelessExecutionArray& execArray) const final;
VTKM_CONT void CopyFromControl(const void* src, void* dst, vtkm::UInt64 bytes) const final;

14
vtkm/cont/serial/internal/ArrayManagerExecutionSerial.h
View File

@ -24,6 +24,7 @@
#include <vtkm/cont/internal/ArrayManagerExecution.h>
#include <vtkm/cont/internal/ArrayManagerExecutionShareWithControl.h>
#include <vtkm/cont/serial/internal/DeviceAdapterTagSerial.h>
#include <vtkm/cont/serial/internal/ExecutionArrayInterfaceBasicSerial.h>
namespace vtkm
{
@ -62,19 +63,6 @@ struct ExecutionPortalFactoryBasic<T, DeviceAdapterTagSerial>
using Superclass::CreatePortalConst;
};
template <>
struct VTKM_CONT_EXPORT ExecutionArrayInterfaceBasic<DeviceAdapterTagSerial>
: public ExecutionArrayInterfaceBasicShareWithControl
{
using Superclass = ExecutionArrayInterfaceBasicShareWithControl;
VTKM_CONT
ExecutionArrayInterfaceBasic(StorageBasicBase& storage);
VTKM_CONT
virtual DeviceAdapterId GetDeviceId() const final { return VTKM_DEVICE_ADAPTER_SERIAL; }
};
} // namespace internal
#ifndef vtk_m_cont_serial_internal_ArrayManagerExecutionSerial_cxx

2
vtkm/cont/serial/internal/CMakeLists.txt
View File

@ -22,6 +22,7 @@ set(headers
ArrayManagerExecutionSerial.h
DeviceAdapterAlgorithmSerial.h
DeviceAdapterTagSerial.h
ExecutionArrayInterfaceBasicSerial.h
VirtualObjectTransferSerial.h
)
vtkm_declare_headers(${headers})
@ -29,6 +30,7 @@ vtkm_declare_headers(${headers})
add_library(vtkm_cont_serial OBJECT
ArrayManagerExecutionSerial.cxx
DeviceAdapterAlgorithmSerial.cxx
ExecutionArrayInterfaceBasicSerial.cxx
)

36
vtkm/cont/serial/internal/ExecutionArrayInterfaceBasicSerial.cxx Normal file
View File

@ -0,0 +1,36 @@
//============================================================================
// 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 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// 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.
//============================================================================
#include <vtkm/cont/serial/internal/ExecutionArrayInterfaceBasicSerial.h>
namespace vtkm
{
namespace cont
{
namespace internal
{
vtkm::cont::DeviceAdapterId ExecutionArrayInterfaceBasic<DeviceAdapterTagSerial>::GetDeviceId()
const
{
return VTKM_DEVICE_ADAPTER_SERIAL;
}
} // namespace internal
}
} // namespace vtkm::cont

50
vtkm/cont/serial/internal/ExecutionArrayInterfaceBasicSerial.h Normal file
View File

@ -0,0 +1,50 @@
//============================================================================
// 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 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// 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_cont_serial_internal_ExecutionArrayInterfaceBasicSerial_h
#define vtk_m_cont_serial_internal_ExecutionArrayInterfaceBasicSerial_h
#include <vtkm/cont/internal/ArrayManagerExecutionShareWithControl.h>
#include <vtkm/cont/serial/internal/DeviceAdapterTagSerial.h>
namespace vtkm
{
namespace cont
{
namespace internal
{
template <>
struct VTKM_CONT_EXPORT ExecutionArrayInterfaceBasic<DeviceAdapterTagSerial>
: public ExecutionArrayInterfaceBasicShareWithControl
{
using Superclass = ExecutionArrayInterfaceBasicShareWithControl;
VTKM_CONT
ExecutionArrayInterfaceBasic(StorageBasicBase& storage);
VTKM_CONT
DeviceAdapterId GetDeviceId() const final;
};
} // namespace internal
}
} // namespace vtkm::cont
#endif //vtk_m_cont_serial_internal_ExecutionArrayInterfaceBasicSerial_h

18
vtkm/cont/tbb/internal/ArrayManagerExecutionTBB.h
View File

@ -20,12 +20,13 @@
#ifndef vtk_m_cont_tbb_internal_ArrayManagerExecutionTBB_h
#define vtk_m_cont_tbb_internal_ArrayManagerExecutionTBB_h
#include <vtkm/cont/tbb/internal/DeviceAdapterTagTBB.h>
#include <vtkm/cont/internal/ArrayExportMacros.h>
#include <vtkm/cont/internal/ArrayManagerExecution.h>
#include <vtkm/cont/internal/ArrayManagerExecutionShareWithControl.h>
#include <vtkm/cont/tbb/internal/DeviceAdapterTagTBB.h>
#include <vtkm/cont/tbb/internal/ExecutionArrayInterfaceBasicTBB.h>
// These must be placed in the vtkm::cont::internal namespace so that
// the template can be found.
@ -86,19 +87,6 @@ struct ExecutionPortalFactoryBasic<T, DeviceAdapterTagTBB>
using Superclass::CreatePortalConst;
};
template <>
struct VTKM_CONT_EXPORT ExecutionArrayInterfaceBasic<DeviceAdapterTagTBB>
: public ExecutionArrayInterfaceBasicShareWithControl
{
using Superclass = ExecutionArrayInterfaceBasicShareWithControl;
VTKM_CONT
ExecutionArrayInterfaceBasic(StorageBasicBase& storage);
VTKM_CONT
virtual DeviceAdapterId GetDeviceId() const final { return VTKM_DEVICE_ADAPTER_TBB; }
};
} // namespace internal
#ifndef vtk_m_cont_tbb_internal_ArrayManagerExecutionTBB_cxx
VTKM_EXPORT_ARRAYHANDLES_FOR_DEVICE_ADAPTER(DeviceAdapterTagTBB)

2
vtkm/cont/tbb/internal/CMakeLists.txt
View File

@ -22,6 +22,7 @@ set(headers
ArrayManagerExecutionTBB.h
DeviceAdapterAlgorithmTBB.h
DeviceAdapterTagTBB.h
ExecutionArrayInterfaceBasicTBB.h
FunctorsTBB.h
ParallelSortTBB.h
VirtualObjectTransferTBB.h
@ -48,6 +49,7 @@ endif()
add_library(vtkm_cont_tbb OBJECT
ArrayManagerExecutionTBB.cxx
DeviceAdapterAlgorithmTBB.cxx
ExecutionArrayInterfaceBasicTBB.cxx
ParallelSortTBB.cxx
)

35
vtkm/cont/tbb/internal/ExecutionArrayInterfaceBasicTBB.cxx Normal file
View File

@ -0,0 +1,35 @@
//============================================================================
// 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 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// 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.
//============================================================================
#include <vtkm/cont/tbb/internal/ExecutionArrayInterfaceBasicTBB.h>
namespace vtkm
{
namespace cont
{
namespace internal
{
vtkm::cont::DeviceAdapterId ExecutionArrayInterfaceBasic<DeviceAdapterTagTBB>::GetDeviceId() const
{
return VTKM_DEVICE_ADAPTER_TBB;
}
} // namespace internal
}
} // namespace vtkm::cont

53
vtkm/cont/tbb/internal/ExecutionArrayInterfaceBasicTBB.h Normal file
View File

@ -0,0 +1,53 @@
//============================================================================
// 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 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2014 UT-Battelle, LLC.
// Copyright 2014 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// 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_cont_tbb_internal_ExecutionArrayInterfaceBasicTBB_h
#define vtk_m_cont_tbb_internal_ExecutionArrayInterfaceBasicTBB_h
#include <vtkm/cont/internal/ArrayManagerExecutionShareWithControl.h>
#include <vtkm/cont/tbb/internal/DeviceAdapterTagTBB.h>
// These must be placed in the vtkm::cont::internal namespace so that
// the template can be found.
namespace vtkm
{
namespace cont
{
namespace internal
{
template <>
struct VTKM_CONT_EXPORT ExecutionArrayInterfaceBasic<DeviceAdapterTagTBB>
: public ExecutionArrayInterfaceBasicShareWithControl
{
using Superclass = ExecutionArrayInterfaceBasicShareWithControl;
VTKM_CONT
ExecutionArrayInterfaceBasic(StorageBasicBase& storage);
VTKM_CONT
DeviceAdapterId GetDeviceId() const final;
};
} // namespace internal
}
} // namespace vtkm::cont
#endif //vtk_m_cont_tbb_internal_ExecutionArrayInterfaceBasicTBB_h

2
vtkm/cont/testing/UnitTestDeviceAdapterAlgorithmGeneral.cxx
View File

@ -126,7 +126,7 @@ struct ExecutionArrayInterfaceBasic<DeviceAdapterTagTestAlgorithmGeneral>
}
VTKM_CONT
virtual DeviceAdapterId GetDeviceId() const final { return -3; }
DeviceAdapterId GetDeviceId() const final { return -3; }
};
}
}

62
vtkm/cont/testing/UnitTestStorageBasic.cxx
View File

@ -26,16 +26,6 @@
#include <vtkm/VecTraits.h>
#include <vtkm/cont/testing/Testing.h>
// We use these to check if the aligned allocator provided by
// StorageBasic can be used with all STL containers
#include <deque>
#include <list>
#include <map>
#include <queue>
#include <set>
#include <stack>
#include <vector>
namespace
{
@ -72,53 +62,6 @@ struct TemplatedTests
typename vtkm::VecTraits<ValueType>::ComponentType STOLEN_ARRAY_VALUE() { return 29; }
void TestAlignedAllocatorSTL()
{
using Allocator = typename StorageType::AllocatorType;
std::vector<ValueType, Allocator> vec(ARRAY_SIZE, ValueType());
StorageType store(&vec[0], ARRAY_SIZE);
}
// This test checks that we can compile and use the allocator with all
// STL containers
void CompileSTLAllocator()
{
using Allocator = typename StorageType::AllocatorType;
using PairAllocator =
typename StorageType::AllocatorType::template rebind<std::pair<ValueType, ValueType>>::other;
std::vector<ValueType, Allocator> v;
ValueType value = vtkm::TypeTraits<ValueType>::ZeroInitialization();
v.push_back(value);
std::deque<ValueType, Allocator> d;
d.push_front(value);
std::list<ValueType, Allocator> l;
l.push_front(value);
std::set<ValueType, std::less<ValueType>, Allocator> set;
set.insert(value);
std::map<ValueType, ValueType, std::less<ValueType>, PairAllocator> m;
m[value] = value;
std::multiset<ValueType, std::less<ValueType>, Allocator> ms;
ms.insert(value);
std::multimap<ValueType, ValueType, std::less<ValueType>, PairAllocator> mm;
mm.insert(std::pair<ValueType, ValueType>(value, value));
std::stack<ValueType, std::deque<ValueType, Allocator>> stack;
stack.push(value);
std::queue<ValueType, std::deque<ValueType, Allocator>> queue;
queue.push(value);
std::priority_queue<ValueType, std::vector<ValueType, Allocator>> pqueue;
pqueue.push(value);
}
/// Returned value should later be passed to StealArray2. It is best to
/// put as much between the two test parts to maximize the chance of a
/// deallocated array being overridden (and thus detected).
@ -153,7 +96,7 @@ struct TemplatedTests
"Stolen array did not retain values.");
}
typename StorageType::AllocatorType allocator;
allocator.deallocate(stolenArray, ARRAY_SIZE);
allocator.deallocate(stolenArray);
}
void BasicAllocation()
@ -199,9 +142,6 @@ struct TemplatedTests
BasicAllocation();
StealArray2(stolenArray);
TestAlignedAllocatorSTL();
CompileSTLAllocator();
}
};

5
vtkm/internal/Configure.h.in
View File

@ -113,6 +113,11 @@
# endif
#endif
// Defines the cache line size in bytes to align allocations to
#ifndef VTKM_ALLOCATION_ALIGNMENT
#define VTKM_ALLOCATION_ALIGNMENT 64
#endif
#ifdef VTKM_USE_64BIT_IDS
# ifndef VTKM_SIZE_ID
# define VTKM_SIZE_ID 8

5
vtkm/interop/testing/TestingOpenGLInterop.h
View File

@ -177,7 +177,6 @@ private:
template <typename T>
void operator()(const T t) const
{
std::cout << "TransferFunctor" << std::endl;
const std::size_t Size = 10;
GLuint GLHandle;
//verify that T is able to be transfer to openGL.
@ -197,7 +196,7 @@ private:
//verify the results match what is in the array handle
temp.SyncControlArray();
T* expectedValues = temp.Internals->ControlArray.StealArray();
T* expectedValues = temp.GetStorage().GetArray();
for (std::size_t i = 0; i < Size; ++i)
{
@ -214,7 +213,7 @@ private:
returnedValues = CopyGLBuffer(GLHandle, t);
//verify the results match what is in the array handle
temp.SyncControlArray();
expectedValues = temp.Internals->ControlArray.StealArray();
expectedValues = temp.GetStorage().GetArray();
for (std::size_t i = 0; i < Size * 2; ++i)
{

13
vtkm/worklet/Clip.h
View File

@ -134,9 +134,9 @@ private:
struct ClipStats
{
vtkm::Id NumberOfCells;
vtkm::Id NumberOfIndices;
vtkm::Id NumberOfNewPoints;
vtkm::Id NumberOfCells = 0;
vtkm::Id NumberOfIndices = 0;
vtkm::Id NumberOfNewPoints = 0;
struct SumOp
{
@ -154,8 +154,9 @@ struct ClipStats
struct EdgeInterpolation
{
vtkm::Id Vertex1, Vertex2;
vtkm::Float64 Weight;
vtkm::Id Vertex1 = -1;
vtkm::Id Vertex2 = -1;
vtkm::Float64 Weight = 0;
struct LessThanOp
{
@ -449,7 +450,7 @@ public:
.Invoke(cellSet, scalars, clipTableIdxs, stats);
// compute offsets for each invocation
ClipStats zero = { 0, 0, 0 };
ClipStats zero;
vtkm::cont::ArrayHandle<ClipStats> cellSetIndices;
ClipStats total = Algorithm::ScanExclusive(stats, cellSetIndices, ClipStats::SumOp(), zero);
stats.ReleaseResources();