Merge topic 'propagate_id3_to_shared_indices'

8de216c0 Propagate vtkm::Id3 scheduling down to the ThreadIndex classes.

Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Kenneth Moreland <kmorel@sandia.gov>
Merge-request: !240

vtkm/cont/cuda/internal/DeviceAdapterAlgorithmThrust.h

@@ -31,9 +31,9 @@
 #include <vtkm/cont/cuda/internal/MakeThrustIterator.h>
 #include <vtkm/cont/cuda/internal/ThrustExceptionHandler.h>
 
+#include <vtkm/exec/cuda/internal/WrappedOperators.h>
 #include <vtkm/exec/internal/ErrorMessageBuffer.h>
 #include <vtkm/exec/internal/WorkletInvokeFunctor.h>
-#include <vtkm/exec/cuda/internal/WrappedOperators.h>
 
 // Disable warnings we check vtkm for but Thrust does not.
 VTKM_THIRDPARTY_PRE_INCLUDE
@@ -98,18 +98,16 @@ template<class FunctorType>
 __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<typename T, typename BinaryOperationType >

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<typename T>
+    VTKM_EXEC_EXPORT void operator()(const T& index) const
     {
       this->Functor(index);
     }
@@ -309,17 +310,18 @@ public:
 
     DeviceAdapterAlgorithm<Device>::ScheduleKernel<Functor> 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];
         }
       }
 

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 <vtkm/cont/internal/IteratorFromArrayPortal.h>
-#include <vtkm/cont/tbb/internal/DeviceAdapterTagTBB.h>
-#include <vtkm/cont/tbb/internal/ArrayManagerExecutionTBB.h>
-#include <vtkm/exec/internal/ErrorMessageBuffer.h>
-#include <vtkm/Extent.h>
 #include <vtkm/cont/ArrayHandle.h>
 #include <vtkm/cont/ArrayHandleIndex.h>
 #include <vtkm/cont/ArrayHandleZip.h>
 #include <vtkm/cont/DeviceAdapterAlgorithm.h>
 #include <vtkm/cont/ErrorExecution.h>
 #include <vtkm/cont/internal/DeviceAdapterAlgorithmGeneral.h>
+#include <vtkm/cont/internal/IteratorFromArrayPortal.h>
+#include <vtkm/cont/tbb/internal/ArrayManagerExecutionTBB.h>
+#include <vtkm/cont/tbb/internal/DeviceAdapterTagTBB.h>
+#include <vtkm/exec/internal/ErrorMessageBuffer.h>
+#include <vtkm/Extent.h>
 
 VTKM_THIRDPARTY_PRE_INCLUDE
 #include <boost/type_traits/remove_reference.hpp>
@@ -466,17 +465,18 @@ private:
     void operator()(const ::tbb::blocked_range3d<vtkm::Id> &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];
             }
           }
         }

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;

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<typename FromTopology,
          vtkm::IdComponent Dimension>
 class ThreadIndicesTopologyMap<
     vtkm::exec::ConnectivityStructured<FromTopology,ToTopology,Dimension> >
-    : public vtkm::exec::arg::ThreadIndicesBasic
 {
   typedef vtkm::exec::arg::ThreadIndicesBasic Superclass;
   typedef vtkm::exec::ConnectivityStructured<FromTopology,ToTopology,Dimension>
@@ -160,7 +171,6 @@ public:
   template<typename Invocation>
   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<typename Invocation>
+  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;

vtkm/exec/internal/WorkletInvokeFunctor.h

@@ -50,8 +50,9 @@ public:
   }
 
   VTKM_SUPPRESS_EXEC_WARNINGS
+  template<typename T>
   VTKM_EXEC_EXPORT
-  void operator()(vtkm::Id index) const
+  void operator()(T index) const
   {
     detail::DoWorkletInvokeFunctor(this->Worklet,
                                    this->Invocation,

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:

vtkm/worklet/WorkletMapTopology.h

@@ -142,10 +142,10 @@ public:
 
   /// Topology map worklets use topology map indices.
   ///
-  template<typename Invocation>
+  template<typename T, typename Invocation>
   VTKM_EXEC_EXPORT
   vtkm::exec::arg::ThreadIndicesTopologyMap<typename Invocation::InputDomainType>
-  GetThreadIndices(vtkm::Id threadIndex, const Invocation &invocation) const
+  GetThreadIndices(const T& threadIndex, const Invocation &invocation) const
   {
     return vtkm::exec::arg::ThreadIndicesTopologyMap<typename Invocation::InputDomainType>(
           threadIndex, invocation);

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<typename T, typename Invocation>
+  VTKM_EXEC_EXPORT
+  vtkm::exec::arg::ThreadIndicesBasic
+  GetThreadIndices(const T& threadIndex, const Invocation &invocation) const
+  {
+    return vtkm::exec::arg::ThreadIndicesBasic(threadIndex, invocation);
+  }
 };
 
 }