Merge branch 'datamodel-design' of public.kitware.com:vtkm into datamodel-design
This commit is contained in:
commit
e8ee229bc8
@ -89,7 +89,7 @@ function (get_year var)
|
||||
set(${var} "${result}" PARENT_SCOPE)
|
||||
endfunction (get_year)
|
||||
|
||||
set(copyright_file_year 2011)
|
||||
set(copyright_file_year 2014)
|
||||
get_year(current_year)
|
||||
|
||||
# Escapes ';' characters (list delimiters) and splits the given string into
|
||||
|
@ -24,6 +24,7 @@
|
||||
|
||||
#include <vtkm/cont/Assert.h>
|
||||
#include <vtkm/cont/ErrorControlBadValue.h>
|
||||
#include <vtkm/cont/ErrorControlInternal.h>
|
||||
#include <vtkm/cont/Storage.h>
|
||||
|
||||
#include <vtkm/cont/internal/ArrayHandleExecutionManager.h>
|
||||
@ -173,12 +174,13 @@ public:
|
||||
// If the user writes into the iterator we return, then the execution
|
||||
// array will become invalid. Play it safe and release the execution
|
||||
// resources. (Use the const version to preserve the execution array.)
|
||||
this->ReleaseResourcesExecution();
|
||||
this->ReleaseResourcesExecutionInternal();
|
||||
return this->Internals->ControlArray.GetPortal();
|
||||
}
|
||||
else
|
||||
{
|
||||
throw vtkm::cont::ErrorControlBadValue("ArrayHandle contains no data.");
|
||||
throw vtkm::cont::ErrorControlInternal(
|
||||
"ArrayHandle::SyncControlArray did not make control array valid.");
|
||||
}
|
||||
}
|
||||
|
||||
@ -197,7 +199,8 @@ public:
|
||||
}
|
||||
else
|
||||
{
|
||||
throw vtkm::cont::ErrorControlBadValue("ArrayHandle contains no data.");
|
||||
throw vtkm::cont::ErrorControlInternal(
|
||||
"ArrayHandle::SyncControlArray did not make control array valid.");
|
||||
}
|
||||
}
|
||||
|
||||
@ -224,6 +227,23 @@ public:
|
||||
}
|
||||
}
|
||||
|
||||
/// \brief Allocates an array large enough to hold the given number of values.
|
||||
///
|
||||
/// The allocation may be done on an already existing array, but can wipe out
|
||||
/// any data already in the array. This method can throw
|
||||
/// ErrorControlOutOfMemory if the array cannot be allocated or
|
||||
/// ErrorControlBadValue if the allocation is not feasible (for example, the
|
||||
/// array storage is read-only).
|
||||
///
|
||||
VTKM_CONT_EXPORT
|
||||
void Allocate(vtkm::Id numberOfValues)
|
||||
{
|
||||
this->ReleaseResourcesExecutionInternal();
|
||||
this->Internals->UserPortalValid = false;
|
||||
this->Internals->ControlArray.Allocate(numberOfValues);
|
||||
this->Internals->ControlArrayValid = true;
|
||||
}
|
||||
|
||||
/// \brief Reduces the size of the array without changing its values.
|
||||
///
|
||||
/// This method allows you to resize the array without reallocating it. The
|
||||
@ -270,18 +290,18 @@ public:
|
||||
///
|
||||
VTKM_CONT_EXPORT void ReleaseResourcesExecution()
|
||||
{
|
||||
if (this->Internals->ExecutionArrayValid)
|
||||
{
|
||||
this->Internals->ExecutionArray->ReleaseResources();
|
||||
this->Internals->ExecutionArrayValid = false;
|
||||
}
|
||||
// Save any data in the execution environment by making sure it is synced
|
||||
// with the control environment.
|
||||
this->SyncControlArray();
|
||||
|
||||
this->ReleaseResourcesExecutionInternal();
|
||||
}
|
||||
|
||||
/// Releases all resources in both the control and execution environments.
|
||||
///
|
||||
VTKM_CONT_EXPORT void ReleaseResources()
|
||||
{
|
||||
this->ReleaseResourcesExecution();
|
||||
this->ReleaseResourcesExecutionInternal();
|
||||
|
||||
// Forget about any user iterators.
|
||||
this->Internals->UserPortalValid = false;
|
||||
@ -507,15 +527,26 @@ public:
|
||||
///
|
||||
VTKM_CONT_EXPORT void SyncControlArray() const
|
||||
{
|
||||
if ( !this->Internals->UserPortalValid
|
||||
&& !this->Internals->ControlArrayValid)
|
||||
if (!this->Internals->UserPortalValid
|
||||
&& !this->Internals->ControlArrayValid)
|
||||
{
|
||||
// Need to change some state that does not change the logical state from
|
||||
// an external point of view.
|
||||
InternalStruct *internals
|
||||
= const_cast<InternalStruct*>(this->Internals.get());
|
||||
internals->ExecutionArray->RetrieveOutputData(internals->ControlArray);
|
||||
internals->ControlArrayValid = true;
|
||||
if (this->Internals->ExecutionArrayValid)
|
||||
{
|
||||
internals->ExecutionArray->RetrieveOutputData(internals->ControlArray);
|
||||
internals->ControlArrayValid = true;
|
||||
}
|
||||
else
|
||||
{
|
||||
// This array is in the null state (there is nothing allocated), but
|
||||
// the calling function wants to do something with the array. Put this
|
||||
// class into a valid state by allocating an array of size 0.
|
||||
internals->ControlArray.Allocate(0);
|
||||
internals->ControlArrayValid = true;
|
||||
}
|
||||
}
|
||||
else
|
||||
{
|
||||
@ -527,6 +558,16 @@ public:
|
||||
}
|
||||
}
|
||||
|
||||
VTKM_CONT_EXPORT
|
||||
void ReleaseResourcesExecutionInternal()
|
||||
{
|
||||
if (this->Internals->ExecutionArrayValid)
|
||||
{
|
||||
this->Internals->ExecutionArray->ReleaseResources();
|
||||
this->Internals->ExecutionArrayValid = false;
|
||||
}
|
||||
}
|
||||
|
||||
boost::shared_ptr<InternalStruct> Internals;
|
||||
};
|
||||
|
||||
|
@ -34,6 +34,7 @@ set(headers
|
||||
DeviceAdapterSerial.h
|
||||
DynamicArrayHandle.h
|
||||
DynamicPointCoordinates.h
|
||||
ExplicitConnectivity.h
|
||||
Error.h
|
||||
ErrorControl.h
|
||||
ErrorControlAssert.h
|
||||
@ -45,6 +46,7 @@ set(headers
|
||||
PointCoordinatesArray.h
|
||||
PointCoordinatesListTag.h
|
||||
PointCoordinatesUniform.h
|
||||
RegularConnectivity.h
|
||||
Storage.h
|
||||
StorageBasic.h
|
||||
StorageImplicit.h
|
||||
|
@ -6,123 +6,12 @@
|
||||
#include <vtkm/cont/Field.h>
|
||||
#include <vtkm/cont/DynamicArrayHandle.h>
|
||||
#include <vtkm/cont/DeviceAdapterAlgorithm.h>
|
||||
#include <vtkm/cont/ExplicitConnectivity.h>
|
||||
#include <vtkm/cont/RegularConnectivity.h>
|
||||
|
||||
namespace vtkm {
|
||||
namespace cont {
|
||||
|
||||
class ExplicitConnectivity
|
||||
{
|
||||
public:
|
||||
ExplicitConnectivity() {}
|
||||
|
||||
vtkm::Id GetNumberOfElements()
|
||||
{
|
||||
return Shapes.GetNumberOfValues();
|
||||
}
|
||||
vtkm::Id GetNumberOfIndices(vtkm::Id index)
|
||||
{
|
||||
return NumIndices.GetPortalControl().Get(index);
|
||||
}
|
||||
vtkm::Id GetElementShapeType(vtkm::Id index)
|
||||
{
|
||||
return Shapes.GetPortalControl().Get(index);
|
||||
}
|
||||
template <vtkm::IdComponent ItemTupleLength>
|
||||
void GetIndices(vtkm::Id index, vtkm::Vec<vtkm::Id,ItemTupleLength> &ids)
|
||||
{
|
||||
int n = GetNumberOfIndices(index);
|
||||
int start = MapCellToConnectivityIndex.GetPortalControl().Get(index);
|
||||
for (int i=0; i<n && i<ItemTupleLength; i++)
|
||||
ids[i] = Connectivity.GetPortalControl().Get(start+i);
|
||||
}
|
||||
template <vtkm::IdComponent ItemTupleLength>
|
||||
void AddShape(vtkm::CellType cellType, int numVertices, vtkm::Vec<vtkm::Id,ItemTupleLength> &ids)
|
||||
{
|
||||
///\todo: how do I modify an array handle?
|
||||
}
|
||||
|
||||
vtkm::cont::ArrayHandle<vtkm::Id, vtkm::cont::StorageTagBasic> Shapes;
|
||||
vtkm::cont::ArrayHandle<vtkm::Id, vtkm::cont::StorageTagBasic> NumIndices;
|
||||
vtkm::cont::ArrayHandle<vtkm::Id, vtkm::cont::StorageTagBasic> Connectivity;
|
||||
vtkm::cont::ArrayHandle<vtkm::Id, vtkm::cont::StorageTagBasic> MapCellToConnectivityIndex;
|
||||
};
|
||||
|
||||
class RegularConnectivity3D
|
||||
{
|
||||
public:
|
||||
void SetNodeDimension3D(int node_i, int node_j, int node_k)
|
||||
{
|
||||
cellDims[0] = node_i-1;
|
||||
cellDims[1] = node_j-1;
|
||||
cellDims[2] = node_k-1;
|
||||
nodeDims[0] = node_i;
|
||||
nodeDims[1] = node_j;
|
||||
nodeDims[2] = node_k;
|
||||
}
|
||||
|
||||
vtkm::Id GetNumberOfElements()
|
||||
{
|
||||
return cellDims[0]*cellDims[1]*cellDims[2];
|
||||
}
|
||||
vtkm::Id GetNumberOfIndices(vtkm::Id)
|
||||
{
|
||||
return 8;
|
||||
}
|
||||
vtkm::Id GetElementShapeType(vtkm::Id)
|
||||
{
|
||||
return VTKM_VOXEL;
|
||||
}
|
||||
template <vtkm::IdComponent ItemTupleLength>
|
||||
void GetIndices(vtkm::Id index, vtkm::Vec<vtkm::Id,ItemTupleLength> &ids)
|
||||
{
|
||||
int i,j,k;
|
||||
CalculateLogicalCellIndices3D(index, i,j,k);
|
||||
///\todo: assert ItemTupleLength >= 8, or return early?
|
||||
ids[0] = CalculateNodeIndex3D(i, j, k);
|
||||
if (ItemTupleLength <= 1) return;
|
||||
ids[1] = ids[0] + 1;
|
||||
if (ItemTupleLength <= 2) return;
|
||||
ids[2] = ids[0] + nodeDims[0];
|
||||
if (ItemTupleLength <= 3) return;
|
||||
ids[3] = ids[2] + 1;
|
||||
if (ItemTupleLength <= 4) return;
|
||||
ids[4] = ids[0] + nodeDims[0]*nodeDims[1];
|
||||
if (ItemTupleLength <= 5) return;
|
||||
ids[5] = ids[4] + 1;
|
||||
if (ItemTupleLength <= 6) return;
|
||||
ids[6] = ids[4] + nodeDims[0];
|
||||
if (ItemTupleLength <= 7) return;
|
||||
ids[7] = ids[6] + 1;
|
||||
}
|
||||
private:
|
||||
int cellDims[3];
|
||||
int nodeDims[3];
|
||||
int CalculateCellIndex3D(int i, int j, int k)
|
||||
{
|
||||
return (k * cellDims[1] + j) * cellDims[0] + i;
|
||||
}
|
||||
int CalculateNodeIndex3D(int i, int j, int k)
|
||||
{
|
||||
return (k * nodeDims[1] + j) * nodeDims[0] + i;
|
||||
}
|
||||
void CalculateLogicalCellIndices3D(int index, int &i, int &j, int &k)
|
||||
{
|
||||
int cellDims01 = cellDims[0] * cellDims[1];
|
||||
k = index / cellDims01;
|
||||
int indexij = index % cellDims01;
|
||||
j = indexij / cellDims[0];
|
||||
i = indexij % cellDims[0];
|
||||
}
|
||||
void CalculateLogicalNodeIndices3D(int index, int &i, int &j, int &k)
|
||||
{
|
||||
int nodeDims01 = nodeDims[0] * nodeDims[1];
|
||||
k = index / nodeDims01;
|
||||
int indexij = index % nodeDims01;
|
||||
j = indexij / nodeDims[0];
|
||||
i = indexij % nodeDims[0];
|
||||
}
|
||||
};
|
||||
|
||||
class DataSet
|
||||
{
|
||||
public:
|
||||
@ -151,6 +40,7 @@ public:
|
||||
|
||||
ExplicitConnectivity conn;
|
||||
RegularConnectivity3D reg;
|
||||
//TODO: Logical structure: vtkm::Extents? Use EAVL logicalStructure?
|
||||
|
||||
//traditional data-model
|
||||
vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::FloatDefault,3> > Points;
|
||||
|
53
vtkm/cont/ExplicitConnectivity.h
Normal file
53
vtkm/cont/ExplicitConnectivity.h
Normal file
@ -0,0 +1,53 @@
|
||||
#ifndef vtk_m_cont_ExplicitConnectivity_h
|
||||
#define vtk_m_cont_ExplicitConnectivity_h
|
||||
|
||||
#include <vtkm/CellType.h>
|
||||
#include <vtkm/cont/ArrayHandle.h>
|
||||
#include <vtkm/cont/Field.h>
|
||||
#include <vtkm/cont/DynamicArrayHandle.h>
|
||||
#include <vtkm/cont/DeviceAdapterAlgorithm.h>
|
||||
|
||||
namespace vtkm {
|
||||
namespace cont {
|
||||
|
||||
class ExplicitConnectivity
|
||||
{
|
||||
public:
|
||||
ExplicitConnectivity() {}
|
||||
|
||||
vtkm::Id GetNumberOfElements()
|
||||
{
|
||||
return Shapes.GetNumberOfValues();
|
||||
}
|
||||
vtkm::Id GetNumberOfIndices(vtkm::Id index)
|
||||
{
|
||||
return NumIndices.GetPortalControl().Get(index);
|
||||
}
|
||||
vtkm::Id GetElementShapeType(vtkm::Id index)
|
||||
{
|
||||
return Shapes.GetPortalControl().Get(index);
|
||||
}
|
||||
template <vtkm::IdComponent ItemTupleLength>
|
||||
void GetIndices(vtkm::Id index, vtkm::Vec<vtkm::Id,ItemTupleLength> &ids)
|
||||
{
|
||||
int n = GetNumberOfIndices(index);
|
||||
int start = MapCellToConnectivityIndex.GetPortalControl().Get(index);
|
||||
for (int i=0; i<n && i<ItemTupleLength; i++)
|
||||
ids[i] = Connectivity.GetPortalControl().Get(start+i);
|
||||
}
|
||||
template <vtkm::IdComponent ItemTupleLength>
|
||||
void AddShape(vtkm::CellType cellType, int numVertices, vtkm::Vec<vtkm::Id,ItemTupleLength> &ids)
|
||||
{
|
||||
///\todo: how do I modify an array handle?
|
||||
}
|
||||
|
||||
vtkm::cont::ArrayHandle<vtkm::Id, vtkm::cont::StorageTagBasic> Shapes;
|
||||
vtkm::cont::ArrayHandle<vtkm::Id, vtkm::cont::StorageTagBasic> NumIndices;
|
||||
vtkm::cont::ArrayHandle<vtkm::Id, vtkm::cont::StorageTagBasic> Connectivity;
|
||||
vtkm::cont::ArrayHandle<vtkm::Id, vtkm::cont::StorageTagBasic> MapCellToConnectivityIndex;
|
||||
};
|
||||
|
||||
}
|
||||
} // namespace vtkm::cont
|
||||
|
||||
#endif //vtk_m_cont_ExplicitConnectivity_h
|
95
vtkm/cont/RegularConnectivity.h
Normal file
95
vtkm/cont/RegularConnectivity.h
Normal file
@ -0,0 +1,95 @@
|
||||
#ifndef vtk_m_cont_RegularConnectivity_h
|
||||
#define vtk_m_cont_RegularConnectivity_h
|
||||
|
||||
#include <vtkm/CellType.h>
|
||||
#include <vtkm/cont/ArrayHandle.h>
|
||||
#include <vtkm/cont/Field.h>
|
||||
#include <vtkm/cont/DynamicArrayHandle.h>
|
||||
#include <vtkm/cont/DeviceAdapterAlgorithm.h>
|
||||
|
||||
namespace vtkm {
|
||||
namespace cont {
|
||||
|
||||
class RegularConnectivity
|
||||
{
|
||||
public:
|
||||
void SetNodeDimension(int node_i, int node_j, int node_k)
|
||||
{
|
||||
cellDims[0] = node_i-1;
|
||||
cellDims[1] = node_j-1;
|
||||
cellDims[2] = node_k-1;
|
||||
nodeDims[0] = node_i;
|
||||
nodeDims[1] = node_j;
|
||||
nodeDims[2] = node_k;
|
||||
}
|
||||
|
||||
vtkm::Id GetNumberOfElements()
|
||||
{
|
||||
return cellDims[0]*cellDims[1]*cellDims[2];
|
||||
}
|
||||
vtkm::Id GetNumberOfIndices(vtkm::Id)
|
||||
{
|
||||
return 8;
|
||||
}
|
||||
vtkm::Id GetElementShapeType(vtkm::Id)
|
||||
{
|
||||
return VTKM_VOXEL;
|
||||
}
|
||||
template <vtkm::IdComponent ItemTupleLength>
|
||||
void GetIndices(vtkm::Id index, vtkm::Vec<vtkm::Id,ItemTupleLength> &ids)
|
||||
{
|
||||
int i,j,k;
|
||||
CalculateLogicalCellIndices3D(index, i,j,k);
|
||||
///\todo: assert ItemTupleLength >= 8, or return early?
|
||||
ids[0] = CalculateNodeIndex3D(i, j, k);
|
||||
if (ItemTupleLength <= 1) return;
|
||||
ids[1] = ids[0] + 1;
|
||||
if (ItemTupleLength <= 2) return;
|
||||
ids[2] = ids[0] + nodeDims[0];
|
||||
if (ItemTupleLength <= 3) return;
|
||||
ids[3] = ids[2] + 1;
|
||||
if (ItemTupleLength <= 4) return;
|
||||
ids[4] = ids[0] + nodeDims[0]*nodeDims[1];
|
||||
if (ItemTupleLength <= 5) return;
|
||||
ids[5] = ids[4] + 1;
|
||||
if (ItemTupleLength <= 6) return;
|
||||
ids[6] = ids[4] + nodeDims[0];
|
||||
if (ItemTupleLength <= 7) return;
|
||||
ids[7] = ids[6] + 1;
|
||||
}
|
||||
private:
|
||||
int cellDims[3];
|
||||
int nodeDims[3];
|
||||
int CalculateCellIndex(int i, int j, int k)
|
||||
{
|
||||
return (k * cellDims[1] + j) * cellDims[0] + i;
|
||||
}
|
||||
int CalculateNodeIndex(int i, int j, int k)
|
||||
{
|
||||
return (k * nodeDims[1] + j) * nodeDims[0] + i;
|
||||
}
|
||||
void CalculateLogicalCellIndices(int index, int &i, int &j, int &k)
|
||||
{
|
||||
int cellDims01 = cellDims[0] * cellDims[1];
|
||||
k = index / cellDims01;
|
||||
int indexij = index % cellDims01;
|
||||
j = indexij / cellDims[0];
|
||||
i = indexij % cellDims[0];
|
||||
}
|
||||
void CalculateLogicalNodeIndices(int index, int &i, int &j, int &k)
|
||||
{
|
||||
int nodeDims01 = nodeDims[0] * nodeDims[1];
|
||||
k = index / nodeDims01;
|
||||
int indexij = index % nodeDims01;
|
||||
j = indexij / nodeDims[0];
|
||||
i = indexij % nodeDims[0];
|
||||
}
|
||||
};
|
||||
|
||||
//TODO:
|
||||
//Add specialized 1D and 2D versions.
|
||||
|
||||
}
|
||||
} // namespace vtkm::cont
|
||||
|
||||
#endif //vtk_m_cont_RegularConnectivity_h
|
@ -129,7 +129,9 @@ public:
|
||||
///
|
||||
/// The allocation may be done on an already existing array, but can wipe out
|
||||
/// any data already in the array. This method can throw
|
||||
/// ErrorControlOutOfMemory if the array cannot be allocated.
|
||||
/// ErrorControlOutOfMemory if the array cannot be allocated or
|
||||
/// ErrorControlBadValue if the allocation is not feasible (for example, the
|
||||
/// array storage is read-only).
|
||||
///
|
||||
VTKM_CONT_EXPORT
|
||||
void Allocate(vtkm::Id numberOfValues);
|
||||
|
@ -48,6 +48,8 @@ struct Transport<vtkm::cont::arg::TransportTagTopologyIn, ContObjectType, Device
|
||||
VTKM_CONT_EXPORT
|
||||
ExecObjectType operator()(const ContObjectType &object, vtkm::Id) const
|
||||
{
|
||||
//DRP: object.PrepareForinput(Device()); //create CUDA version of connectivity array.
|
||||
//make an execution version of the connectivity array.
|
||||
return object;
|
||||
}
|
||||
};
|
||||
|
@ -23,6 +23,8 @@
|
||||
#include <vtkm/cont/Storage.h>
|
||||
#include <vtkm/cont/ErrorControlOutOfMemory.h>
|
||||
|
||||
#include <iostream>
|
||||
|
||||
// Disable GCC warnings we check vtkmfor but Thrust does not.
|
||||
#if defined(__GNUC__) && !defined(VTKM_CUDA)
|
||||
#if (__GNUC__ >= 4) && (__GNUC_MINOR__ >= 6)
|
||||
@ -34,7 +36,7 @@
|
||||
#endif // gcc version >= 4.2
|
||||
#endif // gcc && !CUDA
|
||||
|
||||
#include <thrust/system/cuda/memory.h>
|
||||
#include <thrust/system/cuda/vector.h>
|
||||
#include <thrust/copy.h>
|
||||
|
||||
#if defined(__GNUC__) && !defined(VTKM_CUDA)
|
||||
@ -44,30 +46,12 @@
|
||||
#endif // gcc && !CUDA
|
||||
|
||||
#include <vtkm/exec/cuda/internal/ArrayPortalFromThrust.h>
|
||||
#include <vtkm/exec/cuda/internal/ArrayPortalFromTexture.h>
|
||||
|
||||
#include <boost/utility/enable_if.hpp>
|
||||
|
||||
namespace vtkm {
|
||||
namespace cont {
|
||||
namespace cuda {
|
||||
namespace internal {
|
||||
|
||||
template<typename T> struct UseTexturePortal {typedef boost::false_type type;};
|
||||
|
||||
//Currently disabled as we are still tracking down issues with Texture
|
||||
//Memory. The major issue is that in testing it is slower than classic arrays
|
||||
#ifdef VTKM_USE_TEXTURE_MEM
|
||||
template<> struct UseTexturePortal<vtkm::Int8> {typedef boost::true_type type; };
|
||||
template<> struct UseTexturePortal<vtkm::UInt8> {typedef boost::true_type type; };
|
||||
template<> struct UseTexturePortal<vtkm::Int16> {typedef boost::true_type type; };
|
||||
template<> struct UseTexturePortal<vtkm::UInt16> {typedef boost::true_type type; };
|
||||
template<> struct UseTexturePortal<vtkm::Int32> {typedef boost::true_type type; };
|
||||
template<> struct UseTexturePortal<vtkm::UInt32> {typedef boost::true_type type; };
|
||||
template<> struct UseTexturePortal<vtkm::Float32> {typedef boost::true_type type; };
|
||||
#endif
|
||||
|
||||
|
||||
/// \c ArrayManagerExecutionThrustDevice provides an implementation for a \c
|
||||
/// ArrayManagerExecution class for a thrust device adapter that is designed
|
||||
/// for the cuda backend which has separate memory spaces for host and device. This
|
||||
@ -77,35 +61,32 @@ template<> struct UseTexturePortal<vtkm::Float32> {typedef boost::true_type type
|
||||
/// This array manager should only be used with the cuda device adapter,
|
||||
/// since in the future it will take advantage of texture memory and
|
||||
/// the unique memory access patterns of cuda systems.
|
||||
template<typename T, class StorageTag, typename Enable= void>
|
||||
template<typename T, class StorageTag>
|
||||
class ArrayManagerExecutionThrustDevice
|
||||
{
|
||||
//we need a way to detect that we are using FERMI or lower and disable
|
||||
//the usage of texture iterator. The __CUDA_ARCH__ define is only around
|
||||
//for device code so that can't be used. I expect that we will have to devise
|
||||
//some form of Try/Compile with CUDA or just offer this as an advanced CMake
|
||||
//option. We could also try and see if a runtime switch is possible.
|
||||
|
||||
public:
|
||||
typedef T ValueType;
|
||||
|
||||
typedef vtkm::cont::internal::Storage<ValueType, StorageTag> ContainerType;
|
||||
|
||||
typedef vtkm::exec::cuda::internal::ArrayPortalFromThrust< T > PortalType;
|
||||
typedef vtkm::exec::cuda::internal::ConstArrayPortalFromThrust< T > PortalConstType;
|
||||
typedef vtkm::exec::cuda::internal::ConstArrayPortalFromThrust< const T > PortalConstType;
|
||||
|
||||
VTKM_CONT_EXPORT ArrayManagerExecutionThrustDevice():
|
||||
NumberOfValues(0),
|
||||
ArrayBegin(),
|
||||
ArrayEnd()
|
||||
Array()
|
||||
{
|
||||
|
||||
}
|
||||
|
||||
~ArrayManagerExecutionThrustDevice()
|
||||
{
|
||||
this->ReleaseResources();
|
||||
}
|
||||
|
||||
/// Returns the size of the array.
|
||||
///
|
||||
VTKM_CONT_EXPORT vtkm::Id GetNumberOfValues() const {
|
||||
return this->NumberOfValues;
|
||||
return this->Array.size();
|
||||
}
|
||||
|
||||
/// Allocates the appropriate size of the array and copies the given data
|
||||
@ -119,13 +100,8 @@ public:
|
||||
//calling get portal const
|
||||
try
|
||||
{
|
||||
this->NumberOfValues = arrayPortal.GetNumberOfValues();
|
||||
this->ArrayBegin = ::thrust::system::cuda::malloc<T>( static_cast<std::size_t>(this->NumberOfValues) );
|
||||
this->ArrayEnd = this->ArrayBegin + this->NumberOfValues;
|
||||
|
||||
::thrust::copy(arrayPortal.GetRawIterator(),
|
||||
arrayPortal.GetRawIterator() + this->NumberOfValues,
|
||||
this->ArrayBegin);
|
||||
this->Array.assign(arrayPortal.GetRawIterator(),
|
||||
arrayPortal.GetRawIterator() + arrayPortal.GetNumberOfValues());
|
||||
}
|
||||
catch (std::bad_alloc error)
|
||||
{
|
||||
@ -148,13 +124,16 @@ public:
|
||||
ContainerType &vtkmNotUsed(container),
|
||||
vtkm::Id numberOfValues)
|
||||
{
|
||||
if(this->NumberOfValues > 0)
|
||||
try
|
||||
{
|
||||
::thrust::system::cuda::free( this->ArrayBegin );
|
||||
this->Array.resize(numberOfValues);
|
||||
}
|
||||
this->NumberOfValues = numberOfValues;
|
||||
this->ArrayBegin = ::thrust::system::cuda::malloc<T>( this->NumberOfValues );
|
||||
this->ArrayEnd = this->ArrayBegin + numberOfValues;
|
||||
catch (std::bad_alloc error)
|
||||
{
|
||||
throw vtkm::cont::ErrorControlOutOfMemory(error.what());
|
||||
}
|
||||
|
||||
|
||||
}
|
||||
|
||||
/// Allocates enough space in \c controlArray and copies the data in the
|
||||
@ -162,9 +141,9 @@ public:
|
||||
///
|
||||
VTKM_CONT_EXPORT void RetrieveOutputData(ContainerType &controlArray) const
|
||||
{
|
||||
controlArray.Allocate(this->NumberOfValues);
|
||||
::thrust::copy(this->ArrayBegin,
|
||||
this->ArrayEnd,
|
||||
controlArray.Allocate(this->Array.size());
|
||||
::thrust::copy( this->Array.data(),
|
||||
this->Array.data() + this->Array.size(),
|
||||
controlArray.GetPortal().GetRawIterator());
|
||||
}
|
||||
|
||||
@ -174,19 +153,21 @@ public:
|
||||
{
|
||||
// The operation will succeed even if this assertion fails, but this
|
||||
// is still supposed to be a precondition to Shrink.
|
||||
VTKM_ASSERT_CONT(numberOfValues <= this->NumberOfValues);
|
||||
this->NumberOfValues = numberOfValues;
|
||||
this->ArrayEnd = this->ArrayBegin + this->NumberOfValues;
|
||||
VTKM_ASSERT_CONT(numberOfValues <= this->Array.size());
|
||||
|
||||
this->Array.resize(numberOfValues);
|
||||
}
|
||||
|
||||
VTKM_CONT_EXPORT PortalType GetPortal()
|
||||
{
|
||||
return PortalType(this->ArrayBegin, this->ArrayEnd);
|
||||
return PortalType( this->Array.data(),
|
||||
this->Array.data() + this->Array.size());
|
||||
}
|
||||
|
||||
VTKM_CONT_EXPORT PortalConstType GetPortalConst() const
|
||||
{
|
||||
return PortalConstType(this->ArrayBegin, this->ArrayEnd);
|
||||
return PortalConstType( this->Array.data(),
|
||||
this->Array.data() + this->Array.size());
|
||||
}
|
||||
|
||||
|
||||
@ -194,9 +175,8 @@ public:
|
||||
///
|
||||
VTKM_CONT_EXPORT void ReleaseResources()
|
||||
{
|
||||
::thrust::system::cuda::free( this->ArrayBegin );
|
||||
this->ArrayBegin = ::thrust::system::cuda::pointer<ValueType>();
|
||||
this->ArrayEnd = ::thrust::system::cuda::pointer<ValueType>();
|
||||
this->Array.clear();
|
||||
this->Array.shrink_to_fit();
|
||||
}
|
||||
|
||||
private:
|
||||
@ -206,171 +186,7 @@ private:
|
||||
void operator=(
|
||||
ArrayManagerExecutionThrustDevice<T, StorageTag> &);
|
||||
|
||||
vtkm::Id NumberOfValues;
|
||||
::thrust::system::cuda::pointer<ValueType> ArrayBegin;
|
||||
::thrust::system::cuda::pointer<ValueType> ArrayEnd;
|
||||
};
|
||||
|
||||
|
||||
/// This is a specialization that is used to enable texture memory iterators
|
||||
template<typename T, class StorageTag>
|
||||
class ArrayManagerExecutionThrustDevice<T, StorageTag,
|
||||
typename ::boost::enable_if< typename UseTexturePortal<T>::type >::type >
|
||||
{
|
||||
//we need a way to detect that we are using FERMI or lower and disable
|
||||
//the usage of texture iterator. The __CUDA_ARCH__ define is only around
|
||||
//for device code so that can't be used. I expect that we will have to devise
|
||||
//some form of Try/Compile with CUDA or just offer this as an advanced CMake
|
||||
//option. We could also try and see if a runtime switch is possible.
|
||||
|
||||
public:
|
||||
typedef T ValueType;
|
||||
|
||||
typedef vtkm::cont::internal::Storage<ValueType, StorageTag> ContainerType;
|
||||
|
||||
typedef vtkm::exec::cuda::internal::ArrayPortalFromThrust< T > PortalType;
|
||||
typedef ::vtkm::exec::cuda::internal::DaxTexObjInputIterator<T> TextureIteratorType;
|
||||
typedef ::vtkm::exec::cuda::internal::ConstArrayPortalFromTexture< TextureIteratorType > PortalConstType;
|
||||
|
||||
VTKM_CONT_EXPORT ArrayManagerExecutionThrustDevice():
|
||||
NumberOfValues(0),
|
||||
ArrayBegin(),
|
||||
ArrayEnd(),
|
||||
HaveTextureBound(false),
|
||||
InputArrayIterator()
|
||||
{
|
||||
|
||||
}
|
||||
|
||||
~ArrayManagerExecutionThrustDevice()
|
||||
{
|
||||
if(this->HaveTextureBound)
|
||||
{
|
||||
this->HaveTextureBound = false;
|
||||
this->InputArrayIterator.UnbindTexture();
|
||||
}
|
||||
}
|
||||
|
||||
/// Returns the size of the array.
|
||||
///
|
||||
VTKM_CONT_EXPORT vtkm::Id GetNumberOfValues() const {
|
||||
return this->NumberOfValues;
|
||||
}
|
||||
|
||||
/// Allocates the appropriate size of the array and copies the given data
|
||||
/// into the array.
|
||||
///
|
||||
template<class PortalControl>
|
||||
VTKM_CONT_EXPORT void LoadDataForInput(PortalControl arrayPortal)
|
||||
{
|
||||
//don't bind to the texture yet, as we could have allocate the array
|
||||
//on a previous call with AllocateArrayForOutput and now are directly
|
||||
//calling get portal const
|
||||
try
|
||||
{
|
||||
this->NumberOfValues = arrayPortal.GetNumberOfValues();
|
||||
this->ArrayBegin = ::thrust::system::cuda::malloc<T>( static_cast<std::size_t>(this->NumberOfValues) );
|
||||
this->ArrayEnd = this->ArrayBegin + this->NumberOfValues;
|
||||
|
||||
::thrust::copy(arrayPortal.GetRawIterator(),
|
||||
arrayPortal.GetRawIterator() + this->NumberOfValues,
|
||||
this->ArrayBegin);
|
||||
}
|
||||
catch (std::bad_alloc error)
|
||||
{
|
||||
throw vtkm::cont::ErrorControlOutOfMemory(error.what());
|
||||
}
|
||||
}
|
||||
|
||||
/// Allocates the appropriate size of the array and copies the given data
|
||||
/// into the array.
|
||||
///
|
||||
template<class PortalControl>
|
||||
VTKM_CONT_EXPORT void LoadDataForInPlace(PortalControl arrayPortal)
|
||||
{
|
||||
this->LoadDataForInput(arrayPortal);
|
||||
}
|
||||
|
||||
/// Allocates the array to the given size.
|
||||
///
|
||||
VTKM_CONT_EXPORT void AllocateArrayForOutput(
|
||||
ContainerType &vtkmNotUsed(container),
|
||||
vtkm::Id numberOfValues)
|
||||
{
|
||||
if(this->NumberOfValues > 0)
|
||||
{
|
||||
::thrust::system::cuda::free( this->ArrayBegin );
|
||||
}
|
||||
this->NumberOfValues = numberOfValues;
|
||||
this->ArrayBegin = ::thrust::system::cuda::malloc<T>( this->NumberOfValues );
|
||||
this->ArrayEnd = this->ArrayBegin + numberOfValues;
|
||||
}
|
||||
|
||||
/// Allocates enough space in \c controlArray and copies the data in the
|
||||
/// device vector into it.
|
||||
///
|
||||
VTKM_CONT_EXPORT void RetrieveOutputData(ContainerType &controlArray) const
|
||||
{
|
||||
controlArray.Allocate(this->NumberOfValues);
|
||||
::thrust::copy(this->ArrayBegin,
|
||||
this->ArrayEnd,
|
||||
controlArray.GetPortal().GetRawIterator());
|
||||
}
|
||||
|
||||
/// Resizes the device vector.
|
||||
///
|
||||
VTKM_CONT_EXPORT void Shrink(vtkm::Id numberOfValues)
|
||||
{
|
||||
// The operation will succeed even if this assertion fails, but this
|
||||
// is still supposed to be a precondition to Shrink.
|
||||
VTKM_ASSERT_CONT(numberOfValues <= this->NumberOfValues);
|
||||
this->NumberOfValues = numberOfValues;
|
||||
this->ArrayEnd = this->ArrayBegin + this->NumberOfValues;
|
||||
}
|
||||
|
||||
VTKM_CONT_EXPORT PortalType GetPortal()
|
||||
{
|
||||
return PortalType(this->ArrayBegin, this->ArrayEnd);
|
||||
}
|
||||
|
||||
VTKM_CONT_EXPORT PortalConstType GetPortalConst() const
|
||||
{
|
||||
if(!this->HaveTextureBound)
|
||||
{
|
||||
this->HaveTextureBound = true;
|
||||
this->InputArrayIterator.BindTexture(ArrayBegin,this->NumberOfValues);
|
||||
}
|
||||
|
||||
//if we have a texture iterator bound use that
|
||||
return PortalConstType(this->InputArrayIterator, this->NumberOfValues);
|
||||
}
|
||||
|
||||
|
||||
/// Frees all memory.
|
||||
///
|
||||
VTKM_CONT_EXPORT void ReleaseResources() {
|
||||
if(this->HaveTextureBound)
|
||||
{
|
||||
this->HaveTextureBound = false;
|
||||
this->InputArrayIterator.UnbindTexture();
|
||||
}
|
||||
::thrust::system::cuda::free( this->ArrayBegin );
|
||||
this->ArrayBegin = ::thrust::system::cuda::pointer<ValueType>();
|
||||
this->ArrayEnd = ::thrust::system::cuda::pointer<ValueType>();
|
||||
}
|
||||
|
||||
private:
|
||||
// Not implemented
|
||||
ArrayManagerExecutionThrustDevice(
|
||||
ArrayManagerExecutionThrustDevice<T, StorageTag> &);
|
||||
void operator=(
|
||||
ArrayManagerExecutionThrustDevice<T, StorageTag> &);
|
||||
|
||||
vtkm::Id NumberOfValues;
|
||||
::thrust::system::cuda::pointer<ValueType> ArrayBegin;
|
||||
::thrust::system::cuda::pointer<ValueType> ArrayEnd;
|
||||
mutable bool HaveTextureBound;
|
||||
mutable TextureIteratorType InputArrayIterator;
|
||||
::thrust::system::cuda::vector<ValueType> Array;
|
||||
};
|
||||
|
||||
|
||||
|
@ -24,7 +24,6 @@
|
||||
#include <vtkm/internal/ExportMacros.h>
|
||||
|
||||
#include <vtkm/exec/cuda/internal/ArrayPortalFromThrust.h>
|
||||
#include <vtkm/exec/cuda/internal/ArrayPortalFromTexture.h>
|
||||
|
||||
// Disable GCC warnings we check vtkmfor but Thrust does not.
|
||||
#if defined(__GNUC__) && !defined(VTKM_CUDA)
|
||||
@ -59,7 +58,6 @@ namespace detail {
|
||||
// Tags to specify what type of thrust iterator to use.
|
||||
struct ThrustIteratorTransformTag { };
|
||||
struct ThrustIteratorDevicePtrTag { };
|
||||
struct ThrustIteratorDeviceTextureTag { };
|
||||
|
||||
// Traits to help classify what thrust iterators will be used.
|
||||
template<class IteratorType>
|
||||
@ -149,13 +147,6 @@ struct IteratorTraits
|
||||
typedef typename IteratorChooser<PortalType, Tag>::Type IteratorType;
|
||||
};
|
||||
|
||||
template<typename T>
|
||||
struct IteratorTraits< vtkm::exec::cuda::internal::ConstArrayPortalFromTexture< T > >
|
||||
{
|
||||
typedef vtkm::exec::cuda::internal::ConstArrayPortalFromTexture< T > PortalType;
|
||||
typedef ThrustIteratorDeviceTextureTag Tag;
|
||||
typedef typename PortalType::IteratorType IteratorType;
|
||||
};
|
||||
|
||||
template<typename T>
|
||||
VTKM_CONT_EXPORT static
|
||||
@ -190,14 +181,6 @@ MakeIteratorBegin(PortalType portal, detail::ThrustIteratorDevicePtrTag)
|
||||
return MakeDevicePtr(portal.GetIteratorBegin());
|
||||
}
|
||||
|
||||
template<class PortalType>
|
||||
VTKM_CONT_EXPORT static
|
||||
typename IteratorTraits<PortalType>::IteratorType
|
||||
MakeIteratorBegin(PortalType portal, detail::ThrustIteratorDeviceTextureTag)
|
||||
{
|
||||
return portal.GetIteratorBegin();
|
||||
}
|
||||
|
||||
} // namespace detail
|
||||
|
||||
|
||||
|
@ -113,6 +113,12 @@ struct TryArrayHandleType
|
||||
"Array size did not shrink correctly.");
|
||||
CheckArray(arrayHandle);
|
||||
|
||||
std::cout << "Try reallocating array." << std::endl;
|
||||
arrayHandle.Allocate(ARRAY_SIZE*2);
|
||||
VTKM_TEST_ASSERT(arrayHandle.GetNumberOfValues() == ARRAY_SIZE*2,
|
||||
"Array size did not allocate correctly.");
|
||||
// No point in checking values. This method can invalidate them.
|
||||
|
||||
std::cout << "Try in place operation." << std::endl;
|
||||
{
|
||||
typedef typename vtkm::cont::ArrayHandle<T>::template
|
||||
@ -135,6 +141,25 @@ struct TryArrayHandleType
|
||||
TestValue(index, T()) + T(1)),
|
||||
"Did not get result from in place operation.");
|
||||
}
|
||||
|
||||
std::cout << "Try operations on empty arrays." << std::endl;
|
||||
// After each operation, reinitialize array in case something gets
|
||||
// allocated.
|
||||
arrayHandle = vtkm::cont::ArrayHandle<T>();
|
||||
VTKM_TEST_ASSERT(arrayHandle.GetNumberOfValues() == 0,
|
||||
"Uninitialized array does not report zero values.");
|
||||
arrayHandle = vtkm::cont::ArrayHandle<T>();
|
||||
VTKM_TEST_ASSERT(
|
||||
arrayHandle.GetPortalConstControl().GetNumberOfValues() == 0,
|
||||
"Uninitialized array does not give portal with zero values.");
|
||||
arrayHandle = vtkm::cont::ArrayHandle<T>();
|
||||
arrayHandle.Shrink(0);
|
||||
arrayHandle = vtkm::cont::ArrayHandle<T>();
|
||||
arrayHandle.ReleaseResourcesExecution();
|
||||
arrayHandle = vtkm::cont::ArrayHandle<T>();
|
||||
arrayHandle.ReleaseResources();
|
||||
arrayHandle = vtkm::cont::ArrayHandle<T>();
|
||||
arrayHandle.PrepareForOutput(ARRAY_SIZE, VTKM_DEFAULT_DEVICE_ADAPTER_TAG());
|
||||
}
|
||||
};
|
||||
|
||||
|
@ -18,6 +18,10 @@
|
||||
// this software.
|
||||
//============================================================================
|
||||
|
||||
//#define VTKM_DEVICE_ADAPTER VTKM_DEVICE_ADAPTER_SERIAL
|
||||
//#define VTKM_DEVICE_ADAPTER VTKM_DEVICE_ADAPTER_CUDA
|
||||
//#include <vtkm/cont/DeviceAdapter.h>
|
||||
|
||||
#include <vtkm/cont/testing/Testing.h>
|
||||
#include <vtkm/cont/DataSet.h>
|
||||
#include <vtkm/cont/DeviceAdapterAlgorithm.h>
|
||||
@ -27,6 +31,13 @@
|
||||
#include <vtkm/exec/arg/TopologyIdCount.h>
|
||||
#include <vtkm/exec/arg/TopologyElementType.h>
|
||||
|
||||
/*
|
||||
call notes.
|
||||
wrap execution portal with restrictors. (abandone the fixed length array).
|
||||
compile time polymorphic types.
|
||||
|
||||
*/
|
||||
|
||||
static const int LEN_IDS = 6;
|
||||
|
||||
class CellType : public vtkm::worklet::WorkletMapTopology
|
||||
|
@ -1,490 +0,0 @@
|
||||
//============================================================================
|
||||
// Copyright (c) Kitware, Inc.
|
||||
// All rights reserved.
|
||||
// See LICENSE.txt for details.
|
||||
// This software is distributed WITHOUT ANY WARRANTY; without even
|
||||
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
|
||||
// PURPOSE. See the above copyright notice for more information.
|
||||
//
|
||||
// Copyright 2014 Sandia Corporation.
|
||||
// Copyright 2014 UT-Battelle, LLC.
|
||||
// Copyright 2014. Los Alamos National Security
|
||||
//
|
||||
// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
|
||||
// the U.S. Government retains certain rights in this software.
|
||||
//
|
||||
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
|
||||
// Laboratory (LANL), the U.S. Government retains certain rights in
|
||||
// this software.
|
||||
//============================================================================
|
||||
|
||||
/******************************************************************************
|
||||
* Copyright (c) 2011, Duane Merrill. All rights reserved.
|
||||
* Copyright (c) 2011-2014, 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 TORT
|
||||
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
|
||||
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
|
||||
*
|
||||
******************************************************************************/
|
||||
|
||||
#ifndef vtk_m_exec_cuda_internal_ArrayPortalFromTexture_h
|
||||
#define vtk_m_exec_cuda_internal_ArrayPortalFromTexture_h
|
||||
|
||||
#include <vtkm/Types.h>
|
||||
|
||||
#include <iterator>
|
||||
|
||||
// Disable GCC warnings we check vtkmfor but Thrust does not.
|
||||
#if defined(__GNUC__) && !defined(VTKM_CUDA)
|
||||
#if (__GNUC__ >= 4) && (__GNUC_MINOR__ >= 6)
|
||||
#pragma GCC diagnostic push
|
||||
#endif // gcc version >= 4.6
|
||||
#if (__GNUC__ >= 4) && (__GNUC_MINOR__ >= 2)
|
||||
#pragma GCC diagnostic ignored "-Wshadow"
|
||||
#pragma GCC diagnostic ignored "-Wunused-parameter"
|
||||
#endif // gcc version >= 4.2
|
||||
#endif // gcc && !CUDA
|
||||
|
||||
#include <thrust/system/cuda/memory.h>
|
||||
#include <thrust/iterator/iterator_facade.h>
|
||||
|
||||
#if defined(__GNUC__) && !defined(VTKM_CUDA)
|
||||
#if (__GNUC__ >= 4) && (__GNUC_MINOR__ >= 6)
|
||||
#pragma GCC diagnostic pop
|
||||
#endif // gcc version >= 4.6
|
||||
#endif // gcc && !CUDA
|
||||
|
||||
|
||||
namespace
|
||||
{
|
||||
|
||||
/**
|
||||
* \brief Type selection (<tt>IF ? ThenType : ElseType</tt>)
|
||||
*/
|
||||
template <bool IF, typename ThenType, typename ElseType>
|
||||
struct If
|
||||
{
|
||||
/// Conditional type result
|
||||
typedef ThenType Type; // true
|
||||
};
|
||||
|
||||
template <typename ThenType, typename ElseType>
|
||||
struct If<false, ThenType, ElseType>
|
||||
{
|
||||
typedef ElseType Type; // false
|
||||
};
|
||||
|
||||
/******************************************************************************
|
||||
* Size and alignment
|
||||
******************************************************************************/
|
||||
|
||||
/// Structure alignment
|
||||
template <typename T>
|
||||
struct AlignBytes
|
||||
{
|
||||
struct Pad
|
||||
{
|
||||
T val;
|
||||
char byte;
|
||||
};
|
||||
|
||||
enum
|
||||
{
|
||||
/// The alignment of T in bytes
|
||||
ALIGN_BYTES = sizeof(Pad) - sizeof(T)
|
||||
};
|
||||
};
|
||||
|
||||
// Specializations where host C++ compilers (e.g., Windows) may disagree with device C++ compilers (EDG)
|
||||
|
||||
template <> struct AlignBytes<short4> { enum { ALIGN_BYTES = 8 }; };
|
||||
template <> struct AlignBytes<ushort4> { enum { ALIGN_BYTES = 8 }; };
|
||||
template <> struct AlignBytes<int2> { enum { ALIGN_BYTES = 8 }; };
|
||||
template <> struct AlignBytes<uint2> { enum { ALIGN_BYTES = 8 }; };
|
||||
#ifdef _WIN32
|
||||
template <> struct AlignBytes<long2> { enum { ALIGN_BYTES = 8 }; };
|
||||
template <> struct AlignBytes<ulong2> { enum { ALIGN_BYTES = 8 }; };
|
||||
#endif
|
||||
template <> struct AlignBytes<long long> { enum { ALIGN_BYTES = 8 }; };
|
||||
template <> struct AlignBytes<unsigned long long> { enum { ALIGN_BYTES = 8 }; };
|
||||
template <> struct AlignBytes<float2> { enum { ALIGN_BYTES = 8 }; };
|
||||
template <> struct AlignBytes<double> { enum { ALIGN_BYTES = 8 }; };
|
||||
|
||||
template <> struct AlignBytes<int4> { enum { ALIGN_BYTES = 16 }; };
|
||||
template <> struct AlignBytes<uint4> { enum { ALIGN_BYTES = 16 }; };
|
||||
template <> struct AlignBytes<float4> { enum { ALIGN_BYTES = 16 }; };
|
||||
#ifndef _WIN32
|
||||
template <> struct AlignBytes<long2> { enum { ALIGN_BYTES = 16 }; };
|
||||
template <> struct AlignBytes<ulong2> { enum { ALIGN_BYTES = 16 }; };
|
||||
#endif
|
||||
template <> struct AlignBytes<long4> { enum { ALIGN_BYTES = 16 }; };
|
||||
template <> struct AlignBytes<ulong4> { enum { ALIGN_BYTES = 16 }; };
|
||||
template <> struct AlignBytes<longlong2> { enum { ALIGN_BYTES = 16 }; };
|
||||
template <> struct AlignBytes<ulonglong2> { enum { ALIGN_BYTES = 16 }; };
|
||||
template <> struct AlignBytes<double2> { enum { ALIGN_BYTES = 16 }; };
|
||||
template <> struct AlignBytes<longlong4> { enum { ALIGN_BYTES = 16 }; };
|
||||
template <> struct AlignBytes<ulonglong4> { enum { ALIGN_BYTES = 16 }; };
|
||||
template <> struct AlignBytes<double4> { enum { ALIGN_BYTES = 16 }; };
|
||||
|
||||
|
||||
/// Unit-words of data movement
|
||||
template <typename T>
|
||||
struct UnitWord
|
||||
{
|
||||
enum {
|
||||
ALIGN_BYTES = AlignBytes<T>::ALIGN_BYTES
|
||||
};
|
||||
|
||||
template <typename Unit>
|
||||
struct IsMultiple
|
||||
{
|
||||
enum {
|
||||
UNIT_ALIGN_BYTES = AlignBytes<Unit>::ALIGN_BYTES,
|
||||
IS_MULTIPLE = (sizeof(T) % sizeof(Unit) == 0) && (ALIGN_BYTES % UNIT_ALIGN_BYTES == 0)
|
||||
};
|
||||
};
|
||||
|
||||
/// Biggest shuffle word that T is a whole multiple of and is not larger than the alignment of T
|
||||
typedef typename If<IsMultiple<int>::IS_MULTIPLE,
|
||||
unsigned int,
|
||||
typename If<IsMultiple<short>::IS_MULTIPLE,
|
||||
unsigned short,
|
||||
unsigned char>::Type>::Type ShuffleWord;
|
||||
|
||||
/// Biggest volatile word that T is a whole multiple of and is not larger than the alignment of T
|
||||
typedef typename If<IsMultiple<long long>::IS_MULTIPLE,
|
||||
unsigned long long,
|
||||
ShuffleWord>::Type VolatileWord;
|
||||
|
||||
/// Biggest memory-access word that T is a whole multiple of and is not larger than the alignment of T
|
||||
typedef typename If<IsMultiple<longlong2>::IS_MULTIPLE,
|
||||
ulonglong2,
|
||||
VolatileWord>::Type DeviceWord;
|
||||
|
||||
/// Biggest texture reference word that T is a whole multiple of and is not larger than the alignment of T
|
||||
typedef typename If<IsMultiple<int4>::IS_MULTIPLE,
|
||||
uint4,
|
||||
typename If<IsMultiple<int2>::IS_MULTIPLE,
|
||||
uint2,
|
||||
ShuffleWord>::Type>::Type TextureWord;
|
||||
};
|
||||
|
||||
}
|
||||
|
||||
|
||||
namespace vtkm {
|
||||
namespace exec {
|
||||
namespace cuda {
|
||||
namespace internal {
|
||||
|
||||
template <
|
||||
typename T,
|
||||
typename Offset = ptrdiff_t>
|
||||
class DaxTexObjInputIterator
|
||||
{
|
||||
public:
|
||||
|
||||
// Required iterator traits
|
||||
typedef DaxTexObjInputIterator self_type; ///< My own type
|
||||
typedef Offset difference_type; ///< Type to express the result of subtracting one iterator from another
|
||||
typedef T value_type; ///< The type of the element the iterator can point to
|
||||
typedef T* pointer; ///< The type of a pointer to an element the iterator can point to
|
||||
typedef T reference; ///< The type of a reference to an element the iterator can point to
|
||||
|
||||
// Use Thrust's iterator categories so we can use these iterators in Thrust 1.7 (or newer) methods
|
||||
typedef typename ::thrust::detail::iterator_facade_category<
|
||||
::thrust::device_system_tag,
|
||||
::thrust::random_access_traversal_tag,
|
||||
value_type,
|
||||
reference
|
||||
>::type iterator_category; ///< The iterator category
|
||||
|
||||
private:
|
||||
|
||||
// Largest texture word we can use in device
|
||||
typedef typename UnitWord<T>::TextureWord TextureWord;
|
||||
|
||||
// Number of texture words per T
|
||||
enum { TEXTURE_MULTIPLE = sizeof(T) / sizeof(TextureWord) };
|
||||
|
||||
private:
|
||||
|
||||
const T* ptr;
|
||||
difference_type tex_offset;
|
||||
cudaTextureObject_t tex_obj;
|
||||
|
||||
public:
|
||||
|
||||
/// Constructor
|
||||
__host__ __device__ __forceinline__ DaxTexObjInputIterator()
|
||||
:
|
||||
ptr(NULL),
|
||||
tex_offset(0),
|
||||
tex_obj(0)
|
||||
{}
|
||||
|
||||
/// Use this iterator to bind \p ptr with a texture reference
|
||||
cudaError_t BindTexture(
|
||||
const ::thrust::system::cuda::pointer<T> ptr, ///< Native pointer to wrap that is aligned to cudaDeviceProp::textureAlignment
|
||||
size_t numElements, ///< Number of elements in the range
|
||||
size_t tex_offset = 0) ///< Offset (in items) from \p ptr denoting the position of the iterator
|
||||
{
|
||||
this->ptr = ptr.get();
|
||||
this->tex_offset = tex_offset;
|
||||
|
||||
cudaChannelFormatDesc channel_desc = cudaCreateChannelDesc<TextureWord>();
|
||||
cudaResourceDesc res_desc;
|
||||
cudaTextureDesc tex_desc;
|
||||
memset(&res_desc, 0, sizeof(cudaResourceDesc));
|
||||
memset(&tex_desc, 0, sizeof(cudaTextureDesc));
|
||||
res_desc.resType = cudaResourceTypeLinear;
|
||||
res_desc.res.linear.devPtr = (void*)ptr.get();
|
||||
res_desc.res.linear.desc = channel_desc;
|
||||
res_desc.res.linear.sizeInBytes = numElements * sizeof(T);
|
||||
tex_desc.readMode = cudaReadModeElementType;
|
||||
|
||||
return cudaCreateTextureObject(&tex_obj, &res_desc, &tex_desc, NULL);
|
||||
}
|
||||
|
||||
/// Unbind this iterator from its texture reference
|
||||
cudaError_t UnbindTexture()
|
||||
{
|
||||
return cudaDestroyTextureObject(tex_obj);
|
||||
}
|
||||
|
||||
/// Postfix increment
|
||||
__host__ __device__ __forceinline__ self_type operator++(int)
|
||||
{
|
||||
self_type retval = *this;
|
||||
tex_offset++;
|
||||
return retval;
|
||||
}
|
||||
|
||||
/// Prefix increment
|
||||
__host__ __device__ __forceinline__ self_type operator++()
|
||||
{
|
||||
tex_offset++;
|
||||
return *this;
|
||||
}
|
||||
|
||||
/// Postfix decrement
|
||||
__host__ __device__ __forceinline__ self_type operator--(int)
|
||||
{
|
||||
self_type retval = *this;
|
||||
tex_offset--;
|
||||
return retval;
|
||||
}
|
||||
|
||||
/// Prefix decrement
|
||||
__host__ __device__ __forceinline__ self_type operator--()
|
||||
{
|
||||
tex_offset--;
|
||||
return *this;
|
||||
}
|
||||
|
||||
/// Indirection
|
||||
__host__ __device__ __forceinline__ reference operator*() const
|
||||
{
|
||||
#ifndef DAX_CUDA_COMPILATION
|
||||
// Simply dereference the pointer on the host
|
||||
return ptr[tex_offset];
|
||||
#else
|
||||
// Move array of uninitialized words, then alias and assign to return value
|
||||
TextureWord words[TEXTURE_MULTIPLE];
|
||||
|
||||
#pragma unroll
|
||||
for (int i = 0; i < TEXTURE_MULTIPLE; ++i)
|
||||
{
|
||||
words[i] = tex1Dfetch<TextureWord>(
|
||||
tex_obj,
|
||||
(tex_offset * TEXTURE_MULTIPLE) + i);
|
||||
}
|
||||
|
||||
// Load from words
|
||||
return *reinterpret_cast<T*>(words);
|
||||
#endif
|
||||
}
|
||||
|
||||
/// Addition
|
||||
template <typename Distance>
|
||||
__host__ __device__ __forceinline__ self_type operator+(Distance n) const
|
||||
{
|
||||
self_type retval;
|
||||
retval.ptr = ptr;
|
||||
retval.tex_obj = tex_obj;
|
||||
retval.tex_offset = tex_offset + n;
|
||||
return retval;
|
||||
}
|
||||
|
||||
/// Addition assignment
|
||||
template <typename Distance>
|
||||
__host__ __device__ __forceinline__ self_type& operator+=(Distance n)
|
||||
{
|
||||
tex_offset += n;
|
||||
return *this;
|
||||
}
|
||||
|
||||
/// Subtraction
|
||||
template <typename Distance>
|
||||
__host__ __device__ __forceinline__ self_type operator-(Distance n) const
|
||||
{
|
||||
self_type retval;
|
||||
retval.ptr = ptr;
|
||||
retval.tex_obj = tex_obj;
|
||||
retval.tex_offset = tex_offset - n;
|
||||
return retval;
|
||||
}
|
||||
|
||||
/// Subtraction assignment
|
||||
template <typename Distance>
|
||||
__host__ __device__ __forceinline__ self_type& operator-=(Distance n)
|
||||
{
|
||||
tex_offset -= n;
|
||||
return *this;
|
||||
}
|
||||
|
||||
/// Distance
|
||||
__host__ __device__ __forceinline__ difference_type operator-(self_type other) const
|
||||
{
|
||||
return tex_offset - other.tex_offset;
|
||||
}
|
||||
|
||||
/// Array subscript
|
||||
template <typename Distance>
|
||||
__host__ __device__ __forceinline__ reference operator[](Distance n) const
|
||||
{
|
||||
return *(*this + n);
|
||||
}
|
||||
|
||||
/// Structure dereference
|
||||
__host__ __device__ __forceinline__ pointer operator->()
|
||||
{
|
||||
return &(*(*this));
|
||||
}
|
||||
|
||||
/// Equal to
|
||||
__host__ __device__ __forceinline__ bool operator==(const self_type& rhs) const
|
||||
{
|
||||
return ((ptr == rhs.ptr) && (tex_offset == rhs.tex_offset) && (tex_obj == rhs.tex_obj));
|
||||
}
|
||||
|
||||
/// Not equal to
|
||||
__host__ __device__ __forceinline__ bool operator!=(const self_type& rhs) const
|
||||
{
|
||||
return ((ptr != rhs.ptr) || (tex_offset != rhs.tex_offset) || (tex_obj != rhs.tex_obj));
|
||||
}
|
||||
|
||||
/// less than
|
||||
__host__ __device__ __forceinline__ bool operator<(const self_type& rhs)
|
||||
{
|
||||
return (tex_offset < rhs.tex_offset);
|
||||
}
|
||||
|
||||
/// ostream operator
|
||||
friend std::ostream& operator<<(std::ostream& os, const self_type& itr)
|
||||
{
|
||||
return os;
|
||||
}
|
||||
|
||||
};
|
||||
|
||||
|
||||
template<class TextureIterator>
|
||||
class ConstArrayPortalFromTexture
|
||||
{
|
||||
public:
|
||||
|
||||
typedef typename TextureIterator::value_type ValueType;
|
||||
typedef TextureIterator IteratorType;
|
||||
|
||||
VTKM_EXEC_CONT_EXPORT ConstArrayPortalFromTexture() { }
|
||||
|
||||
VTKM_CONT_EXPORT
|
||||
ConstArrayPortalFromTexture(IteratorType begin, ptrdiff_t size)
|
||||
: Length(size),
|
||||
BeginIterator(begin),
|
||||
EndIterator(begin+size)
|
||||
{ }
|
||||
|
||||
/// Copy constructor for any other ConstArrayPortalFromTexture with an iterator
|
||||
/// type that can be copied to this iterator type. This allows us to do any
|
||||
/// type casting that the iterators do (like the non-const to const cast).
|
||||
///
|
||||
template<typename OtherIteratorT>
|
||||
VTKM_EXEC_CONT_EXPORT
|
||||
ConstArrayPortalFromTexture(const ConstArrayPortalFromTexture<OtherIteratorT> &src)
|
||||
: Length(src.Length),
|
||||
BeginIterator(src.BeginIterator),
|
||||
EndIterator(src.EndIterator)
|
||||
{ }
|
||||
|
||||
template<typename OtherIteratorT>
|
||||
VTKM_EXEC_CONT_EXPORT
|
||||
ConstArrayPortalFromTexture<IteratorType> &operator=(
|
||||
const ConstArrayPortalFromTexture<OtherIteratorT> &src)
|
||||
{
|
||||
this->Length = src.Length;
|
||||
this->BeginIterator = src.BeginIterator;
|
||||
this->EndIterator = src.EndIterator;
|
||||
return *this;
|
||||
}
|
||||
|
||||
VTKM_EXEC_CONT_EXPORT
|
||||
vtkm::Id GetNumberOfValues() const {
|
||||
return static_cast<vtkm::Id>(this->Length);
|
||||
}
|
||||
|
||||
VTKM_EXEC_EXPORT
|
||||
ValueType Get(vtkm::Id index) const {
|
||||
return *this->IteratorAt(index);
|
||||
}
|
||||
|
||||
VTKM_EXEC_EXPORT
|
||||
void Set(vtkm::Id index, ValueType value) const {
|
||||
*this->IteratorAt(index) = value;
|
||||
}
|
||||
|
||||
VTKM_CONT_EXPORT
|
||||
IteratorType GetIteratorBegin() const { return this->BeginIterator; }
|
||||
|
||||
VTKM_CONT_EXPORT
|
||||
IteratorType GetIteratorEnd() const { return this->EndIterator; }
|
||||
|
||||
private:
|
||||
ptrdiff_t Length;
|
||||
IteratorType BeginIterator;
|
||||
IteratorType EndIterator;
|
||||
|
||||
VTKM_EXEC_EXPORT
|
||||
IteratorType IteratorAt(vtkm::Id index) const {
|
||||
// Not using std::advance because on CUDA it cannot be used on a device.
|
||||
return (this->BeginIterator + index);
|
||||
}
|
||||
};
|
||||
|
||||
}
|
||||
}
|
||||
}
|
||||
} // namespace vtkm::exec::cuda::internal
|
||||
|
||||
|
||||
#endif //vtk_m_exec_cuda_internal_ArrayPortalFromTexture_h
|
@ -42,11 +42,63 @@
|
||||
#pragma GCC diagnostic pop
|
||||
#endif // gcc version >= 4.6
|
||||
#endif // gcc && !CUDA
|
||||
|
||||
#include <boost/utility/enable_if.hpp>
|
||||
|
||||
namespace vtkm {
|
||||
namespace exec {
|
||||
namespace cuda {
|
||||
namespace internal {
|
||||
|
||||
template<typename T> struct UseTextureLoad {typedef boost::false_type type;};
|
||||
|
||||
template<> struct UseTextureLoad<vtkm::Int8*> {typedef boost::true_type type; };
|
||||
template<> struct UseTextureLoad<vtkm::UInt8*> {typedef boost::true_type type; };
|
||||
template<> struct UseTextureLoad<vtkm::Int16*> {typedef boost::true_type type; };
|
||||
template<> struct UseTextureLoad<vtkm::UInt16*> {typedef boost::true_type type; };
|
||||
template<> struct UseTextureLoad<vtkm::Int32*> {typedef boost::true_type type; };
|
||||
template<> struct UseTextureLoad<vtkm::UInt32*> {typedef boost::true_type type; };
|
||||
|
||||
template<> struct UseTextureLoad<vtkm::Vec<vtkm::Int32,2>* > {typedef boost::true_type type; };
|
||||
template<> struct UseTextureLoad<vtkm::Vec<vtkm::UInt32,2>* > {typedef boost::true_type type; };
|
||||
template<> struct UseTextureLoad<vtkm::Vec<vtkm::Int32,4>* > {typedef boost::true_type type; };
|
||||
template<> struct UseTextureLoad<vtkm::Vec<vtkm::UInt32,4>* > {typedef boost::true_type type; };
|
||||
|
||||
template<> struct UseTextureLoad<vtkm::Float32* > {typedef boost::true_type type; };
|
||||
template<> struct UseTextureLoad<vtkm::Float64* > {typedef boost::true_type type; };
|
||||
|
||||
template<> struct UseTextureLoad<vtkm::Vec<vtkm::Float32,2>* > {typedef boost::true_type type; };
|
||||
template<> struct UseTextureLoad<vtkm::Vec<vtkm::Float32,4>* > {typedef boost::true_type type; };
|
||||
template<> struct UseTextureLoad<vtkm::Vec<vtkm::Float64,2>* > {typedef boost::true_type type; };
|
||||
|
||||
//this T type is not one that is valid to be loaded through texture memory
|
||||
template<typename T, typename Enable = void>
|
||||
struct load_through_texture
|
||||
{
|
||||
VTKM_EXEC_EXPORT
|
||||
static T get(const thrust::system::cuda::pointer<T> data)
|
||||
{
|
||||
return *(data.get());
|
||||
}
|
||||
};
|
||||
|
||||
//this T type is valid to be loaded through texture memory
|
||||
template<typename T>
|
||||
struct load_through_texture<T, typename ::boost::enable_if< typename UseTextureLoad<T>::type >::type >
|
||||
{
|
||||
VTKM_EXEC_EXPORT
|
||||
static T get(const thrust::system::cuda::pointer<T> data)
|
||||
{
|
||||
//only load through a texture if we have sm 35 support
|
||||
#if __CUDA_ARCH__ >= 350
|
||||
return __ldg(data.get());
|
||||
#else
|
||||
return *(data.get());
|
||||
#endif
|
||||
}
|
||||
};
|
||||
|
||||
|
||||
class ArrayPortalFromThrustBase {};
|
||||
|
||||
/// This templated implementation of an ArrayPortal allows you to adapt a pair
|
||||
@ -168,7 +220,7 @@ public:
|
||||
|
||||
VTKM_EXEC_EXPORT
|
||||
ValueType Get(vtkm::Id index) const {
|
||||
return *this->IteratorAt(index);
|
||||
return vtkm::exec::cuda::internal::load_through_texture<ValueType>::get( this->IteratorAt(index) );
|
||||
}
|
||||
|
||||
VTKM_EXEC_EXPORT
|
||||
|
Loading…
Reference in New Issue
Block a user