diff --git a/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmThrust.h b/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmThrust.h index d8bbf1e57..5ba4a35f6 100644 --- a/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmThrust.h +++ b/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmThrust.h @@ -31,9 +31,9 @@ #include #include +#include #include #include -#include // Disable warnings we check vtkm for but Thrust does not. VTKM_THIRDPARTY_PRE_INCLUDE @@ -98,18 +98,16 @@ template __global__ void Schedule3DIndexKernel(FunctorType functor, dim3 size) { - const vtkm::Id x = blockIdx.x*blockDim.x + threadIdx.x; - const vtkm::Id y = blockIdx.y*blockDim.y + threadIdx.y; - const vtkm::Id z = blockIdx.z*blockDim.z + threadIdx.z; - - if (x >= size.x || y >= size.y || z >= size.z) + const vtkm::Id3 index( + blockIdx.x*blockDim.x + threadIdx.x, + blockIdx.y*blockDim.y + threadIdx.y, + blockIdx.z*blockDim.z + threadIdx.z + ); + if (index[0] >= size.x || index[1] >= size.y || index[2] >= size.z) { return; } - - //now convert back to flat memory - const int idx = x + size.x*(y + size.y*z); - functor( idx ); + functor( index ); } template diff --git a/vtkm/cont/internal/DeviceAdapterAlgorithmSerial.h b/vtkm/cont/internal/DeviceAdapterAlgorithmSerial.h index a36f31a5e..5b631e372 100644 --- a/vtkm/cont/internal/DeviceAdapterAlgorithmSerial.h +++ b/vtkm/cont/internal/DeviceAdapterAlgorithmSerial.h @@ -258,7 +258,8 @@ private: : Functor(functor) { } //needed for when calling from schedule on a range - VTKM_EXEC_EXPORT void operator()(vtkm::Id index) const + template + VTKM_EXEC_EXPORT void operator()(const T& index) const { this->Functor(index); } @@ -309,17 +310,18 @@ public: DeviceAdapterAlgorithm::ScheduleKernel kernel(functor); - //use a const variable to hint to compiler this doesn't change + vtkm::Id3 index; for(vtkm::Id k=0; k < rangeMax[2]; ++k) { - vtkm::Id index = k * rangeMax[1] * rangeMax[0]; + index[2] = k; for(vtkm::Id j=0; j < rangeMax[1]; ++j) { + index[1] = j; for(vtkm::Id i=0; i < rangeMax[0]; ++i) { - kernel( index + i ); + index[0] = i; + kernel( index ); } - index += rangeMax[0]; } } diff --git a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h index 9d0f93a11..6b78c1e9a 100644 --- a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h +++ b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h @@ -20,18 +20,17 @@ #ifndef vtk_m_cont_tbb_internal_DeviceAdapterAlgorithmTBB_h #define vtk_m_cont_tbb_internal_DeviceAdapterAlgorithmTBB_h - -#include -#include -#include -#include -#include #include #include #include #include #include #include +#include +#include +#include +#include +#include VTKM_THIRDPARTY_PRE_INCLUDE #include @@ -466,17 +465,18 @@ private: void operator()(const ::tbb::blocked_range3d &range) const { try { + vtkm::Id3 index; for( vtkm::Id k=range.pages().begin(); k!=range.pages().end(); ++k) { - vtkm::Id index = k * this->Dims[1] * this->Dims[0]; - index += range.rows().begin() * this->Dims[0]; + index[2] = k; for( vtkm::Id j=range.rows().begin(); j!=range.rows().end(); ++j) { + index[1] = j; for( vtkm::Id i=range.cols().begin(); i!=range.cols().end(); ++i) { - this->Functor(index + i); + index[0] = i; + this->Functor( index ); } - index += this->Dims[0]; } } } diff --git a/vtkm/cont/testing/TestingDeviceAdapter.h b/vtkm/cont/testing/TestingDeviceAdapter.h index ce814a91c..dfd7fae6e 100644 --- a/vtkm/cont/testing/TestingDeviceAdapter.h +++ b/vtkm/cont/testing/TestingDeviceAdapter.h @@ -129,17 +129,30 @@ public: struct ClearArrayKernel { VTKM_CONT_EXPORT - ClearArrayKernel(const IdPortalType &array) : Array(array) { } + ClearArrayKernel(const IdPortalType &array) : Array(array), Dims() { } + + VTKM_CONT_EXPORT + ClearArrayKernel(const IdPortalType &array, + const vtkm::Id3& dims) : Array(array), Dims(dims) { } VTKM_EXEC_EXPORT void operator()(vtkm::Id index) const { this->Array.Set(index, OFFSET); } + VTKM_EXEC_EXPORT void operator()(vtkm::Id3 index) const + { + //convert from id3 to id + vtkm::Id flatIndex = + index[0]+ this->Dims[0]*(index[1]+ this->Dims[1]*index[2]); + this->operator()(flatIndex); + } + VTKM_CONT_EXPORT void SetErrorMessageBuffer( const vtkm::exec::internal::ErrorMessageBuffer &) { } IdPortalType Array; + vtkm::Id3 Dims; }; struct ClearArrayMapKernel //: public vtkm::exec::WorkletMapField @@ -158,17 +171,31 @@ public: struct AddArrayKernel { VTKM_CONT_EXPORT - AddArrayKernel(const IdPortalType &array) : Array(array) { } + AddArrayKernel(const IdPortalType &array) : Array(array), Dims() { } + + VTKM_CONT_EXPORT + AddArrayKernel(const IdPortalType &array, + const vtkm::Id3& dims) : Array(array), Dims(dims) { } + VTKM_EXEC_EXPORT void operator()(vtkm::Id index) const { this->Array.Set(index, this->Array.Get(index) + index); } + VTKM_EXEC_EXPORT void operator()(vtkm::Id3 index) const + { + //convert from id3 to id + vtkm::Id flatIndex = + index[0]+ this->Dims[0]*(index[1]+ this->Dims[1]*index[2]); + this->operator()(flatIndex); + } + VTKM_CONT_EXPORT void SetErrorMessageBuffer( const vtkm::exec::internal::ErrorMessageBuffer &) { } IdPortalType Array; + vtkm::Id3 Dims; }; struct OneErrorKernel @@ -486,11 +513,11 @@ private: std::cout << "Running clear." << std::endl; Algorithm::Schedule( ClearArrayKernel(manager.PrepareForOutput( - DIM_SIZE * DIM_SIZE * DIM_SIZE)), + DIM_SIZE * DIM_SIZE * DIM_SIZE), maxRange), maxRange); std::cout << "Running add." << std::endl; - Algorithm::Schedule(AddArrayKernel(manager.PrepareForInPlace(false)), + Algorithm::Schedule(AddArrayKernel(manager.PrepareForInPlace(false), maxRange), maxRange); std::cout << "Checking results." << std::endl; diff --git a/vtkm/exec/arg/ThreadIndicesTopologyMap.h b/vtkm/exec/arg/ThreadIndicesTopologyMap.h index 06633e2df..7a18c0d32 100644 --- a/vtkm/exec/arg/ThreadIndicesTopologyMap.h +++ b/vtkm/exec/arg/ThreadIndicesTopologyMap.h @@ -138,6 +138,18 @@ vtkm::Id3 InflateTo3D(vtkm::Id index) return vtkm::Id3(index, 0, 0); } +VTKM_EXEC_EXPORT +vtkm::Id3 Deflate(const vtkm::Id3& index, vtkm::Id3) +{ + return index; +} + +VTKM_EXEC_EXPORT +vtkm::Id2 Deflate(const vtkm::Id3& index, vtkm::Id2) +{ + return vtkm::Id2(index[0], index[1]); +} + } // namespace detail // Specialization for structured connectivity types. @@ -146,7 +158,6 @@ template class ThreadIndicesTopologyMap< vtkm::exec::ConnectivityStructured > - : public vtkm::exec::arg::ThreadIndicesBasic { typedef vtkm::exec::arg::ThreadIndicesBasic Superclass; typedef vtkm::exec::ConnectivityStructured @@ -160,7 +171,6 @@ public: template VTKM_EXEC_EXPORT ThreadIndicesTopologyMap(vtkm::Id threadIndex, const Invocation &invocation) - : Superclass(threadIndex, invocation) { // The connectivity is stored in the invocation parameter at the given // input domain index. If this class is being used correctly, the type @@ -169,9 +179,33 @@ public: // set its input domain incorrectly. const ConnectivityType &connectivity = invocation.GetInputDomain(); - this->LogicalIndex = connectivity.FlatToLogicalToIndex(this->GetIndex()); - this->IndicesFrom = connectivity.GetIndices(this->LogicalIndex); - this->CellShape = connectivity.GetCellShape(this->GetIndex()); + const LogicalIndexType logicalIndex = connectivity.FlatToLogicalToIndex(threadIndex); + + this->Index = threadIndex; + this->LogicalIndex = logicalIndex; + this->IndicesFrom = connectivity.GetIndices(logicalIndex); + this->CellShape = connectivity.GetCellShape(threadIndex); + } + + template + VTKM_EXEC_EXPORT + ThreadIndicesTopologyMap(vtkm::Id3 threadIndex, const Invocation &invocation) + { + // The connectivity is stored in the invocation parameter at the given + // input domain index. If this class is being used correctly, the type + // of the domain will match the connectivity type used here. If there is + // a compile error here about a type mismatch, chances are a worklet has + // set its input domain incorrectly. + const ConnectivityType &connectivity = invocation.GetInputDomain(); + + const LogicalIndexType logicalIndex = detail::Deflate(threadIndex, LogicalIndexType()); + const vtkm::Id index = connectivity.LogicalToFlatToIndex(logicalIndex); + + + this->Index = index; + this->LogicalIndex = logicalIndex; + this->IndicesFrom = connectivity.GetIndices(logicalIndex); + this->CellShape = connectivity.GetCellShape(index); } /// \brief The logical index into the input domain. @@ -185,6 +219,18 @@ public: return this->LogicalIndex; } + /// \brief The index into the input domain. + /// + /// This index refers to the input element (array value, cell, etc.) that + /// this thread is being invoked for. This is the typical index used during + /// fetches. + /// + VTKM_EXEC_EXPORT + vtkm::Id GetIndex() const + { + return this->Index; + } + /// \brief The 3D index into the input domain. /// /// Overloads the implementation in the base class to return the 3D index @@ -217,6 +263,7 @@ public: CellShapeTag GetCellShape() const { return this->CellShape; } private: + vtkm::Id Index; LogicalIndexType LogicalIndex; IndicesFromType IndicesFrom; CellShapeTag CellShape; diff --git a/vtkm/exec/internal/WorkletInvokeFunctor.h b/vtkm/exec/internal/WorkletInvokeFunctor.h index ce0ef9726..4cb7e8f11 100644 --- a/vtkm/exec/internal/WorkletInvokeFunctor.h +++ b/vtkm/exec/internal/WorkletInvokeFunctor.h @@ -50,8 +50,9 @@ public: } VTKM_SUPPRESS_EXEC_WARNINGS + template VTKM_EXEC_EXPORT - void operator()(vtkm::Id index) const + void operator()(T index) const { detail::DoWorkletInvokeFunctor(this->Worklet, this->Invocation, diff --git a/vtkm/internal/ConnectivityStructuredInternals.h b/vtkm/internal/ConnectivityStructuredInternals.h index a206cadfd..94279cbab 100644 --- a/vtkm/internal/ConnectivityStructuredInternals.h +++ b/vtkm/internal/ConnectivityStructuredInternals.h @@ -340,15 +340,15 @@ public: } VTKM_EXEC_CONT_EXPORT - const vtkm::Id3 GetPointDimensions() const + const vtkm::Id3& GetPointDimensions() const { return this->PointDimensions; } VTKM_EXEC_CONT_EXPORT - vtkm::Id3 GetCellDimensions() const + const vtkm::Id3& GetCellDimensions() const { - return this->PointDimensions - vtkm::Id3(1); + return this->CellDimensions; } VTKM_EXEC_CONT_EXPORT @@ -359,11 +359,11 @@ public: //returns an id3 to signal what kind of scheduling to use VTKM_EXEC_CONT_EXPORT - vtkm::Id3 GetSchedulingRange(vtkm::TopologyElementTagCell) const { + const vtkm::Id3& GetSchedulingRange(vtkm::TopologyElementTagCell) const { return this->GetCellDimensions(); } VTKM_EXEC_CONT_EXPORT - vtkm::Id3 GetSchedulingRange(vtkm::TopologyElementTagPoint) const { + const vtkm::Id3& GetSchedulingRange(vtkm::TopologyElementTagPoint) const { return this->GetPointDimensions(); } @@ -535,10 +535,9 @@ public: VTKM_EXEC_CONT_EXPORT vtkm::Id LogicalToFlatCellIndex(const vtkm::Id3 &logicalCellIndex) const { - vtkm::Id3 cellDimensions = this->GetCellDimensions(); return logicalCellIndex[0] - + cellDimensions[0]*(logicalCellIndex[1] - + cellDimensions[1]*logicalCellIndex[2]); + + this->CellDimensions[0]*(logicalCellIndex[1] + + this->CellDimensions[1]*logicalCellIndex[2]); } private: diff --git a/vtkm/worklet/WorkletMapTopology.h b/vtkm/worklet/WorkletMapTopology.h index 5fbef73c0..531f133ca 100644 --- a/vtkm/worklet/WorkletMapTopology.h +++ b/vtkm/worklet/WorkletMapTopology.h @@ -142,10 +142,10 @@ public: /// Topology map worklets use topology map indices. /// - template + template VTKM_EXEC_EXPORT vtkm::exec::arg::ThreadIndicesTopologyMap - GetThreadIndices(vtkm::Id threadIndex, const Invocation &invocation) const + GetThreadIndices(const T& threadIndex, const Invocation &invocation) const { return vtkm::exec::arg::ThreadIndicesTopologyMap( threadIndex, invocation); diff --git a/vtkm/worklet/internal/WorkletBase.h b/vtkm/worklet/internal/WorkletBase.h index 41f0143b5..01c79c1e9 100644 --- a/vtkm/worklet/internal/WorkletBase.h +++ b/vtkm/worklet/internal/WorkletBase.h @@ -167,6 +167,19 @@ public: { return vtkm::exec::arg::ThreadIndicesBasic(threadIndex, invocation); } + + /// \brief Creates a \c ThreadIndices object. + /// + /// Worklet types can add additional indices by returning different object + /// types. + /// + template + VTKM_EXEC_EXPORT + vtkm::exec::arg::ThreadIndicesBasic + GetThreadIndices(const T& threadIndex, const Invocation &invocation) const + { + return vtkm::exec::arg::ThreadIndicesBasic(threadIndex, invocation); + } }; }