mirror of
https://gitlab.kitware.com/vtk/vtk-m
synced 2024-09-16 17:22:55 +00:00
Add Buffer class
The buffer class encapsulates the movement of raw C arrays between host and devices. The `Buffer` class itself is not associated with any device. Instead, `Buffer` is used in conjunction with a new templated class named `DeviceAdapterMemoryManager` that can allocate data on a given device and transfer data as necessary. `DeviceAdapterMemoryManager` will eventually replace the more complicated device adapter classes that manage data on a device. The code in `DeviceAdapterMemoryManager` is actually enclosed in virtual methods. This allows us to limit the number of classes that need to be compiled for a device. Rather, the implementation of `DeviceAdapterMemoryManager` is compiled once with whatever compiler is necessary, and then the `RuntimeDeviceInformation` is used to get the correct object instance.
This commit is contained in:
parent
a338612e95
commit
8f7b0d18be
@ -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
|
||||
|
@ -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 <vtkm/cont/cuda/DeviceAdapterCuda.h>
|
||||
@ -23,6 +23,7 @@
|
||||
#include <vtkm/cont/DeviceAdapterAlgorithm.h>
|
||||
#include <vtkm/cont/DeviceAdapterTag.h>
|
||||
#include <vtkm/cont/internal/ArrayManagerExecution.h>
|
||||
#include <vtkm/cont/internal/DeviceAdapterMemoryManager.h>
|
||||
|
||||
// clang-format on
|
||||
|
||||
|
@ -13,6 +13,7 @@
|
||||
#include <vtkm/cont/DeviceAdapter.h>
|
||||
#include <vtkm/cont/DeviceAdapterList.h>
|
||||
#include <vtkm/cont/DeviceAdapterTag.h>
|
||||
#include <vtkm/cont/ErrorBadDevice.h>
|
||||
|
||||
//Bring in each device adapters runtime class
|
||||
#include <vtkm/cont/cuda/internal/DeviceAdapterRuntimeDetectorCuda.h>
|
||||
@ -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<vtkm::cont::internal::BufferInfo> 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<vtkm::cont::internal::BufferInfo> ManageArray(
|
||||
std::shared_ptr<vtkm::UInt8>,
|
||||
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::cont::internal::BufferInfo>,
|
||||
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<vtkm::cont::internal::BufferInfo> CopyHostToDevice(
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfoHost>) 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<vtkm::cont::internal::BufferInfoHost> CopyDeviceToHost(
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfo>) 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<vtkm::cont::internal::BufferInfo> CopyDeviceToDevice(
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfo>) 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<vtkm::cont::internal::DeviceAdapterMemoryManagerBase>* Managers;
|
||||
|
||||
VTKM_CONT
|
||||
InitializeDeviceMemoryManagers(
|
||||
std::unique_ptr<vtkm::cont::internal::DeviceAdapterMemoryManagerBase>* managers)
|
||||
: Managers(managers)
|
||||
{
|
||||
}
|
||||
|
||||
template <typename Device>
|
||||
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<Device>::GetName();
|
||||
this->Managers[id].reset(new vtkm::cont::internal::DeviceAdapterMemoryManager<Device>);
|
||||
}
|
||||
}
|
||||
|
||||
template <typename Device>
|
||||
VTKM_CONT void CreateManager(Device, std::false_type)
|
||||
{
|
||||
// No manager for invalid devices.
|
||||
}
|
||||
|
||||
template <typename Device>
|
||||
VTKM_CONT void operator()(Device device)
|
||||
{
|
||||
this->CreateManager(device, std::integral_constant<bool, device.IsEnabled>{});
|
||||
}
|
||||
};
|
||||
|
||||
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<vtkm::cont::internal::DeviceAdapterMemoryManagerBase>
|
||||
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
|
||||
|
@ -11,6 +11,7 @@
|
||||
#define vtk_m_cont_RuntimeDeviceInformation_h
|
||||
|
||||
#include <vtkm/cont/DeviceAdapterTag.h>
|
||||
#include <vtkm/cont/internal/DeviceAdapterMemoryManager.h>
|
||||
#include <vtkm/internal/ExportMacros.h>
|
||||
|
||||
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
|
||||
|
@ -20,6 +20,7 @@
|
||||
|
||||
#include <vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.h>
|
||||
#include <vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h>
|
||||
#include <vtkm/cont/cuda/internal/DeviceAdapterMemoryManagerCuda.h>
|
||||
#include <vtkm/cont/cuda/internal/VirtualObjectTransferCuda.h>
|
||||
#endif
|
||||
|
||||
|
@ -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
|
||||
|
200
vtkm/cont/cuda/internal/DeviceAdapterMemoryManagerCuda.cu
Normal file
200
vtkm/cont/cuda/internal/DeviceAdapterMemoryManagerCuda.cu
Normal file
@ -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 <vtkm/cont/cuda/DeviceAdapterCuda.h>
|
||||
#include <vtkm/cont/cuda/ErrorCuda.h>
|
||||
#include <vtkm/cont/cuda/internal/CudaAllocator.h>
|
||||
#include <vtkm/cont/cuda/internal/DeviceAdapterMemoryManagerCuda.h>
|
||||
|
||||
#include <vtkm/cont/ErrorBadAllocation.h>
|
||||
|
||||
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<vtkm::UInt8> CudaBuffer;
|
||||
vtkm::BufferSizeType Size;
|
||||
|
||||
VTKM_CONT BufferInfoCuda(const std::shared_ptr<vtkm::UInt8> 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::BufferInfo> vtkm::cont::internal::DeviceAdapterMemoryManager<
|
||||
vtkm::cont::DeviceAdapterTagCuda>::Allocate(vtkm::BufferSizeType size)
|
||||
{
|
||||
try
|
||||
{
|
||||
vtkm::UInt8* buffer = static_cast<vtkm::UInt8*>(
|
||||
vtkm::cont::cuda::internal::CudaAllocator::Allocate(static_cast<std::size_t>(size)));
|
||||
std::shared_ptr<vtkm::UInt8> bufferPtr(buffer, CudaDeleter{});
|
||||
return std::shared_ptr<vtkm::cont::internal::BufferInfo>(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::BufferInfo> vtkm::cont::internal::DeviceAdapterMemoryManager<
|
||||
vtkm::cont::DeviceAdapterTagCuda>::ManageArray(std::shared_ptr<vtkm::UInt8> buffer,
|
||||
vtkm::BufferSizeType size)
|
||||
{
|
||||
return std::shared_ptr<vtkm::cont::internal::BufferInfo>(new BufferInfoCuda(buffer, size));
|
||||
}
|
||||
|
||||
void vtkm::cont::internal::DeviceAdapterMemoryManager<vtkm::cont::DeviceAdapterTagCuda>::Reallocate(
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfo> b,
|
||||
vtkm::BufferSizeType newSize)
|
||||
{
|
||||
BufferInfoCuda* buffer = dynamic_cast<BufferInfoCuda*>(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<vtkm::cont::internal::BufferInfo> newBufferInfo = this->Allocate(newSize);
|
||||
BufferInfoCuda* newBuffer = dynamic_cast<BufferInfoCuda*>(newBufferInfo.get());
|
||||
VTKM_ASSERT(newBuffer != nullptr);
|
||||
|
||||
// Copy the data to the new buffer
|
||||
VTKM_CUDA_CALL(cudaMemcpyAsync(newBuffer->GetPointer(),
|
||||
buffer->GetPointer(),
|
||||
static_cast<std::size_t>(buffer->Size),
|
||||
cudaMemcpyDeviceToDevice,
|
||||
cudaStreamPerThread));
|
||||
|
||||
// Reset the buffer in the passed in info
|
||||
*buffer = *newBuffer;
|
||||
}
|
||||
}
|
||||
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfo> vtkm::cont::internal::
|
||||
DeviceAdapterMemoryManager<vtkm::cont::DeviceAdapterTagCuda>::CopyHostToDevice(
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfoHost> 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<std::size_t>(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<vtkm::cont::internal::BufferInfo>(
|
||||
new BufferInfoCuda(src->GetSharedPointer(), src->GetSize()));
|
||||
}
|
||||
else
|
||||
{
|
||||
// Make a new buffer
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfo> destInfo = this->Allocate(src->GetSize());
|
||||
BufferInfoCuda* dest = dynamic_cast<BufferInfoCuda*>(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<std::size_t>(src->GetSize())).c_str(),
|
||||
src->GetSize());
|
||||
|
||||
VTKM_CUDA_CALL(cudaMemcpyAsync(dest->GetPointer(),
|
||||
src->GetPointer(),
|
||||
static_cast<std::size_t>(src->GetSize()),
|
||||
cudaMemcpyHostToDevice,
|
||||
cudaStreamPerThread));
|
||||
|
||||
return destInfo;
|
||||
}
|
||||
}
|
||||
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfoHost> vtkm::cont::internal::
|
||||
DeviceAdapterMemoryManager<vtkm::cont::DeviceAdapterTagCuda>::CopyDeviceToHost(
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfo> src_)
|
||||
{
|
||||
using vtkm::cont::cuda::internal::CudaAllocator;
|
||||
BufferInfoCuda* src = dynamic_cast<BufferInfoCuda*>(src_.get());
|
||||
VTKM_ASSERT(src != nullptr);
|
||||
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfoHost> 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<std::size_t>(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<vtkm::UInt64>(src->Size)).c_str(),
|
||||
src->Size);
|
||||
|
||||
VTKM_CUDA_CALL(cudaMemcpyAsync(dest->GetPointer(),
|
||||
src->GetPointer(),
|
||||
static_cast<std::size_t>(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<DeviceAdapterTagCuda>::Synchronize();
|
||||
|
||||
return dest;
|
||||
}
|
||||
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfo> vtkm::cont::internal::
|
||||
DeviceAdapterMemoryManager<vtkm::cont::DeviceAdapterTagCuda>::CopyDeviceToDevice(
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfo> src)
|
||||
{
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfo> dest = this->Allocate(src->GetSize());
|
||||
VTKM_CUDA_CALL(cudaMemcpyAsync(dest->GetPointer(),
|
||||
src->GetPointer(),
|
||||
static_cast<std::size_t>(src->GetSize()),
|
||||
cudaMemcpyDeviceToDevice,
|
||||
cudaStreamPerThread));
|
||||
|
||||
return dest;
|
||||
}
|
50
vtkm/cont/cuda/internal/DeviceAdapterMemoryManagerCuda.h
Normal file
50
vtkm/cont/cuda/internal/DeviceAdapterMemoryManagerCuda.h
Normal 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.
|
||||
//============================================================================
|
||||
#ifndef vtk_m_cont_cuda_internal_DeviceAdapterMemoryManagerCuda_h
|
||||
#define vtk_m_cont_cuda_internal_DeviceAdapterMemoryManagerCuda_h
|
||||
|
||||
#include <vtkm/cont/cuda/internal/DeviceAdapterTagCuda.h>
|
||||
|
||||
#include <vtkm/cont/internal/DeviceAdapterMemoryManager.h>
|
||||
|
||||
namespace vtkm
|
||||
{
|
||||
namespace cont
|
||||
{
|
||||
namespace internal
|
||||
{
|
||||
|
||||
template <>
|
||||
class VTKM_CONT_EXPORT DeviceAdapterMemoryManager<vtkm::cont::DeviceAdapterTagCuda>
|
||||
: public DeviceAdapterMemoryManagerBase
|
||||
{
|
||||
public:
|
||||
VTKM_CONT std::shared_ptr<BufferInfo> Allocate(vtkm::BufferSizeType size) override;
|
||||
|
||||
VTKM_CONT std::shared_ptr<BufferInfo> ManageArray(std::shared_ptr<vtkm::UInt8> buffer,
|
||||
vtkm::BufferSizeType size) override;
|
||||
|
||||
VTKM_CONT void Reallocate(std::shared_ptr<vtkm::cont::internal::BufferInfo> buffer,
|
||||
vtkm::BufferSizeType newSize) override;
|
||||
|
||||
VTKM_CONT std::shared_ptr<vtkm::cont::internal::BufferInfo> CopyHostToDevice(
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfoHost> src) override;
|
||||
|
||||
VTKM_CONT std::shared_ptr<vtkm::cont::internal::BufferInfoHost> CopyDeviceToHost(
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfo> src) override;
|
||||
|
||||
VTKM_CONT std::shared_ptr<vtkm::cont::internal::BufferInfo> CopyDeviceToDevice(
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfo> src) override;
|
||||
};
|
||||
}
|
||||
}
|
||||
} // namespace vtkm::cont::internal
|
||||
|
||||
#endif //vtk_m_cont_cuda_internal_DeviceAdapterMemoryManagerCuda_h
|
508
vtkm/cont/internal/Buffer.cxx
Normal file
508
vtkm/cont/internal/Buffer.cxx
Normal file
@ -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 <vtkm/internal/Assume.h>
|
||||
|
||||
#include <vtkm/cont/RuntimeDeviceInformation.h>
|
||||
|
||||
#include <vtkm/cont/internal/Buffer.h>
|
||||
#include <vtkm/cont/internal/DeviceAdapterMemoryManager.h>
|
||||
|
||||
#include <condition_variable>
|
||||
#include <map>
|
||||
|
||||
using LockType = std::unique_lock<std::mutex>;
|
||||
|
||||
namespace vtkm
|
||||
{
|
||||
namespace cont
|
||||
{
|
||||
namespace internal
|
||||
{
|
||||
|
||||
class Buffer::InternalsStruct
|
||||
{
|
||||
public:
|
||||
using DeviceBufferMap =
|
||||
std::map<vtkm::cont::DeviceAdapterId, std::shared_ptr<vtkm::cont::internal::BufferInfo>>;
|
||||
using HostBufferPointer = std::shared_ptr<vtkm::cont::internal::BufferInfoHost>;
|
||||
|
||||
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<std::mutex>& 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<std::mutex>& 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<const vtkm::UInt8*>(srcBuffer->GetPointer()),
|
||||
reinterpret_cast<const vtkm::UInt8*>(srcBuffer->GetPointer()) + srcBuffer->GetSize(),
|
||||
reinterpret_cast<vtkm::UInt8*>(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<vtkm::UInt8> 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<vtkm::UInt8> 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<vtkm::cont::internal::Buffer>::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<const vtkm::UInt8*>(obj.ReadPointerHost(token));
|
||||
vtkmdiy::save(bb, data, static_cast<std::size_t>(size));
|
||||
}
|
||||
|
||||
void Serialization<vtkm::cont::internal::Buffer>::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<vtkm::UInt8*>(obj.WritePointerHost(token));
|
||||
vtkmdiy::load(bb, data, static_cast<std::size_t>(size));
|
||||
}
|
||||
|
||||
} // namespace diy
|
236
vtkm/cont/internal/Buffer.h
Normal file
236
vtkm/cont/internal/Buffer.h
Normal file
@ -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 <vtkm/cont/vtkm_cont_export.h>
|
||||
|
||||
#include <vtkm/cont/DeviceAdapterTag.h>
|
||||
#include <vtkm/cont/Serialization.h>
|
||||
#include <vtkm/cont/Token.h>
|
||||
|
||||
#include <vtkm/cont/internal/DeviceAdapterMemoryManager.h>
|
||||
|
||||
#include <memory>
|
||||
#include <mutex>
|
||||
|
||||
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<InternalsStruct> 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<vtkm::UInt8> 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<vtkm::UInt8> 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 <typename T>
|
||||
VTKM_CONT void Reset(T* buffer, vtkm::BufferSizeType numberOfBytes)
|
||||
{
|
||||
this->Reset(buffer, numberOfBytes, std::default_delete<T[]>{});
|
||||
}
|
||||
|
||||
/// \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 <typename T, typename Deleter>
|
||||
VTKM_CONT void Reset(T* buffer, vtkm::BufferSizeType numberOfBytes, Deleter deleter)
|
||||
{
|
||||
std::shared_ptr<vtkm::UInt8> sharedP(reinterpret_cast<vtkm::UInt8*>(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 <typename T, typename Deleter>
|
||||
VTKM_CONT void Reset(T* buffer,
|
||||
vtkm::BufferSizeType numberOfBytes,
|
||||
Deleter deleter,
|
||||
vtkm::cont::DeviceAdapterId device)
|
||||
{
|
||||
std::shared_ptr<vtkm::UInt8> sharedP(reinterpret_cast<vtkm::UInt8*>(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<vtkm::cont::internal::Buffer>
|
||||
{
|
||||
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 <vtkm::IdComponent N>
|
||||
struct Serialization<vtkm::Vec<vtkm::cont::internal::Buffer, N>>
|
||||
{
|
||||
static VTKM_CONT void save(BinaryBuffer& bb,
|
||||
const vtkm::Vec<vtkm::cont::internal::Buffer, N>& obj)
|
||||
{
|
||||
for (vtkm::IdComponent index = 0; index < N; ++index)
|
||||
{
|
||||
vtkmdiy::save(bb, obj[index]);
|
||||
}
|
||||
}
|
||||
|
||||
static VTKM_CONT void load(BinaryBuffer& bb, vtkm::Vec<vtkm::cont::internal::Buffer, N>& 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
|
@ -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
|
||||
|
153
vtkm/cont/internal/DeviceAdapterMemoryManager.cxx
Normal file
153
vtkm/cont/internal/DeviceAdapterMemoryManager.cxx
Normal file
@ -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 <vtkm/cont/internal/DeviceAdapterMemoryManager.h>
|
||||
|
||||
#include <algorithm>
|
||||
|
||||
//----------------------------------------------------------------------------------------
|
||||
// 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 <stdlib.h>
|
||||
#elif defined(VTKM_MEMALIGN_WIN)
|
||||
#include <malloc.h>
|
||||
#elif defined(VTKM_MEMALIGN_SSE)
|
||||
#include <xmmintrin.h>
|
||||
#else
|
||||
#include <malloc.h>
|
||||
#endif
|
||||
|
||||
#include <cstddef>
|
||||
#include <cstdlib>
|
||||
|
||||
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<vtkm::UInt8> HostAllocate(vtkm::BufferSizeType numBytes)
|
||||
{
|
||||
const std::size_t size = static_cast<std::size_t>(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<vtkm::UInt8>(reinterpret_cast<vtkm::UInt8*>(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<UInt8>& 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<vtkm::UInt8> 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::UInt8> vtkm::cont::internal::BufferInfoHost::GetSharedPointer() const
|
||||
{
|
||||
return this->Buffer;
|
||||
}
|
||||
|
||||
vtkm::cont::internal::DeviceAdapterMemoryManagerBase::~DeviceAdapterMemoryManagerBase()
|
||||
{
|
||||
}
|
121
vtkm/cont/internal/DeviceAdapterMemoryManager.h
Normal file
121
vtkm/cont/internal/DeviceAdapterMemoryManager.h
Normal file
@ -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 <vtkm/cont/vtkm_cont_export.h>
|
||||
|
||||
#include <vtkm/Flags.h>
|
||||
#include <vtkm/Types.h>
|
||||
|
||||
#include <vtkm/cont/DeviceAdapterTag.h>
|
||||
|
||||
#include <memory>
|
||||
|
||||
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<vtkm::UInt8> Buffer;
|
||||
vtkm::BufferSizeType Size;
|
||||
|
||||
public:
|
||||
VTKM_CONT BufferInfoHost();
|
||||
VTKM_CONT BufferInfoHost(const std::shared_ptr<vtkm::UInt8>& 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<vtkm::UInt8> 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<BufferInfo> 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<BufferInfo> ManageArray(std::shared_ptr<vtkm::UInt8> 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<vtkm::cont::internal::BufferInfo> 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<vtkm::cont::internal::BufferInfo> CopyHostToDevice(
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfoHost> 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<vtkm::cont::internal::BufferInfoHost> CopyDeviceToHost(
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfo> 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<vtkm::cont::internal::BufferInfo> CopyDeviceToDevice(
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfo> 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 <typename DeviceAdapterTag>
|
||||
class DeviceAdapterMemoryManager;
|
||||
}
|
||||
}
|
||||
} // namespace vtkm::cont::internal
|
||||
|
||||
#endif //vtk_m_cont_internal_DeviceAdapterMemoryManager_h
|
101
vtkm/cont/internal/DeviceAdapterMemoryManagerShared.cxx
Normal file
101
vtkm/cont/internal/DeviceAdapterMemoryManagerShared.cxx
Normal file
@ -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 <vtkm/cont/internal/DeviceAdapterMemoryManagerShared.h>
|
||||
|
||||
#include <algorithm>
|
||||
|
||||
namespace
|
||||
{
|
||||
|
||||
class BufferInfoShared : public vtkm::cont::internal::BufferInfo
|
||||
{
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfoHost> HostBuffer;
|
||||
|
||||
public:
|
||||
VTKM_CONT BufferInfoShared()
|
||||
: HostBuffer(new vtkm::cont::internal::BufferInfoHost)
|
||||
{
|
||||
}
|
||||
|
||||
VTKM_CONT BufferInfoShared(std::shared_ptr<vtkm::cont::internal::BufferInfoHost> hostBuffer)
|
||||
: HostBuffer(hostBuffer)
|
||||
{
|
||||
}
|
||||
|
||||
template <typename... Ts>
|
||||
VTKM_CONT BufferInfoShared(Ts&&... bufferInfoHostArgs)
|
||||
: HostBuffer(new vtkm::cont::internal::BufferInfoHost(std::forward<Ts>(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<vtkm::cont::internal::BufferInfoHost> GetHostBuffer() const
|
||||
{
|
||||
return this->HostBuffer;
|
||||
}
|
||||
};
|
||||
|
||||
} // anonymous namespace
|
||||
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfo>
|
||||
vtkm::cont::internal::DeviceAdapterMemoryManagerShared::Allocate(vtkm::BufferSizeType size)
|
||||
{
|
||||
return std::shared_ptr<vtkm::cont::internal::BufferInfo>(new BufferInfoShared(size));
|
||||
}
|
||||
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfo>
|
||||
vtkm::cont::internal::DeviceAdapterMemoryManagerShared::ManageArray(
|
||||
std::shared_ptr<vtkm::UInt8> buffer,
|
||||
vtkm::BufferSizeType size)
|
||||
{
|
||||
return std::shared_ptr<vtkm::cont::internal::BufferInfo>(new BufferInfoShared(buffer, size));
|
||||
}
|
||||
|
||||
void vtkm::cont::internal::DeviceAdapterMemoryManagerShared::Reallocate(
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfo> b,
|
||||
vtkm::BufferSizeType newSize)
|
||||
{
|
||||
BufferInfoShared* buffer = dynamic_cast<BufferInfoShared*>(b.get());
|
||||
VTKM_ASSERT(buffer != nullptr);
|
||||
|
||||
buffer->GetHostBuffer()->Allocate(newSize, vtkm::CopyFlag::On);
|
||||
}
|
||||
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfo>
|
||||
vtkm::cont::internal::DeviceAdapterMemoryManagerShared::CopyHostToDevice(
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfoHost> src)
|
||||
{
|
||||
return std::shared_ptr<vtkm::cont::internal::BufferInfo>(new BufferInfoShared(src));
|
||||
}
|
||||
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfoHost>
|
||||
vtkm::cont::internal::DeviceAdapterMemoryManagerShared::CopyDeviceToHost(
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfo> src)
|
||||
{
|
||||
BufferInfoShared* buffer = dynamic_cast<BufferInfoShared*>(src.get());
|
||||
VTKM_ASSERT(buffer != nullptr);
|
||||
|
||||
return buffer->GetHostBuffer();
|
||||
}
|
||||
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfo>
|
||||
vtkm::cont::internal::DeviceAdapterMemoryManagerShared::CopyDeviceToDevice(
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfo> src)
|
||||
{
|
||||
BufferInfoShared* dest = new BufferInfoShared(src->GetSize());
|
||||
vtkm::UInt8* srcP = reinterpret_cast<vtkm::UInt8*>(src->GetPointer());
|
||||
vtkm::UInt8* destP = reinterpret_cast<vtkm::UInt8*>(dest->GetPointer());
|
||||
std::copy(srcP, srcP + src->GetSize(), destP);
|
||||
return std::shared_ptr<vtkm::cont::internal::BufferInfo>(dest);
|
||||
}
|
52
vtkm/cont/internal/DeviceAdapterMemoryManagerShared.h
Normal file
52
vtkm/cont/internal/DeviceAdapterMemoryManagerShared.h
Normal file
@ -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 <vtkm/cont/internal/DeviceAdapterMemoryManager.h>
|
||||
|
||||
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<BufferInfo> Allocate(vtkm::BufferSizeType size) override;
|
||||
|
||||
VTKM_CONT std::shared_ptr<BufferInfo> ManageArray(std::shared_ptr<vtkm::UInt8> buffer,
|
||||
vtkm::BufferSizeType size) override;
|
||||
|
||||
VTKM_CONT void Reallocate(std::shared_ptr<vtkm::cont::internal::BufferInfo> buffer,
|
||||
vtkm::BufferSizeType newSize) override;
|
||||
|
||||
VTKM_CONT std::shared_ptr<vtkm::cont::internal::BufferInfo> CopyHostToDevice(
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfoHost> src) override;
|
||||
|
||||
VTKM_CONT std::shared_ptr<vtkm::cont::internal::BufferInfoHost> CopyDeviceToHost(
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfo> src) override;
|
||||
|
||||
VTKM_CONT std::shared_ptr<vtkm::cont::internal::BufferInfo> CopyDeviceToDevice(
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfo> src) override;
|
||||
};
|
||||
}
|
||||
}
|
||||
} // namespace vtkm::cont::internal
|
||||
|
||||
#endif //vtk_m_cont_internal_DeviceAdapterMemoryManagerShared_h
|
@ -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)
|
||||
|
139
vtkm/cont/internal/testing/UnitTestBuffer.cxx
Normal file
139
vtkm/cont/internal/testing/UnitTestBuffer.cxx
Normal file
@ -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 <vtkm/cont/internal/Buffer.h>
|
||||
|
||||
#include <vtkm/cont/serial/DeviceAdapterSerial.h>
|
||||
|
||||
#include <vtkm/cont/testing/Testing.h>
|
||||
|
||||
namespace
|
||||
{
|
||||
|
||||
using T = vtkm::FloatDefault;
|
||||
constexpr vtkm::Id ARRAY_SIZE = 20;
|
||||
|
||||
using PortalType = vtkm::cont::internal::ArrayPortalFromIterators<T*>;
|
||||
using PortalTypeConst = vtkm::cont::internal::ArrayPortalFromIterators<const T*>;
|
||||
|
||||
PortalType MakePortal(void* buffer, vtkm::Id numValues)
|
||||
{
|
||||
return PortalType(static_cast<T*>(buffer),
|
||||
static_cast<T*>(buffer) + static_cast<std::size_t>(numValues));
|
||||
};
|
||||
|
||||
PortalTypeConst MakePortal(const void* buffer, vtkm::Id numValues)
|
||||
{
|
||||
return PortalTypeConst(static_cast<const T*>(buffer),
|
||||
static_cast<const T*>(buffer) + static_cast<std::size_t>(numValues));
|
||||
};
|
||||
|
||||
struct VectorDeleter
|
||||
{
|
||||
std::shared_ptr<std::vector<T>> Data;
|
||||
|
||||
VectorDeleter(vtkm::Id numValues)
|
||||
: Data(new std::vector<T>(static_cast<size_t>(numValues)))
|
||||
{
|
||||
}
|
||||
|
||||
template <typename U>
|
||||
void operator()(U* p)
|
||||
{
|
||||
if (this->Data)
|
||||
{
|
||||
VTKM_TEST_ASSERT(reinterpret_cast<T*>(p) == &this->Data->front());
|
||||
this->Data.reset();
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
|
||||
void DoTest()
|
||||
{
|
||||
constexpr vtkm::Id BUFFER_SIZE = ARRAY_SIZE * static_cast<vtkm::Id>(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);
|
||||
}
|
@ -18,6 +18,7 @@
|
||||
#include <vtkm/cont/openmp/internal/ArrayManagerExecutionOpenMP.h>
|
||||
#include <vtkm/cont/openmp/internal/AtomicInterfaceExecutionOpenMP.h>
|
||||
#include <vtkm/cont/openmp/internal/DeviceAdapterAlgorithmOpenMP.h>
|
||||
#include <vtkm/cont/openmp/internal/DeviceAdapterMemoryManagerOpenMP.h>
|
||||
#include <vtkm/cont/openmp/internal/VirtualObjectTransferOpenMP.h>
|
||||
#endif
|
||||
|
||||
|
@ -11,6 +11,7 @@
|
||||
set(headers
|
||||
ArrayManagerExecutionOpenMP.h
|
||||
DeviceAdapterAlgorithmOpenMP.h
|
||||
DeviceAdapterMemoryManagerOpenMP.h
|
||||
DeviceAdapterRuntimeDetectorOpenMP.h
|
||||
DeviceAdapterTagOpenMP.h
|
||||
AtomicInterfaceExecutionOpenMP.h
|
||||
|
33
vtkm/cont/openmp/internal/DeviceAdapterMemoryManagerOpenMP.h
Normal file
33
vtkm/cont/openmp/internal/DeviceAdapterMemoryManagerOpenMP.h
Normal file
@ -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 <vtkm/cont/openmp/internal/DeviceAdapterTagOpenMP.h>
|
||||
|
||||
#include <vtkm/cont/internal/DeviceAdapterMemoryManagerShared.h>
|
||||
|
||||
namespace vtkm
|
||||
{
|
||||
namespace cont
|
||||
{
|
||||
namespace internal
|
||||
{
|
||||
|
||||
template <>
|
||||
class DeviceAdapterMemoryManager<vtkm::cont::DeviceAdapterTagOpenMP>
|
||||
: public vtkm::cont::internal::DeviceAdapterMemoryManagerShared
|
||||
{
|
||||
};
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
#endif //vtk_m_cont_openmp_internal_DeviceAdapterMemoryManagerOpenMP_h
|
@ -14,6 +14,7 @@
|
||||
// clang-format off
|
||||
#include <vtkm/cont/serial/internal/DeviceAdapterTagSerial.h>
|
||||
#include <vtkm/cont/serial/internal/DeviceAdapterRuntimeDetectorSerial.h>
|
||||
#include <vtkm/cont/serial/internal/DeviceAdapterMemoryManagerSerial.h>
|
||||
#include <vtkm/cont/serial/internal/AtomicInterfaceExecutionSerial.h>
|
||||
#include <vtkm/cont/serial/internal/ArrayManagerExecutionSerial.h>
|
||||
#include <vtkm/cont/serial/internal/DeviceAdapterAlgorithmSerial.h>
|
||||
|
@ -12,6 +12,7 @@ set(headers
|
||||
ArrayManagerExecutionSerial.h
|
||||
AtomicInterfaceExecutionSerial.h
|
||||
DeviceAdapterAlgorithmSerial.h
|
||||
DeviceAdapterMemoryManagerSerial.h
|
||||
DeviceAdapterRuntimeDetectorSerial.h
|
||||
DeviceAdapterTagSerial.h
|
||||
ExecutionArrayInterfaceBasicSerial.h
|
||||
|
33
vtkm/cont/serial/internal/DeviceAdapterMemoryManagerSerial.h
Normal file
33
vtkm/cont/serial/internal/DeviceAdapterMemoryManagerSerial.h
Normal file
@ -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 <vtkm/cont/serial/internal/DeviceAdapterTagSerial.h>
|
||||
|
||||
#include <vtkm/cont/internal/DeviceAdapterMemoryManagerShared.h>
|
||||
|
||||
namespace vtkm
|
||||
{
|
||||
namespace cont
|
||||
{
|
||||
namespace internal
|
||||
{
|
||||
|
||||
template <>
|
||||
class DeviceAdapterMemoryManager<vtkm::cont::DeviceAdapterTagSerial>
|
||||
: public vtkm::cont::internal::DeviceAdapterMemoryManagerShared
|
||||
{
|
||||
};
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
#endif //vtk_m_cont_serial_internal_DeviceAdapterMemoryManagerSerial_h
|
@ -17,6 +17,7 @@
|
||||
#include <vtkm/cont/tbb/internal/ArrayManagerExecutionTBB.h>
|
||||
#include <vtkm/cont/tbb/internal/AtomicInterfaceExecutionTBB.h>
|
||||
#include <vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h>
|
||||
#include <vtkm/cont/tbb/internal/DeviceAdapterMemoryManagerTBB.h>
|
||||
#include <vtkm/cont/tbb/internal/VirtualObjectTransferTBB.h>
|
||||
#endif
|
||||
|
||||
|
@ -12,6 +12,7 @@ set(headers
|
||||
ArrayManagerExecutionTBB.h
|
||||
AtomicInterfaceExecutionTBB.h
|
||||
DeviceAdapterAlgorithmTBB.h
|
||||
DeviceAdapterMemoryManagerTBB.h
|
||||
DeviceAdapterRuntimeDetectorTBB.h
|
||||
DeviceAdapterTagTBB.h
|
||||
ExecutionArrayInterfaceBasicTBB.h
|
||||
|
33
vtkm/cont/tbb/internal/DeviceAdapterMemoryManagerTBB.h
Normal file
33
vtkm/cont/tbb/internal/DeviceAdapterMemoryManagerTBB.h
Normal file
@ -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 <vtkm/cont/tbb/internal/DeviceAdapterTagTBB.h>
|
||||
|
||||
#include <vtkm/cont/internal/DeviceAdapterMemoryManagerShared.h>
|
||||
|
||||
namespace vtkm
|
||||
{
|
||||
namespace cont
|
||||
{
|
||||
namespace internal
|
||||
{
|
||||
|
||||
template <>
|
||||
class DeviceAdapterMemoryManager<vtkm::cont::DeviceAdapterTagTBB>
|
||||
: public vtkm::cont::internal::DeviceAdapterMemoryManagerShared
|
||||
{
|
||||
};
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
#endif //vtk_m_cont_tbb_internal_DeviceAdapterMemoryManagerTBB_h
|
@ -157,17 +157,18 @@ public:
|
||||
}
|
||||
};
|
||||
|
||||
template <typename PortalType>
|
||||
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 <typename PortalType>
|
||||
static VTKM_CONT AddArrayKernel<PortalType> MakeAddArrayKernel(
|
||||
PortalType portal,
|
||||
const vtkm::Id3& dims = vtkm::Id3{})
|
||||
{
|
||||
return AddArrayKernel<PortalType>(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<T*>;
|
||||
auto makePortal = [](const vtkm::cont::internal::BufferInfo& buffer) {
|
||||
return PortalType(static_cast<T*>(buffer.GetPointer()),
|
||||
static_cast<T*>(buffer.GetPointer()) +
|
||||
static_cast<std::size_t>(buffer.GetSize()) / sizeof(T));
|
||||
};
|
||||
|
||||
constexpr vtkm::BufferSizeType BUFFER_SIZE =
|
||||
ARRAY_SIZE * static_cast<vtkm::BufferSizeType>(sizeof(T));
|
||||
|
||||
// Set up buffer on host.
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfoHost> hostBufferSrc(
|
||||
new vtkm::cont::internal::BufferInfoHost(BUFFER_SIZE));
|
||||
VTKM_TEST_ASSERT(hostBufferSrc->GetSize() == BUFFER_SIZE);
|
||||
SetPortal(makePortal(*hostBufferSrc));
|
||||
|
||||
vtkm::cont::internal::DeviceAdapterMemoryManager<DeviceAdapterTag> memoryManager;
|
||||
|
||||
std::cout << "Allocate a buffer." << std::endl;
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfo> 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<vtkm::cont::internal::BufferInfo> workingMemory =
|
||||
memoryManager.CopyDeviceToDevice(allocatedMemory);
|
||||
VTKM_TEST_ASSERT(workingMemory->GetSize() == BUFFER_SIZE);
|
||||
|
||||
std::cout << "Copy data back to host." << std::endl;
|
||||
std::shared_ptr<vtkm::cont::internal::BufferInfoHost> 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();
|
||||
|
@ -71,6 +71,12 @@ public:
|
||||
namespace internal
|
||||
{
|
||||
|
||||
template <>
|
||||
class DeviceAdapterMemoryManager<vtkm::cont::DeviceAdapterTagTestAlgorithmGeneral>
|
||||
: public vtkm::cont::internal::DeviceAdapterMemoryManagerShared
|
||||
{
|
||||
};
|
||||
|
||||
template <typename T, class StorageTag>
|
||||
class ArrayManagerExecution<T, StorageTag, vtkm::cont::DeviceAdapterTagTestAlgorithmGeneral>
|
||||
: public vtkm::cont::internal::ArrayManagerExecution<T,
|
||||
|
Loading…
Reference in New Issue
Block a user