Commit 1145b108 authored by Sujin Philip's avatar Sujin Philip
Browse files

Add DynamicImplicitFunction

parent 2481dd62
Pipeline #51518 passed with stage
......@@ -201,13 +201,13 @@ public:
{ }
VTKM_EXEC_CONT
FloatDefault operator()(const vtkm::Vec<FloatDefault, 3> x) const
vtkm::Vec<FloatDefault, 3> operator()(const vtkm::Vec<FloatDefault, 3> x) const
{
return this->Function.Gradient(x);
}
VTKM_EXEC_CONT
FloatDefault operator()(FloatDefault x, FloatDefault y, FloatDefault z) const
vtkm::Vec<FloatDefault, 3> operator()(FloatDefault x, FloatDefault y, FloatDefault z) const
{
return this->Function.Gradient(x, y, z);
}
......
......@@ -57,6 +57,7 @@ set(headers
DeviceAdapterListTag.h
DynamicArrayHandle.h
DynamicCellSet.h
DynamicImplicitFunction.h
Error.h
ErrorBadAllocation.h
ErrorBadType.h
......
//============================================================================
// 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_cont_DynamicImplicitFunction_h
#define vtk_m_cont_DynamicImplicitFunction_h
#include <vtkm/cont/DeviceAdapterListTag.h>
#include <vtkm/cont/ErrorBadType.h>
#include <vtkm/cont/internal/DynamicImplicitFunctionManager.h>
#include <vtkm/ListTag.h>
#include <vtkm/Types.h>
#include <functional>
#include <map>
#include <memory>
#include <sstream>
#include <type_traits>
namespace vtkm {
namespace cont {
class DynamicImplicitFunction
{
public:
VTKM_EXEC
FloatDefault Value(FloatDefault x, FloatDefault y, FloatDefault z) const
{
return this->ValueCaller(this->Function, x, y, z);
}
VTKM_EXEC
FloatDefault Value(const vtkm::Vec<FloatDefault, 3> &x) const
{
return this->ValueCaller(this->Function, x[0], x[1], x[2]);
}
VTKM_EXEC
vtkm::Vec<FloatDefault, 3> Gradient(FloatDefault x, FloatDefault y, FloatDefault z) const
{
return this->GradientCaller(this->Function, x, y, z);
}
VTKM_EXEC
vtkm::Vec<FloatDefault, 3> Gradient(const vtkm::Vec<FloatDefault, 3> &x) const
{
return this->GradientCaller(this->Function, x[0], x[1], x[2]);
}
private:
VTKM_CONT
DynamicImplicitFunction() = default;
void *Function;
internal::ValueCallerType *ValueCaller;
internal::GradientCallerType *GradientCaller;
friend class DynamicImplicitFunctionHandle;
};
class DynamicImplicitFunctionHandle
{
public:
VTKM_CONT
template <typename ImplicitFunction,
typename DeviceAdapterList = VTKM_DEFAULT_DEVICE_ADAPTER_LIST_TAG>
DynamicImplicitFunctionHandle(const ImplicitFunction &f,
DeviceAdapterList devices = DeviceAdapterList())
{
ImplicitFunction *copy = new ImplicitFunction(f);
CreateManager<ImplicitFunction> create(copy, &this->Managers);
vtkm::ListForEach(create, devices);
this->ImplicitFunctionInstance = copy;
this->Deleter = [copy] { delete copy; };
}
~DynamicImplicitFunctionHandle()
{
this->Deleter();
}
VTKM_CONT
template <typename DeviceAdapter>
DynamicImplicitFunction PrepareForExecution(DeviceAdapter)
{
DynamicImplicitFunction func;
auto it = this->Managers.find(vtkm::cont::DeviceAdapterTraits<DeviceAdapter>::GetId());
if (it == this->Managers.end())
{
std::stringstream ss;
ss << DeviceAdapterTraits<DeviceAdapter>::GetName()
<< " was not in the list specified during construction";
throw vtkm::cont::ErrorBadType(ss.str());
}
it->second->Get(func.Function, func.ValueCaller, func.GradientCaller);
return func;
}
private:
using ManagersMap =
std::map<vtkm::cont::DeviceAdapterId,
std::unique_ptr<internal::DynamicImplicitFunctionManagerBase>>;
template <typename ImplicitFunction>
class CreateManager
{
public:
CreateManager(ImplicitFunction *func, ManagersMap *managers)
: Function(func), Managers(managers)
{ }
template<typename DeviceAdapter>
typename std::enable_if<vtkm::cont::DeviceAdapterTraits<DeviceAdapter>::Valid>::type
operator()(DeviceAdapter)
{
using ManagerType =
internal::DynamicImplicitFunctionManager<ImplicitFunction, DeviceAdapter>;
this->Managers->emplace(
vtkm::cont::DeviceAdapterTraits<DeviceAdapter>::GetId(),
typename ManagersMap::mapped_type(new ManagerType(this->Function)));
}
template<typename DeviceAdapter>
typename std::enable_if<!vtkm::cont::DeviceAdapterTraits<DeviceAdapter>::Valid>::type
operator()(DeviceAdapter) {}
private:
ImplicitFunction *Function;
ManagersMap *Managers;
};
void *ImplicitFunctionInstance;
std::function<void(void)> Deleter;
ManagersMap Managers;
};
}
} // vtkm::cont
#endif // vtk_m_cont_DynamicImplicitFunction_h
......@@ -30,6 +30,7 @@
#include <vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.h>
#include <vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h>
#include <vtkm/cont/cuda/internal/DynamicImplicitFunctionManagerCuda.h>
#endif
#endif //vtk_m_cont_cuda_DeviceAdapterCuda_h
......@@ -24,6 +24,7 @@ set(headers
DeviceAdapterAlgorithmCuda.h
DeviceAdapterAlgorithmThrust.h
DeviceAdapterTagCuda.h
DynamicImplicitFunctionManagerCuda.h
MakeThrustIterator.h
ThrustExceptionHandler.h
)
......
//============================================================================
// 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_cont_cuda_internal_DynamicImplicitFunctionManagerCuda_h
#define vtk_m_cont_cuda_internal_DynamicImplicitFunctionManagerCuda_h
#include <vtkm/cont/cuda/internal/DeviceAdapterTagCuda.h>
#include <vtkm/cont/internal/DynamicImplicitFunctionManager.h>
namespace vtkm {
namespace cont {
namespace internal {
namespace detail {
template<typename ImplicitFunction>
struct DeviceData
{
DeviceData()
: Instance(nullptr), ValueCaller(nullptr), GradientCaller(nullptr)
{ }
ImplicitFunction *Instance;
ValueCallerType *ValueCaller;
GradientCallerType *GradientCaller;
};
template <typename ImplicitFunction>
VTKM_EXEC
static FloatDefault CallValue(void *func, FloatDefault x, FloatDefault y, FloatDefault z)
{
return static_cast<ImplicitFunction*>(func)->Value(x, y, z);
}
template <typename ImplicitFunction>
VTKM_EXEC
static vtkm::Vec<FloatDefault, 3>
CallGradient(void *func, FloatDefault x, FloatDefault y, FloatDefault z)
{
return static_cast<ImplicitFunction*>(func)->Gradient(x, y, z);
}
template <typename ImplicitFunction>
__global__
void CreateKernel(DeviceData<ImplicitFunction> *ddata, ImplicitFunction func)
{
ddata->Instance = new ImplicitFunction(func);
ddata->ValueCaller = &CallValue<ImplicitFunction>;
ddata->GradientCaller = &CallGradient<ImplicitFunction>;
}
template <typename ImplicitFunction>
__global__
void DeleteKernel(ImplicitFunction *deviceFunc)
{
delete deviceFunc;
}
} // detail
template <typename ImplicitFunction>
class DynamicImplicitFunctionManager<ImplicitFunction, vtkm::cont::DeviceAdapterTagCuda>
: public vtkm::cont::internal::DynamicImplicitFunctionManagerBase
{
private:
using DeviceData = detail::DeviceData<ImplicitFunction>;
public:
VTKM_CONT
explicit DynamicImplicitFunctionManager(ImplicitFunction *func)
: HostInstance(func)
{ }
VTKM_CONT
~DynamicImplicitFunctionManager()
{
if (this->Device.Instance)
{
detail::DeleteKernel<<<1, 1>>>(this->Device.Instance);
}
}
VTKM_CONT
void Get(void* &instance,
ValueCallerType* &vcaller,
GradientCallerType* &gcaller) override
{
if (!this->Device.Instance)
{
DeviceData *ddata;
cudaMalloc(&ddata, sizeof(DeviceData));
detail::CreateKernel<<<1, 1>>>(ddata, *this->HostInstance);
cudaMemcpy(&this->Device, ddata, sizeof(DeviceData), cudaMemcpyDeviceToHost);
cudaFree(ddata);
}
instance = this->Device.Instance;
vcaller = this->Device.ValueCaller;
gcaller = this->Device.GradientCaller;
}
DynamicImplicitFunctionManager(const DynamicImplicitFunctionManager&) = delete;
DynamicImplicitFunctionManager& operator=(const DynamicImplicitFunctionManager&) = delete;
private:
ImplicitFunction *HostInstance;
DeviceData Device;
};
}
}
} // vtkm::cont::internal
#endif // vtk_m_cont_cuda_internal_DynamicImplicitFunctionManagerCuda_h
......@@ -25,6 +25,7 @@ set(unit_tests
UnitTestCudaDataSetExplicit.cu
UnitTestCudaDataSetSingleType.cu
UnitTestCudaDeviceAdapter.cu
UnitTestCudaDynamicImplicitFunction.cu
UnitTestCudaMath.cu
)
vtkm_unit_tests(CUDA SOURCES ${unit_tests})
//============================================================================
// 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/testing/TestingDynamicImplicitFunction.h>
namespace {
void TestDynamicImplicitFunction()
{
TestingDynamicImplicitFunction testing;
testing.Run(vtkm::cont::DeviceAdapterTagSerial());
testing.Run(vtkm::cont::DeviceAdapterTagCuda());
}
} // anonymous namespace
int UnitTestCudaDynamicImplicitFunction(int, char *[])
{
return vtkm::cont::testing::Testing::Run(TestDynamicImplicitFunction);
}
......@@ -30,6 +30,7 @@ set(headers
DeviceAdapterDefaultSelection.h
DeviceAdapterError.h
DeviceAdapterTag.h
DynamicImplicitFunctionManager.h
DynamicTransform.h
FunctorsGeneral.h
IteratorFromArrayPortal.h
......
//============================================================================
// 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_cont_internal_DynamicImplicitFunctionManager_h
#define vtk_m_cont_internal_DynamicImplicitFunctionManager_h
#include <vtkm/Types.h>
namespace vtkm {
namespace cont {
namespace internal {
using ValueCallerType = FloatDefault(void*, FloatDefault, FloatDefault, FloatDefault);
using GradientCallerType = vtkm::Vec<FloatDefault, 3>(void*, FloatDefault, FloatDefault, FloatDefault);
class DynamicImplicitFunctionManagerBase
{
public:
virtual ~DynamicImplicitFunctionManagerBase() { }
virtual void Get(void* &, ValueCallerType* &, GradientCallerType* &) = 0;
};
template <typename ImplicitFunction, typename DeviceAdapter>
class DynamicImplicitFunctionManager;
}
}
} // vtkm::cont::internal
#endif // vtk_m_cont_internal_DynamicImplicitFunctionManager_h
......@@ -23,5 +23,6 @@
#include <vtkm/cont/serial/internal/DeviceAdapterTagSerial.h>
#include <vtkm/cont/serial/internal/ArrayManagerExecutionSerial.h>
#include <vtkm/cont/serial/internal/DeviceAdapterAlgorithmSerial.h>
#include <vtkm/cont/serial/internal/DynamicImplicitFunctionManagerSerial.h>
#endif //vtk_m_cont_serial_DeviceAdapterSerial_h
......@@ -22,6 +22,7 @@ set(headers
ArrayManagerExecutionSerial.h
DeviceAdapterAlgorithmSerial.h
DeviceAdapterTagSerial.h
DynamicImplicitFunctionManagerSerial.h
)
vtkm_declare_headers(${headers})
//============================================================================
// 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_cont_serial_internal_DynamicImplicitFunctionManagerSerial_h
#define vtk_m_cont_serial_internal_DynamicImplicitFunctionManagerSerial_h
#include <vtkm/cont/serial/internal/DeviceAdapterTagSerial.h>
#include <vtkm/cont/internal/DynamicImplicitFunctionManager.h>
namespace vtkm {
namespace cont {
namespace internal {
template <typename ImplicitFunction>
class DynamicImplicitFunctionManager<ImplicitFunction, vtkm::cont::DeviceAdapterTagSerial>
: public vtkm::cont::internal::DynamicImplicitFunctionManagerBase
{
public:
VTKM_CONT
explicit DynamicImplicitFunctionManager(ImplicitFunction *func)
: HostInstance(func)
{ }
VTKM_CONT
void Get(void* &instance, ValueCallerType* &vcaller, GradientCallerType* &gcaller) override
{
instance = HostInstance;
vcaller = &CallValue;
gcaller = &CallGradient;
}
DynamicImplicitFunctionManager(const DynamicImplicitFunctionManager&) = delete;
DynamicImplicitFunctionManager& operator=(const DynamicImplicitFunctionManager&) = delete;
private:
VTKM_EXEC
static FloatDefault CallValue(void *func, FloatDefault x, FloatDefault y, FloatDefault z)
{
return static_cast<ImplicitFunction*>(func)->Value(x, y, z);
}
VTKM_EXEC
static vtkm::Vec<FloatDefault, 3>
CallGradient(void *func, FloatDefault x, FloatDefault y, FloatDefault z)
{
return static_cast<ImplicitFunction*>(func)->Gradient(x, y, z);
}
ImplicitFunction *HostInstance;
};
}
}
} // vtkm::cont::internal
#endif // vtk_m_cont_serial_internal_DynamicImplicitFunctionManagerSerial_h
......@@ -25,5 +25,6 @@ set(unit_tests
UnitTestSerialDataSetExplicit.cxx
UnitTestSerialDataSetSingleType.cxx
UnitTestSerialDeviceAdapter.cxx
UnitTestSerialDynamicImplicitFunction.cxx
)
vtkm_unit_tests(TBB SOURCES ${unit_tests})
//============================================================================
// 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/testing/TestingDynamicImplicitFunction.h>
namespace {
void TestDynamicImplicitFunction()
{
TestingDynamicImplicitFunction testing;
testing.Run(vtkm::cont::DeviceAdapterTagSerial());
}
} // anonymous namespace
int UnitTestSerialDynamicImplicitFunction(int, char *[])
{
return vtkm::cont::testing::Testing::Run(TestDynamicImplicitFunction);
}
......@@ -25,6 +25,7 @@
#ifdef VTKM_ENABLE_TBB
#include <vtkm/cont/tbb/internal/ArrayManagerExecutionTBB.h>
#include <vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h>
#include <vtkm/cont/tbb/internal/DynamicImplicitFunctionManagerTBB.h>
#endif
#endif //vtk_m_cont_tbb_DeviceAdapterTBB_h
......@@ -22,6 +22,7 @@ set(headers
ArrayManagerExecutionTBB.h
DeviceAdapterAlgorithmTBB.h
DeviceAdapterTagTBB.h
DynamicImplicitFunctionManagerTBB.h
FunctorsTBB.h
)
......
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
// This software is distributed WITHOUT ANY WARRANTY; without even