diff --git a/vtkm/cont/CMakeLists.txt b/vtkm/cont/CMakeLists.txt index 0e07b874a..c2a27539b 100644 --- a/vtkm/cont/CMakeLists.txt +++ b/vtkm/cont/CMakeLists.txt @@ -146,6 +146,9 @@ set(sources ErrorBadType.cxx internal/ArrayHandleBasicImpl.cxx internal/ArrayManagerExecutionShareWithControl.cxx + internal/Buffer.cxx + internal/DeviceAdapterMemoryManager.cxx + internal/DeviceAdapterMemoryManagerShared.cxx internal/TransferInfo.cxx internal/VirtualObjectTransfer.cxx Initialize.cxx diff --git a/vtkm/cont/DeviceAdapter.h b/vtkm/cont/DeviceAdapter.h index ffc631823..d66458451 100644 --- a/vtkm/cont/DeviceAdapter.h +++ b/vtkm/cont/DeviceAdapter.h @@ -12,7 +12,7 @@ // These are listed in non-alphabetical order because this is the conceptual // order in which the sub-files are loaded. (But the compile should still -// succeed of the order is changed.) Turn off formatting to keep the order. +// succeed if the order is changed.) Turn off formatting to keep the order. // clang-format off #include @@ -23,6 +23,7 @@ #include #include #include +#include // clang-format on diff --git a/vtkm/cont/RuntimeDeviceInformation.cxx b/vtkm/cont/RuntimeDeviceInformation.cxx index 888cad222..1051f1809 100644 --- a/vtkm/cont/RuntimeDeviceInformation.cxx +++ b/vtkm/cont/RuntimeDeviceInformation.cxx @@ -13,6 +13,7 @@ #include #include #include +#include //Bring in each device adapters runtime class #include @@ -24,6 +25,62 @@ namespace { +class DeviceAdapterMemoryManagerInvalid + : public vtkm::cont::internal::DeviceAdapterMemoryManagerBase +{ +public: + VTKM_CONT virtual ~DeviceAdapterMemoryManagerInvalid() override {} + + /// Allocates a buffer of the specified size in bytes and returns a BufferInfo object + /// containing information about it. + VTKM_CONT virtual std::shared_ptr Allocate( + vtkm::BufferSizeType) override + { + throw vtkm::cont::ErrorBadDevice("Tried to manage memory on an invalid device."); + } + + /// Manages the provided array. Returns a `BufferInfo` object that contains the data. + /// The deleter in the `shared_ptr` is expected to correctly free the data. + VTKM_CONT virtual std::shared_ptr ManageArray( + std::shared_ptr, + vtkm::BufferSizeType) override + { + throw vtkm::cont::ErrorBadDevice("Tried to manage memory on an invalid device."); + } + + /// Reallocates the provided buffer to a new size. The passed in `BufferInfo` should be + /// modified to reflect the changes. + VTKM_CONT virtual void Reallocate(std::shared_ptr, + vtkm::BufferSizeType) override + { + throw vtkm::cont::ErrorBadDevice("Tried to manage memory on an invalid device."); + } + + /// Copies data from the host buffer provided onto the device and returns a buffer info + /// object holding the pointer for the device. + VTKM_CONT virtual std::shared_ptr CopyHostToDevice( + std::shared_ptr) override + { + throw vtkm::cont::ErrorBadDevice("Tried to manage memory on an invalid device."); + } + + /// Copies data from the device buffer provided to the host. The passed in `BufferInfo` object + /// was created by a previous call to this object. + VTKM_CONT virtual std::shared_ptr CopyDeviceToHost( + std::shared_ptr) override + { + throw vtkm::cont::ErrorBadDevice("Tried to manage memory on an invalid device."); + } + + /// Copies data from one device buffer to another device buffer. The passed in `BufferInfo` object + /// was created by a previous call to this object. + VTKM_CONT virtual std::shared_ptr CopyDeviceToDevice( + std::shared_ptr) override + { + throw vtkm::cont::ErrorBadDevice("Tried to manage memory on an invalid device."); + } +}; + struct VTKM_NEVER_EXPORT InitializeDeviceNames { vtkm::cont::DeviceAdapterNameType* Names; @@ -58,6 +115,42 @@ struct VTKM_NEVER_EXPORT InitializeDeviceNames } }; +struct VTKM_NEVER_EXPORT InitializeDeviceMemoryManagers +{ + std::unique_ptr* Managers; + + VTKM_CONT + InitializeDeviceMemoryManagers( + std::unique_ptr* managers) + : Managers(managers) + { + } + + template + VTKM_CONT void CreateManager(Device device, std::true_type) + { + auto id = device.GetValue(); + + if (id > 0 && id < VTKM_MAX_DEVICE_ADAPTER_ID) + { + auto name = vtkm::cont::DeviceAdapterTraits::GetName(); + this->Managers[id].reset(new vtkm::cont::internal::DeviceAdapterMemoryManager); + } + } + + template + VTKM_CONT void CreateManager(Device, std::false_type) + { + // No manager for invalid devices. + } + + template + VTKM_CONT void operator()(Device device) + { + this->CreateManager(device, std::integral_constant{}); + } +}; + struct VTKM_NEVER_EXPORT RuntimeDeviceInformationFunctor { bool Exists = false; @@ -70,24 +163,16 @@ struct VTKM_NEVER_EXPORT RuntimeDeviceInformationFunctor } } }; -} - -namespace vtkm -{ -namespace cont -{ -namespace detail -{ class RuntimeDeviceNames { public: - static const DeviceAdapterNameType& GetDeviceName(vtkm::Int8 id) + static const vtkm::cont::DeviceAdapterNameType& GetDeviceName(vtkm::Int8 id) { return Instance().DeviceNames[id]; } - static const DeviceAdapterNameType& GetLowerCaseDeviceName(vtkm::Int8 id) + static const vtkm::cont::DeviceAdapterNameType& GetLowerCaseDeviceName(vtkm::Int8 id) { return Instance().LowerCaseDeviceNames[id]; } @@ -107,9 +192,63 @@ private: friend struct InitializeDeviceNames; - DeviceAdapterNameType DeviceNames[VTKM_MAX_DEVICE_ADAPTER_ID]; - DeviceAdapterNameType LowerCaseDeviceNames[VTKM_MAX_DEVICE_ADAPTER_ID]; + vtkm::cont::DeviceAdapterNameType DeviceNames[VTKM_MAX_DEVICE_ADAPTER_ID]; + vtkm::cont::DeviceAdapterNameType LowerCaseDeviceNames[VTKM_MAX_DEVICE_ADAPTER_ID]; }; + +class RuntimeDeviceMemoryManagers +{ +public: + static vtkm::cont::internal::DeviceAdapterMemoryManagerBase& GetDeviceMemoryManager( + vtkm::cont::DeviceAdapterId device) + { + const auto id = device.GetValue(); + + if (device.IsValueValid()) + { + auto&& manager = Instance().DeviceMemoryManagers[id]; + if (manager) + { + return *manager.get(); + } + else + { + return Instance().InvalidManager; + } + } + else + { + return Instance().InvalidManager; + } + } + +private: + static RuntimeDeviceMemoryManagers& Instance() + { + static RuntimeDeviceMemoryManagers instance; + return instance; + } + + RuntimeDeviceMemoryManagers() + { + InitializeDeviceMemoryManagers functor(this->DeviceMemoryManagers); + vtkm::ListForEach(functor, VTKM_DEFAULT_DEVICE_ADAPTER_LIST()); + } + + friend struct InitializeDeviceMemoryManagers; + + std::unique_ptr + DeviceMemoryManagers[VTKM_MAX_DEVICE_ADAPTER_ID]; + DeviceAdapterMemoryManagerInvalid InvalidManager; +}; +} + +namespace vtkm +{ +namespace cont +{ +namespace detail +{ } VTKM_CONT @@ -119,7 +258,7 @@ DeviceAdapterNameType RuntimeDeviceInformation::GetName(DeviceAdapterId device) if (device.IsValueValid()) { - return detail::RuntimeDeviceNames::GetDeviceName(id); + return RuntimeDeviceNames::GetDeviceName(id); } else if (id == VTKM_DEVICE_ADAPTER_UNDEFINED) { @@ -131,7 +270,7 @@ DeviceAdapterNameType RuntimeDeviceInformation::GetName(DeviceAdapterId device) } // Deviceis invalid: - return detail::RuntimeDeviceNames::GetDeviceName(0); + return RuntimeDeviceNames::GetDeviceName(0); } VTKM_CONT @@ -156,7 +295,7 @@ DeviceAdapterId RuntimeDeviceInformation::GetId(DeviceAdapterNameType name) cons for (vtkm::Int8 id = 0; id < VTKM_MAX_DEVICE_ADAPTER_ID; ++id) { - if (name == detail::RuntimeDeviceNames::GetLowerCaseDeviceName(id)) + if (name == RuntimeDeviceNames::GetLowerCaseDeviceName(id)) { return vtkm::cont::make_DeviceAdapterId(id); } @@ -178,5 +317,19 @@ bool RuntimeDeviceInformation::Exists(DeviceAdapterId id) const vtkm::ListForEach(functor, VTKM_DEFAULT_DEVICE_ADAPTER_LIST(), id); return functor.Exists; } + +VTKM_CONT vtkm::cont::internal::DeviceAdapterMemoryManagerBase& +RuntimeDeviceInformation::GetMemoryManager(DeviceAdapterId device) const +{ + if (device.IsValueValid()) + { + return RuntimeDeviceMemoryManagers::GetDeviceMemoryManager(device); + } + else + { + throw vtkm::cont::ErrorBadValue( + "Attempted to get a DeviceAdapterMemoryManager for an invalid device."); + } +} } } // namespace vtkm::cont diff --git a/vtkm/cont/RuntimeDeviceInformation.h b/vtkm/cont/RuntimeDeviceInformation.h index c3ee1282c..49af37d85 100644 --- a/vtkm/cont/RuntimeDeviceInformation.h +++ b/vtkm/cont/RuntimeDeviceInformation.h @@ -11,6 +11,7 @@ #define vtk_m_cont_RuntimeDeviceInformation_h #include +#include #include namespace vtkm @@ -43,6 +44,15 @@ public: /// VTKM_CONT bool Exists(DeviceAdapterId id) const; + + /// Returns a reference to a `DeviceAdapterMemoryManager` that will work with the + /// given device. This method will throw an exception if the device id is not a + /// real device (for example `DeviceAdapterTagAny`). If the device in question is + /// not valid, a `DeviceAdapterMemoryManager` will be returned, but attempting to + /// call any of the methods will result in a runtime exception. + /// + VTKM_CONT + vtkm::cont::internal::DeviceAdapterMemoryManagerBase& GetMemoryManager(DeviceAdapterId id) const; }; } } // namespace vtkm::cont diff --git a/vtkm/cont/cuda/DeviceAdapterCuda.h b/vtkm/cont/cuda/DeviceAdapterCuda.h index e8e3382c3..b1491ef33 100644 --- a/vtkm/cont/cuda/DeviceAdapterCuda.h +++ b/vtkm/cont/cuda/DeviceAdapterCuda.h @@ -20,6 +20,7 @@ #include #include +#include #include #endif diff --git a/vtkm/cont/cuda/internal/CMakeLists.txt b/vtkm/cont/cuda/internal/CMakeLists.txt index f859a7e1b..9345f7825 100644 --- a/vtkm/cont/cuda/internal/CMakeLists.txt +++ b/vtkm/cont/cuda/internal/CMakeLists.txt @@ -13,6 +13,7 @@ set(headers AtomicInterfaceExecutionCuda.h CudaAllocator.h DeviceAdapterAlgorithmCuda.h + DeviceAdapterMemoryManagerCuda.h DeviceAdapterRuntimeDetectorCuda.h DeviceAdapterTagCuda.h DeviceAdapterTimerImplementationCuda.h @@ -30,6 +31,7 @@ if (TARGET vtkm::cuda) ${CMAKE_CURRENT_SOURCE_DIR}/ArrayManagerExecutionCuda.cu ${CMAKE_CURRENT_SOURCE_DIR}/CudaAllocator.cu ${CMAKE_CURRENT_SOURCE_DIR}/DeviceAdapterAlgorithmCuda.cu + ${CMAKE_CURRENT_SOURCE_DIR}/DeviceAdapterMemoryManagerCuda.cu ${CMAKE_CURRENT_SOURCE_DIR}/DeviceAdapterRuntimeDetectorCuda.cu ${CMAKE_CURRENT_SOURCE_DIR}/DeviceAdapterTimerImplementationCuda.cu ${CMAKE_CURRENT_SOURCE_DIR}/ExecutionArrayInterfaceBasicCuda.cu diff --git a/vtkm/cont/cuda/internal/DeviceAdapterMemoryManagerCuda.cu b/vtkm/cont/cuda/internal/DeviceAdapterMemoryManagerCuda.cu new file mode 100644 index 000000000..ca8d4680e --- /dev/null +++ b/vtkm/cont/cuda/internal/DeviceAdapterMemoryManagerCuda.cu @@ -0,0 +1,200 @@ +//============================================================================ +// 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. +//============================================================================ + +#include +#include +#include +#include + +#include + +namespace +{ + +/// Deletes a pointer declared with the CUDA API +struct CudaDeleter +{ + VTKM_CONT void operator()(void* buffer) const + { + VTKM_ASSERT(buffer != nullptr); + vtkm::cont::cuda::internal::CudaAllocator::Free(buffer); + } +}; + +struct BufferInfoCuda : vtkm::cont::internal::BufferInfo +{ + std::shared_ptr CudaBuffer; + vtkm::BufferSizeType Size; + + VTKM_CONT BufferInfoCuda(const std::shared_ptr buffer, vtkm::BufferSizeType size) + : CudaBuffer(buffer) + , Size(size) + { + } + + VTKM_CONT void* GetPointer() const override { return this->CudaBuffer.get(); } + + VTKM_CONT vtkm::BufferSizeType GetSize() const override { return this->Size; } +}; + +} // anonymous namespace + +std::shared_ptr vtkm::cont::internal::DeviceAdapterMemoryManager< + vtkm::cont::DeviceAdapterTagCuda>::Allocate(vtkm::BufferSizeType size) +{ + try + { + vtkm::UInt8* buffer = static_cast( + vtkm::cont::cuda::internal::CudaAllocator::Allocate(static_cast(size))); + std::shared_ptr bufferPtr(buffer, CudaDeleter{}); + return std::shared_ptr(new BufferInfoCuda(bufferPtr, size)); + } + catch (const std::exception& error) + { + std::ostringstream err; + err << "Failed to allocate " << size << " bytes on device: " << error.what(); + throw vtkm::cont::ErrorBadAllocation(err.str()); + } +} + +std::shared_ptr vtkm::cont::internal::DeviceAdapterMemoryManager< + vtkm::cont::DeviceAdapterTagCuda>::ManageArray(std::shared_ptr buffer, + vtkm::BufferSizeType size) +{ + return std::shared_ptr(new BufferInfoCuda(buffer, size)); +} + +void vtkm::cont::internal::DeviceAdapterMemoryManager::Reallocate( + std::shared_ptr b, + vtkm::BufferSizeType newSize) +{ + BufferInfoCuda* buffer = dynamic_cast(b.get()); + VTKM_ASSERT(buffer); + + if (newSize <= buffer->Size) + { + // Just reuse the buffer. (Would be nice to free up memory.) + buffer->Size = newSize; + } + else + { + // Make a new buffer + std::shared_ptr newBufferInfo = this->Allocate(newSize); + BufferInfoCuda* newBuffer = dynamic_cast(newBufferInfo.get()); + VTKM_ASSERT(newBuffer != nullptr); + + // Copy the data to the new buffer + VTKM_CUDA_CALL(cudaMemcpyAsync(newBuffer->GetPointer(), + buffer->GetPointer(), + static_cast(buffer->Size), + cudaMemcpyDeviceToDevice, + cudaStreamPerThread)); + + // Reset the buffer in the passed in info + *buffer = *newBuffer; + } +} + +std::shared_ptr vtkm::cont::internal:: + DeviceAdapterMemoryManager::CopyHostToDevice( + std::shared_ptr src) +{ + using vtkm::cont::cuda::internal::CudaAllocator; + if (CudaAllocator::IsManagedPointer(src->GetPointer())) + { + // In the current code structure, we don't know whether this buffer is going to be used + // for input or output. (Currently, I don't think there is any difference.) + CudaAllocator::PrepareForOutput(src->GetPointer(), static_cast(src->GetSize())); + + // The provided control pointer is already cuda managed and can be accessed on the device + // via unified memory. Just shallow copy the pointer. + return std::shared_ptr( + new BufferInfoCuda(src->GetSharedPointer(), src->GetSize())); + } + else + { + // Make a new buffer + std::shared_ptr destInfo = this->Allocate(src->GetSize()); + BufferInfoCuda* dest = dynamic_cast(destInfo.get()); + VTKM_ASSERT(dest != nullptr); + + // Copy the data to the new buffer + VTKM_LOG_F(vtkm::cont::LogLevel::MemTransfer, + "Copying host --> CUDA dev: %s (%lld bytes)", + vtkm::cont::GetHumanReadableSize(static_cast(src->GetSize())).c_str(), + src->GetSize()); + + VTKM_CUDA_CALL(cudaMemcpyAsync(dest->GetPointer(), + src->GetPointer(), + static_cast(src->GetSize()), + cudaMemcpyHostToDevice, + cudaStreamPerThread)); + + return destInfo; + } +} + +std::shared_ptr vtkm::cont::internal:: + DeviceAdapterMemoryManager::CopyDeviceToHost( + std::shared_ptr src_) +{ + using vtkm::cont::cuda::internal::CudaAllocator; + BufferInfoCuda* src = dynamic_cast(src_.get()); + VTKM_ASSERT(src != nullptr); + + std::shared_ptr dest; + + if (CudaAllocator::IsManagedPointer(src->GetPointer())) + { + // The provided control pointer is already cuda managed and can be accessed on the host + // via unified memory. Just shallow copy the pointer. + CudaAllocator::PrepareForControl(src->GetPointer(), static_cast(src->Size)); + dest.reset(new vtkm::cont::internal::BufferInfoHost(src->CudaBuffer, src->Size)); + } + else + { + // Make a new buffer + dest.reset(new vtkm::cont::internal::BufferInfoHost(src->Size)); + + // Copy the data to the new buffer + VTKM_LOG_F(vtkm::cont::LogLevel::MemTransfer, + "Copying CUDA dev --> host: %s (%lld bytes)", + vtkm::cont::GetHumanReadableSize(static_cast(src->Size)).c_str(), + src->Size); + + VTKM_CUDA_CALL(cudaMemcpyAsync(dest->GetPointer(), + src->GetPointer(), + static_cast(src->Size), + cudaMemcpyDeviceToHost, + cudaStreamPerThread)); + } + + //In all cases we have possibly multiple async calls queued up in + //our stream. We need to block on the copy back to control since + //we don't wanting it accessing memory that hasn't finished + //being used by the GPU + vtkm::cont::DeviceAdapterAlgorithm::Synchronize(); + + return dest; +} + +std::shared_ptr vtkm::cont::internal:: + DeviceAdapterMemoryManager::CopyDeviceToDevice( + std::shared_ptr src) +{ + std::shared_ptr dest = this->Allocate(src->GetSize()); + VTKM_CUDA_CALL(cudaMemcpyAsync(dest->GetPointer(), + src->GetPointer(), + static_cast(src->GetSize()), + cudaMemcpyDeviceToDevice, + cudaStreamPerThread)); + + return dest; +} diff --git a/vtkm/cont/cuda/internal/DeviceAdapterMemoryManagerCuda.h b/vtkm/cont/cuda/internal/DeviceAdapterMemoryManagerCuda.h new file mode 100644 index 000000000..c8b5295c1 --- /dev/null +++ b/vtkm/cont/cuda/internal/DeviceAdapterMemoryManagerCuda.h @@ -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. +//============================================================================ +#ifndef vtk_m_cont_cuda_internal_DeviceAdapterMemoryManagerCuda_h +#define vtk_m_cont_cuda_internal_DeviceAdapterMemoryManagerCuda_h + +#include + +#include + +namespace vtkm +{ +namespace cont +{ +namespace internal +{ + +template <> +class VTKM_CONT_EXPORT DeviceAdapterMemoryManager + : public DeviceAdapterMemoryManagerBase +{ +public: + VTKM_CONT std::shared_ptr Allocate(vtkm::BufferSizeType size) override; + + VTKM_CONT std::shared_ptr ManageArray(std::shared_ptr buffer, + vtkm::BufferSizeType size) override; + + VTKM_CONT void Reallocate(std::shared_ptr buffer, + vtkm::BufferSizeType newSize) override; + + VTKM_CONT std::shared_ptr CopyHostToDevice( + std::shared_ptr src) override; + + VTKM_CONT std::shared_ptr CopyDeviceToHost( + std::shared_ptr src) override; + + VTKM_CONT std::shared_ptr CopyDeviceToDevice( + std::shared_ptr src) override; +}; +} +} +} // namespace vtkm::cont::internal + +#endif //vtk_m_cont_cuda_internal_DeviceAdapterMemoryManagerCuda_h diff --git a/vtkm/cont/internal/Buffer.cxx b/vtkm/cont/internal/Buffer.cxx new file mode 100644 index 000000000..2774a6b81 --- /dev/null +++ b/vtkm/cont/internal/Buffer.cxx @@ -0,0 +1,508 @@ +//============================================================================ +// 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. +//============================================================================ + +#include + +#include + +#include +#include + +#include +#include + +using LockType = std::unique_lock; + +namespace vtkm +{ +namespace cont +{ +namespace internal +{ + +class Buffer::InternalsStruct +{ +public: + using DeviceBufferMap = + std::map>; + using HostBufferPointer = std::shared_ptr; + +private: + vtkm::cont::Token::ReferenceCount ReadCount = 0; + vtkm::cont::Token::ReferenceCount WriteCount = 0; + + VTKM_CONT void CheckLock(const LockType& lock) const + { + VTKM_ASSERT((lock.mutex() == &this->Mutex) && (lock.owns_lock())); + } + + // If this number disagrees with the size of the buffers, then they should be resized and + // data preserved. + vtkm::BufferSizeType NumberOfBytes = 0; + + DeviceBufferMap DeviceBuffers; + HostBufferPointer HostBuffer; + +public: + std::mutex Mutex; + std::condition_variable ConditionVariable; + + LockType GetLock() { return LockType(this->Mutex); } + + VTKM_CONT vtkm::cont::Token::ReferenceCount* GetReadCount(const LockType& lock) + { + this->CheckLock(lock); + return &this->ReadCount; + } + VTKM_CONT vtkm::cont::Token::ReferenceCount* GetWriteCount(const LockType& lock) + { + this->CheckLock(lock); + return &this->WriteCount; + } + + VTKM_CONT DeviceBufferMap& GetDeviceBuffers(const LockType& lock) + { + this->CheckLock(lock); + return this->DeviceBuffers; + } + + VTKM_CONT HostBufferPointer& GetHostBuffer(const LockType& lock) + { + this->CheckLock(lock); + return this->HostBuffer; + } + + VTKM_CONT vtkm::BufferSizeType GetNumberOfBytes(const LockType& lock) + { + this->CheckLock(lock); + return this->NumberOfBytes; + } + VTKM_CONT void SetNumberOfBytes(const LockType& lock, vtkm::BufferSizeType numberOfBytes) + { + this->CheckLock(lock); + this->NumberOfBytes = numberOfBytes; + } +}; + +namespace detail +{ + +struct BufferHelper +{ + static bool CanRead(Buffer::InternalsStruct* internals, + const LockType& lock, + const vtkm::cont::Token& token) + { + return ((*internals->GetWriteCount(lock) < 1) || + (token.IsAttached(internals->GetWriteCount(lock)))); + } + + static bool CanWrite(Buffer::InternalsStruct* internals, + const LockType& lock, + const vtkm::cont::Token& token) + { + return ( + ((*internals->GetWriteCount(lock) < 1) || + (token.IsAttached(internals->GetWriteCount(lock)))) && + ((*internals->GetReadCount(lock) < 1) || + ((*internals->GetReadCount(lock) == 1) && token.IsAttached(internals->GetReadCount(lock))))); + } + + static void WaitToRead(Buffer::InternalsStruct* internals, + LockType& lock, + const vtkm::cont::Token& token) + { + // Note that if you deadlocked here, that means that you are trying to do a read operation on an + // array where an object is writing to it. This could happen on the same thread. For example, if + // you call `WritePortal()` then no other operation that can result in reading or writing + // data in the array can happen while the resulting portal is still in scope. + internals->ConditionVariable.wait( + lock, [&lock, &token, internals] { return CanRead(internals, lock, token); }); + } + + static void WaitToWrite(Buffer::InternalsStruct* internals, + LockType& lock, + const vtkm::cont::Token& token) + { + // Note that if you deadlocked here, that means that you are trying to do a write operation on + // an array where an object is reading or writing to it. This could happen on the same thread. + // For example, if you call `WritePortal()` then no other operation that can result in reading + // or writing data in the array can happen while the resulting portal is still in scope. + internals->ConditionVariable.wait( + lock, [&lock, &token, internals] { return CanWrite(internals, lock, token); }); + } + + static void AllocateOnHost(Buffer::InternalsStruct* internals, + std::unique_lock& lock, + vtkm::cont::Token& token) + { + WaitToRead(internals, lock, token); + Buffer::InternalsStruct::HostBufferPointer& hostBuffer = internals->GetHostBuffer(lock); + vtkm::BufferSizeType targetSize = internals->GetNumberOfBytes(lock); + if (hostBuffer.get() != nullptr) + { + // Buffer already exists on the host. Make sure it is the right size. + if (hostBuffer->GetSize() != targetSize) + { + hostBuffer->Allocate(targetSize, vtkm::CopyFlag::On); + } + return; + } + + // Buffer does not exist on host. See if we can find data on a device. + for (auto&& deviceBuffer : internals->GetDeviceBuffers(lock)) + { + if (deviceBuffer.second.get() == nullptr) + { + continue; + } + + vtkm::cont::internal::DeviceAdapterMemoryManagerBase& memoryManager = + vtkm::cont::RuntimeDeviceInformation().GetMemoryManager(deviceBuffer.first); + + if (deviceBuffer.second->GetSize() > targetSize) + { + // Device buffer too large. Resize. + memoryManager.Reallocate(deviceBuffer.second, targetSize); + } + + hostBuffer = memoryManager.CopyDeviceToHost(deviceBuffer.second); + + if (hostBuffer->GetSize() != targetSize) + { + hostBuffer->Allocate(targetSize, vtkm::CopyFlag::On); + } + + return; + } + + // Buffer not allocated on host or any device, so just allocate a buffer. + hostBuffer.reset(new vtkm::cont::internal::BufferInfoHost); + hostBuffer->Allocate(targetSize, vtkm::CopyFlag::Off); + } + + static void AllocateOnDevice(Buffer::InternalsStruct* internals, + std::unique_lock& lock, + vtkm::cont::Token& token, + vtkm::cont::DeviceAdapterId device) + { + WaitToRead(internals, lock, token); + Buffer::InternalsStruct::DeviceBufferMap& deviceBuffers = internals->GetDeviceBuffers(lock); + vtkm::BufferSizeType targetSize = internals->GetNumberOfBytes(lock); + vtkm::cont::internal::DeviceAdapterMemoryManagerBase& memoryManager = + vtkm::cont::RuntimeDeviceInformation().GetMemoryManager(device); + + if (deviceBuffers[device].get() != nullptr) + { + // Buffer already exists on the device. Make sure it is the right size. + if (deviceBuffers[device]->GetSize() != targetSize) + { + memoryManager.Reallocate(deviceBuffers[device], targetSize); + } + VTKM_ASSERT(deviceBuffers[device]->GetSize() == targetSize); + return; + } + + // Buffer does not exist on device. Check to see if it is on another device but not a host. + // We currently do not support device-to-device transfers, so the data has to go to the + // host first. + if (internals->GetHostBuffer(lock).get() == nullptr) + { + for (auto&& deviceBuffer : deviceBuffers) + { + if (deviceBuffer.second.get() != nullptr) + { + // Copy data to host. + AllocateOnHost(internals, lock, token); + break; + } + } + } + + // If the buffer is now on the host, copy it to the device. + Buffer::InternalsStruct::HostBufferPointer& hostBuffer = internals->GetHostBuffer(lock); + if (hostBuffer.get() != nullptr) + { + if (hostBuffer->GetSize() > targetSize) + { + // Host buffer too large. Resize. + hostBuffer->Allocate(targetSize, vtkm::CopyFlag::On); + } + + deviceBuffers[device] = memoryManager.CopyHostToDevice(hostBuffer); + + if (deviceBuffers[device]->GetSize() != targetSize) + { + memoryManager.Reallocate(deviceBuffers[device], targetSize); + } + VTKM_ASSERT(deviceBuffers[device]->GetSize() == targetSize); + + return; + } + + // Buffer not allocated anywhere, so just allocate a buffer. + deviceBuffers[device] = memoryManager.Allocate(targetSize); + } + + static void CopyOnHost(vtkm::cont::internal::Buffer::InternalsStruct* srcInternals, + LockType& srcLock, + vtkm::cont::internal::Buffer::InternalsStruct* destInternals, + LockType& destLock, + vtkm::cont::Token& token) + { + WaitToRead(srcInternals, srcLock, token); + WaitToWrite(destInternals, destLock, token); + + // Any current buffers in destination can be (and should be) deleted. + destInternals->GetDeviceBuffers(destLock).clear(); + + // Do the copy + Buffer::InternalsStruct::HostBufferPointer& destBuffer = destInternals->GetHostBuffer(destLock); + Buffer::InternalsStruct::HostBufferPointer& srcBuffer = srcInternals->GetHostBuffer(srcLock); + + destBuffer.reset(new vtkm::cont::internal::BufferInfoHost); + destBuffer->Allocate(srcBuffer->GetSize(), vtkm::CopyFlag::Off); + + std::copy(reinterpret_cast(srcBuffer->GetPointer()), + reinterpret_cast(srcBuffer->GetPointer()) + srcBuffer->GetSize(), + reinterpret_cast(destBuffer->GetPointer())); + + destInternals->SetNumberOfBytes(destLock, srcInternals->GetNumberOfBytes(srcLock)); + } + + static void CopyOnDevice(vtkm::cont::DeviceAdapterId device, + vtkm::cont::internal::Buffer::InternalsStruct* srcInternals, + LockType& srcLock, + vtkm::cont::internal::Buffer::InternalsStruct* destInternals, + LockType& destLock, + vtkm::cont::Token& token) + { + WaitToRead(srcInternals, srcLock, token); + WaitToWrite(destInternals, destLock, token); + + // Any current buffers in destination can be (and should be) deleted. + destInternals->GetHostBuffer(destLock).reset(); + destInternals->GetDeviceBuffers(destLock).clear(); + + // Do the copy + vtkm::cont::internal::DeviceAdapterMemoryManagerBase& memoryManager = + vtkm::cont::RuntimeDeviceInformation().GetMemoryManager(device); + + destInternals->GetDeviceBuffers(destLock)[device] = + memoryManager.CopyDeviceToDevice(srcInternals->GetDeviceBuffers(srcLock)[device]); + + destInternals->SetNumberOfBytes(destLock, srcInternals->GetNumberOfBytes(srcLock)); + } +}; + +} // namespace detail + +Buffer::Buffer() + : Internals(new InternalsStruct) +{ +} + +vtkm::BufferSizeType Buffer::GetNumberOfBytes() const +{ + LockType lock = this->Internals->GetLock(); + return this->Internals->GetNumberOfBytes(lock); +} + +void Buffer::SetNumberOfBytes(vtkm::BufferSizeType numberOfBytes, vtkm::CopyFlag preserve) +{ + VTKM_ASSUME(numberOfBytes >= 0); + LockType lock = this->Internals->GetLock(); + if (this->Internals->GetNumberOfBytes(lock) == numberOfBytes) + { + // Allocation has not changed. Just return. + // Note, if you set the size to the old size and then try to get the buffer on a different + // place, a copy might happen. + return; + } + + // We are altering the array, so make sure we can write to it. + vtkm::cont::Token token; + detail::BufferHelper::WaitToWrite(this->Internals.get(), lock, token); + + this->Internals->SetNumberOfBytes(lock, numberOfBytes); + if ((preserve == vtkm::CopyFlag::Off) || (numberOfBytes == 0)) + { + // No longer need these buffers. Just delete them. + this->Internals->GetHostBuffer(lock).reset(); + this->Internals->GetDeviceBuffers(lock).clear(); + } + else + { + // Do nothing (other than resetting numberOfBytes). Buffers will get resized when you get the + // pointer. + } +} + +bool Buffer::IsAllocatedOnHost() const +{ + LockType lock = this->Internals->GetLock(); + return (this->Internals->GetHostBuffer(lock).get() != nullptr); +} + +bool Buffer::IsAllocatedOnDevice(vtkm::cont::DeviceAdapterId device) const +{ + LockType lock = this->Internals->GetLock(); + return (this->Internals->GetDeviceBuffers(lock)[device].get() != nullptr); +} + +const void* Buffer::ReadPointerHost(vtkm::cont::Token& token) const +{ + LockType lock = this->Internals->GetLock(); + detail::BufferHelper::WaitToRead(this->Internals.get(), lock, token); + detail::BufferHelper::AllocateOnHost(this->Internals.get(), lock, token); + return this->Internals->GetHostBuffer(lock)->GetPointer(); +} + +const void* Buffer::ReadPointerDevice(vtkm::cont::DeviceAdapterId device, + vtkm::cont::Token& token) const +{ + LockType lock = this->Internals->GetLock(); + detail::BufferHelper::WaitToRead(this->Internals.get(), lock, token); + detail::BufferHelper::AllocateOnDevice(this->Internals.get(), lock, token, device); + return this->Internals->GetDeviceBuffers(lock)[device]->GetPointer(); +} + +void* Buffer::WritePointerHost(vtkm::cont::Token& token) const +{ + LockType lock = this->Internals->GetLock(); + detail::BufferHelper::WaitToWrite(this->Internals.get(), lock, token); + detail::BufferHelper::AllocateOnHost(this->Internals.get(), lock, token); + + // Array is being written on host. All other buffers invalidated, so delete them. + this->Internals->GetDeviceBuffers(lock).clear(); + + return this->Internals->GetHostBuffer(lock)->GetPointer(); +} + +void* Buffer::WritePointerDevice(vtkm::cont::DeviceAdapterId device, vtkm::cont::Token& token) const +{ + LockType lock = this->Internals->GetLock(); + detail::BufferHelper::WaitToWrite(this->Internals.get(), lock, token); + detail::BufferHelper::AllocateOnDevice(this->Internals.get(), lock, token, device); + + // Array is being written on this device. All other buffers invalided, so delete them. + this->Internals->GetHostBuffer(lock).reset(); + InternalsStruct::DeviceBufferMap& deviceBuffers = this->Internals->GetDeviceBuffers(lock); + auto iter = deviceBuffers.find(device); + deviceBuffers.erase(deviceBuffers.begin(), iter); + iter = deviceBuffers.find(device); + std::advance(iter, 1); + deviceBuffers.erase(iter, deviceBuffers.end()); + + return this->Internals->GetDeviceBuffers(lock)[device]->GetPointer(); +} + +void Buffer::DeepCopy(vtkm::cont::internal::Buffer& dest) const +{ + LockType srcLock = this->Internals->GetLock(); + LockType destLock = dest.Internals->GetLock(); + vtkm::cont::Token token; + + detail::BufferHelper::WaitToRead(this->Internals.get(), srcLock, token); + + if (!this->Internals->GetDeviceBuffers(srcLock).empty()) + { + // If we are on a device, copy there. + detail::BufferHelper::CopyOnDevice(this->Internals->GetDeviceBuffers(srcLock).begin()->first, + this->Internals.get(), + srcLock, + dest.Internals.get(), + destLock, + token); + } + else if (this->Internals->GetHostBuffer(srcLock) != nullptr) + { + // If we are on a host, copy there + detail::BufferHelper::CopyOnHost( + this->Internals.get(), srcLock, dest.Internals.get(), destLock, token); + } + else + { + // Nothing actually allocated. Just unallocate everything on destination. + dest.Internals->GetHostBuffer(destLock).reset(); + dest.Internals->GetDeviceBuffers(destLock).clear(); + dest.Internals->SetNumberOfBytes(destLock, this->Internals->GetNumberOfBytes(srcLock)); + } +} + +void Buffer::DeepCopy(vtkm::cont::internal::Buffer& dest, vtkm::cont::DeviceAdapterId device) const +{ + LockType srcLock = this->Internals->GetLock(); + LockType destLock = dest.Internals->GetLock(); + vtkm::cont::Token token; + detail::BufferHelper::CopyOnDevice( + device, this->Internals.get(), srcLock, dest.Internals.get(), destLock, token); +} + +void Buffer::Reset(const std::shared_ptr buffer, vtkm::BufferSizeType numberOfBytes) +{ + LockType lock = this->Internals->GetLock(); + vtkm::cont::Token token; + detail::BufferHelper::WaitToWrite(this->Internals.get(), lock, token); + + this->Internals->GetDeviceBuffers(lock).clear(); + + this->Internals->GetHostBuffer(lock).reset( + new vtkm::cont::internal::BufferInfoHost(buffer, numberOfBytes)); + this->Internals->SetNumberOfBytes(lock, numberOfBytes); +} + +void Buffer::Reset(const std::shared_ptr buffer, + vtkm::BufferSizeType numberOfBytes, + vtkm::cont::DeviceAdapterId device) +{ + LockType lock = this->Internals->GetLock(); + vtkm::cont::Token token; + detail::BufferHelper::WaitToWrite(this->Internals.get(), lock, token); + + this->Internals->GetDeviceBuffers(lock).clear(); + this->Internals->GetHostBuffer(lock).reset(); + + this->Internals->GetDeviceBuffers(lock)[device] = + vtkm::cont::RuntimeDeviceInformation().GetMemoryManager(device).ManageArray(buffer, + numberOfBytes); + this->Internals->SetNumberOfBytes(lock, numberOfBytes); +} +} +} +} // namespace vtkm::cont::internal + +namespace mangled_diy_namespace +{ + +void Serialization::save(BinaryBuffer& bb, + const vtkm::cont::internal::Buffer& obj) +{ + vtkm::BufferSizeType size = obj.GetNumberOfBytes(); + vtkmdiy::save(bb, size); + + vtkm::cont::Token token; + const vtkm::UInt8* data = reinterpret_cast(obj.ReadPointerHost(token)); + vtkmdiy::save(bb, data, static_cast(size)); +} + +void Serialization::load(BinaryBuffer& bb, + vtkm::cont::internal::Buffer& obj) +{ + vtkm::BufferSizeType size; + vtkmdiy::load(bb, size); + + obj.SetNumberOfBytes(size, vtkm::CopyFlag::Off); + vtkm::cont::Token token; + vtkm::UInt8* data = reinterpret_cast(obj.WritePointerHost(token)); + vtkmdiy::load(bb, data, static_cast(size)); +} + +} // namespace diy diff --git a/vtkm/cont/internal/Buffer.h b/vtkm/cont/internal/Buffer.h new file mode 100644 index 000000000..12f91fbba --- /dev/null +++ b/vtkm/cont/internal/Buffer.h @@ -0,0 +1,236 @@ +//============================================================================ +// 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. +//============================================================================ +#ifndef vtk_m_vtkm_cont_internal_Buffer_h +#define vtk_m_vtkm_cont_internal_Buffer_h + +#include + +#include +#include +#include + +#include + +#include +#include + +namespace vtkm +{ +namespace cont +{ +namespace internal +{ + +namespace detail +{ + +struct BufferHelper; + +} // namespace detail + +/// \brief Manages a buffer data among the host and various devices. +/// +/// The `Buffer` class defines a contiguous section of memory of a specified number of bytes. The +/// data in this buffer is managed in the host and across the devices supported by VTK-m. `Buffer` +/// will allocate memory and transfer data as necessary. +/// +class VTKM_CONT_EXPORT Buffer +{ + class InternalsStruct; + std::shared_ptr Internals; + + friend struct vtkm::cont::internal::detail::BufferHelper; + +public: + /// \brief Create an empty `Buffer`. + /// + VTKM_CONT Buffer(); + + /// \brief Returns the number of bytes held by the buffer. + /// + /// Note that `Buffer` allocates memory lazily. So there might not actually be any memory + /// allocated anywhere. It is also possible that memory is simultaneously allocated on multiple + /// devices. + /// + VTKM_CONT vtkm::BufferSizeType GetNumberOfBytes() const; + + /// \brief Changes the size of the buffer. + /// + /// Note that `Buffer` alloates memory lazily. So there might not be any memory allocated at + /// the return of the call. (However, later calls to retrieve pointers will allocate memory + /// as necessary.) + /// + /// The `preserve` argument flags whether any existing data in the buffer is preserved. + /// Preserving data might cost more time or memory. + /// + VTKM_CONT void SetNumberOfBytes(vtkm::BufferSizeType numberOfBytes, vtkm::CopyFlag preserve); + + /// \brief Returns `true` if the buffer is allocated on the host. + /// + VTKM_CONT bool IsAllocatedOnHost() const; + + /// \brief Returns `true` if the buffer is allocated on the given device. + /// + VTKM_CONT bool IsAllocatedOnDevice(vtkm::cont::DeviceAdapterId device) const; + + /// \brief Returns a readable host (control environment) pointer to the buffer. + /// + /// Memory will be allocated and data will be copied as necessary. The memory at the pointer will + /// be valid as long as `token` is still in scope. Any write operation to this buffer will be + /// blocked until the `token` goes out of scope. + /// + VTKM_CONT const void* ReadPointerHost(vtkm::cont::Token& token) const; + + /// \brief Returns a readable device pointer to the buffer. + /// + /// Memory will be allocated and data will be copied as necessary. The memory at the pointer will + /// be valid as long as `token` is still in scope. Any write operation to this buffer will be + /// blocked until the `token` goes out of scope. + /// + VTKM_CONT const void* ReadPointerDevice(vtkm::cont::DeviceAdapterId device, + vtkm::cont::Token& token) const; + + /// \brief Returns a writable host (control environment) pointer to the buffer. + /// + /// Memory will be allocated and data will be copied as necessary. The memory at the pointer will + /// be valid as long as `token` is still in scope. Any read or write operation to this buffer + /// will be blocked until the `token` goes out of scope. + /// + VTKM_CONT void* WritePointerHost(vtkm::cont::Token& token) const; + + /// \brief Returns a writable device pointer to the buffer. + /// + /// Memory will be allocated and data will be copied as necessary. The memory at the pointer will + /// be valid as long as `token` is still in scope. Any read or write operation to this buffer + /// will be blocked until the `token` goes out of scope. + /// + VTKM_CONT void* WritePointerDevice(vtkm::cont::DeviceAdapterId device, + vtkm::cont::Token& token) const; + + /// @{ + /// \brief Copies the data from this buffer to the target buffer. + /// + /// If a device is given, then the copy will be preferred for that device. Otherwise, a device + /// already containing the data will be used for the copy. If no such device exists, the host + /// will be used. + /// + VTKM_CONT void DeepCopy(vtkm::cont::internal::Buffer& dest) const; + VTKM_CONT void DeepCopy(vtkm::cont::internal::Buffer& dest, + vtkm::cont::DeviceAdapterId device) const; + /// @} + + /// \brief Resets the `Buffer` to the memory allocated at the given pointer. + /// + /// The `Buffer` is initialized to a state that contains the given `buffer` of data. The + /// provided `shared_ptr` should have a deleter that appropriately deletes the buffer when + /// it is no longer used. + /// + VTKM_CONT void Reset(const std::shared_ptr buffer, + vtkm::BufferSizeType numberOfBytes); + + /// \brief Resets the `Buffer` to the memory allocated at the given pointer on a device. + /// + /// The `Buffer` is initialized to a state that contains the given `buffer` of data. The pointer + /// is assumed to be memory in the given `device`. The provided `shared_ptr` should have a + /// deleter that appropriately deletes the buffer when it is no longer used. This is likely a + /// device-specific function. + /// + VTKM_CONT void Reset(const std::shared_ptr buffer, + vtkm::BufferSizeType numberOfBytes, + vtkm::cont::DeviceAdapterId device); + + /// \brief Resets the `Buffer` to the memory allocated at the given pointer. + /// + /// The `Buffer` is initialized to a state that contains the given `buffer` of data. The memory + /// is assumed to be deleted with the standard `delete[]` keyword. + /// + /// Note that the size passed in is the number of bytes, not the number of values. + /// + template + VTKM_CONT void Reset(T* buffer, vtkm::BufferSizeType numberOfBytes) + { + this->Reset(buffer, numberOfBytes, std::default_delete{}); + } + + /// \brief Resets the `Buffer` to the memory allocated at the given pointer. + /// + /// The `Buffer` is initialized to a state that contains the given `buffer` of data. The + /// `deleter` is an object with an `operator()(void*)` that will properly delete the provided + /// buffer. + /// + /// Note that the size passed in is the number of bytes, not the number of values. + /// + template + VTKM_CONT void Reset(T* buffer, vtkm::BufferSizeType numberOfBytes, Deleter deleter) + { + std::shared_ptr sharedP(reinterpret_cast(buffer), deleter); + this->Reset(sharedP, numberOfBytes); + } + + /// \brief Resets the `Buffer` to the memory allocated at the given pointer on a device. + /// + /// The `Buffer` is initialized to a state that contains the given `buffer` of data. The + /// `deleter` is an object with an `operator()(void*)` that will properly delete the provided + /// buffer. This is likely a device-specific function. + /// + /// Note that the size passed in is the number of bytes, not the number of values. + /// + template + VTKM_CONT void Reset(T* buffer, + vtkm::BufferSizeType numberOfBytes, + Deleter deleter, + vtkm::cont::DeviceAdapterId device) + { + std::shared_ptr sharedP(reinterpret_cast(buffer), deleter); + this->Reset(sharedP, numberOfBytes, device); + } +}; +} +} +} // namespace vtkm::cont::internal + +//============================================================================= +// Specializations of serialization related classes +/// @cond SERIALIZATION +namespace mangled_diy_namespace +{ + +template <> +struct VTKM_CONT_EXPORT Serialization +{ + static VTKM_CONT void save(BinaryBuffer& bb, const vtkm::cont::internal::Buffer& obj); + static VTKM_CONT void load(BinaryBuffer& bb, vtkm::cont::internal::Buffer& obj); +}; + +template +struct Serialization> +{ + static VTKM_CONT void save(BinaryBuffer& bb, + const vtkm::Vec& obj) + { + for (vtkm::IdComponent index = 0; index < N; ++index) + { + vtkmdiy::save(bb, obj[index]); + } + } + + static VTKM_CONT void load(BinaryBuffer& bb, vtkm::Vec& obj) + { + for (vtkm::IdComponent index = 0; index < N; ++index) + { + vtkmdiy::load(bb, obj[index]); + } + } +}; + +} // diy +/// @endcond SERIALIZATION + +#endif //vtk_m_vtkm_cont_internal_Buffer_h diff --git a/vtkm/cont/internal/CMakeLists.txt b/vtkm/cont/internal/CMakeLists.txt index e322a90ea..081fbdf42 100644 --- a/vtkm/cont/internal/CMakeLists.txt +++ b/vtkm/cont/internal/CMakeLists.txt @@ -20,9 +20,12 @@ set(headers ArrayTransfer.h AtomicInterfaceControl.h AtomicInterfaceExecution.h + Buffer.h CastInvalidValue.h ConnectivityExplicitInternals.h DeviceAdapterAlgorithmGeneral.h + DeviceAdapterMemoryManager.h + DeviceAdapterMemoryManagerShared.h DeviceAdapterListHelpers.h FunctorsGeneral.h IteratorFromArrayPortal.h diff --git a/vtkm/cont/internal/DeviceAdapterMemoryManager.cxx b/vtkm/cont/internal/DeviceAdapterMemoryManager.cxx new file mode 100644 index 000000000..b96d29ba8 --- /dev/null +++ b/vtkm/cont/internal/DeviceAdapterMemoryManager.cxx @@ -0,0 +1,153 @@ +//============================================================================ +// 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. +//============================================================================ + +#include + +#include + +//---------------------------------------------------------------------------------------- +// Special allocation/deallocation code + +#if defined(VTKM_POSIX) +#define VTKM_MEMALIGN_POSIX +#elif defined(_WIN32) +#define VTKM_MEMALIGN_WIN +#elif defined(__SSE__) +#define VTKM_MEMALIGN_SSE +#else +#define VTKM_MEMALIGN_NONE +#endif + +#if defined(VTKM_MEMALIGN_POSIX) +#include +#elif defined(VTKM_MEMALIGN_WIN) +#include +#elif defined(VTKM_MEMALIGN_SSE) +#include +#else +#include +#endif + +#include +#include + +namespace +{ + +/// A deleter object that can be used with our aligned mallocs +struct HostDeleter +{ + void operator()(void* memory) + { +#if defined(VTKM_MEMALIGN_POSIX) + free(memory); +#elif defined(VTKM_MEMALIGN_WIN) + _aligned_free(memory); +#elif defined(VTKM_MEMALIGN_SSE) + _mm_free(memory); +#else + free(memory); +#endif + } +}; + +/// Allocates a buffer of a specified size using VTK-m's preferred memory alignment. +/// Returns as a shared_ptr with the proper deleter. +std::shared_ptr HostAllocate(vtkm::BufferSizeType numBytes) +{ + const std::size_t size = static_cast(numBytes); + constexpr std::size_t align = VTKM_ALLOCATION_ALIGNMENT; + +#if defined(VTKM_MEMALIGN_POSIX) + void* memory = nullptr; + if (posix_memalign(&memory, align, size) != 0) + { + memory = nullptr; + } +#elif defined(VTKM_MEMALIGN_WIN) + void* memory = _aligned_malloc(size, align); +#elif defined(VTKM_MEMALIGN_SSE) + void* memory = _mm_malloc(size, align); +#else + void* memory = malloc(size); +#endif + + return std::shared_ptr(reinterpret_cast(memory), HostDeleter{}); +} + +} // anonymous namespace + +//---------------------------------------------------------------------------------------- +vtkm::cont::internal::BufferInfo::~BufferInfo() +{ +} + +vtkm::cont::internal::BufferInfoHost::BufferInfoHost() + : Buffer() + , Size(0) +{ +} + +vtkm::cont::internal::BufferInfoHost::BufferInfoHost(const std::shared_ptr& buffer, + vtkm::BufferSizeType size) + : Buffer(buffer) + , Size(size) +{ +} + +vtkm::cont::internal::BufferInfoHost::BufferInfoHost(vtkm::BufferSizeType size) + : Buffer() + , Size(0) +{ + this->Allocate(size, vtkm::CopyFlag::Off); +} + +void* vtkm::cont::internal::BufferInfoHost::GetPointer() const +{ + return this->Buffer.get(); +} + +vtkm::BufferSizeType vtkm::cont::internal::BufferInfoHost::GetSize() const +{ + return this->Size; +} + +void vtkm::cont::internal::BufferInfoHost::Allocate(vtkm::BufferSizeType size, + vtkm::CopyFlag preserve) +{ + if (size < 1) + { + // Compilers apparently have a problem calling std::shared_ptr::reset with nullptr. + vtkm::UInt8* empty = nullptr; + this->Buffer.reset(empty); + this->Size = 0; + } + else if (preserve == vtkm::CopyFlag::Off) + { + this->Buffer = HostAllocate(size); + this->Size = size; + } + else // preserve == vtkm::CopyFlag::On + { + std::shared_ptr newBuffer = HostAllocate(size); + std::copy(this->Buffer.get(), this->Buffer.get() + std::min(size, this->Size), newBuffer.get()); + this->Buffer = newBuffer; + this->Size = size; + } +} + +std::shared_ptr vtkm::cont::internal::BufferInfoHost::GetSharedPointer() const +{ + return this->Buffer; +} + +vtkm::cont::internal::DeviceAdapterMemoryManagerBase::~DeviceAdapterMemoryManagerBase() +{ +} diff --git a/vtkm/cont/internal/DeviceAdapterMemoryManager.h b/vtkm/cont/internal/DeviceAdapterMemoryManager.h new file mode 100644 index 000000000..b841161d8 --- /dev/null +++ b/vtkm/cont/internal/DeviceAdapterMemoryManager.h @@ -0,0 +1,121 @@ +//============================================================================ +// 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. +//============================================================================ +#ifndef vtk_m_cont_internal_DeviceAdapterMemoryManager_h +#define vtk_m_cont_internal_DeviceAdapterMemoryManager_h + +#include + +#include +#include + +#include + +#include + +namespace vtkm +{ + +using BufferSizeType = vtkm::Int64; + +namespace cont +{ +namespace internal +{ + +class VTKM_CONT_EXPORT BufferInfo +{ +public: + /// Returns a pointer to the memory that is allocated. This pointer may only be referenced on + /// the associated device. + /// + VTKM_CONT virtual void* GetPointer() const = 0; + + /// Returns the size of the buffer in bytes. + /// + VTKM_CONT virtual vtkm::BufferSizeType GetSize() const = 0; + + VTKM_CONT virtual ~BufferInfo(); + +protected: + BufferInfo() = default; +}; + +class VTKM_CONT_EXPORT BufferInfoHost : public BufferInfo +{ + std::shared_ptr Buffer; + vtkm::BufferSizeType Size; + +public: + VTKM_CONT BufferInfoHost(); + VTKM_CONT BufferInfoHost(const std::shared_ptr& buffer, vtkm::BufferSizeType size); + VTKM_CONT BufferInfoHost(vtkm::BufferSizeType size); + + VTKM_CONT virtual void* GetPointer() const override; + + VTKM_CONT virtual vtkm::BufferSizeType GetSize() const override; + + VTKM_CONT void Allocate(vtkm::BufferSizeType size, vtkm::CopyFlag preserve); + + VTKM_CONT std::shared_ptr GetSharedPointer() const; +}; + +/// \brief The base class for device adapter memory managers. +/// +/// Every device adapter is expected to define a specialization of `DeviceAdapterMemoryManager`, +/// and they are all expected to subclass this base class. +/// +class VTKM_CONT_EXPORT DeviceAdapterMemoryManagerBase +{ +public: + VTKM_CONT virtual ~DeviceAdapterMemoryManagerBase(); + + /// Allocates a buffer of the specified size in bytes and returns a BufferInfo object + /// containing information about it. + VTKM_CONT virtual std::shared_ptr Allocate(vtkm::BufferSizeType size) = 0; + + /// Manages the provided array. Returns a `BufferInfo` object that contains the data. + /// The deleter in the `shared_ptr` is expected to correctly free the data. + VTKM_CONT virtual std::shared_ptr ManageArray(std::shared_ptr buffer, + vtkm::BufferSizeType size) = 0; + + /// Reallocates the provided buffer to a new size. The passed in `BufferInfo` should be + /// modified to reflect the changes. + VTKM_CONT virtual void Reallocate(std::shared_ptr buffer, + vtkm::BufferSizeType newSize) = 0; + + /// Copies data from the host buffer provided onto the device and returns a buffer info + /// object holding the pointer for the device. + VTKM_CONT virtual std::shared_ptr CopyHostToDevice( + std::shared_ptr src) = 0; + + /// Copies data from the device buffer provided to the host. The passed in `BufferInfo` object + /// was created by a previous call to this object. + VTKM_CONT virtual std::shared_ptr CopyDeviceToHost( + std::shared_ptr src) = 0; + + /// Copies data from one device buffer to another device buffer. The passed in `BufferInfo` object + /// was created by a previous call to this object. + VTKM_CONT virtual std::shared_ptr CopyDeviceToDevice( + std::shared_ptr src) = 0; +}; + +/// \brief The device adapter memory manager. +/// +/// Every device adapter is expected to define a specialization of `DeviceAdapterMemoryManager`. +/// This class must be a (perhaps indirect) subclass of `DeviceAdapterMemoryManagerBase`. All +/// abstract methods must be implemented. +/// +template +class DeviceAdapterMemoryManager; +} +} +} // namespace vtkm::cont::internal + +#endif //vtk_m_cont_internal_DeviceAdapterMemoryManager_h diff --git a/vtkm/cont/internal/DeviceAdapterMemoryManagerShared.cxx b/vtkm/cont/internal/DeviceAdapterMemoryManagerShared.cxx new file mode 100644 index 000000000..f3b9f7895 --- /dev/null +++ b/vtkm/cont/internal/DeviceAdapterMemoryManagerShared.cxx @@ -0,0 +1,101 @@ +//============================================================================ +// 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. +//============================================================================ + +#include + +#include + +namespace +{ + +class BufferInfoShared : public vtkm::cont::internal::BufferInfo +{ + std::shared_ptr HostBuffer; + +public: + VTKM_CONT BufferInfoShared() + : HostBuffer(new vtkm::cont::internal::BufferInfoHost) + { + } + + VTKM_CONT BufferInfoShared(std::shared_ptr hostBuffer) + : HostBuffer(hostBuffer) + { + } + + template + VTKM_CONT BufferInfoShared(Ts&&... bufferInfoHostArgs) + : HostBuffer(new vtkm::cont::internal::BufferInfoHost(std::forward(bufferInfoHostArgs)...)) + { + } + + VTKM_CONT void* GetPointer() const override { return this->HostBuffer->GetPointer(); } + + VTKM_CONT vtkm::BufferSizeType GetSize() const override { return this->HostBuffer->GetSize(); } + + VTKM_CONT std::shared_ptr GetHostBuffer() const + { + return this->HostBuffer; + } +}; + +} // anonymous namespace + +std::shared_ptr +vtkm::cont::internal::DeviceAdapterMemoryManagerShared::Allocate(vtkm::BufferSizeType size) +{ + return std::shared_ptr(new BufferInfoShared(size)); +} + +std::shared_ptr +vtkm::cont::internal::DeviceAdapterMemoryManagerShared::ManageArray( + std::shared_ptr buffer, + vtkm::BufferSizeType size) +{ + return std::shared_ptr(new BufferInfoShared(buffer, size)); +} + +void vtkm::cont::internal::DeviceAdapterMemoryManagerShared::Reallocate( + std::shared_ptr b, + vtkm::BufferSizeType newSize) +{ + BufferInfoShared* buffer = dynamic_cast(b.get()); + VTKM_ASSERT(buffer != nullptr); + + buffer->GetHostBuffer()->Allocate(newSize, vtkm::CopyFlag::On); +} + +std::shared_ptr +vtkm::cont::internal::DeviceAdapterMemoryManagerShared::CopyHostToDevice( + std::shared_ptr src) +{ + return std::shared_ptr(new BufferInfoShared(src)); +} + +std::shared_ptr +vtkm::cont::internal::DeviceAdapterMemoryManagerShared::CopyDeviceToHost( + std::shared_ptr src) +{ + BufferInfoShared* buffer = dynamic_cast(src.get()); + VTKM_ASSERT(buffer != nullptr); + + return buffer->GetHostBuffer(); +} + +std::shared_ptr +vtkm::cont::internal::DeviceAdapterMemoryManagerShared::CopyDeviceToDevice( + std::shared_ptr src) +{ + BufferInfoShared* dest = new BufferInfoShared(src->GetSize()); + vtkm::UInt8* srcP = reinterpret_cast(src->GetPointer()); + vtkm::UInt8* destP = reinterpret_cast(dest->GetPointer()); + std::copy(srcP, srcP + src->GetSize(), destP); + return std::shared_ptr(dest); +} diff --git a/vtkm/cont/internal/DeviceAdapterMemoryManagerShared.h b/vtkm/cont/internal/DeviceAdapterMemoryManagerShared.h new file mode 100644 index 000000000..8abd9808c --- /dev/null +++ b/vtkm/cont/internal/DeviceAdapterMemoryManagerShared.h @@ -0,0 +1,52 @@ +//============================================================================ +// 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. +//============================================================================ +#ifndef vtk_m_cont_internal_DeviceAdapterMemoryManagerShared_h +#define vtk_m_cont_internal_DeviceAdapterMemoryManagerShared_h + +#include + +namespace vtkm +{ +namespace cont +{ +namespace internal +{ + +/// \brief An implementation of DeviceAdapterMemoryManager for devices that share memory with the +/// host. +/// +/// Device adapters that have shared memory with the host can implement their +/// `DeviceAdapterMemoryManager` as a simple subclass of this. +/// +class VTKM_CONT_EXPORT DeviceAdapterMemoryManagerShared : public DeviceAdapterMemoryManagerBase +{ +public: + VTKM_CONT std::shared_ptr Allocate(vtkm::BufferSizeType size) override; + + VTKM_CONT std::shared_ptr ManageArray(std::shared_ptr buffer, + vtkm::BufferSizeType size) override; + + VTKM_CONT void Reallocate(std::shared_ptr buffer, + vtkm::BufferSizeType newSize) override; + + VTKM_CONT std::shared_ptr CopyHostToDevice( + std::shared_ptr src) override; + + VTKM_CONT std::shared_ptr CopyDeviceToHost( + std::shared_ptr src) override; + + VTKM_CONT std::shared_ptr CopyDeviceToDevice( + std::shared_ptr src) override; +}; +} +} +} // namespace vtkm::cont::internal + +#endif //vtk_m_cont_internal_DeviceAdapterMemoryManagerShared_h diff --git a/vtkm/cont/internal/testing/CMakeLists.txt b/vtkm/cont/internal/testing/CMakeLists.txt index d7b7c3835..5e0d34d1a 100644 --- a/vtkm/cont/internal/testing/CMakeLists.txt +++ b/vtkm/cont/internal/testing/CMakeLists.txt @@ -12,6 +12,7 @@ set(unit_tests UnitTestArrayManagerExecutionShareWithControl.cxx UnitTestArrayPortalFromIterators.cxx + UnitTestBuffer.cxx UnitTestIteratorFromArrayPortal.cxx ) vtkm_unit_tests(SOURCES ${unit_tests} DEFINES VTKM_NO_ERROR_ON_MIXED_CUDA_CXX_TAG) diff --git a/vtkm/cont/internal/testing/UnitTestBuffer.cxx b/vtkm/cont/internal/testing/UnitTestBuffer.cxx new file mode 100644 index 000000000..291bb85be --- /dev/null +++ b/vtkm/cont/internal/testing/UnitTestBuffer.cxx @@ -0,0 +1,139 @@ +//============================================================================ +// 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. +//============================================================================ + +#include + +#include + +#include + +namespace +{ + +using T = vtkm::FloatDefault; +constexpr vtkm::Id ARRAY_SIZE = 20; + +using PortalType = vtkm::cont::internal::ArrayPortalFromIterators; +using PortalTypeConst = vtkm::cont::internal::ArrayPortalFromIterators; + +PortalType MakePortal(void* buffer, vtkm::Id numValues) +{ + return PortalType(static_cast(buffer), + static_cast(buffer) + static_cast(numValues)); +}; + +PortalTypeConst MakePortal(const void* buffer, vtkm::Id numValues) +{ + return PortalTypeConst(static_cast(buffer), + static_cast(buffer) + static_cast(numValues)); +}; + +struct VectorDeleter +{ + std::shared_ptr> Data; + + VectorDeleter(vtkm::Id numValues) + : Data(new std::vector(static_cast(numValues))) + { + } + + template + void operator()(U* p) + { + if (this->Data) + { + VTKM_TEST_ASSERT(reinterpret_cast(p) == &this->Data->front()); + this->Data.reset(); + } + } +}; + + +void DoTest() +{ + constexpr vtkm::Id BUFFER_SIZE = ARRAY_SIZE * static_cast(sizeof(T)); + constexpr vtkm::cont::DeviceAdapterTagSerial device; + + std::cout << "Initialize buffer" << std::endl; + vtkm::cont::internal::Buffer buffer; + buffer.SetNumberOfBytes(BUFFER_SIZE, vtkm::CopyFlag::Off); + + VTKM_TEST_ASSERT(buffer.GetNumberOfBytes() == BUFFER_SIZE); + + std::cout << "Fill up values on host" << std::endl; + { + vtkm::cont::Token token; + SetPortal(MakePortal(buffer.WritePointerHost(token), ARRAY_SIZE)); + } + VTKM_TEST_ASSERT(buffer.IsAllocatedOnHost()); + VTKM_TEST_ASSERT(!buffer.IsAllocatedOnDevice(device)); + + std::cout << "Check values" << std::endl; + { + vtkm::cont::Token token; + std::cout << " Host" << std::endl; + CheckPortal(MakePortal(buffer.ReadPointerHost(token), ARRAY_SIZE)); + std::cout << " Device" << std::endl; + CheckPortal(MakePortal(buffer.ReadPointerDevice(device, token), ARRAY_SIZE)); + } + VTKM_TEST_ASSERT(buffer.IsAllocatedOnHost()); + VTKM_TEST_ASSERT(buffer.IsAllocatedOnDevice(device)); + + std::cout << "Resize array and access write on device" << std::endl; + buffer.SetNumberOfBytes(BUFFER_SIZE / 2, vtkm::CopyFlag::On); + VTKM_TEST_ASSERT(buffer.GetNumberOfBytes() == BUFFER_SIZE / 2); + { + vtkm::cont::Token token; + CheckPortal(MakePortal(buffer.WritePointerDevice(device, token), ARRAY_SIZE / 2)); + } + VTKM_TEST_ASSERT(!buffer.IsAllocatedOnHost()); + VTKM_TEST_ASSERT(buffer.IsAllocatedOnDevice(device)); + + std::cout << "Resize array and access write on host" << std::endl; + // Note that this is a weird corner case where the array was resized while saving the data + // and then requested on another device. + buffer.SetNumberOfBytes(BUFFER_SIZE * 2, vtkm::CopyFlag::On); + VTKM_TEST_ASSERT(buffer.GetNumberOfBytes() == BUFFER_SIZE * 2); + { + vtkm::cont::Token token; + // Although the array is twice ARRAY_SIZE, the valid values are only ARRAY_SIZE/2 + CheckPortal(MakePortal(buffer.WritePointerHost(token), ARRAY_SIZE / 2)); + } + VTKM_TEST_ASSERT(buffer.IsAllocatedOnHost()); + VTKM_TEST_ASSERT(!buffer.IsAllocatedOnDevice(device)); + + std::cout << "Reset with device data" << std::endl; + VectorDeleter vectorDeleter(ARRAY_SIZE); + void* devicePointer = &vectorDeleter.Data->front(); + SetPortal(MakePortal(devicePointer, ARRAY_SIZE)); + buffer.Reset(devicePointer, BUFFER_SIZE, std::move(vectorDeleter), device); + VTKM_TEST_ASSERT(buffer.GetNumberOfBytes() == BUFFER_SIZE); + VTKM_TEST_ASSERT(!buffer.IsAllocatedOnHost()); + VTKM_TEST_ASSERT(buffer.IsAllocatedOnDevice(device)); + + std::cout << "Make sure device pointer is as expected" << std::endl; + { + vtkm::cont::Token token; + VTKM_TEST_ASSERT(buffer.WritePointerDevice(device, token) == devicePointer); + } + + std::cout << "Pull data to host" << std::endl; + { + vtkm::cont::Token token; + CheckPortal(MakePortal(buffer.ReadPointerHost(token), ARRAY_SIZE)); + } +} + +} // anonymous namespace + +int UnitTestBuffer(int argc, char* argv[]) +{ + return vtkm::cont::testing::Testing::Run(DoTest, argc, argv); +} diff --git a/vtkm/cont/openmp/DeviceAdapterOpenMP.h b/vtkm/cont/openmp/DeviceAdapterOpenMP.h index 9204eea54..0203e9bed 100644 --- a/vtkm/cont/openmp/DeviceAdapterOpenMP.h +++ b/vtkm/cont/openmp/DeviceAdapterOpenMP.h @@ -18,6 +18,7 @@ #include #include #include +#include #include #endif diff --git a/vtkm/cont/openmp/internal/CMakeLists.txt b/vtkm/cont/openmp/internal/CMakeLists.txt index 907452a0b..58f550378 100644 --- a/vtkm/cont/openmp/internal/CMakeLists.txt +++ b/vtkm/cont/openmp/internal/CMakeLists.txt @@ -11,6 +11,7 @@ set(headers ArrayManagerExecutionOpenMP.h DeviceAdapterAlgorithmOpenMP.h + DeviceAdapterMemoryManagerOpenMP.h DeviceAdapterRuntimeDetectorOpenMP.h DeviceAdapterTagOpenMP.h AtomicInterfaceExecutionOpenMP.h diff --git a/vtkm/cont/openmp/internal/DeviceAdapterMemoryManagerOpenMP.h b/vtkm/cont/openmp/internal/DeviceAdapterMemoryManagerOpenMP.h new file mode 100644 index 000000000..933fe5caf --- /dev/null +++ b/vtkm/cont/openmp/internal/DeviceAdapterMemoryManagerOpenMP.h @@ -0,0 +1,33 @@ +//============================================================================ +// 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. +//============================================================================ +#ifndef vtk_m_cont_openmp_internal_DeviceAdapterMemoryManagerOpenMP_h +#define vtk_m_cont_openmp_internal_DeviceAdapterMemoryManagerOpenMP_h + +#include + +#include + +namespace vtkm +{ +namespace cont +{ +namespace internal +{ + +template <> +class DeviceAdapterMemoryManager + : public vtkm::cont::internal::DeviceAdapterMemoryManagerShared +{ +}; +} +} +} + +#endif //vtk_m_cont_openmp_internal_DeviceAdapterMemoryManagerOpenMP_h diff --git a/vtkm/cont/serial/DeviceAdapterSerial.h b/vtkm/cont/serial/DeviceAdapterSerial.h index c000bef68..01f6d424d 100644 --- a/vtkm/cont/serial/DeviceAdapterSerial.h +++ b/vtkm/cont/serial/DeviceAdapterSerial.h @@ -14,6 +14,7 @@ // clang-format off #include #include +#include #include #include #include diff --git a/vtkm/cont/serial/internal/CMakeLists.txt b/vtkm/cont/serial/internal/CMakeLists.txt index 28995ac52..df3e344b8 100644 --- a/vtkm/cont/serial/internal/CMakeLists.txt +++ b/vtkm/cont/serial/internal/CMakeLists.txt @@ -12,6 +12,7 @@ set(headers ArrayManagerExecutionSerial.h AtomicInterfaceExecutionSerial.h DeviceAdapterAlgorithmSerial.h + DeviceAdapterMemoryManagerSerial.h DeviceAdapterRuntimeDetectorSerial.h DeviceAdapterTagSerial.h ExecutionArrayInterfaceBasicSerial.h diff --git a/vtkm/cont/serial/internal/DeviceAdapterMemoryManagerSerial.h b/vtkm/cont/serial/internal/DeviceAdapterMemoryManagerSerial.h new file mode 100644 index 000000000..d695954ff --- /dev/null +++ b/vtkm/cont/serial/internal/DeviceAdapterMemoryManagerSerial.h @@ -0,0 +1,33 @@ +//============================================================================ +// 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. +//============================================================================ +#ifndef vtk_m_cont_serial_internal_DeviceAdapterMemoryManagerSerial_h +#define vtk_m_cont_serial_internal_DeviceAdapterMemoryManagerSerial_h + +#include + +#include + +namespace vtkm +{ +namespace cont +{ +namespace internal +{ + +template <> +class DeviceAdapterMemoryManager + : public vtkm::cont::internal::DeviceAdapterMemoryManagerShared +{ +}; +} +} +} + +#endif //vtk_m_cont_serial_internal_DeviceAdapterMemoryManagerSerial_h diff --git a/vtkm/cont/tbb/DeviceAdapterTBB.h b/vtkm/cont/tbb/DeviceAdapterTBB.h index f870fe37f..b3852b3d2 100644 --- a/vtkm/cont/tbb/DeviceAdapterTBB.h +++ b/vtkm/cont/tbb/DeviceAdapterTBB.h @@ -17,6 +17,7 @@ #include #include #include +#include #include #endif diff --git a/vtkm/cont/tbb/internal/CMakeLists.txt b/vtkm/cont/tbb/internal/CMakeLists.txt index 32289294d..5469a0535 100644 --- a/vtkm/cont/tbb/internal/CMakeLists.txt +++ b/vtkm/cont/tbb/internal/CMakeLists.txt @@ -12,6 +12,7 @@ set(headers ArrayManagerExecutionTBB.h AtomicInterfaceExecutionTBB.h DeviceAdapterAlgorithmTBB.h + DeviceAdapterMemoryManagerTBB.h DeviceAdapterRuntimeDetectorTBB.h DeviceAdapterTagTBB.h ExecutionArrayInterfaceBasicTBB.h diff --git a/vtkm/cont/tbb/internal/DeviceAdapterMemoryManagerTBB.h b/vtkm/cont/tbb/internal/DeviceAdapterMemoryManagerTBB.h new file mode 100644 index 000000000..d03e9464a --- /dev/null +++ b/vtkm/cont/tbb/internal/DeviceAdapterMemoryManagerTBB.h @@ -0,0 +1,33 @@ +//============================================================================ +// 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. +//============================================================================ +#ifndef vtk_m_cont_tbb_internal_DeviceAdapterMemoryManagerTBB_h +#define vtk_m_cont_tbb_internal_DeviceAdapterMemoryManagerTBB_h + +#include + +#include + +namespace vtkm +{ +namespace cont +{ +namespace internal +{ + +template <> +class DeviceAdapterMemoryManager + : public vtkm::cont::internal::DeviceAdapterMemoryManagerShared +{ +}; +} +} +} + +#endif //vtk_m_cont_tbb_internal_DeviceAdapterMemoryManagerTBB_h diff --git a/vtkm/cont/testing/TestingDeviceAdapter.h b/vtkm/cont/testing/TestingDeviceAdapter.h index a50831d1e..f881b9aee 100644 --- a/vtkm/cont/testing/TestingDeviceAdapter.h +++ b/vtkm/cont/testing/TestingDeviceAdapter.h @@ -157,17 +157,18 @@ public: } }; + template struct AddArrayKernel { VTKM_CONT - AddArrayKernel(const IdPortalType& array) + AddArrayKernel(const PortalType& array) : Array(array) , Dims() { } VTKM_CONT - AddArrayKernel(const IdPortalType& array, const vtkm::Id3& dims) + AddArrayKernel(const PortalType& array, const vtkm::Id3& dims) : Array(array) , Dims(dims) { @@ -187,10 +188,18 @@ public: VTKM_CONT void SetErrorMessageBuffer(const vtkm::exec::internal::ErrorMessageBuffer&) {} - IdPortalType Array; + PortalType Array; vtkm::Id3 Dims; }; + template + static VTKM_CONT AddArrayKernel MakeAddArrayKernel( + PortalType portal, + const vtkm::Id3& dims = vtkm::Id3{}) + { + return AddArrayKernel(portal, dims); + } + // Checks that each instance is only visited once: struct OverlapKernel { @@ -503,10 +512,79 @@ private: "Device adapter Name does not equal itself."); } - // Note: this test does not actually test to make sure the data is available - // in the execution environment. It tests to make sure data gets to the array - // and back, but it is possible that the data is not available in the - // execution environment. + static VTKM_CONT void TestMemoryTransfer() + { + std::cout << "-------------------------------------------" << std::endl; + std::cout << "Testing Memory Transfer" << std::endl; + + using T = vtkm::Id; + using PortalType = vtkm::cont::internal::ArrayPortalFromIterators; + auto makePortal = [](const vtkm::cont::internal::BufferInfo& buffer) { + return PortalType(static_cast(buffer.GetPointer()), + static_cast(buffer.GetPointer()) + + static_cast(buffer.GetSize()) / sizeof(T)); + }; + + constexpr vtkm::BufferSizeType BUFFER_SIZE = + ARRAY_SIZE * static_cast(sizeof(T)); + + // Set up buffer on host. + std::shared_ptr hostBufferSrc( + new vtkm::cont::internal::BufferInfoHost(BUFFER_SIZE)); + VTKM_TEST_ASSERT(hostBufferSrc->GetSize() == BUFFER_SIZE); + SetPortal(makePortal(*hostBufferSrc)); + + vtkm::cont::internal::DeviceAdapterMemoryManager memoryManager; + + std::cout << "Allocate a buffer." << std::endl; + std::shared_ptr allocatedMemory = + memoryManager.Allocate(BUFFER_SIZE); + VTKM_TEST_ASSERT(allocatedMemory->GetSize() == BUFFER_SIZE); + + std::cout << "Copy data from host to device." << std::endl; + allocatedMemory = memoryManager.CopyHostToDevice(hostBufferSrc); + + std::cout << "Copy data within device." << std::endl; + std::shared_ptr workingMemory = + memoryManager.CopyDeviceToDevice(allocatedMemory); + VTKM_TEST_ASSERT(workingMemory->GetSize() == BUFFER_SIZE); + + std::cout << "Copy data back to host." << std::endl; + std::shared_ptr hostBufferDest = + memoryManager.CopyDeviceToHost(workingMemory); + VTKM_TEST_ASSERT(hostBufferDest->GetSize() == BUFFER_SIZE); + CheckPortal(makePortal(*hostBufferDest)); + + std::cout << "Shrink a buffer (and preserve memory)" << std::endl; + memoryManager.Reallocate(workingMemory, BUFFER_SIZE / 2); + hostBufferDest = memoryManager.CopyDeviceToHost(workingMemory); + VTKM_TEST_ASSERT(hostBufferDest->GetSize() == BUFFER_SIZE / 2); + CheckPortal(makePortal(*hostBufferDest)); + + std::cout << "Grow a buffer (and preserve memory)" << std::endl; + memoryManager.Reallocate(workingMemory, BUFFER_SIZE * 2); + hostBufferDest = memoryManager.CopyDeviceToHost(workingMemory); + VTKM_TEST_ASSERT(hostBufferDest->GetSize() == BUFFER_SIZE * 2); + hostBufferDest->Allocate(BUFFER_SIZE / 2, vtkm::CopyFlag::On); + CheckPortal(makePortal(*hostBufferDest)); + + std::cout << "Make sure data is actually available on the device." << std::endl; + // This actually requires running schedule. + workingMemory = memoryManager.CopyDeviceToDevice(allocatedMemory); + Algorithm::Schedule(MakeAddArrayKernel(makePortal(*workingMemory)), ARRAY_SIZE); + + hostBufferDest = memoryManager.CopyDeviceToHost(workingMemory); + + PortalType portal = makePortal(*hostBufferDest); + VTKM_TEST_ASSERT(portal.GetNumberOfValues() == ARRAY_SIZE); + for (vtkm::Id index = 0; index < ARRAY_SIZE; ++index) + { + T expected = TestValue(index, T()) + T(index); + T computed = portal.Get(index); + VTKM_TEST_ASSERT(test_equal(expected, computed), expected, " != ", computed); + } + } + static VTKM_CONT void TestArrayTransfer() { std::cout << "-------------------------------------------" << std::endl; @@ -534,7 +612,7 @@ private: // because we are about to shrink. { vtkm::cont::Token token; - Algorithm::Schedule(AddArrayKernel(handle.PrepareForInPlace(DeviceAdapterTag{}, token)), + Algorithm::Schedule(MakeAddArrayKernel(handle.PrepareForInPlace(DeviceAdapterTag{}, token)), ARRAY_SIZE); } @@ -669,7 +747,8 @@ private: std::cout << "Running add." << std::endl; { vtkm::cont::Token token; - Algorithm::Schedule(AddArrayKernel(handle.PrepareForInPlace(DeviceAdapterTag{}, token)), 1); + Algorithm::Schedule(MakeAddArrayKernel(handle.PrepareForInPlace(DeviceAdapterTag{}, token)), + 1); } std::cout << "Checking results." << std::endl; @@ -700,7 +779,7 @@ private: std::cout << "Running add." << std::endl; { vtkm::cont::Token token; - Algorithm::Schedule(AddArrayKernel(handle.PrepareForInPlace(DeviceAdapterTag{}, token)), + Algorithm::Schedule(MakeAddArrayKernel(handle.PrepareForInPlace(DeviceAdapterTag{}, token)), ARRAY_SIZE); } @@ -734,7 +813,7 @@ private: std::cout << "Running add." << std::endl; { vtkm::cont::Token token; - Algorithm::Schedule(AddArrayKernel(handle.PrepareForInPlace(DeviceAdapterTag{}, token)), + Algorithm::Schedule(MakeAddArrayKernel(handle.PrepareForInPlace(DeviceAdapterTag{}, token)), size); } @@ -776,7 +855,8 @@ private: { vtkm::cont::Token token; Algorithm::Schedule( - AddArrayKernel(handle.PrepareForInPlace(DeviceAdapterTag{}, token), maxRange), maxRange); + MakeAddArrayKernel(handle.PrepareForInPlace(DeviceAdapterTag{}, token), maxRange), + maxRange); } std::cout << "Checking results." << std::endl; @@ -2071,7 +2151,8 @@ private: ClearArrayKernel(array.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag(), token)), ARRAY_SIZE); Algorithm::Schedule( - AddArrayKernel(array.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag(), token)), ARRAY_SIZE); + MakeAddArrayKernel(array.PrepareForOutput(ARRAY_SIZE, DeviceAdapterTag(), token)), + ARRAY_SIZE); } //we know have an array whose sum is equal to OFFSET * ARRAY_SIZE, @@ -2366,7 +2447,7 @@ private: Algorithm::Schedule(OneErrorKernel(), ARRAY_SIZE); for (; nkernels < 100; ++nkernels) { - Algorithm::Schedule(AddArrayKernel(portal), ARRAY_SIZE); + Algorithm::Schedule(MakeAddArrayKernel(portal), ARRAY_SIZE); std::this_thread::sleep_for(std::chrono::milliseconds(20)); } Algorithm::Synchronize(); @@ -3097,6 +3178,7 @@ private: { std::cout << "Doing DeviceAdapter tests" << std::endl; + TestMemoryTransfer(); TestArrayTransfer(); TestOutOfMemory(); TestTimer(); diff --git a/vtkm/cont/testing/UnitTestDeviceAdapterAlgorithmGeneral.cxx b/vtkm/cont/testing/UnitTestDeviceAdapterAlgorithmGeneral.cxx index ec5def2bc..112ee2246 100644 --- a/vtkm/cont/testing/UnitTestDeviceAdapterAlgorithmGeneral.cxx +++ b/vtkm/cont/testing/UnitTestDeviceAdapterAlgorithmGeneral.cxx @@ -71,6 +71,12 @@ public: namespace internal { +template <> +class DeviceAdapterMemoryManager + : public vtkm::cont::internal::DeviceAdapterMemoryManagerShared +{ +}; + template class ArrayManagerExecution : public vtkm::cont::internal::ArrayManagerExecution