Shuenhoy/vtk-m
1
0
Fork 0
mirror of https://gitlab.kitware.com/vtk/vtk-m synced 2024-09-08 21:33:55 +00:00
Code Issues 2 Actions Packages Projects Releases Wiki Activity

Initial infrastructure to allow WorkletMapField to have 3D scheduling

Browse Source
This commit is contained in:
Robert Maynard 2020-01-06 17:06:51 -05:00
parent 37d55a4ada
commit 1f1688483e
20 changed files with 509 additions and 63 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

10
vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h
View File

@ -182,7 +182,7 @@ __global__ void TaskStrided1DLaunch(TaskType task, vtkm::Id size)
}
template <typename TaskType>
__global__ void TaskStrided3DLaunch(TaskType task, dim3 size)
__global__ void TaskStrided3DLaunch(TaskType task, vtkm::Id3 size)
{
//This is the 3D version of executing in a grid-stride manner
const dim3 start(blockIdx.x * blockDim.x + threadIdx.x,
@ -190,11 +190,11 @@ __global__ void TaskStrided3DLaunch(TaskType task, dim3 size)
blockIdx.z * blockDim.z + threadIdx.z);
const dim3 inc(blockDim.x * gridDim.x, blockDim.y * gridDim.y, blockDim.z * gridDim.z);
for (vtkm::Id k = start.z; k < size.z; k += inc.z)
for (vtkm::Id k = start.z; k < size[2]; k += inc.z)
{
for (vtkm::Id j = start.y; j < size.y; j += inc.y)
for (vtkm::Id j = start.y; j < size[1]; j += inc.y)
{
task(start.x, size.x, inc.x, j, k);
task(size, start.x, size[0], inc.x, j, k);
}
}
}
@ -1664,7 +1664,7 @@ public:
#endif
cuda::internal::TaskStrided3DLaunch<<<blocks, threadsPerBlock, 0, cudaStreamPerThread>>>(
functor, ranges);
functor, rangeMax);
}
template <class Functor>

2
vtkm/cont/openmp/internal/DeviceAdapterAlgorithmOpenMP.cxx
View File

@ -132,7 +132,7 @@ void DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagOpenMP>::ScheduleTask(
{
for (vtkm::Id j = startIJK[1]; j < endIJK[1]; ++j)
{
functor(startIJK[0], endIJK[0], j, k);
functor(size, startIJK[0], endIJK[0], j, k);
}
}
}

2
vtkm/cont/serial/internal/DeviceAdapterAlgorithmSerial.cxx
View File

@ -58,7 +58,7 @@ void DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagSerial>::ScheduleTask(
{
for (vtkm::Id j = 0; j < size[1]; ++j)
{
functor(0, size[0], j, k);
functor(size, 0, size[0], j, k);
}
}

2
vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.cxx
View File

@ -65,7 +65,7 @@ void DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagTBB>::ScheduleTask(
{
const vtkm::Id start = r.cols().begin();
const vtkm::Id end = r.cols().end();
functor(start, end, j, k);
functor(size, start, end, j, k);
}
}
});

1
vtkm/exec/arg/CMakeLists.txt
View File

