mirror of
https://gitlab.kitware.com/vtk/vtk-m
synced 2024-09-08 13:23:51 +00:00
Add TaskTiling1D/3D which use faux virtuals to reduce binary size.
Redesigns the TBB and Serial backends and the vtkm::exec::Task concept so that we can re-use the same launching logic for all Worklets, instead of generating per worlet code. To keep the performance the same the TilingTask now is past a range of indices to work on, rather than a single index. Binary size reduction: WorkletTests_SERIAL old - 19MB WorkletTests_SERIAL new - 18MB WorkletTests_TBB old - 39MB WorkletTests_TBB new - 18MB libvtkAcceleratorsVTKm old - 48MB libvtkAcceleratorsVTKm new - 19MB
This commit is contained in:
parent
4655f5af4d
commit
60a405ef65
@ -103,15 +103,20 @@ set(device_sources
|
||||
RuntimeDeviceTracker.cxx
|
||||
)
|
||||
|
||||
vtkm_declare_headers(${headers})
|
||||
|
||||
#-----------------------------------------------------------------------------
|
||||
add_subdirectory(internal)
|
||||
add_subdirectory(arg)
|
||||
|
||||
vtkm_declare_headers(${headers})
|
||||
|
||||
add_subdirectory(serial)
|
||||
set(sources ${sources} $<TARGET_OBJECTS:vtkm_cont_serial>)
|
||||
|
||||
add_subdirectory(tbb)
|
||||
if (VTKm_ENABLE_TBB)
|
||||
set(sources ${sources} $<TARGET_OBJECTS:vtkm_cont_tbb>)
|
||||
endif()
|
||||
|
||||
add_subdirectory(cuda)
|
||||
if (VTKm_ENABLE_CUDA)
|
||||
get_property(vtkm_cont_cuda_object_files GLOBAL
|
||||
@ -122,15 +127,8 @@ if (VTKm_ENABLE_CUDA)
|
||||
PROPERTIES GENERATED TRUE)
|
||||
|
||||
set(sources ${sources} ${vtkm_cont_cuda_object_files})
|
||||
|
||||
endif()
|
||||
|
||||
add_subdirectory(tbb)
|
||||
if (VTKm_ENABLE_TBB)
|
||||
set(sources ${sources} $<TARGET_OBJECTS:vtkm_cont_tbb>)
|
||||
endif()
|
||||
|
||||
|
||||
vtkm_library(
|
||||
SOURCES ${sources}
|
||||
WRAP_FOR_CUDA ${device_sources}
|
||||
|
@ -583,6 +583,22 @@ public:
|
||||
///
|
||||
template <typename T, typename DeviceTag>
|
||||
class DeviceAdapterAtomicArrayImplementation;
|
||||
|
||||
/// \brief Class providing a device-specific support for selecting the optimal
|
||||
/// Task type for a given worklet.
|
||||
///
|
||||
/// When worklets are launched inside the execution enviornment we need to
|
||||
/// ask the device adapter what is the preferred execution style, be it
|
||||
/// a tiled iteration pattern, or strided. This class
|
||||
///
|
||||
/// By default if not specialized for a device adapter the default
|
||||
/// is to use vtkm::exec::internal::TaskSingular
|
||||
///
|
||||
/// The class provide the actual implementation used by
|
||||
/// vtkm::cont::DeviceTaskTypes.
|
||||
///
|
||||
template <typename DeviceTag>
|
||||
class DeviceTaskTypes;
|
||||
}
|
||||
} // namespace vtkm::cont
|
||||
|
||||
|
@ -33,7 +33,7 @@ namespace
|
||||
static const vtkm::Id ARRAY_SIZE = 10;
|
||||
|
||||
template <typename PortalType>
|
||||
struct TestKernel : public vtkm::exec::FunctorBase
|
||||
struct TestKernelIn : public vtkm::exec::FunctorBase
|
||||
{
|
||||
PortalType Portal;
|
||||
|
||||
@ -68,7 +68,7 @@ struct TryArrayInType
|
||||
vtkm::cont::arg::Transport<vtkm::cont::arg::TransportTagArrayIn, ArrayHandleType, Device>
|
||||
transport;
|
||||
|
||||
TestKernel<PortalType> kernel;
|
||||
TestKernelIn<PortalType> kernel;
|
||||
kernel.Portal = transport(handle, handle, ARRAY_SIZE, ARRAY_SIZE);
|
||||
|
||||
vtkm::cont::DeviceAdapterAlgorithm<Device>::Schedule(kernel, ARRAY_SIZE);
|
||||
|
@ -33,7 +33,7 @@ namespace
|
||||
static const vtkm::Id ARRAY_SIZE = 10;
|
||||
|
||||
template <typename PortalType>
|
||||
struct TestKernel : public vtkm::exec::FunctorBase
|
||||
struct TestKernelInOut : public vtkm::exec::FunctorBase
|
||||
{
|
||||
PortalType Portal;
|
||||
|
||||
@ -66,7 +66,7 @@ struct TryArrayInOutType
|
||||
vtkm::cont::arg::Transport<vtkm::cont::arg::TransportTagArrayInOut, ArrayHandleType, Device>
|
||||
transport;
|
||||
|
||||
TestKernel<PortalType> kernel;
|
||||
TestKernelInOut<PortalType> kernel;
|
||||
kernel.Portal = transport(handle, handle, ARRAY_SIZE, ARRAY_SIZE);
|
||||
|
||||
vtkm::cont::DeviceAdapterAlgorithm<Device>::Schedule(kernel, ARRAY_SIZE);
|
||||
|
@ -34,7 +34,7 @@ namespace
|
||||
static const vtkm::Id ARRAY_SIZE = 10;
|
||||
|
||||
template <typename PortalType>
|
||||
struct TestKernel : public vtkm::exec::FunctorBase
|
||||
struct TestKernelOut : public vtkm::exec::FunctorBase
|
||||
{
|
||||
PortalType Portal;
|
||||
|
||||
@ -60,7 +60,7 @@ struct TryArrayOutType
|
||||
vtkm::cont::arg::Transport<vtkm::cont::arg::TransportTagArrayOut, ArrayHandleType, Device>
|
||||
transport;
|
||||
|
||||
TestKernel<PortalType> kernel;
|
||||
TestKernelOut<PortalType> kernel;
|
||||
kernel.Portal =
|
||||
transport(handle, vtkm::cont::ArrayHandleIndex(ARRAY_SIZE), ARRAY_SIZE, ARRAY_SIZE);
|
||||
|
||||
|
@ -241,6 +241,29 @@ private:
|
||||
(unsigned long long int)newValue);
|
||||
}
|
||||
};
|
||||
|
||||
template <>
|
||||
class DeviceTaskTypes<vtkm::cont::DeviceAdapterTagCuda>
|
||||
{
|
||||
public:
|
||||
template <typename WorkletType, typename InvocationType>
|
||||
static vtkm::exec::internal::TaskSingular<WorkletType, InvocationType> MakeTask(
|
||||
const WorkletType& worklet, const InvocationType& invocation, vtkm::Id,
|
||||
vtkm::Id globalIndexOffset = 0)
|
||||
{
|
||||
using Task = vtkm::exec::internal::TaskSingular<WorkletType, InvocationType>;
|
||||
return Task(worklet, invocation, globalIndexOffset);
|
||||
}
|
||||
|
||||
template <typename WorkletType, typename InvocationType>
|
||||
static vtkm::exec::internal::TaskSingular<WorkletType, InvocationType> MakeTask(
|
||||
const WorkletType& worklet, const InvocationType& invocation, vtkm::Id3,
|
||||
vtkm::Id globalIndexOffset = 0)
|
||||
{
|
||||
using Task = vtkm::exec::internal::TaskSingular<WorkletType, InvocationType>;
|
||||
return Task(worklet, invocation, globalIndexOffset);
|
||||
}
|
||||
};
|
||||
}
|
||||
} // namespace vtkm::cont
|
||||
|
||||
|
@ -1187,6 +1187,14 @@ private:
|
||||
}
|
||||
|
||||
public:
|
||||
template <class TaskType, typename RangeType>
|
||||
VTKM_CONT static void ScheduleTask(TaskType task, RangeType size)
|
||||
{ //for now defer to the schedule api, we need to do a significant
|
||||
//amount of build infrastructure work to implement type erasure tasks
|
||||
//for cuda
|
||||
Schedule(task, size);
|
||||
}
|
||||
|
||||
template <class Functor>
|
||||
VTKM_CONT static void Schedule(Functor functor, vtkm::Id numInstances)
|
||||
{
|
||||
|
@ -29,6 +29,7 @@
|
||||
#include <vtkm/cont/internal/FunctorsGeneral.h>
|
||||
|
||||
#include <vtkm/exec/internal/ErrorMessageBuffer.h>
|
||||
#include <vtkm/exec/internal/TaskSingular.h>
|
||||
|
||||
#include <vtkm/TypeTraits.h>
|
||||
|
||||
@ -1029,6 +1030,39 @@ private:
|
||||
|
||||
#endif
|
||||
};
|
||||
|
||||
/// \brief Class providing a device-specific support for selecting the optimal
|
||||
/// Task type for a given worklet.
|
||||
///
|
||||
/// When worklets are launched inside the execution enviornment we need to
|
||||
/// ask the device adapter what is the preferred execution style, be it
|
||||
/// a tiled iteration pattern, or strided. This class
|
||||
///
|
||||
/// By default if not specialized for a device adapter the default
|
||||
/// is to use vtkm::exec::internal::TaskSingular
|
||||
///
|
||||
template <typename DeviceTag>
|
||||
class DeviceTaskTypes
|
||||
{
|
||||
public:
|
||||
template <typename WorkletType, typename InvocationType>
|
||||
static vtkm::exec::internal::TaskSingular<WorkletType, InvocationType> MakeTask(
|
||||
const WorkletType& worklet, const InvocationType& invocation, vtkm::Id,
|
||||
vtkm::Id globalIndexOffset = 0)
|
||||
{
|
||||
using Task = vtkm::exec::internal::TaskSingular<WorkletType, InvocationType>;
|
||||
return Task(worklet, invocation, globalIndexOffset);
|
||||
}
|
||||
|
||||
template <typename WorkletType, typename InvocationType>
|
||||
static vtkm::exec::internal::TaskSingular<WorkletType, InvocationType> MakeTask(
|
||||
const WorkletType& worklet, const InvocationType& invocation, vtkm::Id3,
|
||||
vtkm::Id globalIndexOffset = 0)
|
||||
{
|
||||
using Task = vtkm::exec::internal::TaskSingular<WorkletType, InvocationType>;
|
||||
return Task(worklet, invocation, globalIndexOffset);
|
||||
}
|
||||
};
|
||||
}
|
||||
} // namespace vtkm::cont
|
||||
|
||||
|
@ -25,22 +25,19 @@ set(headers
|
||||
VirtualObjectTransferSerial.h
|
||||
)
|
||||
|
||||
vtkm_declare_headers(${headers})
|
||||
|
||||
add_library(vtkm_cont_serial OBJECT ArrayManagerExecutionSerial.cxx)
|
||||
add_library(vtkm_cont_serial OBJECT
|
||||
ArrayManagerExecutionSerial.cxx
|
||||
DeviceAdapterAlgorithmSerial.cxx
|
||||
)
|
||||
|
||||
target_compile_features(vtkm_cont_serial PRIVATE cxx_auto_type)
|
||||
target_compile_definitions(vtkm_cont_serial PRIVATE vtkm_cont_EXPORTS)
|
||||
if(BUILD_SHARED_LIBS)
|
||||
set_property(TARGET vtkm_cont_serial PROPERTY POSITION_INDEPENDENT_CODE ON)
|
||||
endif()
|
||||
|
||||
target_include_directories(vtkm_cont_serial PRIVATE
|
||||
${VTKm_BACKEND_INCLUDE_DIRS}
|
||||
${VTKm_SOURCE_DIR}
|
||||
${VTKm_BINARY_DIR}/include
|
||||
)
|
||||
|
||||
if (BUILD_SHARED_LIBS)
|
||||
set_property(TARGET vtkm_cont_serial
|
||||
PROPERTY POSITION_INDEPENDENT_CODE ON
|
||||
)
|
||||
endif()
|
||||
|
||||
|
75
vtkm/cont/serial/internal/DeviceAdapterAlgorithmSerial.cxx
Normal file
75
vtkm/cont/serial/internal/DeviceAdapterAlgorithmSerial.cxx
Normal file
@ -0,0 +1,75 @@
|
||||
//============================================================================
|
||||
// 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.
|
||||
//============================================================================
|
||||
|
||||
#include <vtkm/cont/serial/internal/DeviceAdapterAlgorithmSerial.h>
|
||||
|
||||
namespace vtkm
|
||||
{
|
||||
namespace cont
|
||||
{
|
||||
|
||||
void DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagSerial>::ScheduleTask(
|
||||
vtkm::exec::serial::internal::TaskTiling1D& functor, vtkm::Id size)
|
||||
{
|
||||
const vtkm::Id MESSAGE_SIZE = 1024;
|
||||
char errorString[MESSAGE_SIZE];
|
||||
errorString[0] = '\0';
|
||||
vtkm::exec::internal::ErrorMessageBuffer errorMessage(errorString, MESSAGE_SIZE);
|
||||
functor.SetErrorMessageBuffer(errorMessage);
|
||||
|
||||
const vtkm::Id iterations = size / 1024;
|
||||
vtkm::Id index = 0;
|
||||
for (vtkm::Id i = 0; i < iterations; ++i)
|
||||
{
|
||||
functor(index, index + 1024);
|
||||
index += 1024;
|
||||
}
|
||||
functor(index, size);
|
||||
|
||||
if (errorMessage.IsErrorRaised())
|
||||
{
|
||||
throw vtkm::cont::ErrorExecution(errorString);
|
||||
}
|
||||
}
|
||||
|
||||
void DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagSerial>::ScheduleTask(
|
||||
vtkm::exec::serial::internal::TaskTiling3D& functor, vtkm::Id3 size)
|
||||
{
|
||||
const vtkm::Id MESSAGE_SIZE = 1024;
|
||||
char errorString[MESSAGE_SIZE];
|
||||
errorString[0] = '\0';
|
||||
vtkm::exec::internal::ErrorMessageBuffer errorMessage(errorString, MESSAGE_SIZE);
|
||||
functor.SetErrorMessageBuffer(errorMessage);
|
||||
|
||||
for (vtkm::Id k = 0; k < size[2]; ++k)
|
||||
{
|
||||
for (vtkm::Id j = 0; j < size[1]; ++j)
|
||||
{
|
||||
functor(0, size[0], j, k);
|
||||
}
|
||||
}
|
||||
|
||||
if (errorMessage.IsErrorRaised())
|
||||
{
|
||||
throw vtkm::cont::ErrorExecution(errorString);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
@ -30,7 +30,7 @@
|
||||
|
||||
#include <vtkm/BinaryOperators.h>
|
||||
|
||||
#include <vtkm/exec/internal/ErrorMessageBuffer.h>
|
||||
#include <vtkm/exec/serial/internal/TaskTiling.h>
|
||||
|
||||
#include <algorithm>
|
||||
#include <numeric>
|
||||
@ -237,86 +237,23 @@ public:
|
||||
return ScanExclusive(input, output, vtkm::Sum(), vtkm::TypeTraits<T>::ZeroInitialization());
|
||||
}
|
||||
|
||||
private:
|
||||
// This runs in the execution environment.
|
||||
VTKM_CONT_EXPORT static void ScheduleTask(vtkm::exec::serial::internal::TaskTiling1D& functor,
|
||||
vtkm::Id size);
|
||||
VTKM_CONT_EXPORT static void ScheduleTask(vtkm::exec::serial::internal::TaskTiling3D& functor,
|
||||
vtkm::Id3 size);
|
||||
|
||||
template <class FunctorType>
|
||||
class ScheduleKernel
|
||||
VTKM_CONT static inline void Schedule(FunctorType functor, vtkm::Id size)
|
||||
{
|
||||
public:
|
||||
ScheduleKernel(const FunctorType& functor)
|
||||
: Functor(functor)
|
||||
{
|
||||
}
|
||||
|
||||
//needed for when calling from schedule on a range
|
||||
template <typename T>
|
||||
VTKM_EXEC void operator()(const T& index) const
|
||||
{
|
||||
this->Functor(index);
|
||||
}
|
||||
|
||||
private:
|
||||
const FunctorType Functor;
|
||||
|
||||
void operator=(const ScheduleKernel<FunctorType>&) = delete;
|
||||
};
|
||||
|
||||
public:
|
||||
template <class Functor>
|
||||
VTKM_CONT static void Schedule(Functor functor, vtkm::Id numInstances)
|
||||
{
|
||||
const vtkm::Id MESSAGE_SIZE = 1024;
|
||||
char errorString[MESSAGE_SIZE];
|
||||
errorString[0] = '\0';
|
||||
vtkm::exec::internal::ErrorMessageBuffer errorMessage(errorString, MESSAGE_SIZE);
|
||||
|
||||
functor.SetErrorMessageBuffer(errorMessage);
|
||||
|
||||
DeviceAdapterAlgorithm<Device>::ScheduleKernel<Functor> kernel(functor);
|
||||
|
||||
const vtkm::Id size = numInstances;
|
||||
|
||||
VTKM_VECTORIZATION_PRE_LOOP
|
||||
for (vtkm::Id i = 0; i < size; ++i)
|
||||
{
|
||||
VTKM_VECTORIZATION_IN_LOOP
|
||||
kernel(i);
|
||||
}
|
||||
|
||||
if (errorMessage.IsErrorRaised())
|
||||
{
|
||||
throw vtkm::cont::ErrorExecution(errorString);
|
||||
}
|
||||
vtkm::exec::serial::internal::TaskTiling1D kernel(functor);
|
||||
ScheduleTask(kernel, size);
|
||||
}
|
||||
|
||||
template <class Functor>
|
||||
VTKM_CONT static void Schedule(Functor functor, vtkm::Id3 rangeMax)
|
||||
template <class FunctorType>
|
||||
VTKM_CONT static inline void Schedule(FunctorType functor, vtkm::Id3 size)
|
||||
{
|
||||
const vtkm::Id MESSAGE_SIZE = 1024;
|
||||
char errorString[MESSAGE_SIZE];
|
||||
errorString[0] = '\0';
|
||||
vtkm::exec::internal::ErrorMessageBuffer errorMessage(errorString, MESSAGE_SIZE);
|
||||
|
||||
functor.SetErrorMessageBuffer(errorMessage);
|
||||
|
||||
DeviceAdapterAlgorithm<Device>::ScheduleKernel<Functor> kernel(functor);
|
||||
for (vtkm::Id k = 0; k < rangeMax[2]; ++k)
|
||||
{
|
||||
for (vtkm::Id j = 0; j < rangeMax[1]; ++j)
|
||||
{
|
||||
VTKM_VECTORIZATION_PRE_LOOP
|
||||
for (vtkm::Id i = 0; i < rangeMax[0]; ++i)
|
||||
{
|
||||
VTKM_VECTORIZATION_IN_LOOP
|
||||
kernel(vtkm::Id3(i, j, k));
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
if (errorMessage.IsErrorRaised())
|
||||
{
|
||||
throw vtkm::cont::ErrorExecution(errorString);
|
||||
}
|
||||
vtkm::exec::serial::internal::TaskTiling3D kernel(functor);
|
||||
ScheduleTask(kernel, size);
|
||||
}
|
||||
|
||||
private:
|
||||
@ -427,6 +364,29 @@ public:
|
||||
// Nothing to do. This device is serial and has no asynchronous operations.
|
||||
}
|
||||
};
|
||||
|
||||
template <>
|
||||
class DeviceTaskTypes<vtkm::cont::DeviceAdapterTagSerial>
|
||||
{
|
||||
public:
|
||||
template <typename WorkletType, typename InvocationType>
|
||||
static vtkm::exec::serial::internal::TaskTiling1D MakeTask(const WorkletType& worklet,
|
||||
const InvocationType& invocation,
|
||||
vtkm::Id,
|
||||
vtkm::Id globalIndexOffset = 0)
|
||||
{
|
||||
return vtkm::exec::serial::internal::TaskTiling1D(worklet, invocation, globalIndexOffset);
|
||||
}
|
||||
|
||||
template <typename WorkletType, typename InvocationType>
|
||||
static vtkm::exec::serial::internal::TaskTiling3D MakeTask(const WorkletType& worklet,
|
||||
const InvocationType& invocation,
|
||||
vtkm::Id3,
|
||||
vtkm::Id globalIndexOffset = 0)
|
||||
{
|
||||
return vtkm::exec::serial::internal::TaskTiling3D(worklet, invocation, globalIndexOffset);
|
||||
}
|
||||
};
|
||||
}
|
||||
} // namespace vtkm::cont
|
||||
|
||||
|
@ -38,16 +38,23 @@ vtkm_declare_headers(parallel_sort.h TESTABLE OFF)
|
||||
vtkm_declare_headers(${headers} TESTABLE ${VTKm_ENABLE_TBB})
|
||||
|
||||
#-----------------------------------------------------------------------------
|
||||
if (VTKm_ENABLE_TBB)
|
||||
add_library(vtkm_cont_tbb OBJECT ArrayManagerExecutionTBB.cxx)
|
||||
target_compile_features(vtkm_cont_tbb PRIVATE cxx_auto_type)
|
||||
target_compile_definitions(vtkm_cont_tbb PRIVATE vtkm_cont_EXPORTS)
|
||||
target_include_directories(vtkm_cont_tbb PRIVATE
|
||||
"${VTKm_SOURCE_DIR}"
|
||||
"${VTKm_BINARY_DIR}/include"
|
||||
${VTKm_BACKEND_INCLUDE_DIRS}
|
||||
)
|
||||
if (BUILD_SHARED_LIBS)
|
||||
set_property(TARGET vtkm_cont_tbb PROPERTY POSITION_INDEPENDENT_CODE ON)
|
||||
endif()
|
||||
if (NOT VTKm_ENABLE_TBB)
|
||||
return()
|
||||
endif()
|
||||
|
||||
add_library(vtkm_cont_tbb OBJECT
|
||||
ArrayManagerExecutionTBB.cxx
|
||||
DeviceAdapterAlgorithmTBB.cxx
|
||||
)
|
||||
|
||||
target_compile_features(vtkm_cont_tbb PRIVATE cxx_auto_type)
|
||||
target_compile_definitions(vtkm_cont_tbb PRIVATE vtkm_cont_EXPORTS)
|
||||
if(BUILD_SHARED_LIBS)
|
||||
set_property(TARGET vtkm_cont_tbb PROPERTY POSITION_INDEPENDENT_CODE ON)
|
||||
endif()
|
||||
|
||||
target_include_directories(vtkm_cont_tbb PRIVATE
|
||||
"${VTKm_SOURCE_DIR}"
|
||||
"${VTKm_BINARY_DIR}/include"
|
||||
${VTKm_BACKEND_INCLUDE_DIRS}
|
||||
)
|
||||
|
80
vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.cxx
Normal file
80
vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.cxx
Normal file
@ -0,0 +1,80 @@
|
||||
//============================================================================
|
||||
// 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.
|
||||
//============================================================================
|
||||
|
||||
#include <vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h>
|
||||
|
||||
namespace vtkm
|
||||
{
|
||||
namespace cont
|
||||
{
|
||||
|
||||
void DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagTBB>::ScheduleTask(
|
||||
vtkm::exec::tbb::internal::TaskTiling1D& functor, vtkm::Id size)
|
||||
{
|
||||
const vtkm::Id MESSAGE_SIZE = 1024;
|
||||
char errorString[MESSAGE_SIZE];
|
||||
errorString[0] = '\0';
|
||||
vtkm::exec::internal::ErrorMessageBuffer errorMessage(errorString, MESSAGE_SIZE);
|
||||
functor.SetErrorMessageBuffer(errorMessage);
|
||||
|
||||
::tbb::blocked_range<vtkm::Id> range(0, size);
|
||||
|
||||
::tbb::parallel_for(
|
||||
range, [&](const ::tbb::blocked_range<vtkm::Id>& r) { functor(r.begin(), r.end()); });
|
||||
|
||||
if (errorMessage.IsErrorRaised())
|
||||
{
|
||||
throw vtkm::cont::ErrorExecution(errorString);
|
||||
}
|
||||
}
|
||||
|
||||
void DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagTBB>::ScheduleTask(
|
||||
vtkm::exec::tbb::internal::TaskTiling3D& functor, vtkm::Id3 size)
|
||||
{
|
||||
static const vtkm::UInt32 TBB_GRAIN_SIZE_3D[3] = { 1, 4, 256 };
|
||||
const vtkm::Id MESSAGE_SIZE = 1024;
|
||||
char errorString[MESSAGE_SIZE];
|
||||
errorString[0] = '\0';
|
||||
vtkm::exec::internal::ErrorMessageBuffer errorMessage(errorString, MESSAGE_SIZE);
|
||||
functor.SetErrorMessageBuffer(errorMessage);
|
||||
|
||||
//memory is generally setup in a way that iterating the first range
|
||||
//in the tightest loop has the best cache coherence.
|
||||
::tbb::blocked_range3d<vtkm::Id> range(0, size[2], TBB_GRAIN_SIZE_3D[0], 0, size[1],
|
||||
TBB_GRAIN_SIZE_3D[1], 0, size[0], TBB_GRAIN_SIZE_3D[2]);
|
||||
::tbb::parallel_for(range, [&](const ::tbb::blocked_range3d<vtkm::Id>& r) {
|
||||
for (vtkm::Id k = r.pages().begin(); k != r.pages().end(); ++k)
|
||||
{
|
||||
for (vtkm::Id j = r.rows().begin(); j != r.rows().end(); ++j)
|
||||
{
|
||||
const vtkm::Id start = r.cols().begin();
|
||||
const vtkm::Id end = r.cols().end();
|
||||
functor(start, end, j, k);
|
||||
}
|
||||
}
|
||||
});
|
||||
|
||||
if (errorMessage.IsErrorRaised())
|
||||
{
|
||||
throw vtkm::cont::ErrorExecution(errorString);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
@ -30,7 +30,8 @@
|
||||
#include <vtkm/cont/tbb/internal/ArrayManagerExecutionTBB.h>
|
||||
#include <vtkm/cont/tbb/internal/DeviceAdapterTagTBB.h>
|
||||
#include <vtkm/cont/tbb/internal/FunctorsTBB.h>
|
||||
#include <vtkm/exec/internal/ErrorMessageBuffer.h>
|
||||
|
||||
#include <vtkm/exec/tbb/internal/TaskTiling.h>
|
||||
|
||||
namespace vtkm
|
||||
{
|
||||
@ -99,53 +100,23 @@ public:
|
||||
binary_functor, initialValue);
|
||||
}
|
||||
|
||||
VTKM_CONT_EXPORT static void ScheduleTask(vtkm::exec::tbb::internal::TaskTiling1D& functor,
|
||||
vtkm::Id size);
|
||||
VTKM_CONT_EXPORT static void ScheduleTask(vtkm::exec::tbb::internal::TaskTiling3D& functor,
|
||||
vtkm::Id3 size);
|
||||
|
||||
template <class FunctorType>
|
||||
VTKM_CONT static void Schedule(FunctorType functor, vtkm::Id numInstances)
|
||||
VTKM_CONT static inline void Schedule(FunctorType functor, vtkm::Id numInstances)
|
||||
{
|
||||
const vtkm::Id MESSAGE_SIZE = 1024;
|
||||
char errorString[MESSAGE_SIZE];
|
||||
errorString[0] = '\0';
|
||||
vtkm::exec::internal::ErrorMessageBuffer errorMessage(errorString, MESSAGE_SIZE);
|
||||
|
||||
tbb::ScheduleKernel<FunctorType> kernel(functor);
|
||||
kernel.SetErrorMessageBuffer(errorMessage);
|
||||
|
||||
::tbb::blocked_range<vtkm::Id> range(0, numInstances, tbb::TBB_GRAIN_SIZE);
|
||||
|
||||
::tbb::parallel_for(range, kernel);
|
||||
|
||||
if (errorMessage.IsErrorRaised())
|
||||
{
|
||||
throw vtkm::cont::ErrorExecution(errorString);
|
||||
}
|
||||
vtkm::exec::tbb::internal::TaskTiling1D kernel(functor);
|
||||
ScheduleTask(kernel, numInstances);
|
||||
}
|
||||
|
||||
template <class FunctorType>
|
||||
VTKM_CONT static void Schedule(FunctorType functor, vtkm::Id3 rangeMax)
|
||||
VTKM_CONT static inline void Schedule(FunctorType functor, vtkm::Id3 rangeMax)
|
||||
{
|
||||
static const vtkm::UInt32 TBB_GRAIN_SIZE_3D[3] = { 1, 4, 256 };
|
||||
|
||||
//we need to extract from the functor that uniform grid information
|
||||
const vtkm::Id MESSAGE_SIZE = 1024;
|
||||
char errorString[MESSAGE_SIZE];
|
||||
errorString[0] = '\0';
|
||||
vtkm::exec::internal::ErrorMessageBuffer errorMessage(errorString, MESSAGE_SIZE);
|
||||
|
||||
//memory is generally setup in a way that iterating the first range
|
||||
//in the tightest loop has the best cache coherence.
|
||||
::tbb::blocked_range3d<vtkm::Id> range(0, rangeMax[2], TBB_GRAIN_SIZE_3D[0], 0, rangeMax[1],
|
||||
TBB_GRAIN_SIZE_3D[1], 0, rangeMax[0],
|
||||
TBB_GRAIN_SIZE_3D[2]);
|
||||
|
||||
tbb::ScheduleKernelId3<FunctorType> kernel(functor);
|
||||
kernel.SetErrorMessageBuffer(errorMessage);
|
||||
|
||||
::tbb::parallel_for(range, kernel);
|
||||
|
||||
if (errorMessage.IsErrorRaised())
|
||||
{
|
||||
throw vtkm::cont::ErrorExecution(errorString);
|
||||
}
|
||||
vtkm::exec::tbb::internal::TaskTiling3D kernel(functor);
|
||||
ScheduleTask(kernel, rangeMax);
|
||||
}
|
||||
|
||||
template <typename T, class Container>
|
||||
@ -249,6 +220,29 @@ public:
|
||||
private:
|
||||
::tbb::tick_count StartTime;
|
||||
};
|
||||
|
||||
template <>
|
||||
class DeviceTaskTypes<vtkm::cont::DeviceAdapterTagTBB>
|
||||
{
|
||||
public:
|
||||
template <typename WorkletType, typename InvocationType>
|
||||
static vtkm::exec::serial::internal::TaskTiling1D MakeTask(const WorkletType& worklet,
|
||||
const InvocationType& invocation,
|
||||
vtkm::Id,
|
||||
vtkm::Id globalIndexOffset = 0)
|
||||
{
|
||||
return vtkm::exec::tbb::internal::TaskTiling1D(worklet, invocation, globalIndexOffset);
|
||||
}
|
||||
|
||||
template <typename WorkletType, typename InvocationType>
|
||||
static vtkm::exec::serial::internal::TaskTiling3D MakeTask(const WorkletType& worklet,
|
||||
const InvocationType& invocation,
|
||||
vtkm::Id3,
|
||||
vtkm::Id globalIndexOffset = 0)
|
||||
{
|
||||
return vtkm::exec::tbb::internal::TaskTiling3D(worklet, invocation, globalIndexOffset);
|
||||
}
|
||||
};
|
||||
}
|
||||
} // namespace vtkm::cont
|
||||
|
||||
|
@ -75,6 +75,12 @@ namespace cont
|
||||
namespace tbb
|
||||
{
|
||||
|
||||
namespace internal
|
||||
{
|
||||
template <typename ResultType, typename Function>
|
||||
using WrappedBinaryOperator = vtkm::cont::internal::WrappedBinaryOperator<ResultType, Function>;
|
||||
}
|
||||
|
||||
// The "grain size" of scheduling with TBB. Not a lot of thought has gone
|
||||
// into picking this size.
|
||||
static const vtkm::Id TBB_GRAIN_SIZE = 1024;
|
||||
@ -423,113 +429,6 @@ ScanExclusivePortals(
|
||||
return body.Sum;
|
||||
}
|
||||
|
||||
template <class FunctorType>
|
||||
class ScheduleKernel
|
||||
{
|
||||
public:
|
||||
VTKM_CONT ScheduleKernel(const FunctorType& functor)
|
||||
: Functor(functor)
|
||||
{
|
||||
}
|
||||
|
||||
VTKM_CONT void SetErrorMessageBuffer(const vtkm::exec::internal::ErrorMessageBuffer& errorMessage)
|
||||
{
|
||||
this->ErrorMessage = errorMessage;
|
||||
this->Functor.SetErrorMessageBuffer(errorMessage);
|
||||
}
|
||||
|
||||
VTKM_CONT
|
||||
void operator()(const ::tbb::blocked_range<vtkm::Id>& range) const
|
||||
{
|
||||
// The TBB device adapter causes array classes to be shared between
|
||||
// control and execution environment. This means that it is possible for
|
||||
// an exception to be thrown even though this is typically not allowed.
|
||||
// Throwing an exception from here is bad because there are several
|
||||
// simultaneous threads running. Get around the problem by catching the
|
||||
// error and setting the message buffer as expected.
|
||||
try
|
||||
{
|
||||
const vtkm::Id start = range.begin();
|
||||
const vtkm::Id end = range.end();
|
||||
VTKM_VECTORIZATION_PRE_LOOP
|
||||
for (vtkm::Id index = start; index != end; index++)
|
||||
{
|
||||
VTKM_VECTORIZATION_IN_LOOP
|
||||
this->Functor(index);
|
||||
}
|
||||
}
|
||||
catch (vtkm::cont::Error& error)
|
||||
{
|
||||
this->ErrorMessage.RaiseError(error.GetMessage().c_str());
|
||||
}
|
||||
catch (...)
|
||||
{
|
||||
this->ErrorMessage.RaiseError("Unexpected error in execution environment.");
|
||||
}
|
||||
}
|
||||
|
||||
private:
|
||||
FunctorType Functor;
|
||||
vtkm::exec::internal::ErrorMessageBuffer ErrorMessage;
|
||||
};
|
||||
|
||||
template <class FunctorType>
|
||||
class ScheduleKernelId3
|
||||
{
|
||||
public:
|
||||
VTKM_CONT ScheduleKernelId3(const FunctorType& functor)
|
||||
: Functor(functor)
|
||||
{
|
||||
}
|
||||
|
||||
VTKM_CONT void SetErrorMessageBuffer(const vtkm::exec::internal::ErrorMessageBuffer& errorMessage)
|
||||
{
|
||||
this->ErrorMessage = errorMessage;
|
||||
this->Functor.SetErrorMessageBuffer(errorMessage);
|
||||
}
|
||||
|
||||
VTKM_CONT
|
||||
void operator()(const ::tbb::blocked_range3d<vtkm::Id>& range) const
|
||||
{
|
||||
try
|
||||
{
|
||||
const vtkm::Id kstart = range.pages().begin();
|
||||
const vtkm::Id kend = range.pages().end();
|
||||
const vtkm::Id jstart = range.rows().begin();
|
||||
const vtkm::Id jend = range.rows().end();
|
||||
const vtkm::Id istart = range.cols().begin();
|
||||
const vtkm::Id iend = range.cols().end();
|
||||
|
||||
vtkm::Id3 index;
|
||||
for (vtkm::Id k = kstart; k != kend; ++k)
|
||||
{
|
||||
index[2] = k;
|
||||
for (vtkm::Id j = jstart; j != jend; ++j)
|
||||
{
|
||||
index[1] = j;
|
||||
for (vtkm::Id i = istart; i != iend; ++i)
|
||||
{
|
||||
index[0] = i;
|
||||
this->Functor(index);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
catch (vtkm::cont::Error& error)
|
||||
{
|
||||
this->ErrorMessage.RaiseError(error.GetMessage().c_str());
|
||||
}
|
||||
catch (...)
|
||||
{
|
||||
this->ErrorMessage.RaiseError("Unexpected error in execution environment.");
|
||||
}
|
||||
}
|
||||
|
||||
private:
|
||||
FunctorType Functor;
|
||||
vtkm::exec::internal::ErrorMessageBuffer ErrorMessage;
|
||||
};
|
||||
|
||||
template <typename InputPortalType, typename IndexPortalType, typename OutputPortalType>
|
||||
class ScatterKernel
|
||||
{
|
||||
|
@ -44,8 +44,9 @@ vtkm_declare_headers(${impl_headers} ${headers})
|
||||
|
||||
|
||||
#-----------------------------------------------------------------------------
|
||||
add_subdirectory(serial)
|
||||
add_subdirectory(tbb)
|
||||
add_subdirectory(cuda)
|
||||
|
||||
|
||||
#-----------------------------------------------------------------------------
|
||||
add_subdirectory(testing)
|
||||
|
@ -31,8 +31,6 @@ namespace exec
|
||||
|
||||
/// Base class for all classes that are used to marshal data from the invocation
|
||||
/// parameters to the user worklets when invoked in the execution environment.
|
||||
/// These classes are generally automatically created by
|
||||
/// vtkm::worklet::internal::DispatcherBase.
|
||||
class TaskBase
|
||||
{
|
||||
};
|
||||
|
@ -40,4 +40,3 @@ set_source_files_properties(ThrustPatches.h
|
||||
if (VTKm_ENABLE_CUDA)
|
||||
add_subdirectory(testing)
|
||||
endif()
|
||||
|
||||
|
@ -22,5 +22,6 @@
|
||||
|
||||
set(unit_tests
|
||||
UnitTestTextureMemorySupport.cu
|
||||
UnitTestTaskSingularCuda.cu
|
||||
)
|
||||
vtkm_unit_tests(CUDA SOURCES ${unit_tests})
|
||||
|
350
vtkm/exec/cuda/internal/testing/UnitTestTaskSingularCuda.cu
Normal file
350
vtkm/exec/cuda/internal/testing/UnitTestTaskSingularCuda.cu
Normal file
@ -0,0 +1,350 @@
|
||||
//============================================================================
|
||||
// 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.
|
||||
//============================================================================
|
||||
|
||||
#include <vtkm/cont/cuda/DeviceAdapterCuda.h>
|
||||
|
||||
#include <vtkm/exec/FunctorBase.h>
|
||||
#include <vtkm/exec/arg/BasicArg.h>
|
||||
#include <vtkm/exec/arg/ThreadIndicesBasic.h>
|
||||
#include <vtkm/exec/internal/TaskSingular.h>
|
||||
|
||||
#include <vtkm/StaticAssert.h>
|
||||
|
||||
#include <vtkm/internal/FunctionInterface.h>
|
||||
#include <vtkm/internal/Invocation.h>
|
||||
|
||||
#include <vtkm/cont/testing/Testing.h>
|
||||
|
||||
namespace
|
||||
{
|
||||
|
||||
struct TestExecObject
|
||||
{
|
||||
VTKM_EXEC_CONT
|
||||
TestExecObject()
|
||||
: Portal()
|
||||
{
|
||||
}
|
||||
|
||||
VTKM_EXEC_CONT
|
||||
TestExecObject(vtkm::exec::cuda::internal::ArrayPortalFromThrust<vtkm::Id> portal)
|
||||
: Portal(portal)
|
||||
{
|
||||
}
|
||||
|
||||
vtkm::exec::cuda::internal::ArrayPortalFromThrust<vtkm::Id> Portal;
|
||||
};
|
||||
|
||||
struct MyOutputToInputMapPortal
|
||||
{
|
||||
typedef vtkm::Id ValueType;
|
||||
VTKM_EXEC_CONT
|
||||
vtkm::Id Get(vtkm::Id index) const { return index; }
|
||||
};
|
||||
|
||||
struct MyVisitArrayPortal
|
||||
{
|
||||
typedef vtkm::IdComponent ValueType;
|
||||
vtkm::IdComponent Get(vtkm::Id) const { return 1; }
|
||||
};
|
||||
|
||||
struct TestFetchTagInput
|
||||
{
|
||||
};
|
||||
struct TestFetchTagOutput
|
||||
{
|
||||
};
|
||||
|
||||
// Missing TransportTag, but we are not testing that so we can leave it out.
|
||||
struct TestControlSignatureTagInput
|
||||
{
|
||||
typedef TestFetchTagInput FetchTag;
|
||||
};
|
||||
struct TestControlSignatureTagOutput
|
||||
{
|
||||
typedef TestFetchTagOutput FetchTag;
|
||||
};
|
||||
|
||||
} // anonymous namespace
|
||||
|
||||
namespace vtkm
|
||||
{
|
||||
namespace exec
|
||||
{
|
||||
namespace arg
|
||||
{
|
||||
|
||||
template <>
|
||||
struct Fetch<TestFetchTagInput, vtkm::exec::arg::AspectTagDefault,
|
||||
vtkm::exec::arg::ThreadIndicesBasic, TestExecObject>
|
||||
{
|
||||
typedef vtkm::Id ValueType;
|
||||
|
||||
VTKM_EXEC
|
||||
ValueType Load(const vtkm::exec::arg::ThreadIndicesBasic& indices,
|
||||
const TestExecObject& execObject) const
|
||||
{
|
||||
return execObject.Portal.Get(indices.GetInputIndex()) + 10 * indices.GetInputIndex();
|
||||
}
|
||||
|
||||
VTKM_EXEC
|
||||
void Store(const vtkm::exec::arg::ThreadIndicesBasic&, const TestExecObject&, ValueType) const
|
||||
{
|
||||
// No-op
|
||||
}
|
||||
};
|
||||
|
||||
template <>
|
||||
struct Fetch<TestFetchTagOutput, vtkm::exec::arg::AspectTagDefault,
|
||||
vtkm::exec::arg::ThreadIndicesBasic, TestExecObject>
|
||||
{
|
||||
typedef vtkm::Id ValueType;
|
||||
|
||||
VTKM_EXEC
|
||||
ValueType Load(const vtkm::exec::arg::ThreadIndicesBasic&, const TestExecObject&) const
|
||||
{
|
||||
// No-op
|
||||
return ValueType();
|
||||
}
|
||||
|
||||
VTKM_EXEC
|
||||
void Store(const vtkm::exec::arg::ThreadIndicesBasic& indices, const TestExecObject& execObject,
|
||||
ValueType value) const
|
||||
{
|
||||
execObject.Portal.Set(indices.GetOutputIndex(), value + 20 * indices.GetOutputIndex());
|
||||
}
|
||||
};
|
||||
}
|
||||
}
|
||||
} // vtkm::exec::arg
|
||||
|
||||
namespace
|
||||
{
|
||||
|
||||
typedef void TestControlSignature(TestControlSignatureTagInput, TestControlSignatureTagOutput);
|
||||
typedef vtkm::internal::FunctionInterface<TestControlSignature> TestControlInterface;
|
||||
|
||||
typedef void TestExecutionSignature1(vtkm::exec::arg::BasicArg<1>, vtkm::exec::arg::BasicArg<2>);
|
||||
typedef vtkm::internal::FunctionInterface<TestExecutionSignature1> TestExecutionInterface1;
|
||||
|
||||
typedef vtkm::exec::arg::BasicArg<2> TestExecutionSignature2(vtkm::exec::arg::BasicArg<1>);
|
||||
typedef vtkm::internal::FunctionInterface<TestExecutionSignature2> TestExecutionInterface2;
|
||||
|
||||
typedef vtkm::internal::FunctionInterface<void(TestExecObject, TestExecObject)>
|
||||
ExecutionParameterInterface;
|
||||
|
||||
typedef vtkm::internal::Invocation<ExecutionParameterInterface, TestControlInterface,
|
||||
TestExecutionInterface1, 1, MyOutputToInputMapPortal,
|
||||
MyVisitArrayPortal>
|
||||
InvocationType1;
|
||||
|
||||
typedef vtkm::internal::Invocation<ExecutionParameterInterface, TestControlInterface,
|
||||
TestExecutionInterface2, 1, MyOutputToInputMapPortal,
|
||||
MyVisitArrayPortal>
|
||||
InvocationType2;
|
||||
|
||||
template <typename TaskType>
|
||||
static __global__ void ScheduleTaskSingular(TaskType task, vtkm::Id start, vtkm::Id end)
|
||||
{
|
||||
|
||||
const vtkm::Id index = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
if (index >= start && index < end)
|
||||
{
|
||||
task(index);
|
||||
}
|
||||
}
|
||||
|
||||
// Not a full worklet, but provides operators that we expect in a worklet.
|
||||
struct TestWorkletProxy : vtkm::exec::FunctorBase
|
||||
{
|
||||
VTKM_EXEC
|
||||
void operator()(vtkm::Id input, vtkm::Id& output) const { output = input + 100; }
|
||||
|
||||
VTKM_EXEC
|
||||
vtkm::Id operator()(vtkm::Id input) const { return input + 200; }
|
||||
|
||||
template <typename T, typename OutToInArrayType, typename VisitArrayType,
|
||||
typename InputDomainType, typename G>
|
||||
VTKM_EXEC vtkm::exec::arg::ThreadIndicesBasic GetThreadIndices(
|
||||
const T& threadIndex, const OutToInArrayType& outToIn, const VisitArrayType& visit,
|
||||
const InputDomainType&, const G& globalThreadIndexOffset) const
|
||||
{
|
||||
return vtkm::exec::arg::ThreadIndicesBasic(threadIndex, outToIn.Get(threadIndex),
|
||||
visit.Get(threadIndex), globalThreadIndexOffset);
|
||||
}
|
||||
};
|
||||
|
||||
#define ERROR_MESSAGE "Expected worklet error."
|
||||
|
||||
// Not a full worklet, but provides operators that we expect in a worklet.
|
||||
struct TestWorkletErrorProxy : vtkm::exec::FunctorBase
|
||||
{
|
||||
VTKM_EXEC
|
||||
void operator()(vtkm::Id, vtkm::Id) const { this->RaiseError(ERROR_MESSAGE); }
|
||||
|
||||
template <typename T, typename OutToInArrayType, typename VisitArrayType,
|
||||
typename InputDomainType, typename G>
|
||||
VTKM_EXEC vtkm::exec::arg::ThreadIndicesBasic GetThreadIndices(
|
||||
const T& threadIndex, const OutToInArrayType& outToIn, const VisitArrayType& visit,
|
||||
const InputDomainType&, const G& globalThreadIndexOffset) const
|
||||
{
|
||||
return vtkm::exec::arg::ThreadIndicesBasic(threadIndex, outToIn.Get(threadIndex),
|
||||
visit.Get(threadIndex), globalThreadIndexOffset);
|
||||
}
|
||||
};
|
||||
|
||||
// Check behavior of InvocationToFetch helper class.
|
||||
|
||||
VTKM_STATIC_ASSERT(
|
||||
(std::is_same<
|
||||
vtkm::exec::internal::detail::InvocationToFetch<vtkm::exec::arg::ThreadIndicesBasic,
|
||||
InvocationType1, 1>::type,
|
||||
vtkm::exec::arg::Fetch<TestFetchTagInput, vtkm::exec::arg::AspectTagDefault,
|
||||
vtkm::exec::arg::ThreadIndicesBasic, TestExecObject>>::type::value));
|
||||
|
||||
VTKM_STATIC_ASSERT(
|
||||
(std::is_same<
|
||||
vtkm::exec::internal::detail::InvocationToFetch<vtkm::exec::arg::ThreadIndicesBasic,
|
||||
InvocationType1, 2>::type,
|
||||
vtkm::exec::arg::Fetch<TestFetchTagOutput, vtkm::exec::arg::AspectTagDefault,
|
||||
vtkm::exec::arg::ThreadIndicesBasic, TestExecObject>>::type::value));
|
||||
|
||||
VTKM_STATIC_ASSERT(
|
||||
(std::is_same<
|
||||
vtkm::exec::internal::detail::InvocationToFetch<vtkm::exec::arg::ThreadIndicesBasic,
|
||||
InvocationType2, 0>::type,
|
||||
vtkm::exec::arg::Fetch<TestFetchTagOutput, vtkm::exec::arg::AspectTagDefault,
|
||||
vtkm::exec::arg::ThreadIndicesBasic, TestExecObject>>::type::value));
|
||||
|
||||
template <typename DeviceAdapter>
|
||||
void TestNormalFunctorInvoke()
|
||||
{
|
||||
std::cout << "Testing normal worklet invoke." << std::endl;
|
||||
|
||||
vtkm::Id inputTestValues[3] = { 5, 5, 6 };
|
||||
vtkm::Id outputTestValues[3] = { static_cast<vtkm::Id>(0xDEADDEAD),
|
||||
static_cast<vtkm::Id>(0xDEADDEAD),
|
||||
static_cast<vtkm::Id>(0xDEADDEAD) };
|
||||
|
||||
vtkm::cont::ArrayHandle<vtkm::Id> input = vtkm::cont::make_ArrayHandle(inputTestValues, 3);
|
||||
vtkm::cont::ArrayHandle<vtkm::Id> output;
|
||||
|
||||
vtkm::internal::FunctionInterface<void(TestExecObject, TestExecObject)> execObjects =
|
||||
vtkm::internal::make_FunctionInterface<void>(
|
||||
TestExecObject(input.PrepareForInPlace(DeviceAdapter())),
|
||||
TestExecObject(output.PrepareForOutput(3, DeviceAdapter())));
|
||||
|
||||
std::cout << " Try void return." << std::endl;
|
||||
typedef vtkm::exec::internal::TaskSingular<TestWorkletProxy, InvocationType1> TaskSingular1;
|
||||
TestWorkletProxy worklet;
|
||||
InvocationType1 invocation1(execObjects);
|
||||
|
||||
using TaskTypes = typename vtkm::cont::DeviceTaskTypes<DeviceAdapter>;
|
||||
auto task1 = TaskTypes::MakeTask(worklet, invocation1, vtkm::Id());
|
||||
|
||||
ScheduleTaskSingular<decltype(task1)><<<32, 256>>>(task1, 1, 2);
|
||||
cudaDeviceSynchronize();
|
||||
input.SyncControlArray();
|
||||
output.SyncControlArray();
|
||||
|
||||
output.CopyInto(outputTestValues, DeviceAdapter());
|
||||
|
||||
VTKM_TEST_ASSERT(inputTestValues[1] == 5, "Input value changed.");
|
||||
VTKM_TEST_ASSERT(outputTestValues[1] == inputTestValues[1] + 100 + 30,
|
||||
"Output value not set right.");
|
||||
|
||||
std::cout << " Try return value." << std::endl;
|
||||
|
||||
execObjects = vtkm::internal::make_FunctionInterface<void>(
|
||||
TestExecObject(input.PrepareForInPlace(DeviceAdapter())),
|
||||
TestExecObject(output.PrepareForOutput(3, DeviceAdapter())));
|
||||
|
||||
typedef vtkm::exec::internal::TaskSingular<TestWorkletProxy, InvocationType2> TaskSingular2;
|
||||
InvocationType2 invocation2(execObjects);
|
||||
|
||||
using TaskTypes = typename vtkm::cont::DeviceTaskTypes<DeviceAdapter>;
|
||||
auto task2 = TaskTypes::MakeTask(worklet, invocation2, vtkm::Id());
|
||||
|
||||
ScheduleTaskSingular<decltype(task2)><<<32, 256>>>(task2, 2, 3);
|
||||
cudaDeviceSynchronize();
|
||||
input.SyncControlArray();
|
||||
output.SyncControlArray();
|
||||
|
||||
output.CopyInto(outputTestValues, DeviceAdapter());
|
||||
|
||||
VTKM_TEST_ASSERT(inputTestValues[2] == 6, "Input value changed.");
|
||||
VTKM_TEST_ASSERT(outputTestValues[2] == inputTestValues[2] + 200 + 30 * 2,
|
||||
"Output value not set right.");
|
||||
}
|
||||
|
||||
template <typename DeviceAdapter>
|
||||
void TestErrorFunctorInvoke()
|
||||
{
|
||||
std::cout << "Testing invoke with an error raised in the worklet." << std::endl;
|
||||
|
||||
vtkm::Id inputTestValue = 5;
|
||||
vtkm::Id outputTestValue = static_cast<vtkm::Id>(0xDEADDEAD);
|
||||
|
||||
vtkm::cont::ArrayHandle<vtkm::Id> input = vtkm::cont::make_ArrayHandle(&inputTestValue, 1);
|
||||
vtkm::cont::ArrayHandle<vtkm::Id> output = vtkm::cont::make_ArrayHandle(&outputTestValue, 1);
|
||||
|
||||
vtkm::internal::FunctionInterface<void(TestExecObject, TestExecObject)> execObjects =
|
||||
vtkm::internal::make_FunctionInterface<void>(
|
||||
TestExecObject(input.PrepareForInPlace(DeviceAdapter())),
|
||||
TestExecObject(output.PrepareForInPlace(DeviceAdapter())));
|
||||
|
||||
typedef vtkm::exec::internal::TaskSingular<TestWorkletErrorProxy, InvocationType1> TaskSingular1;
|
||||
TestWorkletErrorProxy worklet;
|
||||
InvocationType1 invocation(execObjects);
|
||||
|
||||
using TaskTypes = typename vtkm::cont::DeviceTaskTypes<DeviceAdapter>;
|
||||
using Algorithm = vtkm::cont::DeviceAdapterAlgorithm<DeviceAdapter>;
|
||||
|
||||
auto task = TaskTypes::MakeTask(worklet, invocation, vtkm::Id());
|
||||
|
||||
vtkm::Id errorArraySize = 0;
|
||||
char* hostErrorPtr = nullptr;
|
||||
char* deviceErrorPtr = Algorithm::GetPinnedErrorArray(errorArraySize, &hostErrorPtr);
|
||||
|
||||
hostErrorPtr[0] = '\0';
|
||||
vtkm::exec::internal::ErrorMessageBuffer errorMessage(deviceErrorPtr, errorArraySize);
|
||||
task.SetErrorMessageBuffer(errorMessage);
|
||||
|
||||
ScheduleTaskSingular<decltype(task)><<<32, 256>>>(task, 1, 2);
|
||||
cudaDeviceSynchronize();
|
||||
|
||||
VTKM_TEST_ASSERT(errorMessage.IsErrorRaised(), "Error not raised correctly.");
|
||||
VTKM_TEST_ASSERT(hostErrorPtr == std::string(ERROR_MESSAGE), "Got wrong error message.");
|
||||
}
|
||||
|
||||
template <typename DeviceAdapter>
|
||||
void TestTaskSingular()
|
||||
{
|
||||
TestNormalFunctorInvoke<DeviceAdapter>();
|
||||
TestErrorFunctorInvoke<DeviceAdapter>();
|
||||
}
|
||||
|
||||
} // anonymous namespace
|
||||
|
||||
int UnitTestTaskSingularCuda(int, char* [])
|
||||
{
|
||||
return vtkm::cont::testing::Testing::Run(TestTaskSingular<vtkm::cont::DeviceAdapterTagCuda>);
|
||||
}
|
@ -36,6 +36,11 @@ namespace exec
|
||||
namespace internal
|
||||
{
|
||||
|
||||
// TaskSingular represents an execution pattern for a worklet
|
||||
// that is best expressed in terms of single dimension iteration space. Inside
|
||||
// this single dimension no order is preferred.
|
||||
//
|
||||
//
|
||||
template <typename WorkletType, typename InvocationType>
|
||||
class TaskSingular : public vtkm::exec::TaskBase
|
||||
{
|
||||
@ -48,6 +53,7 @@ public:
|
||||
, GlobalIndexOffset(globalIndexOffset)
|
||||
{
|
||||
}
|
||||
|
||||
VTKM_CONT
|
||||
void SetErrorMessageBuffer(const vtkm::exec::internal::ErrorMessageBuffer& buffer)
|
||||
{
|
||||
|
@ -28,8 +28,10 @@
|
||||
#ifndef vtk_m_exec_internal_WorkletInvokeFunctorDetail_h
|
||||
#define vtk_m_exec_internal_WorkletInvokeFunctorDetail_h
|
||||
|
||||
#if !defined(vtk_m_exec_internal_TaskSingular_h) && !defined(VTKM_TEST_HEADER_BUILD)
|
||||
#error WorkletInvokeFunctorDetail.h must be included from TaskSingular.h
|
||||
#if !defined(vtk_m_exec_internal_TaskSingular_h) && \
|
||||
!defined(vtk_m_exec_internal_TaskTiling_h) && \
|
||||
!defined(VTKM_TEST_HEADER_BUILD)
|
||||
#error WorkletInvokeFunctorDetail.h must be included from TaskSingular.h or TaskTiling.h
|
||||
#endif
|
||||
|
||||
#include <vtkm/internal/FunctionInterface.h>
|
||||
|
@ -40,8 +40,10 @@ $# Ignore the following comment. It is meant for the generated file.
|
||||
#ifndef vtk_m_exec_internal_WorkletInvokeFunctorDetail_h
|
||||
#define vtk_m_exec_internal_WorkletInvokeFunctorDetail_h
|
||||
|
||||
#if !defined(vtk_m_exec_internal_TaskSingular_h) && !defined(VTKM_TEST_HEADER_BUILD)
|
||||
#error WorkletInvokeFunctorDetail.h must be included from TaskSingular.h
|
||||
#if !defined(vtk_m_exec_internal_TaskSingular_h) && \
|
||||
!defined(vtk_m_exec_internal_TaskTiling_h) && \
|
||||
!defined(VTKM_TEST_HEADER_BUILD)
|
||||
#error WorkletInvokeFunctorDetail.h must be included from TaskSingular.h or TaskTiling.h
|
||||
#endif
|
||||
|
||||
#include <vtkm/internal/FunctionInterface.h>
|
||||
|
@ -20,6 +20,12 @@
|
||||
##
|
||||
##=============================================================================
|
||||
|
||||
set(headers
|
||||
TestingTaskTiling.h
|
||||
)
|
||||
|
||||
vtkm_declare_headers(${headers})
|
||||
|
||||
set(unit_tests
|
||||
UnitTestErrorMessageBuffer.cxx
|
||||
UnitTestTaskSingular.cxx
|
||||
|
448
vtkm/exec/internal/testing/TestingTaskTiling.h
Normal file
448
vtkm/exec/internal/testing/TestingTaskTiling.h
Normal file
@ -0,0 +1,448 @@
|
||||
//============================================================================
|
||||
// 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.
|
||||
//============================================================================
|
||||
|
||||
#include <vtkm/StaticAssert.h>
|
||||
|
||||
#include <vtkm/cont/DeviceAdapterAlgorithm.h>
|
||||
#include <vtkm/cont/testing/Testing.h>
|
||||
|
||||
#include <vtkm/exec/FunctorBase.h>
|
||||
#include <vtkm/exec/arg/BasicArg.h>
|
||||
#include <vtkm/exec/arg/Fetch.h>
|
||||
#include <vtkm/exec/arg/ThreadIndicesBasic.h>
|
||||
|
||||
#include <vtkm/internal/FunctionInterface.h>
|
||||
#include <vtkm/internal/Invocation.h>
|
||||
|
||||
#include <algorithm>
|
||||
#include <vector>
|
||||
|
||||
namespace vtkm
|
||||
{
|
||||
namespace exec
|
||||
{
|
||||
namespace internal
|
||||
{
|
||||
namespace testing
|
||||
{
|
||||
|
||||
struct TestExecObject
|
||||
{
|
||||
VTKM_EXEC_CONT
|
||||
TestExecObject()
|
||||
: Values(nullptr)
|
||||
{
|
||||
}
|
||||
|
||||
VTKM_EXEC_CONT
|
||||
TestExecObject(std::vector<vtkm::Id>& values)
|
||||
: Values(&values[0])
|
||||
{
|
||||
}
|
||||
|
||||
VTKM_EXEC_CONT
|
||||
TestExecObject(const TestExecObject& other) { Values = other.Values; }
|
||||
|
||||
vtkm::Id* Values;
|
||||
};
|
||||
|
||||
struct MyOutputToInputMapPortal
|
||||
{
|
||||
typedef vtkm::Id ValueType;
|
||||
VTKM_EXEC_CONT
|
||||
vtkm::Id Get(vtkm::Id index) const { return index; }
|
||||
};
|
||||
|
||||
struct MyVisitArrayPortal
|
||||
{
|
||||
typedef vtkm::IdComponent ValueType;
|
||||
vtkm::IdComponent Get(vtkm::Id) const { return 1; }
|
||||
};
|
||||
|
||||
struct TestFetchTagInput
|
||||
{
|
||||
};
|
||||
struct TestFetchTagOutput
|
||||
{
|
||||
};
|
||||
|
||||
// Missing TransportTag, but we are not testing that so we can leave it out.
|
||||
struct TestControlSignatureTagInput
|
||||
{
|
||||
typedef TestFetchTagInput FetchTag;
|
||||
};
|
||||
struct TestControlSignatureTagOutput
|
||||
{
|
||||
typedef TestFetchTagOutput FetchTag;
|
||||
};
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
namespace vtkm
|
||||
{
|
||||
namespace exec
|
||||
{
|
||||
namespace arg
|
||||
{
|
||||
|
||||
using namespace vtkm::exec::internal::testing;
|
||||
|
||||
template <>
|
||||
struct Fetch<TestFetchTagInput, vtkm::exec::arg::AspectTagDefault,
|
||||
vtkm::exec::arg::ThreadIndicesBasic, TestExecObject>
|
||||
{
|
||||
typedef vtkm::Id ValueType;
|
||||
|
||||
VTKM_EXEC
|
||||
ValueType Load(const vtkm::exec::arg::ThreadIndicesBasic& indices,
|
||||
const TestExecObject& execObject) const
|
||||
{
|
||||
return execObject.Values[indices.GetInputIndex()] + 10 * indices.GetInputIndex();
|
||||
}
|
||||
|
||||
VTKM_EXEC
|
||||
void Store(const vtkm::exec::arg::ThreadIndicesBasic&, const TestExecObject&, ValueType) const
|
||||
{
|
||||
// No-op
|
||||
}
|
||||
};
|
||||
|
||||
template <>
|
||||
struct Fetch<TestFetchTagOutput, vtkm::exec::arg::AspectTagDefault,
|
||||
vtkm::exec::arg::ThreadIndicesBasic, TestExecObject>
|
||||
{
|
||||
typedef vtkm::Id ValueType;
|
||||
|
||||
VTKM_EXEC
|
||||
ValueType Load(const vtkm::exec::arg::ThreadIndicesBasic&, const TestExecObject&) const
|
||||
{
|
||||
// No-op
|
||||
return ValueType();
|
||||
}
|
||||
|
||||
VTKM_EXEC
|
||||
void Store(const vtkm::exec::arg::ThreadIndicesBasic& indices, const TestExecObject& execObject,
|
||||
ValueType value) const
|
||||
{
|
||||
execObject.Values[indices.GetOutputIndex()] = value + 20 * indices.GetOutputIndex();
|
||||
}
|
||||
};
|
||||
}
|
||||
}
|
||||
} // vtkm::exec::arg
|
||||
|
||||
namespace vtkm
|
||||
{
|
||||
namespace exec
|
||||
{
|
||||
namespace internal
|
||||
{
|
||||
namespace testing
|
||||
{
|
||||
|
||||
typedef void TestControlSignature(TestControlSignatureTagInput, TestControlSignatureTagOutput);
|
||||
typedef vtkm::internal::FunctionInterface<TestControlSignature> TestControlInterface;
|
||||
|
||||
typedef void TestExecutionSignature1(vtkm::exec::arg::BasicArg<1>, vtkm::exec::arg::BasicArg<2>);
|
||||
typedef vtkm::internal::FunctionInterface<TestExecutionSignature1> TestExecutionInterface1;
|
||||
|
||||
typedef vtkm::exec::arg::BasicArg<2> TestExecutionSignature2(vtkm::exec::arg::BasicArg<1>);
|
||||
typedef vtkm::internal::FunctionInterface<TestExecutionSignature2> TestExecutionInterface2;
|
||||
|
||||
typedef vtkm::internal::FunctionInterface<void(TestExecObject, TestExecObject)>
|
||||
ExecutionParameterInterface;
|
||||
|
||||
typedef vtkm::internal::Invocation<ExecutionParameterInterface, TestControlInterface,
|
||||
TestExecutionInterface1, 1, MyOutputToInputMapPortal,
|
||||
MyVisitArrayPortal>
|
||||
InvocationType1;
|
||||
|
||||
typedef vtkm::internal::Invocation<ExecutionParameterInterface, TestControlInterface,
|
||||
TestExecutionInterface2, 1, MyOutputToInputMapPortal,
|
||||
MyVisitArrayPortal>
|
||||
InvocationType2;
|
||||
|
||||
// Not a full worklet, but provides operators that we expect in a worklet.
|
||||
struct TestWorkletProxy : vtkm::exec::FunctorBase
|
||||
{
|
||||
VTKM_EXEC
|
||||
void operator()(vtkm::Id input, vtkm::Id& output) const { output = input + 100; }
|
||||
|
||||
VTKM_EXEC
|
||||
vtkm::Id operator()(vtkm::Id input) const { return input + 200; }
|
||||
|
||||
template <typename OutToInArrayType, typename VisitArrayType, typename InputDomainType,
|
||||
typename G>
|
||||
VTKM_EXEC vtkm::exec::arg::ThreadIndicesBasic GetThreadIndices(
|
||||
const vtkm::Id& threadIndex, const OutToInArrayType& outToIn, const VisitArrayType& visit,
|
||||
const InputDomainType&, const G& globalThreadIndexOffset) const
|
||||
{
|
||||
return vtkm::exec::arg::ThreadIndicesBasic(threadIndex, outToIn.Get(threadIndex),
|
||||
visit.Get(threadIndex), globalThreadIndexOffset);
|
||||
}
|
||||
|
||||
template <typename OutToInArrayType, typename VisitArrayType, typename InputDomainType,
|
||||
typename G>
|
||||
VTKM_EXEC vtkm::exec::arg::ThreadIndicesBasic GetThreadIndices(
|
||||
const vtkm::Id3& threadIndex, const OutToInArrayType& outToIn, const VisitArrayType& visit,
|
||||
const InputDomainType&, const G& globalThreadIndexOffset) const
|
||||
{
|
||||
const vtkm::Id index = vtkm::dot(threadIndex, vtkm::Id3(1, 8, 64));
|
||||
return vtkm::exec::arg::ThreadIndicesBasic(index, outToIn.Get(index), visit.Get(index),
|
||||
globalThreadIndexOffset);
|
||||
}
|
||||
};
|
||||
|
||||
#define ERROR_MESSAGE "Expected worklet error."
|
||||
|
||||
// Not a full worklet, but provides operators that we expect in a worklet.
|
||||
struct TestWorkletErrorProxy : vtkm::exec::FunctorBase
|
||||
{
|
||||
VTKM_EXEC
|
||||
void operator()(vtkm::Id, vtkm::Id) const { this->RaiseError(ERROR_MESSAGE); }
|
||||
|
||||
template <typename OutToInArrayType, typename VisitArrayType, typename InputDomainType,
|
||||
typename G>
|
||||
VTKM_EXEC vtkm::exec::arg::ThreadIndicesBasic GetThreadIndices(
|
||||
const vtkm::Id& threadIndex, const OutToInArrayType& outToIn, const VisitArrayType& visit,
|
||||
const InputDomainType&, const G& globalThreadIndexOffset) const
|
||||
{
|
||||
return vtkm::exec::arg::ThreadIndicesBasic(threadIndex, outToIn.Get(threadIndex),
|
||||
visit.Get(threadIndex), globalThreadIndexOffset);
|
||||
}
|
||||
|
||||
template <typename OutToInArrayType, typename VisitArrayType, typename InputDomainType,
|
||||
typename G>
|
||||
VTKM_EXEC vtkm::exec::arg::ThreadIndicesBasic GetThreadIndices(
|
||||
const vtkm::Id3& threadIndex, const OutToInArrayType& outToIn, const VisitArrayType& visit,
|
||||
const InputDomainType&, const G& globalThreadIndexOffset) const
|
||||
{
|
||||
const vtkm::Id index = vtkm::dot(threadIndex, vtkm::Id3(1, 8, 64));
|
||||
return vtkm::exec::arg::ThreadIndicesBasic(index, outToIn.Get(index), visit.Get(index),
|
||||
globalThreadIndexOffset);
|
||||
}
|
||||
};
|
||||
|
||||
template <typename DeviceAdapter>
|
||||
void Test1DNormalTaskTilingInvoke()
|
||||
{
|
||||
|
||||
std::cout << "Testing TaskTiling1D." << std::endl;
|
||||
|
||||
std::vector<vtkm::Id> inputTestValues(100, 5);
|
||||
std::vector<vtkm::Id> outputTestValues(100, static_cast<vtkm::Id>(0xDEADDEAD));
|
||||
vtkm::internal::FunctionInterface<void(TestExecObject, TestExecObject)> execObjects =
|
||||
vtkm::internal::make_FunctionInterface<void>(TestExecObject(inputTestValues),
|
||||
TestExecObject(outputTestValues));
|
||||
|
||||
std::cout << " Try void return." << std::endl;
|
||||
TestWorkletProxy worklet;
|
||||
InvocationType1 invocation1(execObjects);
|
||||
|
||||
using TaskTypes = typename vtkm::cont::DeviceTaskTypes<DeviceAdapter>;
|
||||
auto task1 = TaskTypes::MakeTask(worklet, invocation1, vtkm::Id());
|
||||
|
||||
vtkm::exec::internal::ErrorMessageBuffer errorMessage(nullptr, 0);
|
||||
task1.SetErrorMessageBuffer(errorMessage);
|
||||
|
||||
task1(0, 90);
|
||||
task1(90, 99);
|
||||
task1(99, 100); //verify single value ranges work
|
||||
|
||||
for (std::size_t i = 0; i < 100; ++i)
|
||||
{
|
||||
VTKM_TEST_ASSERT(inputTestValues[i] == 5, "Input value changed.");
|
||||
VTKM_TEST_ASSERT(outputTestValues[i] ==
|
||||
inputTestValues[i] + 100 + (30 * static_cast<vtkm::Id>(i)),
|
||||
"Output value not set right.");
|
||||
}
|
||||
|
||||
std::cout << " Try return value." << std::endl;
|
||||
std::fill(inputTestValues.begin(), inputTestValues.end(), 6);
|
||||
std::fill(outputTestValues.begin(), outputTestValues.end(), static_cast<vtkm::Id>(0xDEADDEAD));
|
||||
|
||||
InvocationType2 invocation2(execObjects);
|
||||
|
||||
using TaskTypes = typename vtkm::cont::DeviceTaskTypes<DeviceAdapter>;
|
||||
auto task2 = TaskTypes::MakeTask(worklet, invocation2, vtkm::Id());
|
||||
|
||||
task2.SetErrorMessageBuffer(errorMessage);
|
||||
|
||||
task2(0, 0); //verify zero value ranges work
|
||||
task2(0, 90);
|
||||
task2(90, 100);
|
||||
|
||||
task2(0, 100); //verify that you can invoke worklets multiple times
|
||||
|
||||
for (std::size_t i = 0; i < 100; ++i)
|
||||
{
|
||||
VTKM_TEST_ASSERT(inputTestValues[i] == 6, "Input value changed.");
|
||||
VTKM_TEST_ASSERT(outputTestValues[i] ==
|
||||
inputTestValues[i] + 200 + (30 * static_cast<vtkm::Id>(i)),
|
||||
"Output value not set right.");
|
||||
}
|
||||
}
|
||||
|
||||
template <typename DeviceAdapter>
|
||||
void Test1DErrorTaskTilingInvoke()
|
||||
{
|
||||
|
||||
std::cout << "Testing TaskTiling1D with an error raised in the worklet." << std::endl;
|
||||
|
||||
std::vector<vtkm::Id> inputTestValues(100, 5);
|
||||
std::vector<vtkm::Id> outputTestValues(100, static_cast<vtkm::Id>(0xDEADDEAD));
|
||||
|
||||
TestExecObject arg1(inputTestValues);
|
||||
TestExecObject arg2(outputTestValues);
|
||||
|
||||
vtkm::internal::FunctionInterface<void(TestExecObject, TestExecObject)> execObjects =
|
||||
vtkm::internal::make_FunctionInterface<void>(arg1, arg2);
|
||||
|
||||
TestWorkletErrorProxy worklet;
|
||||
InvocationType1 invocation(execObjects);
|
||||
|
||||
using TaskTypes = typename vtkm::cont::DeviceTaskTypes<DeviceAdapter>;
|
||||
auto task = TaskTypes::MakeTask(worklet, invocation, vtkm::Id());
|
||||
|
||||
char message[1024];
|
||||
message[0] = '\0';
|
||||
vtkm::exec::internal::ErrorMessageBuffer errorMessage(message, 1024);
|
||||
task.SetErrorMessageBuffer(errorMessage);
|
||||
|
||||
task(0, 100);
|
||||
|
||||
VTKM_TEST_ASSERT(errorMessage.IsErrorRaised(), "Error not raised correctly.");
|
||||
VTKM_TEST_ASSERT(message == std::string(ERROR_MESSAGE), "Got wrong error message.");
|
||||
}
|
||||
|
||||
template <typename DeviceAdapter>
|
||||
void Test3DNormalTaskTilingInvoke()
|
||||
{
|
||||
std::cout << "Testing TaskTiling3D." << std::endl;
|
||||
|
||||
std::vector<vtkm::Id> inputTestValues((8 * 8 * 8), 5);
|
||||
std::vector<vtkm::Id> outputTestValues((8 * 8 * 8), static_cast<vtkm::Id>(0xDEADDEAD));
|
||||
vtkm::internal::FunctionInterface<void(TestExecObject, TestExecObject)> execObjects =
|
||||
vtkm::internal::make_FunctionInterface<void>(TestExecObject(inputTestValues),
|
||||
TestExecObject(outputTestValues));
|
||||
|
||||
std::cout << " Try void return." << std::endl;
|
||||
|
||||
TestWorkletProxy worklet;
|
||||
InvocationType1 invocation1(execObjects);
|
||||
|
||||
using TaskTypes = typename vtkm::cont::DeviceTaskTypes<DeviceAdapter>;
|
||||
auto task1 = TaskTypes::MakeTask(worklet, invocation1, vtkm::Id3());
|
||||
for (vtkm::Id k = 0; k < 8; ++k)
|
||||
{
|
||||
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);
|
||||
}
|
||||
}
|
||||
|
||||
for (std::size_t i = 0; i < (8 * 8 * 8); ++i)
|
||||
{
|
||||
VTKM_TEST_ASSERT(inputTestValues[i] == 5, "Input value changed.");
|
||||
VTKM_TEST_ASSERT(outputTestValues[i] ==
|
||||
inputTestValues[i] + 100 + (30 * static_cast<vtkm::Id>(i)),
|
||||
"Output value not set right.");
|
||||
}
|
||||
|
||||
std::cout << " Try return value." << std::endl;
|
||||
std::fill(inputTestValues.begin(), inputTestValues.end(), 6);
|
||||
std::fill(outputTestValues.begin(), outputTestValues.end(), static_cast<vtkm::Id>(0xDEADDEAD));
|
||||
|
||||
InvocationType2 invocation2(execObjects);
|
||||
using TaskTypes = typename vtkm::cont::DeviceTaskTypes<DeviceAdapter>;
|
||||
auto task2 = TaskTypes::MakeTask(worklet, invocation2, vtkm::Id3());
|
||||
|
||||
//verify that linear order of values being processed is not presumed
|
||||
for (vtkm::Id i = 0; i < 8; ++i)
|
||||
{
|
||||
for (vtkm::Id j = 0; j < 8; ++j)
|
||||
{
|
||||
for (vtkm::Id k = 0; k < 8; ++k)
|
||||
{
|
||||
task2(i, i + 1, j, k);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
for (std::size_t i = 0; i < (8 * 8 * 8); ++i)
|
||||
{
|
||||
VTKM_TEST_ASSERT(inputTestValues[i] == 6, "Input value changed.");
|
||||
VTKM_TEST_ASSERT(outputTestValues[i] ==
|
||||
inputTestValues[i] + 200 + (30 * static_cast<vtkm::Id>(i)),
|
||||
"Output value not set right.");
|
||||
}
|
||||
}
|
||||
|
||||
template <typename DeviceAdapter>
|
||||
void Test3DErrorTaskTilingInvoke()
|
||||
{
|
||||
std::cout << "Testing TaskTiling3D with an error raised in the worklet." << std::endl;
|
||||
|
||||
std::vector<vtkm::Id> inputTestValues((8 * 8 * 8), 5);
|
||||
std::vector<vtkm::Id> outputTestValues((8 * 8 * 8), static_cast<vtkm::Id>(0xDEADDEAD));
|
||||
vtkm::internal::FunctionInterface<void(TestExecObject, TestExecObject)> execObjects =
|
||||
vtkm::internal::make_FunctionInterface<void>(TestExecObject(inputTestValues),
|
||||
TestExecObject(outputTestValues));
|
||||
|
||||
TestWorkletErrorProxy worklet;
|
||||
InvocationType1 invocation(execObjects);
|
||||
|
||||
using TaskTypes = typename vtkm::cont::DeviceTaskTypes<DeviceAdapter>;
|
||||
auto task1 = TaskTypes::MakeTask(worklet, invocation, vtkm::Id3());
|
||||
|
||||
char message[1024];
|
||||
message[0] = '\0';
|
||||
vtkm::exec::internal::ErrorMessageBuffer errorMessage(message, 1024);
|
||||
task1.SetErrorMessageBuffer(errorMessage);
|
||||
|
||||
for (vtkm::Id k = 0; k < 8; ++k)
|
||||
{
|
||||
for (vtkm::Id j = 0; j < 8; ++j)
|
||||
{
|
||||
task1(0, 8, j, k);
|
||||
}
|
||||
}
|
||||
|
||||
VTKM_TEST_ASSERT(errorMessage.IsErrorRaised(), "Error not raised correctly.");
|
||||
VTKM_TEST_ASSERT(message == std::string(ERROR_MESSAGE), "Got wrong error message.");
|
||||
}
|
||||
|
||||
template <typename DeviceAdapter>
|
||||
void TestTaskTiling()
|
||||
{
|
||||
Test1DNormalTaskTilingInvoke<DeviceAdapter>();
|
||||
Test1DErrorTaskTilingInvoke<DeviceAdapter>();
|
||||
|
||||
Test3DNormalTaskTilingInvoke<DeviceAdapter>();
|
||||
Test3DErrorTaskTilingInvoke<DeviceAdapter>();
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
@ -235,8 +235,10 @@ void TestNormalFunctorInvoke()
|
||||
inputTestValue = 5;
|
||||
outputTestValue = static_cast<vtkm::Id>(0xDEADDEAD);
|
||||
typedef vtkm::exec::internal::TaskSingular<TestWorkletProxy, InvocationType1> TaskSingular1;
|
||||
TaskSingular1 taskInvokeWorklet1 =
|
||||
TaskSingular1(TestWorkletProxy(), InvocationType1(execObjects));
|
||||
TestWorkletProxy worklet;
|
||||
InvocationType1 invocation1(execObjects);
|
||||
TaskSingular1 taskInvokeWorklet1(worklet, invocation1);
|
||||
|
||||
taskInvokeWorklet1(1);
|
||||
VTKM_TEST_ASSERT(inputTestValue == 5, "Input value changed.");
|
||||
VTKM_TEST_ASSERT(outputTestValue == inputTestValue + 100 + 30, "Output value not set right.");
|
||||
@ -245,8 +247,9 @@ void TestNormalFunctorInvoke()
|
||||
inputTestValue = 6;
|
||||
outputTestValue = static_cast<vtkm::Id>(0xDEADDEAD);
|
||||
typedef vtkm::exec::internal::TaskSingular<TestWorkletProxy, InvocationType2> TaskSingular2;
|
||||
TaskSingular2 taskInvokeWorklet2 =
|
||||
TaskSingular2(TestWorkletProxy(), InvocationType2(execObjects));
|
||||
InvocationType2 invocation2(execObjects);
|
||||
TaskSingular2 taskInvokeWorklet2(worklet, invocation2);
|
||||
|
||||
taskInvokeWorklet2(2);
|
||||
VTKM_TEST_ASSERT(inputTestValue == 6, "Input value changed.");
|
||||
VTKM_TEST_ASSERT(outputTestValue == inputTestValue + 200 + 30 * 2, "Output value not set right.");
|
||||
@ -263,8 +266,10 @@ void TestErrorFunctorInvoke()
|
||||
TestExecObject(&outputTestValue));
|
||||
|
||||
typedef vtkm::exec::internal::TaskSingular<TestWorkletErrorProxy, InvocationType1> TaskSingular1;
|
||||
TaskSingular1 taskInvokeWorklet1 =
|
||||
TaskSingular1(TestWorkletErrorProxy(), InvocationType1(execObjects));
|
||||
TestWorkletErrorProxy worklet;
|
||||
InvocationType1 invocation(execObjects);
|
||||
TaskSingular1 taskInvokeWorklet1 = TaskSingular1(worklet, invocation);
|
||||
|
||||
char message[1024];
|
||||
message[0] = '\0';
|
||||
vtkm::exec::internal::ErrorMessageBuffer errorMessage(message, 1024);
|
||||
|
22
vtkm/exec/serial/CMakeLists.txt
Normal file
22
vtkm/exec/serial/CMakeLists.txt
Normal file
@ -0,0 +1,22 @@
|
||||
##============================================================================
|
||||
## 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.
|
||||
##============================================================================
|
||||
|
||||
#-----------------------------------------------------------------------------
|
||||
add_subdirectory(internal)
|
28
vtkm/exec/serial/internal/CMakeLists.txt
Normal file
28
vtkm/exec/serial/internal/CMakeLists.txt
Normal file
@ -0,0 +1,28 @@
|
||||
##============================================================================
|
||||
## 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.
|
||||
##============================================================================
|
||||
|
||||
set(headers
|
||||
TaskTiling.h
|
||||
)
|
||||
|
||||
vtkm_declare_headers(${headers})
|
||||
|
||||
#-----------------------------------------------------------------------------
|
||||
add_subdirectory(testing)
|
301
vtkm/exec/serial/internal/TaskTiling.h
Normal file
301
vtkm/exec/serial/internal/TaskTiling.h
Normal file
@ -0,0 +1,301 @@
|
||||
//============================================================================
|
||||
// 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.
|
||||
//============================================================================
|
||||
|
||||
#ifndef vtk_m_exec_serial_internal_TaskTiling_h
|
||||
#define vtk_m_exec_serial_internal_TaskTiling_h
|
||||
|
||||
#include <vtkm/exec/TaskBase.h>
|
||||
|
||||
//Todo: rename this header to TaskInvokeWorkletDetail.h
|
||||
#include <vtkm/exec/internal/WorkletInvokeFunctorDetail.h>
|
||||
|
||||
namespace vtkm
|
||||
{
|
||||
namespace exec
|
||||
{
|
||||
namespace serial
|
||||
{
|
||||
namespace internal
|
||||
{
|
||||
|
||||
template <typename WType>
|
||||
void TaskTilingSetErrorBuffer(void* w, const vtkm::exec::internal::ErrorMessageBuffer& buffer)
|
||||
{
|
||||
using WorkletType = typename std::remove_cv<WType>::type;
|
||||
WorkletType* const worklet = static_cast<WorkletType*>(w);
|
||||
worklet->SetErrorMessageBuffer(buffer);
|
||||
}
|
||||
|
||||
template <typename WType, typename IType>
|
||||
void TaskTiling1DExecute(void* w, void* const v, vtkm::Id globalIndexOffset, vtkm::Id start,
|
||||
vtkm::Id end)
|
||||
{
|
||||
using WorkletType = typename std::remove_cv<WType>::type;
|
||||
using InvocationType = typename std::remove_cv<IType>::type;
|
||||
|
||||
WorkletType const* const worklet = static_cast<WorkletType*>(w);
|
||||
InvocationType const* const invocation = static_cast<InvocationType*>(v);
|
||||
|
||||
for (vtkm::Id index = start; index < end; ++index)
|
||||
{
|
||||
//Todo: rename this function to DoTaskInvokeWorklet
|
||||
vtkm::exec::internal::detail::DoWorkletInvokeFunctor(
|
||||
*worklet, *invocation,
|
||||
worklet->GetThreadIndices(index, invocation->OutputToInputMap, invocation->VisitArray,
|
||||
invocation->GetInputDomain(), globalIndexOffset));
|
||||
}
|
||||
}
|
||||
|
||||
template <typename FType>
|
||||
void FunctorTiling1DExecute(void* f, void* const, vtkm::Id, vtkm::Id start, vtkm::Id end)
|
||||
{
|
||||
using FunctorType = typename std::remove_cv<FType>::type;
|
||||
FunctorType const* const functor = static_cast<FunctorType*>(f);
|
||||
|
||||
for (vtkm::Id index = start; index < end; ++index)
|
||||
{
|
||||
functor->operator()(index);
|
||||
}
|
||||
}
|
||||
|
||||
template <typename WType, typename IType>
|
||||
void TaskTiling3DExecute(void* w, void* const v, vtkm::Id globalIndexOffset, vtkm::Id istart,
|
||||
vtkm::Id iend, vtkm::Id j, vtkm::Id k)
|
||||
{
|
||||
using WorkletType = typename std::remove_cv<WType>::type;
|
||||
using InvocationType = typename std::remove_cv<IType>::type;
|
||||
|
||||
WorkletType const* const worklet = static_cast<WorkletType*>(w);
|
||||
InvocationType const* const invocation = static_cast<InvocationType*>(v);
|
||||
|
||||
vtkm::Id3 index(istart, j, k);
|
||||
for (vtkm::Id i = istart; i < iend; ++i)
|
||||
{
|
||||
index[0] = i;
|
||||
//Todo: rename this function to DoTaskInvokeWorklet
|
||||
vtkm::exec::internal::detail::DoWorkletInvokeFunctor(
|
||||
*worklet, *invocation,
|
||||
worklet->GetThreadIndices(index, invocation->OutputToInputMap, invocation->VisitArray,
|
||||
invocation->GetInputDomain(), globalIndexOffset));
|
||||
}
|
||||
}
|
||||
|
||||
template <typename FType>
|
||||
void FunctorTiling3DExecute(void* f, void* const, vtkm::Id, vtkm::Id istart, vtkm::Id iend,
|
||||
vtkm::Id j, vtkm::Id k)
|
||||
{
|
||||
using FunctorType = typename std::remove_cv<FType>::type;
|
||||
FunctorType const* const functor = static_cast<FunctorType*>(f);
|
||||
|
||||
vtkm::Id3 index(istart, j, k);
|
||||
for (vtkm::Id i = istart; i < iend; ++i)
|
||||
{
|
||||
index[0] = i;
|
||||
functor->operator()(index);
|
||||
}
|
||||
}
|
||||
|
||||
// TaskTiling1D represents an execution pattern for a worklet
|
||||
// that is best expressed in terms of single dimension iteration space. TaskTiling1D
|
||||
// also states that for best performance a linear consecutive range of values
|
||||
// should be given to the worklet for optimal performance.
|
||||
//
|
||||
// Note: The worklet and invocation must have a lifetime that is at least
|
||||
// as long as the Task
|
||||
class VTKM_ALWAYS_EXPORT TaskTiling1D : public vtkm::exec::TaskBase
|
||||
{
|
||||
public:
|
||||
TaskTiling1D()
|
||||
: Worklet(nullptr)
|
||||
, Invocation(nullptr)
|
||||
, GlobalIndexOffset(0)
|
||||
{
|
||||
}
|
||||
|
||||
/// This constructor supports general functors that which have a call
|
||||
/// operator with the signature:
|
||||
/// operator()(vtkm::Id)
|
||||
template <typename FunctorType>
|
||||
TaskTiling1D(FunctorType& functor)
|
||||
: Worklet(nullptr)
|
||||
, Invocation(nullptr)
|
||||
, ExecuteFunction(nullptr)
|
||||
, SetErrorBufferFunction(nullptr)
|
||||
, GlobalIndexOffset(0)
|
||||
{
|
||||
//Setup the execute and set error buffer function pointers
|
||||
this->ExecuteFunction = &FunctorTiling1DExecute<FunctorType>;
|
||||
this->SetErrorBufferFunction = &TaskTilingSetErrorBuffer<FunctorType>;
|
||||
|
||||
//Bind the Worklet to void*
|
||||
this->Worklet = (void*)&functor;
|
||||
}
|
||||
|
||||
/// This constructor supports any vtkm worklet and the associated invocation
|
||||
/// parameters that go along with it
|
||||
template <typename WorkletType, typename InvocationType>
|
||||
TaskTiling1D(WorkletType& worklet, const InvocationType& invocation,
|
||||
const vtkm::Id& globalIndexOffset = 0)
|
||||
: Worklet(nullptr)
|
||||
, Invocation(nullptr)
|
||||
, ExecuteFunction(nullptr)
|
||||
, SetErrorBufferFunction(nullptr)
|
||||
, GlobalIndexOffset(globalIndexOffset)
|
||||
{
|
||||
//Setup the execute and set error buffer function pointers
|
||||
this->ExecuteFunction = &TaskTiling1DExecute<WorkletType, InvocationType>;
|
||||
this->SetErrorBufferFunction = &TaskTilingSetErrorBuffer<WorkletType>;
|
||||
|
||||
//Bind the Worklet and Invocation to void*
|
||||
this->Worklet = (void*)&worklet;
|
||||
this->Invocation = (void*)&invocation;
|
||||
}
|
||||
|
||||
/// explicit Copy constructor.
|
||||
/// Note this required so that compilers don't use the templated constructor
|
||||
/// as the copy constructor which will cause compile issues
|
||||
TaskTiling1D(const TaskTiling1D& task)
|
||||
: Worklet(task.Worklet)
|
||||
, Invocation(task.Invocation)
|
||||
, ExecuteFunction(task.ExecuteFunction)
|
||||
, SetErrorBufferFunction(task.SetErrorBufferFunction)
|
||||
, GlobalIndexOffset(task.GlobalIndexOffset)
|
||||
{
|
||||
}
|
||||
|
||||
void SetErrorMessageBuffer(const vtkm::exec::internal::ErrorMessageBuffer& buffer)
|
||||
{
|
||||
this->SetErrorBufferFunction(this->Worklet, buffer);
|
||||
}
|
||||
|
||||
void operator()(vtkm::Id start, vtkm::Id end) const
|
||||
{
|
||||
this->ExecuteFunction(this->Worklet, this->Invocation, this->GlobalIndexOffset, start, end);
|
||||
}
|
||||
|
||||
protected:
|
||||
void* Worklet;
|
||||
void* Invocation;
|
||||
|
||||
typedef void (*ExecuteSignature)(void*, void* const, vtkm::Id, vtkm::Id, vtkm::Id);
|
||||
ExecuteSignature ExecuteFunction;
|
||||
|
||||
typedef void (*SetErrorBufferSignature)(void*, const vtkm::exec::internal::ErrorMessageBuffer&);
|
||||
SetErrorBufferSignature SetErrorBufferFunction;
|
||||
|
||||
const vtkm::Id GlobalIndexOffset;
|
||||
};
|
||||
|
||||
// TaskTiling3D represents an execution pattern for a worklet
|
||||
// that is best expressed in terms of an 3 dimensional iteration space. TaskTiling3D
|
||||
// also states that for best performance a linear consecutive range of values
|
||||
// in the X dimension should be given to the worklet for optimal performance.
|
||||
//
|
||||
// Note: The worklet and invocation must have a lifetime that is at least
|
||||
// as long as the Task
|
||||
class VTKM_ALWAYS_EXPORT TaskTiling3D : public vtkm::exec::TaskBase
|
||||
{
|
||||
public:
|
||||
TaskTiling3D()
|
||||
: Worklet(nullptr)
|
||||
, Invocation(nullptr)
|
||||
, GlobalIndexOffset(0)
|
||||
{
|
||||
}
|
||||
|
||||
/// This constructor supports general functors that which have a call
|
||||
/// operator with the signature:
|
||||
/// operator()(vtkm::Id)
|
||||
template <typename FunctorType>
|
||||
TaskTiling3D(FunctorType& functor)
|
||||
: Worklet(nullptr)
|
||||
, Invocation(nullptr)
|
||||
, ExecuteFunction(nullptr)
|
||||
, SetErrorBufferFunction(nullptr)
|
||||
, GlobalIndexOffset(0)
|
||||
{
|
||||
//Setup the execute and set error buffer function pointers
|
||||
this->ExecuteFunction = &FunctorTiling3DExecute<FunctorType>;
|
||||
this->SetErrorBufferFunction = &TaskTilingSetErrorBuffer<FunctorType>;
|
||||
|
||||
//Bind the Worklet to void*
|
||||
this->Worklet = (void*)&functor;
|
||||
}
|
||||
|
||||
template <typename WorkletType, typename InvocationType>
|
||||
TaskTiling3D(WorkletType& worklet, const InvocationType& invocation,
|
||||
const vtkm::Id& globalIndexOffset = 0)
|
||||
: Worklet(nullptr)
|
||||
, Invocation(nullptr)
|
||||
, ExecuteFunction(nullptr)
|
||||
, SetErrorBufferFunction(nullptr)
|
||||
, GlobalIndexOffset(globalIndexOffset)
|
||||
{
|
||||
// Setup the execute and set error buffer function pointers
|
||||
this->ExecuteFunction = &TaskTiling3DExecute<WorkletType, InvocationType>;
|
||||
this->SetErrorBufferFunction = &TaskTilingSetErrorBuffer<WorkletType>;
|
||||
|
||||
// At this point we bind the Worklet and Invocation to void*
|
||||
this->Worklet = (void*)&worklet;
|
||||
this->Invocation = (void*)&invocation;
|
||||
}
|
||||
|
||||
/// explicit Copy constructor.
|
||||
/// Note this required so that compilers don't use the templated constructor
|
||||
/// as the copy constructor which will cause compile issues
|
||||
TaskTiling3D(const TaskTiling3D& task)
|
||||
: Worklet(task.Worklet)
|
||||
, Invocation(task.Invocation)
|
||||
, ExecuteFunction(task.ExecuteFunction)
|
||||
, SetErrorBufferFunction(task.SetErrorBufferFunction)
|
||||
, GlobalIndexOffset(task.GlobalIndexOffset)
|
||||
{
|
||||
}
|
||||
|
||||
void SetErrorMessageBuffer(const vtkm::exec::internal::ErrorMessageBuffer& buffer)
|
||||
{
|
||||
this->SetErrorBufferFunction(this->Worklet, buffer);
|
||||
}
|
||||
|
||||
void operator()(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);
|
||||
}
|
||||
|
||||
protected:
|
||||
void* Worklet;
|
||||
void* Invocation;
|
||||
|
||||
typedef void (*ExecuteSignature)(void*, void* const, vtkm::Id, vtkm::Id, vtkm::Id, vtkm::Id,
|
||||
vtkm::Id);
|
||||
ExecuteSignature ExecuteFunction;
|
||||
|
||||
typedef void (*SetErrorBufferSignature)(void*, const vtkm::exec::internal::ErrorMessageBuffer&);
|
||||
SetErrorBufferSignature SetErrorBufferFunction;
|
||||
|
||||
const vtkm::Id GlobalIndexOffset;
|
||||
};
|
||||
}
|
||||
}
|
||||
}
|
||||
} // vtkm::exec::serial::internal
|
||||
|
||||
#endif //vtk_m_exec_serial_internal_TaskTiling_h
|
26
vtkm/exec/serial/internal/testing/CMakeLists.txt
Normal file
26
vtkm/exec/serial/internal/testing/CMakeLists.txt
Normal file
@ -0,0 +1,26 @@
|
||||
##=============================================================================
|
||||
##
|
||||
## 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.
|
||||
##
|
||||
##=============================================================================
|
||||
|
||||
set(unit_tests
|
||||
UnitTestTaskTilingSerial.cxx
|
||||
)
|
||||
vtkm_unit_tests(SOURCES ${unit_tests})
|
@ -0,0 +1,29 @@
|
||||
//============================================================================
|
||||
// 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.
|
||||
//============================================================================
|
||||
|
||||
#include <vtkm/cont/serial/DeviceAdapterSerial.h>
|
||||
#include <vtkm/exec/internal/testing/TestingTaskTiling.h>
|
||||
|
||||
int UnitTestTaskTilingSerial(int, char* [])
|
||||
{
|
||||
|
||||
return vtkm::cont::testing::Testing::Run(
|
||||
vtkm::exec::internal::testing::TestTaskTiling<vtkm::cont::DeviceAdapterTagSerial>);
|
||||
}
|
22
vtkm/exec/tbb/CMakeLists.txt
Normal file
22
vtkm/exec/tbb/CMakeLists.txt
Normal file
@ -0,0 +1,22 @@
|
||||
##============================================================================
|
||||
## 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.
|
||||
##============================================================================
|
||||
|
||||
#-----------------------------------------------------------------------------
|
||||
add_subdirectory(internal)
|
30
vtkm/exec/tbb/internal/CMakeLists.txt
Normal file
30
vtkm/exec/tbb/internal/CMakeLists.txt
Normal file
@ -0,0 +1,30 @@
|
||||
##============================================================================
|
||||
## 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.
|
||||
##============================================================================
|
||||
|
||||
set(headers
|
||||
TaskTiling.h
|
||||
)
|
||||
|
||||
vtkm_declare_headers(${headers})
|
||||
|
||||
#-----------------------------------------------------------------------------
|
||||
if (VTKm_ENABLE_TBB)
|
||||
add_subdirectory(testing)
|
||||
endif()
|
41
vtkm/exec/tbb/internal/TaskTiling.h
Normal file
41
vtkm/exec/tbb/internal/TaskTiling.h
Normal file
@ -0,0 +1,41 @@
|
||||
//============================================================================
|
||||
// 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 2017 Sandia Corporation.
|
||||
// Copyright 2017 UT-Battelle, LLC.
|
||||
// Copyright 2017 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.
|
||||
//============================================================================
|
||||
#ifndef vtk_m_exec_tbb_internal_TaskTiling_h
|
||||
#define vtk_m_exec_tbb_internal_TaskTiling_h
|
||||
|
||||
#include <vtkm/exec/serial/internal/TaskTiling.h>
|
||||
|
||||
namespace vtkm
|
||||
{
|
||||
namespace exec
|
||||
{
|
||||
namespace tbb
|
||||
{
|
||||
namespace internal
|
||||
{
|
||||
|
||||
using TaskTiling1D = vtkm::exec::serial::internal::TaskTiling1D;
|
||||
using TaskTiling3D = vtkm::exec::serial::internal::TaskTiling3D;
|
||||
}
|
||||
}
|
||||
}
|
||||
} // namespace vtkm::exec::tbb::internal
|
||||
|
||||
#endif //vtk_m_exec_tbb_internal_TaskTiling_h
|
26
vtkm/exec/tbb/internal/testing/CMakeLists.txt
Normal file
26
vtkm/exec/tbb/internal/testing/CMakeLists.txt
Normal file
@ -0,0 +1,26 @@
|
||||
##=============================================================================
|
||||
##
|
||||
## 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.
|
||||
##
|
||||
##=============================================================================
|
||||
|
||||
set(unit_tests
|
||||
UnitTestTaskTilingTBB.cxx
|
||||
)
|
||||
vtkm_unit_tests(SOURCES ${unit_tests})
|
28
vtkm/exec/tbb/internal/testing/UnitTestTaskTilingTBB.cxx
Normal file
28
vtkm/exec/tbb/internal/testing/UnitTestTaskTilingTBB.cxx
Normal file
@ -0,0 +1,28 @@
|
||||
//============================================================================
|
||||
// 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.
|
||||
//============================================================================
|
||||
|
||||
#include <vtkm/cont/tbb/DeviceAdapterTBB.h>
|
||||
#include <vtkm/exec/internal/testing/TestingTaskTiling.h>
|
||||
|
||||
int UnitTestTaskTilingTBB(int, char* [])
|
||||
{
|
||||
return vtkm::cont::testing::Testing::Run(
|
||||
vtkm::exec::internal::testing::TestTaskTiling<vtkm::cont::DeviceAdapterTagTBB>);
|
||||
}
|
@ -281,14 +281,17 @@ private:
|
||||
RangeType globalIndexOffset, DeviceAdapter) const
|
||||
{
|
||||
using Algorithm = vtkm::cont::DeviceAdapterAlgorithm<DeviceAdapter>;
|
||||
using TaskTypes = typename vtkm::cont::DeviceTaskTypes<DeviceAdapter>;
|
||||
|
||||
// The WorkletInvokeFunctor class handles the magic of fetching values
|
||||
// for each instance and calling the worklet's function. So just create
|
||||
// a WorkletInvokeFunctor and schedule it with the device adapter.
|
||||
using TaskType = vtkm::exec::internal::TaskSingular<WorkletType, Invocation>;
|
||||
TaskType task = TaskType(this->Worklet, invocation, globalIndexOffset);
|
||||
|
||||
Algorithm::Schedule(task, range);
|
||||
// The TaskType class handles the magic of fetching values
|
||||
// for each instance and calling the worklet's function.
|
||||
// The TaskType will evaluate to one of the following classes:
|
||||
//
|
||||
// vtkm::exec::internal::TaskSingular
|
||||
// vtkm::exec::internal::TaskTiling1D
|
||||
// vtkm::exec::internal::TaskTiling3D
|
||||
auto task = TaskTypes::MakeTask(this->Worklet, invocation, range, globalIndexOffset);
|
||||
Algorithm::ScheduleTask(task, range);
|
||||
}
|
||||
|
||||
vtkm::Id NumberOfBlocks;
|
||||
|
@ -26,6 +26,7 @@
|
||||
#include <vtkm/internal/Invocation.h>
|
||||
|
||||
#include <vtkm/cont/DeviceAdapter.h>
|
||||
|
||||
#include <vtkm/cont/ErrorBadType.h>
|
||||
|
||||
#include <vtkm/cont/arg/ControlSignatureTagBase.h>
|
||||
@ -34,7 +35,6 @@
|
||||
#include <vtkm/cont/internal/DynamicTransform.h>
|
||||
|
||||
#include <vtkm/exec/arg/ExecutionSignatureTagBase.h>
|
||||
#include <vtkm/exec/internal/TaskSingular.h>
|
||||
|
||||
#include <vtkm/internal/IntegerSequence.h>
|
||||
#include <vtkm/internal/brigand.hpp>
|
||||
@ -476,15 +476,18 @@ private:
|
||||
template <typename Invocation, typename RangeType, typename DeviceAdapter>
|
||||
VTKM_CONT void InvokeSchedule(const Invocation& invocation, RangeType range, DeviceAdapter) const
|
||||
{
|
||||
// The TaskSingular class handles the magic of fetching values
|
||||
// for each instance and calling the worklet's function. So just create
|
||||
// a TaskSingular and schedule it with the device adapter.
|
||||
typedef vtkm::exec::internal::TaskSingular<WorkletType, Invocation> WorkletInvokeFunctorType;
|
||||
WorkletInvokeFunctorType workletFunctor = WorkletInvokeFunctorType(this->Worklet, invocation);
|
||||
using Algorithm = vtkm::cont::DeviceAdapterAlgorithm<DeviceAdapter>;
|
||||
using TaskTypes = typename vtkm::cont::DeviceTaskTypes<DeviceAdapter>;
|
||||
|
||||
typedef vtkm::cont::DeviceAdapterAlgorithm<DeviceAdapter> Algorithm;
|
||||
|
||||
Algorithm::Schedule(workletFunctor, range);
|
||||
// The TaskType class handles the magic of fetching values
|
||||
// for each instance and calling the worklet's function.
|
||||
// The TaskType will evaluate to one of the following classes:
|
||||
//
|
||||
// vtkm::exec::internal::TaskSingular
|
||||
// vtkm::exec::internal::TaskTiling1D
|
||||
// vtkm::exec::internal::TaskTiling3D
|
||||
auto task = TaskTypes::MakeTask(this->Worklet, invocation, range);
|
||||
Algorithm::ScheduleTask(task, range);
|
||||
}
|
||||
};
|
||||
}
|
||||
|
Loading…
Reference in New Issue
Block a user