host_tensor.h
Go to the documentation of this file.
/***************************************************************************************************
* Copyright (c) 2017-2019, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without modification, are permitted
* provided that the following conditions are met:
* * Redistributions of source code must retain the above copyright notice, this list of
* conditions and the following disclaimer.
* * Redistributions in binary form must reproduce the above copyright notice, this list of
* conditions and the following disclaimer in the documentation and/or other materials
* provided with the distribution.
* * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
* to endorse or promote products derived from this software without specific prior written
* permission.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
* IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
* FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
* FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
* BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
* OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
* STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
* OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
**************************************************************************************************/
#pragma once
#include <vector>
#include "cutlass/cutlass.h"
#include "cutlass/matrix_traits.h"
#include "cutlass/tensor_ref.h"
#include "cutlass/tensor_view.h"
#include "device_memory.h"
namespace cutlass {
template <
typename Element_,
typename Layout_
>
class HostTensor {
public:
using Element = Element_;
using Layout = Layout_;
static int const kRank = Layout::kRank;
using Index = typename Layout::Index;
using LongIndex = typename Layout::LongIndex;
using TensorCoord = typename Layout::TensorCoord;
using Stride = typename Layout::Stride;
using TensorRef = TensorRef<Element, Layout>;
using ConstTensorRef = typename TensorRef::ConstTensorRef;
using TensorView = TensorView<Element, Layout>;
using ConstTensorView = typename TensorView::ConstTensorView;
using Reference = typename TensorRef::Reference;
using ConstReference = typename ConstTensorRef::Reference;
static int const kElementsPerStoredItem = (sizeof_bits<Element>::value < 8 ? sizeof(Element) * 8 / sizeof_bits<Element>::value : 1);
private:
//
// Data members
//
TensorCoord extent_;
Layout layout_;
std::vector<Element> host_;
device_memory::allocation<Element> device_;
public:
//
// Device and Host Methods
//
HostTensor() {}
HostTensor(
TensorCoord const &extent,
bool device_backed = true
) {
this->reset(extent, Layout::packed(extent), device_backed);
}
HostTensor(
TensorCoord const &extent,
Layout const &layout,
bool device_backed = true
) {
this->reset(extent, layout, device_backed);
}
~HostTensor() { }
void reset() {
extent_ = TensorCoord();
layout_ = Layout::packed(extent_);
host_.clear();
device_.reset();
}
void reserve(
size_t count,
bool device_backed_ = true) {
device_.reset();
host_.clear();
count /= kElementsPerStoredItem;
host_.resize(count);
// Allocate memory
Element* device_memory = nullptr;
if (device_backed_) {
device_memory = device_memory::allocate<Element>(count);
}
device_.reset(device_memory, device_backed_ ? count : 0);
}
void reset(
TensorCoord const &extent,
Layout const &layout,
bool device_backed_ = true) {
extent_ = extent;
layout_ = layout;
reserve(size_t(layout_.capacity(extent_)), device_backed_);
}
void reset(
TensorCoord const &extent,
bool device_backed_ = true) {
reset(extent, Layout::packed(extent), device_backed_);
}
void resize(
TensorCoord const &extent,
Layout const &layout,
bool device_backed_ = true) {
extent_ = extent;
layout_ = layout;
LongIndex new_size = size_t(layout_.capacity(extent_));
if (static_cast<decltype(host_.size())>(new_size) > host_.size()) {
reserve(new_size);
}
}
void resize(
TensorCoord const &extent,
bool device_backed_ = true) {
resize(extent, Layout::packed(extent), device_backed_);
}
size_t size() const {
return host_.size() * kElementsPerStoredItem;
}
LongIndex capacity() const {
return layout_.capacity(extent_) * kElementsPerStoredItem;
}
Element * host_data() { return host_.data(); }
Element * host_data_ptr_offset(LongIndex ptr_element_offset) { return host_.data() + ptr_element_offset; }
Reference host_data(LongIndex idx) {
return ReferenceFactory<Element>::get(host_data(), idx);
}
Element const * host_data() const { return host_.data(); }
ConstReference host_data(LongIndex idx) const {
return ReferenceFactory<Element const>::get(host_data(), idx);
}
Element * device_data() { return device_.get(); }
Element * device_data_ptr_offset(LongIndex ptr_element_offset) { return device_.get() + ptr_element_offset; }
Element const * device_data() const { return device_.get(); }
TensorRef host_ref(LongIndex ptr_element_offset=0) { return TensorRef(host_data_ptr_offset(ptr_element_offset), layout_); }
ConstTensorRef host_ref(LongIndex ptr_element_offset=0) const { return ConstTensorRef(host_data_ptr_offset(ptr_element_offset), layout_); }
TensorRef device_ref(LongIndex ptr_element_offset=0) {
return TensorRef(device_data_ptr_offset(ptr_element_offset), layout_);
}
ConstTensorRef device_ref(LongIndex ptr_element_offset=0) const {
return TensorRef(device_data_ptr_offset(ptr_element_offset), layout_);
}
TensorView host_view(LongIndex ptr_element_offset=0) {
return TensorView(host_data_ptr_offset(ptr_element_offset), layout_, extent_);
}
ConstTensorView host_view(LongIndex ptr_element_offset=0) const {
return ConstTensorView(host_data_ptr_offset(ptr_element_offset), layout_, extent_);
}
TensorView device_view(LongIndex ptr_element_offset=0) {
return TensorView(device_data_ptr_offset(ptr_element_offset), layout_, extent_);
}
ConstTensorView device_view(LongIndex ptr_element_offset=0) const {
return ConstTensorView(device_data_ptr_offset(ptr_element_offset), layout_, extent_);
}
bool device_backed() const {
return (device_.get() == nullptr) ? false : true;
}
Layout & layout() {
return layout_;
}
Layout layout() const {
return layout_;
}
Stride stride() const {
return layout_.stride();
}
Stride & stride() {
return layout_.stride();
}
Index stride(int dim) const {
return layout_.stride().at(dim);
}
Index & stride(int dim) {
return layout_.stride().at(dim);
}
LongIndex offset(TensorCoord const& coord) const {
return layout_(coord);
}
Reference at(TensorCoord const& coord) {
return host_data(offset(coord));
}
ConstReference at(TensorCoord const& coord) const {
return host_data(offset(coord));
}
TensorCoord extent() const {
return extent_;
}
TensorCoord & extent() {
return extent_;
}
void sync_host() {
if (device_backed()) {
device_memory::copy_to_host(
host_data(), device_data(), size());
}
}
void sync_device() {
if (device_backed()) {
device_memory::copy_to_device(
device_data(), host_data(), size());
}
}
void copy_in_device_to_host(
Element const* ptr_device,
LongIndex count = -1) {
if (count < 0) {
count = capacity();
}
else {
count = __NV_STD_MIN(capacity(), count);
}
device_memory::copy_to_host(
host_data(), ptr_device, count);
}
void copy_in_device_to_device(
Element const* ptr_device,
LongIndex count = -1) {
if (count < 0) {
count = capacity();
}
else {
count = __NV_STD_MIN(capacity(), count);
}
device_memory::copy_device_to_device(
device_data(), ptr_device, count);
}
void copy_in_host_to_device(
Element const* ptr_host,
LongIndex count = -1) {
if (count < 0) {
count = capacity();
}
else {
count = __NV_STD_MIN(capacity(), count);
}
device_memory::copy_to_device(
device_data(), ptr_host, count);
}
void copy_in_host_to_host(
Element const* ptr_host,
LongIndex count = -1) {
if (count < 0) {
count = capacity();
}
else {
count = __NV_STD_MIN(capacity(), count);
}
device_memory::copy_host_to_host(
host_data(), ptr_host, count);
}
void copy_out_device_to_host(
Element * ptr_host,
LongIndex count = -1) const {
if (count < 0) {
count = capacity();
}
else {
count = __NV_STD_MIN(capacity(), count);
}
device_memory::copy_to_host(
ptr_host, device_data(), count);
}
void copy_out_device_to_device(
Element * ptr_device,
LongIndex count = -1) const {
if (count < 0) {
count = capacity();
}
else {
count = __NV_STD_MIN(capacity(), count);
}
device_memory::copy_device_to_device(
ptr_device, device_data(), count);
}
void copy_out_host_to_device(
Element * ptr_device,
LongIndex count = -1) const {
if (count < 0) {
count = capacity();
}
else {
count = __NV_STD_MIN(capacity(), count);
}
device_memory::copy_to_device(
ptr_device, host_data(), count);
}
void copy_out_host_to_host(
Element * ptr_host,
LongIndex count = -1) const {
if (count < 0) {
count = capacity();
}
else {
count = __NV_STD_MIN(capacity(), count);
}
device_memory::copy_host_to_host(
ptr_host, host_data(), count);
}
};
} // namespace cutlass
cutlass::HostTensor::stride
Stride & stride()
Returns the layout object&#39;s stride vector.
Definition: host_tensor.h:324
cutlass
Definition: aligned_buffer.h:35
cutlass::HostTensor::HostTensor
HostTensor(TensorCoord const &extent, Layout const &layout, bool device_backed=true)
Constructs a tensor given an extent and layout.
Definition: host_tensor.h:140
tensor_ref.h
Defines a structure containing strides, bounds, and a pointer to tensor data.
cutlass::HostTensor::Stride
typename Layout::Stride Stride
Layout&#39;s stride vector.
Definition: host_tensor.h:81
cutlass::HostTensor::copy_in_host_to_device
void copy_in_host_to_device(Element const *ptr_host, LongIndex count=-1)
Copy data from a caller-supplied device pointer into host memory.
Definition: host_tensor.h:410
cutlass::HostTensor::stride
Index & stride(int dim)
Returns the layout object&#39;s stride in a given physical dimension.
Definition: host_tensor.h:334
cutlass::device_memory::copy_to_device
void copy_to_device(T *dst, T const *src, size_t count=1)
Definition: device_memory.h:81
cutlass::HostTensor::device_view
ConstTensorView device_view(LongIndex ptr_element_offset=0) const
Accesses the tensor reference pointing to data.
Definition: host_tensor.h:298
cutlass::HostTensor::LongIndex
typename Layout::LongIndex LongIndex
Long index used for pointer offsets.
Definition: host_tensor.h:75
cutlass::HostTensor::Layout
Layout_ Layout
Mapping function from logical coordinate to linear memory.
Definition: host_tensor.h:66
cutlass::HostTensor::host_view
ConstTensorView host_view(LongIndex ptr_element_offset=0) const
Accesses the tensor reference pointing to data.
Definition: host_tensor.h:288
cutlass::HostTensor::capacity
LongIndex capacity() const
Returns the logical capacity based on extent and layout. May differ from size().
Definition: host_tensor.h:234
cutlass::HostTensor::copy_out_host_to_host
void copy_out_host_to_host(Element *ptr_host, LongIndex count=-1) const
Copy data from a caller-supplied device pointer into host memory.
Definition: host_tensor.h:485
cutlass::HostTensor::sync_host
void sync_host()
Copies data from device to host.
Definition: host_tensor.h:364
cutlass::HostTensor::layout
Layout layout() const
Returns the layout object.
Definition: host_tensor.h:314
cutlass::HostTensor::copy_in_host_to_host
void copy_in_host_to_host(Element const *ptr_host, LongIndex count=-1)
Copy data from a caller-supplied device pointer into host memory.
Definition: host_tensor.h:425
cutlass::HostTensor::device_data
Element * device_data()
Gets pointer to device data.
Definition: host_tensor.h:258
cutlass::device_memory::allocation::get
T * get() const
Returns a pointer to the managed object.
Definition: device_memory.h:162
cutlass::device_memory::allocation::reset
void reset()
Deletes the managed object and resets capacity to zero.
Definition: device_memory.h:171
cutlass::HostTensor::device_data
Element const * device_data() const
Gets pointer to device data.
Definition: host_tensor.h:264
cutlass::HostTensor::kElementsPerStoredItem
static int const kElementsPerStoredItem
Used to handle packing of subbyte elements.
Definition: host_tensor.h:102
tensor_view.h
Defines a structure containing strides and a pointer to tensor data.
device_memory.h
C++ interface to CUDA device memory management functions.
cutlass::HostTensor::~HostTensor
~HostTensor()
Definition: host_tensor.h:149
cutlass::TensorRef< Element, Layout >::ConstTensorRef
TensorRef< typename platform::remove_const< Element >::type const, Layout > ConstTensorRef
TensorRef to constant data.
Definition: tensor_ref.h:179
cutlass::HostTensor::copy_in_device_to_device
void copy_in_device_to_device(Element const *ptr_device, LongIndex count=-1)
Copy data from a caller-supplied device pointer into host memory.
Definition: host_tensor.h:395
cutlass::HostTensor::ConstReference
typename ConstTensorRef::Reference ConstReference
Constant reference to element in tensor.
Definition: host_tensor.h:99
cutlass::HostTensor::stride
Index stride(int dim) const
Returns the layout object&#39;s stride in a given physical dimension.
Definition: host_tensor.h:329
cutlass::HostTensor::sync_device
void sync_device()
Copies data from host to device.
Definition: host_tensor.h:372
cutlass::TensorView< Element, Layout >::ConstTensorView
TensorView< typename platform::remove_const< Element >::type const, Layout > ConstTensorView
TensorView pointing to constant memory.
Definition: tensor_view.h:95
cutlass::HostTensor::TensorView
TensorView< Element, Layout > TensorView
Tensor reference to device memory.
Definition: host_tensor.h:90
cutlass::HostTensor::host_data_ptr_offset
Element * host_data_ptr_offset(LongIndex ptr_element_offset)
Gets pointer to host data with a pointer offset.
Definition: host_tensor.h:242
cutlass::HostTensor
Host tensor.
Definition: host_tensor.h:59
cutlass::HostTensor::device_view
TensorView device_view(LongIndex ptr_element_offset=0)
Accesses the tensor reference pointing to data.
Definition: host_tensor.h:293
cutlass::TensorView< Element, Layout >
cutlass::HostTensor::copy_out_device_to_host
void copy_out_device_to_host(Element *ptr_host, LongIndex count=-1) const
Copy data from a caller-supplied device pointer into host memory.
Definition: host_tensor.h:440
cutlass::HostTensor::copy_out_host_to_device
void copy_out_host_to_device(Element *ptr_device, LongIndex count=-1) const
Copy data from a caller-supplied device pointer into host memory.
Definition: host_tensor.h:470
cutlass::HostTensor::at
ConstReference at(TensorCoord const &coord) const
Returns a const reference to the element at the logical Coord in host memory.
Definition: host_tensor.h:349
cutlass::sizeof_bits
Defines the size of an element in bits.
Definition: numeric_types.h:42
cutlass::device_memory::copy_host_to_host
void copy_host_to_host(T *dst, T const *src, size_t count=1)
Definition: device_memory.h:96
cutlass::HostTensor::host_ref
ConstTensorRef host_ref(LongIndex ptr_element_offset=0) const
Accesses the tensor reference pointing to data.
Definition: host_tensor.h:270
cutlass::HostTensor::extent
TensorCoord extent() const
Returns the extent of the tensor.
Definition: host_tensor.h:354
cutlass::HostTensor::host_data
Element * host_data()
Gets pointer to host data.
Definition: host_tensor.h:239
cutlass::HostTensor::reserve
void reserve(size_t count, bool device_backed_=true)
Resizes internal memory allocations without affecting layout or extent.
Definition: host_tensor.h:161
cutlass::ReferenceFactory
Definition: subbyte_reference.h:557
cutlass::TensorRef< Element, Layout >
__NV_STD_MIN
#define __NV_STD_MIN(a, b)
Select minimum(a, b)
Definition: platform.h:168
cutlass::HostTensor::device_backed
bool device_backed() const
Returns true if device memory is allocated.
Definition: host_tensor.h:303
cutlass::HostTensor::TensorCoord
typename Layout::TensorCoord TensorCoord
Coordinate in logical tensor space.
Definition: host_tensor.h:78
cutlass::TensorRef< Element, Layout >::Reference
typename platform::conditional< sizeof_bits< Element >::value >=8, Element &, SubbyteReference< Element > >::type Reference
Reference type to an element.
Definition: tensor_ref.h:159
cutlass::HostTensor::host_data
Reference host_data(LongIndex idx)
Gets a reference to an element in host memory.
Definition: host_tensor.h:245
cutlass::HostTensor::extent
TensorCoord & extent()
Returns the extent of the tensor.
Definition: host_tensor.h:359
cutlass::HostTensor::HostTensor
HostTensor()
Default constructor.
Definition: host_tensor.h:128
cutlass::HostTensor::ConstTensorView
typename TensorView::ConstTensorView ConstTensorView
Tensor reference to constant device memory.
Definition: host_tensor.h:93
cutlass::HostTensor::resize
void resize(TensorCoord const &extent, bool device_backed_=true)
Definition: host_tensor.h:221
cutlass::device_memory::copy_to_host
void copy_to_host(T *dst, T const *src, size_t count=1)
Definition: device_memory.h:86
cutlass::HostTensor::copy_in_device_to_host
void copy_in_device_to_host(Element const *ptr_device, LongIndex count=-1)
Copy data from a caller-supplied device pointer into host memory.
Definition: host_tensor.h:380
cutlass::HostTensor::offset
LongIndex offset(TensorCoord const &coord) const
Computes the offset of an index from the origin of the tensor.
Definition: host_tensor.h:339
cutlass::HostTensor::host_view
TensorView host_view(LongIndex ptr_element_offset=0)
Accesses the tensor reference pointing to data.
Definition: host_tensor.h:283
cutlass::HostTensor::HostTensor
HostTensor(TensorCoord const &extent, bool device_backed=true)
Constructs a tensor given an extent. Assumes a packed layout.
Definition: host_tensor.h:131
cutlass::HostTensor::device_ref
TensorRef device_ref(LongIndex ptr_element_offset=0)
Accesses the tensor reference pointing to data.
Definition: host_tensor.h:273
cutlass::HostTensor::Element
Element_ Element
Data type of individual access.
Definition: host_tensor.h:63
cutlass::HostTensor::at
Reference at(TensorCoord const &coord)
Returns a reference to the element at the logical Coord in host memory.
Definition: host_tensor.h:344
cutlass::HostTensor::reset
void reset(TensorCoord const &extent, bool device_backed_=true)
Definition: host_tensor.h:195
cutlass::HostTensor::ConstTensorRef
typename TensorRef::ConstTensorRef ConstTensorRef
Tensor reference to constant device memory.
Definition: host_tensor.h:87
cutlass::HostTensor::host_data
ConstReference host_data(LongIndex idx) const
Gets a constant reference to an element in host memory.
Definition: host_tensor.h:253
cutlass::HostTensor::reset
void reset()
Clears the HostTensor allocation to size/capacity = 0.
Definition: host_tensor.h:152
cutlass::HostTensor::kRank
static int const kRank
Logical rank of tensor index space.
Definition: host_tensor.h:69
cutlass::HostTensor::device_data_ptr_offset
Element * device_data_ptr_offset(LongIndex ptr_element_offset)
Gets pointer to device data with a pointer offset.
Definition: host_tensor.h:261
cutlass::HostTensor::size
size_t size() const
Returns the number of elements stored in the host tensor.
Definition: host_tensor.h:229
cutlass::HostTensor::device_ref
ConstTensorRef device_ref(LongIndex ptr_element_offset=0) const
Accesses the tensor reference pointing to data.
Definition: host_tensor.h:278
cutlass::HostTensor::layout
Layout & layout()
Returns the layout object.
Definition: host_tensor.h:309
cutlass::HostTensor::Reference
typename TensorRef::Reference Reference
Reference to element in tensor.
Definition: host_tensor.h:96
cutlass::HostTensor::resize
void resize(TensorCoord const &extent, Layout const &layout, bool device_backed_=true)
Definition: host_tensor.h:204
matrix_traits.h
Defines properties of matrices used to denote layout and operands to GEMM kernels.
cutlass::HostTensor::TensorRef
TensorRef< Element, Layout > TensorRef
Tensor reference to device memory.
Definition: host_tensor.h:84
cutlass::device_memory::allocation< Element >
cutlass::HostTensor::reset
void reset(TensorCoord const &extent, Layout const &layout, bool device_backed_=true)
Definition: host_tensor.h:182
cutlass::HostTensor::stride
Stride stride() const
Returns the layout object&#39;s stride vector.
Definition: host_tensor.h:319
cutlass::HostTensor::Index
typename Layout::Index Index
Index type.
Definition: host_tensor.h:72
cutlass.h
Basic include for CUTLASS.
cutlass::HostTensor::host_ref
TensorRef host_ref(LongIndex ptr_element_offset=0)
Accesses the tensor reference pointing to data.
Definition: host_tensor.h:267
cutlass::HostTensor::copy_out_device_to_device
void copy_out_device_to_device(Element *ptr_device, LongIndex count=-1) const
Copy data from a caller-supplied device pointer into host memory.
Definition: host_tensor.h:455
cutlass::device_memory::copy_device_to_device
void copy_device_to_device(T *dst, T const *src, size_t count=1)
Definition: device_memory.h:91
cutlass::HostTensor::host_data
Element const * host_data() const
Gets pointer to host data.
Definition: host_tensor.h:250