@ -31,6 +31,7 @@ set(headers
OutputIndex.h
ThreadIndices.h
ThreadIndicesBasic.h
ThreadIndicesBasic3D.h
ThreadIndicesExtrude.h
ThreadIndicesPointNeighborhood.h
ThreadIndicesReduceByKey.h

64
vtkm/exec/arg/ThreadIndicesBasic3D.h Normal file
View File

@ -0,0 +1,64 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
//
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//============================================================================
#ifndef vtk_m_exec_arg_ThreadIndicesBasic3D_h
#define vtk_m_exec_arg_ThreadIndicesBasic3D_h
#include <vtkm/exec/arg/ThreadIndicesBasic.h>
namespace vtkm
{
namespace exec
{
namespace arg
{
/// \brief Container for 3D thread indices in a worklet invocation
///
/// During the execution of a worklet function in an execution environment
/// thread, VTK-m has to manage several indices. To simplify this management
/// and to provide a single place to store them (so that they do not have to be
/// recomputed), \c WorkletInvokeFunctor creates a \c ThreadIndices object.
/// This object gets passed to \c Fetch operations to help them load data.
///
///
class ThreadIndicesBasic3D : public vtkm::exec::arg::ThreadIndicesBasic
{
public:
VTKM_EXEC
ThreadIndicesBasic3D(const vtkm::Id3& threadIndex3D,
vtkm::Id threadIndex1D,
vtkm::Id inIndex,
vtkm::IdComponent visitIndex,
vtkm::Id outIndex,
vtkm::Id globalThreadIndexOffset)
: ThreadIndicesBasic(threadIndex1D, inIndex, visitIndex, outIndex, globalThreadIndexOffset)
, ThreadIndex3D(threadIndex3D)
{
}
/// \brief The 3D index into the input domain.
///
/// This index refers to the input element (array value, cell, etc.) that
/// this thread is being invoked for. If the input domain has 2 or 3
/// dimensional indexing, this result will preserve that. If the domain
/// indexing is just one dimensional, the result will have the index in the
/// first component with the remaining components set to 0.
///
VTKM_EXEC
vtkm::Id3 GetInputIndex3D() const { return this->ThreadIndex3D; }
private:
vtkm::Id3 ThreadIndex3D;
};
}
}
} // namespace vtkm::exec::arg
#endif //vtk_m_exec_arg_ThreadIndicesBasic3D_h

24
vtkm/exec/arg/ThreadIndicesExtrude.h
View File

@ -55,19 +55,19 @@ public:
}
VTKM_SUPPRESS_EXEC_WARNINGS
VTKM_EXEC
ThreadIndicesTopologyMap(const vtkm::Id3& threadIndex,
ThreadIndicesTopologyMap(const vtkm::Id3& threadIndex3D,
vtkm::Id threadIndex1D,
const ConnectivityType& connectivity,
vtkm::Id globalThreadIndexOffset = 0)
{
// We currently only support multidimensional indices on one-to-one input-
// to-output mappings. (We don't have a use case otherwise.)
// That is why we treat teh threadIndex as also the inputIndex and outputIndex
const LogicalIndexType logicalIndex = detail::Deflate(threadIndex, LogicalIndexType());
const vtkm::Id index = connectivity.LogicalToFlatToIndex(logicalIndex);
auto logicalIndex = detail::Deflate(threadIndex3D, LogicalIndexType());
this->ThreadIndex = index;
this->InputIndex = index;
this->OutputIndex = index;
this->ThreadIndex = threadIndex1D;
this->InputIndex = threadIndex1D;
this->OutputIndex = threadIndex1D;
this->VisitIndex = 0;
this->LogicalIndex = logicalIndex;
this->IndicesIncident = connectivity.GetIndices(logicalIndex);
@ -196,17 +196,17 @@ public:
this->GlobalThreadIndexOffset = globalThreadIndexOffset;
}
ThreadIndicesTopologyMap(const vtkm::Id3& threadIndex,
ThreadIndicesTopologyMap(const vtkm::Id3& threadIndex3D,
vtkm::Id threadIndex1D,
const ConnectivityType& connectivity,
vtkm::Id globalThreadIndexOffset = 0)
{
const LogicalIndexType logicalIndex = detail::Deflate(threadIndex, LogicalIndexType());
const vtkm::Id index = connectivity.LogicalToFlatToIndex(logicalIndex);
const LogicalIndexType logicalIndex = detail::Deflate(threadIndex3D, LogicalIndexType());
this->ThreadIndex = index;
this->InputIndex = index;
this->OutputIndex = index;
this->ThreadIndex = threadIndex1D;
this->InputIndex = threadIndex1D;
this->OutputIndex = threadIndex1D;
this->VisitIndex = 0;
this->LogicalIndex = logicalIndex;
this->IndicesIncident = connectivity.GetIndices(logicalIndex);

18
vtkm/exec/arg/ThreadIndicesPointNeighborhood.h
View File

@ -57,23 +57,19 @@ class ThreadIndicesPointNeighborhood
public:
template <vtkm::IdComponent Dimension>
VTKM_EXEC ThreadIndicesPointNeighborhood(
const vtkm::Id3& outIndex,
const vtkm::Id3& threadIndex3D,
vtkm::Id threadIndex1D,
const vtkm::exec::ConnectivityStructured<vtkm::TopologyElementTagPoint,
vtkm::TopologyElementTagCell,
Dimension>& connectivity,
vtkm::Id globalThreadIndexOffset = 0)
: State(outIndex, detail::To3D(connectivity.GetPointDimensions()))
: State(threadIndex3D, detail::To3D(connectivity.GetPointDimensions()))
, ThreadIndex(threadIndex1D)
, InputIndex(threadIndex1D)
, OutputIndex(threadIndex1D)
, VisitIndex(0)
, GlobalThreadIndexOffset(globalThreadIndexOffset)
{
using ConnectivityType = vtkm::exec::ConnectivityStructured<vtkm::TopologyElementTagPoint,
vtkm::TopologyElementTagCell,
Dimension>;
using ConnRangeType = typename ConnectivityType::SchedulingRangeType;
const ConnRangeType index = detail::Deflate(outIndex, ConnRangeType());
this->ThreadIndex = connectivity.LogicalToFlatToIndex(index);
this->InputIndex = this->ThreadIndex;
this->VisitIndex = 0;
this->OutputIndex = this->ThreadIndex;
}
template <vtkm::IdComponent Dimension>

15
vtkm/exec/arg/ThreadIndicesTopologyMap.h
View File

@ -174,23 +174,22 @@ public:
this->GlobalThreadIndexOffset = globalThreadIndexOffset;
}
VTKM_EXEC ThreadIndicesTopologyMap(const vtkm::Id3& threadIndex,
VTKM_EXEC ThreadIndicesTopologyMap(const vtkm::Id3& threadIndex3D,
vtkm::Id threadIndex1D,
const ConnectivityType& connectivity,
const vtkm::Id globalThreadIndexOffset = 0)
{
// We currently only support multidimensional indices on one-to-one input-
// to-output mappings. (We don't have a use case otherwise.)
// That is why we treat teh threadIndex as also the inputIndex and outputIndex
const LogicalIndexType logicalIndex = detail::Deflate(threadIndex, LogicalIndexType());
const vtkm::Id index = connectivity.LogicalToFlatToIndex(logicalIndex);
this->ThreadIndex = index;
this->InputIndex = index;
this->OutputIndex = index;
auto logicalIndex = detail::Deflate(threadIndex3D, LogicalIndexType());
this->ThreadIndex = threadIndex1D;
this->InputIndex = threadIndex1D;
this->OutputIndex = threadIndex1D;
this->VisitIndex = 0;
this->LogicalIndex = logicalIndex;
this->IndicesIncident = connectivity.GetIndices(logicalIndex);
this->CellShape = connectivity.GetCellShape(index);
this->CellShape = connectivity.GetCellShape(threadIndex1D);
this->GlobalThreadIndexOffset = globalThreadIndexOffset;
}

1
vtkm/exec/arg/testing/CMakeLists.txt
View File

@ -16,6 +16,7 @@ vtkm_declare_headers(${headers})
set(unit_tests
UnitTestExecutionSignatureTag.cxx
UnitTestFetchArrayDirectIn.cxx
UnitTestFetchArrayDirectIn3d.cxx
UnitTestFetchArrayDirectInOut.cxx
UnitTestFetchArrayDirectOut.cxx
UnitTestFetchArrayNeighborhoodIn.cxx

139
vtkm/exec/arg/testing/UnitTestFetchArrayDirectIn3d.cxx Normal file
View File

@ -0,0 +1,139 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
//
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//============================================================================
#include <vtkm/exec/arg/FetchTagArrayDirectIn.h>
#include <vtkm/exec/arg/ThreadIndicesBasic3D.h>
#include <vtkm/testing/Testing.h>
namespace
{
static constexpr vtkm::Id3 ARRAY_SIZE = { 10, 10, 3 };
template <typename T>
struct TestPortal
{
using ValueType = T;
VTKM_EXEC_CONT
vtkm::Id GetNumberOfValues() const { return vtkm::ReduceProduct(ARRAY_SIZE); }
VTKM_EXEC_CONT
ValueType Get(vtkm::Id3 index) const
{
VTKM_TEST_ASSERT(index[0] >= 0, "Bad portal index.");
VTKM_TEST_ASSERT(index[1] >= 0, "Bad portal index.");
VTKM_TEST_ASSERT(index[2] >= 0, "Bad portal index.");
VTKM_TEST_ASSERT(index[0] < ARRAY_SIZE[0], "Bad portal index.");
VTKM_TEST_ASSERT(index[1] < ARRAY_SIZE[1], "Bad portal index.");
VTKM_TEST_ASSERT(index[2] < ARRAY_SIZE[2], "Bad portal index.");
auto flatIndex = index[0] + ARRAY_SIZE[0] * (index[1] + ARRAY_SIZE[1] * index[2]);
return TestValue(flatIndex, ValueType());
}
};
}
namespace vtkm
{
namespace exec
{
namespace arg
{
// Fetch for ArrayPortalTex3D when being used for Loads
template <typename T>
struct Fetch<vtkm::exec::arg::FetchTagArrayDirectIn,
vtkm::exec::arg::AspectTagDefault,
vtkm::exec::arg::ThreadIndicesBasic3D,
TestPortal<T>>
{
using ValueType = T;
using PortalType = const TestPortal<T>&;
using ThreadIndicesType = vtkm::exec::arg::ThreadIndicesBasic3D;
VTKM_EXEC
ValueType Load(const ThreadIndicesType& indices, PortalType field) const
{
return field.Get(indices.GetInputIndex3D());
}
VTKM_EXEC
void Store(const ThreadIndicesType&, PortalType, ValueType) const {}
};
}
}
}
namespace
{
template <typename T>
struct FetchArrayDirectIn3DTests
{
void operator()()
{
TestPortal<T> execObject;
using FetchType = vtkm::exec::arg::Fetch<vtkm::exec::arg::FetchTagArrayDirectIn,
vtkm::exec::arg::AspectTagDefault,
vtkm::exec::arg::ThreadIndicesBasic3D,
TestPortal<T>>;
FetchType fetch;
vtkm::Id index1d = 0;
vtkm::Id3 index3d = { 0, 0, 0 };
for (vtkm::Id k = 0; k < ARRAY_SIZE[2]; ++k)
{
index3d[2] = k;
for (vtkm::Id j = 0; j < ARRAY_SIZE[1]; ++j)
{
index3d[1] = j;
for (vtkm::Id i = 0; i < ARRAY_SIZE[0]; i++, index1d++)
{
index3d[0] = i;
vtkm::exec::arg::ThreadIndicesBasic3D indices(index3d, index1d, index1d, 0, index1d, 0);
T value = fetch.Load(indices, execObject);
VTKM_TEST_ASSERT(test_equal(value, TestValue(index1d, T())),
"Got invalid value from Load.");
value = T(T(2) * value);
// This should be a no-op, but we should be able to call it.
fetch.Store(indices, execObject, value);
}
}
}
}
};
struct TryType
{
template <typename T>
void operator()(T) const
{
FetchArrayDirectIn3DTests<T>()();
}
};
void TestExecObjectFetch3D()
{
vtkm::testing::Testing::TryTypes(TryType());
}
} // anonymous namespace
int UnitTestFetchArrayDirectIn3d(int argc, char* argv[])
{
return vtkm::testing::Testing::Run(TestExecObjectFetch3D, argc, argv);
}

2
vtkm/exec/arg/testing/UnitTestFetchArrayNeighborhoodIn.cxx
View File

@ -123,7 +123,7 @@ struct FetchArrayNeighborhoodInTests
for (vtkm::Id i = 0; i < POINT_DIMS[0]; i++, index++)
{
index3d[0] = i;
vtkm::exec::arg::ThreadIndicesPointNeighborhood indices(index3d, connectivity);
vtkm::exec::arg::ThreadIndicesPointNeighborhood indices(index3d, index, connectivity);
auto neighbors = fetch.Load(indices, execObject);

20
vtkm/exec/cuda/internal/TaskStrided.h
View File

@ -135,17 +135,24 @@ public:
}
VTKM_EXEC
void operator()(vtkm::Id start, vtkm::Id end, vtkm::Id inc, vtkm::Id j, vtkm::Id k) const
void operator()(const vtkm::Id3& size,
vtkm::Id start,
vtkm::Id end,
vtkm::Id inc,
vtkm::Id j,
vtkm::Id k) const
{
vtkm::Id3 index(start, j, k);
for (vtkm::Id i = start; i < end; i += inc)
auto threadIndex1D = index[0] + size[0] * (index[1] + size[1] * index[2]);
for (vtkm::Id i = start; i < end; i += inc, threadIndex1D += inc)
{
index[0] = i;
//Todo: rename this function to DoTaskInvokeWorklet
vtkm::exec::internal::detail::DoWorkletInvokeFunctor(
this->Worklet,
this->Invocation,
this->Worklet.GetThreadIndices(index,
this->Worklet.GetThreadIndices(threadIndex1D,
index,
this->Invocation.OutputToInputMap,
this->Invocation.VisitArray,
this->Invocation.ThreadToOutputMap,
@ -178,7 +185,12 @@ public:
}
VTKM_EXEC
void operator()(vtkm::Id start, vtkm::Id end, vtkm::Id inc, vtkm::Id j, vtkm::Id k) const
void operator()(const vtkm::Id3& size,
vtkm::Id start,
vtkm::Id end,
vtkm::Id inc,
vtkm::Id j,
vtkm::Id k) const
{
vtkm::Id3 index(start, j, k);
for (vtkm::Id i = start; i < end; i += inc)

10
vtkm/exec/internal/testing/TestingTaskTiling.h
View File

@ -222,6 +222,7 @@ struct TestWorkletProxy : vtkm::exec::FunctorBase
typename InputDomainType,
typename G>
VTKM_EXEC vtkm::exec::arg::ThreadIndicesBasic GetThreadIndices(
const vtkm::Id3& vtkmNotUsed(iterationSpace),
const vtkm::Id3& threadIndex,
const OutToInArrayType& outToIn,
const VisitArrayType& visit,
@ -271,6 +272,7 @@ struct TestWorkletErrorProxy : vtkm::exec::FunctorBase
typename InputDomainType,
typename G>
VTKM_EXEC vtkm::exec::arg::ThreadIndicesBasic GetThreadIndices(
const vtkm::Id3& vtkmNotUsed(iterationSpace),
const vtkm::Id3& threadIndex,
const OutToInArrayType& outToIn,
const VisitArrayType& visit,
@ -403,8 +405,8 @@ void Test3DNormalTaskTilingInvoke()
for (vtkm::Id j = 0; j < 8; j += 2)
{
//verify that order is not required
task1(0, 8, j + 1, k);
task1(0, 8, j, k);
task1(vtkm::Id3{ 8, 8, 8 }, 0, 8, j + 1, k);
task1(vtkm::Id3{ 8, 8, 8 }, 0, 8, j, k);
}
}
@ -431,7 +433,7 @@ void Test3DNormalTaskTilingInvoke()
{
for (vtkm::Id k = 0; k < 8; ++k)
{
task2(i, i + 1, j, k);
task2(vtkm::Id3{ 8, 8, 8 }, i, i + 1, j, k);
}
}
}
@ -471,7 +473,7 @@ void Test3DErrorTaskTilingInvoke()
{
for (vtkm::Id j = 0; j < 8; ++j)
{
task1(0, 8, j, k);
task1(vtkm::Id3{ 8, 8, 8 }, 0, 8, j, k);
}
}

28
vtkm/exec/serial/internal/TaskTiling.h
View File

@ -83,6 +83,7 @@ template <typename WType, typename IType>
VTKM_NEVER_EXPORT void TaskTiling3DExecute(void* w,
void* const v,
vtkm::Id globalIndexOffset,
const vtkm::Id3& maxSize,
vtkm::Id istart,
vtkm::Id iend,
vtkm::Id j,
@ -95,14 +96,16 @@ VTKM_NEVER_EXPORT void TaskTiling3DExecute(void* w,
InvocationType const* const invocation = static_cast<InvocationType*>(v);
vtkm::Id3 index(istart, j, k);
for (vtkm::Id i = istart; i < iend; ++i)
auto threadIndex1D = index[0] + maxSize[0] * (index[1] + maxSize[1] * index[2]);
for (vtkm::Id i = istart; i < iend; ++i, ++threadIndex1D)
{
index[0] = i;
//Todo: rename this function to DoTaskInvokeWorklet
vtkm::exec::internal::detail::DoWorkletInvokeFunctor(
*worklet,
*invocation,
worklet->GetThreadIndices(index,
worklet->GetThreadIndices(threadIndex1D,
index,
invocation->OutputToInputMap,
invocation->VisitArray,
invocation->ThreadToOutputMap,
@ -114,7 +117,8 @@ VTKM_NEVER_EXPORT void TaskTiling3DExecute(void* w,
template <typename FType>
VTKM_NEVER_EXPORT void FunctorTiling3DExecute(void* f,
void* const,
vtkm::Id,
vtkm::Id vtkmNotUsed(globalIndexOffset),
const vtkm::Id3& vtkmNotUsed(maxSize),
vtkm::Id istart,
vtkm::Id iend,
vtkm::Id j,
@ -295,18 +299,28 @@ public:
this->SetErrorBufferFunction(this->Worklet, buffer);
}
void operator()(vtkm::Id istart, vtkm::Id iend, vtkm::Id j, vtkm::Id k) const
void operator()(const vtkm::Id3& maxSize,
vtkm::Id istart,
vtkm::Id iend,
vtkm::Id j,
vtkm::Id k) const
{
this->ExecuteFunction(
this->Worklet, this->Invocation, this->GlobalIndexOffset, istart, iend, j, k);
this->Worklet, this->Invocation, this->GlobalIndexOffset, maxSize, istart, iend, j, k);
}
protected:
void* Worklet;
void* Invocation;
using ExecuteSignature =
void (*)(void*, void* const, vtkm::Id, vtkm::Id, vtkm::Id, vtkm::Id, vtkm::Id);
using ExecuteSignature = void (*)(void*,
void* const,
vtkm::Id,
const vtkm::Id3&,
vtkm::Id,
vtkm::Id,
vtkm::Id,
vtkm::Id);
ExecuteSignature ExecuteFunction;
using SetErrorBufferSignature = void (*)(void*, const vtkm::exec::internal::ErrorMessageBuffer&);

5
vtkm/worklet/WorkletMapTopology.h
View File

@ -185,7 +185,8 @@ public:
typename ThreadToOutArrayType,
typename InputDomainType>
VTKM_EXEC vtkm::exec::arg::ThreadIndicesTopologyMap<InputDomainType> GetThreadIndices(
const vtkm::Id3& threadIndex,
vtkm::Id threadIndex1D,
const vtkm::Id3& threadIndex3D,
const OutToInArrayType& vtkmNotUsed(outToIn),
const VisitArrayType& vtkmNotUsed(visit),
const ThreadToOutArrayType& vtkmNotUsed(threadToOut),
@ -200,7 +201,7 @@ public:
"Scheduling on 3D topologies only works with default MaskNone.");
return vtkm::exec::arg::ThreadIndicesTopologyMap<InputDomainType>(
threadIndex, connectivity, globalThreadIndexOffset);
threadIndex3D, threadIndex1D, connectivity, globalThreadIndexOffset);
}
};

5
vtkm/worklet/WorkletPointNeighborhood.h
View File

@ -203,7 +203,8 @@ public:
typename ThreadToOutArrayType,
typename InputDomainType>
VTKM_EXEC vtkm::exec::arg::ThreadIndicesPointNeighborhood GetThreadIndices(
const vtkm::Id3& threadIndex,
vtkm::Id threadIndex1D,
const vtkm::Id3& threadIndex3D,
const OutToInArrayType& vtkmNotUsed(outToIn),
const VisitArrayType& vtkmNotUsed(visit),
const ThreadToOutArrayType& vtkmNotUsed(threadToOut),
@ -218,7 +219,7 @@ public:
"Scheduling on 3D topologies only works with default MaskNone.");
return vtkm::exec::arg::ThreadIndicesPointNeighborhood(
threadIndex, connectivity, globalThreadIndexOffset);
threadIndex3D, threadIndex1D, connectivity, globalThreadIndexOffset);
}
};
}

36
vtkm/worklet/internal/WorkletBase.h
View File

@ -20,6 +20,7 @@
#include <vtkm/exec/arg/OutputIndex.h>
#include <vtkm/exec/arg/ThreadIndices.h>
#include <vtkm/exec/arg/ThreadIndicesBasic.h>
#include <vtkm/exec/arg/ThreadIndicesBasic3D.h>
#include <vtkm/exec/arg/VisitIndex.h>
#include <vtkm/exec/arg/WorkIndex.h>
@ -256,23 +257,50 @@ public:
/// types.
///
VTKM_SUPPRESS_EXEC_WARNINGS
template <typename T,
typename OutToInArrayType,
template <typename OutToInArrayType,
typename VisitArrayType,
typename ThreadToOutArrayType,
typename InputDomainType>
VTKM_EXEC vtkm::exec::arg::ThreadIndicesBasic GetThreadIndices(
const T& threadIndex,
const vtkm::Id& threadIndex,
const OutToInArrayType& outToIn,
const VisitArrayType& visit,
const ThreadToOutArrayType& threadToOut,
const InputDomainType&,
const T& globalThreadIndexOffset = 0) const
const vtkm::Id& globalThreadIndexOffset = 0) const
{
vtkm::Id outIndex = threadToOut.Get(threadIndex);
return vtkm::exec::arg::ThreadIndicesBasic(
threadIndex, outToIn.Get(outIndex), visit.Get(outIndex), outIndex, globalThreadIndexOffset);
}
/// \brief Creates a \c ThreadIndices object.
///
/// Worklet types can add additional indices by returning different object
/// types.
///
VTKM_SUPPRESS_EXEC_WARNINGS
template <typename OutToInArrayType,
typename VisitArrayType,
typename ThreadToOutArrayType,
typename InputDomainType>
VTKM_EXEC vtkm::exec::arg::ThreadIndicesBasic3D GetThreadIndices(
vtkm::Id threadIndex1D,
const vtkm::Id3& threadIndex3D,
const OutToInArrayType& outToIn,
const VisitArrayType& visit,
const ThreadToOutArrayType& threadToOut,
const InputDomainType&,
const vtkm::Id& globalThreadIndexOffset = 0) const
{
vtkm::Id outIndex = threadToOut.Get(threadIndex1D);
return vtkm::exec::arg::ThreadIndicesBasic3D(threadIndex3D,
threadIndex1D,
outToIn.Get(outIndex),
visit.Get(outIndex),
outIndex,
globalThreadIndexOffset);
}
};
}
}

1
vtkm/worklet/testing/CMakeLists.txt
View File

@ -71,6 +71,7 @@ set(unit_tests
UnitTestTube.cxx
UnitTestWholeCellSetIn.cxx
UnitTestWorkletMapField.cxx
UnitTestWorkletMapField3d.cxx
UnitTestWorkletMapFieldExecArg.cxx
UnitTestWorkletMapFieldWholeArray.cxx
UnitTestWorkletMapFieldWholeArrayAtomic.cxx

187
vtkm/worklet/testing/UnitTestWorkletMapField3d.cxx Normal file
View File

@ -0,0 +1,187 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
//
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//============================================================================
#include <vtkm/cont/ArrayCopy.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/VariantArrayHandle.h>
#include <vtkm/worklet/DispatcherMapField.h>
#include <vtkm/worklet/WorkletMapField.h>
#include <vtkm/cont/testing/Testing.h>
namespace mapfield3d
{
static constexpr vtkm::Id3 SCHEDULE_SIZE = { 10, 10, 10 };
static constexpr vtkm::Id ARRAY_SIZE = SCHEDULE_SIZE[0] * SCHEDULE_SIZE[1] * SCHEDULE_SIZE[2];
template <typename PortalType>
struct ExecutionObject
{
PortalType Portal;
};
template <typename T>
struct ExecutionObjectInterface : public vtkm::cont::ExecutionObjectBase
{
vtkm::cont::ArrayHandle<T> Data;
vtkm::Id3 ScheduleRange;
template <typename Device>
VTKM_CONT auto PrepareForExecution(Device device) const
-> ExecutionObject<decltype(this->Data.PrepareForInput(device))>
{
return ExecutionObject<decltype(this->Data.PrepareForInput(device))>{
this->Data.PrepareForInput(device)
};
}
vtkm::Id3 GetRange3d() const { return this->ScheduleRange; }
};
}
namespace vtkm
{
namespace exec
{
namespace arg
{
// Fetch for ArrayPortalTex3D when being used for Loads
template <typename PType>
struct Fetch<vtkm::exec::arg::FetchTagExecObject,
vtkm::exec::arg::AspectTagDefault,
vtkm::exec::arg::ThreadIndicesBasic3D,
mapfield3d::ExecutionObject<PType>>
{
using ValueType = typename PType::ValueType;
using PortalType = mapfield3d::ExecutionObject<PType>;
using ThreadIndicesType = vtkm::exec::arg::ThreadIndicesBasic3D;
VTKM_EXEC
ValueType Load(const ThreadIndicesType& indices, const PortalType& field) const
{
return field.Portal.Get(indices.GetInputIndex());
}
VTKM_EXEC
void Store(const ThreadIndicesType&, const PortalType&, const ValueType&) const {}
};
}
}
}
namespace mapfield3d
{
class TestMapFieldWorklet : public vtkm::worklet::WorkletMapField
{
public:
using ControlSignature = void(ExecObject, FieldOut, FieldInOut);
using ExecutionSignature = _3(_1, _2, _3, WorkIndex);
template <typename T>
VTKM_EXEC T operator()(const T& in, T& out, T& inout, vtkm::Id workIndex) const
{
auto expected = TestValue(workIndex, T()) + T(100);
if (!test_equal(in, expected))
{
this->RaiseError("Got wrong input value.");
}
out = static_cast<T>(in - T(100));
if (!test_equal(inout, TestValue(workIndex, T()) + T(100)))
{
this->RaiseError("Got wrong in-out value.");
}
// We return the new value of inout. Since _3 is both an arg and return,
// this tests that the return value is set after updating the arg values.
return static_cast<T>(inout - T(100));
}
template <typename T1, typename T2, typename T3>
VTKM_EXEC T3 operator()(const T1&, const T2&, const T3&, vtkm::Id) const
{
this->RaiseError("Cannot call this worklet with different types.");
return vtkm::TypeTraits<T3>::ZeroInitialization();
}
};
template <typename T>
inline vtkm::Id3 SchedulingRange(const ExecutionObjectInterface<T>& inputDomain)
{
return inputDomain.GetRange3d();
}
template <typename T>
inline vtkm::Id3 SchedulingRange(const ExecutionObjectInterface<T>* const inputDomain)
{
return inputDomain->GetRange3d();
}
template <typename WorkletType>
struct DoTestWorklet
{
template <typename T>
VTKM_CONT void operator()(T) const
{
std::cout << "Set up data." << std::endl;
T inputArray[ARRAY_SIZE];
for (vtkm::Id index = 0; index < ARRAY_SIZE; index++)
{
inputArray[index] = static_cast<T>(TestValue(index, T()) + T(100));
}
vtkm::cont::ArrayHandle<T> inputHandle = vtkm::cont::make_ArrayHandle(inputArray, ARRAY_SIZE);
vtkm::cont::ArrayHandle<T> outputHandleAsPtr;
vtkm::cont::ArrayHandle<T> inoutHandleAsPtr;
ExecutionObjectInterface<T> inputExecObject;
inputExecObject.Data = inputHandle;
inputExecObject.ScheduleRange = SCHEDULE_SIZE;
vtkm::cont::ArrayCopy(inputHandle, inoutHandleAsPtr);
std::cout << "Create and run dispatchers." << std::endl;
vtkm::worklet::DispatcherMapField<WorkletType> dispatcher;
dispatcher.Invoke(inputExecObject, &outputHandleAsPtr, &inoutHandleAsPtr);
std::cout << "Check results." << std::endl;
CheckPortal(outputHandleAsPtr.GetPortalConstControl());
CheckPortal(inoutHandleAsPtr.GetPortalConstControl());
}
};
void TestWorkletMapField3d(vtkm::cont::DeviceAdapterId id)
{
using HandleTypesToTest3D =
vtkm::List<vtkm::Id, vtkm::Vec2i_32, vtkm::FloatDefault, vtkm::Vec3f_64>;
std::cout << "Testing Map Field with 3d types on device adapter: " << id.GetName() << std::endl;
//need to test with ExecObject that has 3d range
//need to fetch from ExecObject that has 3d range
vtkm::testing::Testing::TryTypes(mapfield3d::DoTestWorklet<TestMapFieldWorklet>(),
HandleTypesToTest3D());
}
} // mapfield3d namespace
int UnitTestWorkletMapField3d(int argc, char* argv[])
{
return vtkm::cont::testing::Testing::RunOnDevice(mapfield3d::TestWorkletMapField3d, argc, argv);
}