From 968960c1a1a449edd12a18d62688b1cadcf28d97 Mon Sep 17 00:00:00 2001 From: Sujin Philip Date: Wed, 29 Mar 2017 10:48:43 -0400 Subject: [PATCH 1/3] Add Virtual Methods Framework --- vtkm/cont/CMakeLists.txt | 1 + vtkm/cont/VirtualObjectCache.h | 333 ++++++++++++++++++ vtkm/cont/cuda/DeviceAdapterCuda.h | 1 + vtkm/cont/cuda/internal/CMakeLists.txt | 1 + .../cuda/internal/VirtualObjectTransferCuda.h | 78 ++++ vtkm/cont/cuda/testing/CMakeLists.txt | 1 + .../testing/UnitTestCudaVirtualObjectCache.cu | 39 ++ vtkm/cont/internal/CMakeLists.txt | 2 + vtkm/cont/internal/VirtualObjectTransfer.h | 53 +++ .../VirtualObjectTransferShareWithControl.h | 51 +++ vtkm/cont/serial/DeviceAdapterSerial.h | 1 + vtkm/cont/serial/internal/CMakeLists.txt | 1 + .../internal/VirtualObjectTransferSerial.h | 43 +++ vtkm/cont/serial/testing/CMakeLists.txt | 1 + .../UnitTestSerialVirtualObjectCache.cxx | 37 ++ vtkm/cont/tbb/DeviceAdapterTBB.h | 1 + vtkm/cont/tbb/internal/CMakeLists.txt | 1 + .../tbb/internal/VirtualObjectTransferTBB.h | 43 +++ vtkm/cont/tbb/testing/CMakeLists.txt | 1 + .../testing/UnitTestTBBVirtualObjectCache.cxx | 39 ++ vtkm/cont/testing/CMakeLists.txt | 1 + vtkm/cont/testing/TestingDeviceAdapter.h | 90 +++++ vtkm/cont/testing/TestingVirtualObjectCache.h | 220 ++++++++++++ .../UnitTestDeviceAdapterAlgorithmGeneral.cxx | 8 +- 24 files changed, 1046 insertions(+), 1 deletion(-) create mode 100644 vtkm/cont/VirtualObjectCache.h create mode 100644 vtkm/cont/cuda/internal/VirtualObjectTransferCuda.h create mode 100644 vtkm/cont/cuda/testing/UnitTestCudaVirtualObjectCache.cu create mode 100644 vtkm/cont/internal/VirtualObjectTransfer.h create mode 100644 vtkm/cont/internal/VirtualObjectTransferShareWithControl.h create mode 100644 vtkm/cont/serial/internal/VirtualObjectTransferSerial.h create mode 100644 vtkm/cont/serial/testing/UnitTestSerialVirtualObjectCache.cxx create mode 100644 vtkm/cont/tbb/internal/VirtualObjectTransferTBB.h create mode 100644 vtkm/cont/tbb/testing/UnitTestTBBVirtualObjectCache.cxx create mode 100644 vtkm/cont/testing/TestingVirtualObjectCache.h diff --git a/vtkm/cont/CMakeLists.txt b/vtkm/cont/CMakeLists.txt index 4f64f2797..1decadf56 100644 --- a/vtkm/cont/CMakeLists.txt +++ b/vtkm/cont/CMakeLists.txt @@ -72,6 +72,7 @@ set(headers StorageListTag.h Timer.h TryExecute.h + VirtualObjectCache.h ) set(header_impls diff --git a/vtkm/cont/VirtualObjectCache.h b/vtkm/cont/VirtualObjectCache.h new file mode 100644 index 000000000..310f6bc0f --- /dev/null +++ b/vtkm/cont/VirtualObjectCache.h @@ -0,0 +1,333 @@ +//============================================================================ +// 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_cont_VirtualObjectCache_h +#define vtk_m_cont_VirtualObjectCache_h + +#include +#include +#include +#include +#include + +#include +#include + + +#define VTKM_MAX_DEVICE_ADAPTER_ID 8 + +namespace vtkm { +namespace cont { + +/// \brief Implements VTK-m's execution side Virtual Methods +/// functionality. +/// +/// The template parameter \c VirtualObject is the class that acts as the +/// interface. Following is a method for implementing such classes: +/// 1. Create a const void*<\tt> member variable that will hold a +/// reference to the target class. +/// 2. For each virtual-like method: +/// a. Create a typedef for a function with the same signature, +/// except for an extra const void*<\tt> argument. +/// b. Create a function pointer member variable with the type of the +/// associated typedef from 2a. +/// c. Create an implementation of the method, that calls the associated +/// function pointer with the void pointer to the target class as one of +/// the arguments. +/// 3. Create the following template function: +/// +/// template +/// VTKM_EXEC void Bind(const TargetClass *deviceTarget); +/// +/// This function should set the void pointer from 1 to \c deviceTarget, +/// and assign the function pointers from 2b to functions that cast their +/// first argument to const TargetClass*<\tt> and call the corresponding +/// member function on it. +/// +/// Both \c VirtualObject and target class objects should be bitwise copyable. +/// +/// \sa vtkm::exec::ImplicitFunction, vtkm::cont::ImplicitFunction +/// +template +class VirtualObjectCache +{ +public: + VirtualObjectCache() + : Target(nullptr), + CurrentDevice(VTKM_DEVICE_ADAPTER_UNDEFINED), + DeviceState(nullptr), + RefreshFlag(false) + { + } + + ~VirtualObjectCache() + { + this->Reset(); + } + + VirtualObjectCache(const VirtualObjectCache &other) + : Target(other.Target), + Transfers(other.Transfers), + CurrentDevice(VTKM_DEVICE_ADAPTER_UNDEFINED), + DeviceState(nullptr), + RefreshFlag(false) + { + } + + VirtualObjectCache& operator=(const VirtualObjectCache &other) + { + if (this != &other) + { + this->Target = other.Target; + this->Transfers = other.Transfers; + this->CurrentDevice = VTKM_DEVICE_ADAPTER_UNDEFINED; + this->DeviceState = nullptr; + this->RefreshFlag = false; + this->Object = VirtualObject(); + } + return *this; + } + + VirtualObjectCache(VirtualObjectCache &&other) + : Target(other.Target), + Transfers(other.Transfers), + CurrentDevice(other.CurrentDevice), + DeviceState(other.DeviceState), + RefreshFlag(other.RefreshFlag), + Object(other.Object) + { + other.CurrentDevice = VTKM_DEVICE_ADAPTER_UNDEFINED; + other.DeviceState = nullptr; + } + + VirtualObjectCache& operator=(VirtualObjectCache &&other) + { + if (this != &other) + { + this->Target = other.Target; + this->Transfers = std::move(other.Transfers); + this->CurrentDevice = other.CurrentDevice; + this->DeviceState = other.DeviceState; + this->RefreshFlag = other.RefreshFlag; + this->Object = std::move(other.Object); + + other.CurrentDevice = VTKM_DEVICE_ADAPTER_UNDEFINED; + other.DeviceState = nullptr; + } + return *this; + } + + /// Reset to the default constructed state + void Reset() + { + this->ReleaseResources(); + + this->Target = nullptr; + this->Transfers.fill(TransferInterface()); + } + + void ReleaseResources() + { + if (this->CurrentDevice > 0) + { + this->GetCurrentTransfer().Cleanup(this->DeviceState); + this->CurrentDevice = VTKM_DEVICE_ADAPTER_UNDEFINED; + this->DeviceState = nullptr; + this->RefreshFlag = false; + this->Object = VirtualObject(); + } + } + + /// Get if in a valid state (a target is bound) + bool GetValid() const + { + return this->Target != nullptr; + } + + // Set/Get if the cached virtual object should be refreshed to the current + // state of the target + void SetRefreshFlag(bool value) + { + this->RefreshFlag = value; + } + bool GetRefreshFlag() const + { + return this->RefreshFlag; + } + + /// Bind to \c target. The lifetime of target is expected to be managed + /// externally, and should be valid for as long as it is bound. + /// Also accepts a list-tag of device adapters where the virtual + /// object may be used (default = \c VTKM_DEFAULT_DEVICE_ADAPTER_LIST_TAG). + /// + template + void Bind(const TargetClass *target, + DeviceAdapterList devices = DeviceAdapterList()) + { + this->Reset(); + + this->Target = target; + vtkm::ListForEach( + CreateTransferInterface(this->Transfers.data()), + devices); + } + + /// Get a \c VirtualObject for \c DeviceAdapter. + /// VirtualObjectCache and VirtualObject are analogous to ArrayHandle and Portal + /// The returned VirtualObject will be invalidated if: + /// 1. A new VirtualObject is requested for a different DeviceAdapter + /// 2. VirtualObjectCache is destroyed + /// 3. Reset or ReleaseResources is called + /// + template + VirtualObject GetVirtualObject(DeviceAdapter) + { + using DeviceInfo = vtkm::cont::DeviceAdapterTraits; + + if (!this->GetValid()) + { + throw vtkm::cont::ErrorBadValue("No target object bound"); + } + + vtkm::cont::DeviceAdapterId deviceId = DeviceInfo::GetId(); + if (deviceId < 0 || deviceId >= VTKM_MAX_DEVICE_ADAPTER_ID) + { + std::string msg = + "Device '" + DeviceInfo::GetName() + "' has invalid ID of " + + std::to_string(deviceId) + "(VTKM_MAX_DEVICE_ADAPTER_ID = " + + std::to_string(VTKM_MAX_DEVICE_ADAPTER_ID) + ")"; + throw vtkm::cont::ErrorBadType(msg); + } + + if (this->CurrentDevice != deviceId) + { + this->ReleaseResources(); + + std::size_t idx = static_cast(deviceId); + TransferInterface &transfer = this->Transfers[idx]; + if (!TransferInterfaceValid(transfer)) + { + std::string msg = DeviceInfo::GetName() + + " was not in the list specified in Bind"; + throw vtkm::cont::ErrorBadType(msg); + } + this->CurrentDevice = deviceId; + this->DeviceState = transfer.Create(this->Object, this->Target); + } + else if (this->RefreshFlag) + { + this->GetCurrentTransfer().Update(this->DeviceState, this->Target); + } + + this->RefreshFlag = false; + return this->Object; + } + +private: + struct TransferInterface + { + using CreateSig = void*(VirtualObject &, const void *); + using UpdateSig = void(void *, const void *); + using CleanupSig = void(void *); + + CreateSig *Create = nullptr; + UpdateSig *Update = nullptr; + CleanupSig *Cleanup = nullptr; + }; + + static bool TransferInterfaceValid(const TransferInterface &t) + { + return t.Create != nullptr; + } + + TransferInterface& GetCurrentTransfer() + { + return this->Transfers[static_cast(this->CurrentDevice)]; + } + + template + class CreateTransferInterface + { + private: + template + using EnableIfValid = + std::enable_if::Valid>; + + template + using EnableIfInvalid = + std::enable_if::Valid>; + + public: + CreateTransferInterface(TransferInterface *transfers) + : Transfers(transfers) + { } + + // Use SFINAE to create entries for valid device adapters only + template + typename EnableIfValid::type operator()(DeviceAdapter) const + { + using DeviceInfo = vtkm::cont::DeviceAdapterTraits; + + if (DeviceInfo::GetId() >= 0 && + DeviceInfo::GetId() < VTKM_MAX_DEVICE_ADAPTER_ID) + { + using TransferImpl = + internal::VirtualObjectTransfer; + + std::size_t id = static_cast(DeviceInfo::GetId()); + TransferInterface &transfer = this->Transfers[id]; + + transfer.Create = &TransferImpl::Create; + transfer.Update = &TransferImpl::Update; + transfer.Cleanup = &TransferImpl::Cleanup; + } + else + { + std::string msg = + "Device '" + DeviceInfo::GetName() + "' has invalid ID of " + + std::to_string(DeviceInfo::GetId()) + "(VTKM_MAX_DEVICE_ADAPTER_ID = " + + std::to_string(VTKM_MAX_DEVICE_ADAPTER_ID) + ")"; + throw vtkm::cont::ErrorBadType(msg); + } + } + + template + typename EnableIfInvalid::type operator()(DeviceAdapter) const + { } + + private: + TransferInterface *Transfers; + }; + + const void *Target; + std::array Transfers; + + vtkm::cont::DeviceAdapterId CurrentDevice; + void *DeviceState; + bool RefreshFlag; + VirtualObject Object; +}; + +} +} // vtkm::cont + +#undef VTKM_MAX_DEVICE_ADAPTER_ID + +#endif // vtk_m_cont_VirtualObjectCache_h diff --git a/vtkm/cont/cuda/DeviceAdapterCuda.h b/vtkm/cont/cuda/DeviceAdapterCuda.h index 4476dbb71..cf4d06826 100644 --- a/vtkm/cont/cuda/DeviceAdapterCuda.h +++ b/vtkm/cont/cuda/DeviceAdapterCuda.h @@ -30,6 +30,7 @@ #include #include +#include #endif #endif //vtk_m_cont_cuda_DeviceAdapterCuda_h diff --git a/vtkm/cont/cuda/internal/CMakeLists.txt b/vtkm/cont/cuda/internal/CMakeLists.txt index fb9f1acdd..089f4229f 100644 --- a/vtkm/cont/cuda/internal/CMakeLists.txt +++ b/vtkm/cont/cuda/internal/CMakeLists.txt @@ -26,6 +26,7 @@ set(headers DeviceAdapterTagCuda.h MakeThrustIterator.h ThrustExceptionHandler.h + VirtualObjectTransferCuda.h ) vtkm_declare_headers(CUDA ${headers} TESTABLE ${VTKm_ENABLE_CUDA}) diff --git a/vtkm/cont/cuda/internal/VirtualObjectTransferCuda.h b/vtkm/cont/cuda/internal/VirtualObjectTransferCuda.h new file mode 100644 index 000000000..f9e3012eb --- /dev/null +++ b/vtkm/cont/cuda/internal/VirtualObjectTransferCuda.h @@ -0,0 +1,78 @@ +//============================================================================ +// 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_cont_cuda_internal_VirtualObjectTransferCuda_h +#define vtk_m_cont_cuda_internal_VirtualObjectTransferCuda_h + +#include +#include +#include + + +namespace vtkm { +namespace cont { +namespace internal { + +namespace detail { + +template +__global__ +void CreateKernel(VirtualObject *object, const TargetClass *target) +{ + object->Bind(target); +} + +} // detail + +template +struct VirtualObjectTransfer +{ + static void* Create(VirtualObject &object, const void *target) + { + TargetClass *cutarget; + VTKM_CUDA_CALL(cudaMalloc(&cutarget, sizeof(TargetClass))); + VTKM_CUDA_CALL(cudaMemcpy(cutarget, target, sizeof(TargetClass), cudaMemcpyHostToDevice)); + + VirtualObject *cuobject; + VTKM_CUDA_CALL(cudaMalloc(&cuobject, sizeof(VirtualObject))); + detail::CreateKernel<<<1, 1>>>(cuobject, cutarget); + VTKM_CUDA_CHECK_ASYNCHRONOUS_ERROR(); + VTKM_CUDA_CALL(cudaMemcpy(&object, cuobject, sizeof(VirtualObject), cudaMemcpyDeviceToHost)); + VTKM_CUDA_CALL(cudaFree(cuobject)); + + return cutarget; + } + + static void Update(void *deviceState, const void *target) + { + VTKM_CUDA_CALL(cudaMemcpy(deviceState, target, sizeof(TargetClass), cudaMemcpyHostToDevice)); + } + + static void Cleanup(void *deviceState) + { + VTKM_CUDA_CALL(cudaFree(deviceState)); + } +}; + +} +} +} // vtkm::cont::internal + +#endif // vtk_m_cont_cuda_internal_VirtualObjectTransferCuda_h diff --git a/vtkm/cont/cuda/testing/CMakeLists.txt b/vtkm/cont/cuda/testing/CMakeLists.txt index f3db6d9e8..b937fb3c9 100644 --- a/vtkm/cont/cuda/testing/CMakeLists.txt +++ b/vtkm/cont/cuda/testing/CMakeLists.txt @@ -26,5 +26,6 @@ set(unit_tests UnitTestCudaDataSetSingleType.cu UnitTestCudaDeviceAdapter.cu UnitTestCudaMath.cu + UnitTestCudaVirtualObjectCache.cu ) vtkm_unit_tests(CUDA SOURCES ${unit_tests}) diff --git a/vtkm/cont/cuda/testing/UnitTestCudaVirtualObjectCache.cu b/vtkm/cont/cuda/testing/UnitTestCudaVirtualObjectCache.cu new file mode 100644 index 000000000..0b32ee39c --- /dev/null +++ b/vtkm/cont/cuda/testing/UnitTestCudaVirtualObjectCache.cu @@ -0,0 +1,39 @@ +//============================================================================ +// 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. +//============================================================================ +#include + +namespace { + +void TestVirtualObjectCache() +{ + using DeviceAdapterList = + vtkm::ListTagBase; + + vtkm::cont::testing::TestingVirtualObjectCache::Run(); +} + +} // anonymous namespace + + +int UnitTestCudaVirtualObjectCache(int, char *[]) +{ + return vtkm::cont::testing::Testing::Run(TestVirtualObjectCache); +} diff --git a/vtkm/cont/internal/CMakeLists.txt b/vtkm/cont/internal/CMakeLists.txt index 984d6f404..df7e0c107 100644 --- a/vtkm/cont/internal/CMakeLists.txt +++ b/vtkm/cont/internal/CMakeLists.txt @@ -35,6 +35,8 @@ set(headers IteratorFromArrayPortal.h SimplePolymorphicContainer.h StorageError.h + VirtualObjectTransfer.h + VirtualObjectTransferShareWithControl.h ) vtkm_declare_headers(${headers}) diff --git a/vtkm/cont/internal/VirtualObjectTransfer.h b/vtkm/cont/internal/VirtualObjectTransfer.h new file mode 100644 index 000000000..b554289d8 --- /dev/null +++ b/vtkm/cont/internal/VirtualObjectTransfer.h @@ -0,0 +1,53 @@ +//============================================================================ +// 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_cont_internal_VirtualObjectTransfer_h +#define vtk_m_cont_internal_VirtualObjectTransfer_h + + +namespace vtkm { +namespace cont { +namespace internal { + +template +struct VirtualObjectTransfer +#ifdef VTKM_DOXYGEN_ONLY +{ + /// Takes a void* to host copy of the target object, transfers it to the + /// device, binds it to the VirtualObject, and returns a void* to an internal + /// state structure. + /// + static void* Create(VirtualObject &object, const void *hostTarget); + + /// Performs cleanup of the device state used to track the VirtualObject on + /// the device. + /// + static void Cleanup(void *deviceState); + + /// Update the device state with the state of target + static void Update(void *deviceState, const void *target); +} +#endif +; + +} +} +} // vtkm::cont::internal + +#endif // vtkm_cont_internal_VirtualObjectTransfer_h diff --git a/vtkm/cont/internal/VirtualObjectTransferShareWithControl.h b/vtkm/cont/internal/VirtualObjectTransferShareWithControl.h new file mode 100644 index 000000000..153040b60 --- /dev/null +++ b/vtkm/cont/internal/VirtualObjectTransferShareWithControl.h @@ -0,0 +1,51 @@ +//============================================================================ +// 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_cont_internal_VirtualObjectTransferShareWithControl_h +#define vtk_m_cont_internal_VirtualObjectTransferShareWithControl_h + + +namespace vtkm { +namespace cont { +namespace internal { + +template +struct VirtualObjectTransferShareWithControl +{ + static void* Create(VirtualObject &object, const void *target) + { + static int dummyState = 0; + object.Bind(static_cast(target)); + return &dummyState; + } + + static void Update(void*, const void*) + { + } + + static void Cleanup(void*) + { + } +}; + +} +} +} // vtkm::cont::internal + +#endif // vtk_m_cont_internal_VirtualObjectTransferShareWithControl_h diff --git a/vtkm/cont/serial/DeviceAdapterSerial.h b/vtkm/cont/serial/DeviceAdapterSerial.h index 28cc06e35..1c83672bc 100644 --- a/vtkm/cont/serial/DeviceAdapterSerial.h +++ b/vtkm/cont/serial/DeviceAdapterSerial.h @@ -23,5 +23,6 @@ #include #include #include +#include #endif //vtk_m_cont_serial_DeviceAdapterSerial_h diff --git a/vtkm/cont/serial/internal/CMakeLists.txt b/vtkm/cont/serial/internal/CMakeLists.txt index 6e6ea0394..4c86c4fbd 100644 --- a/vtkm/cont/serial/internal/CMakeLists.txt +++ b/vtkm/cont/serial/internal/CMakeLists.txt @@ -22,6 +22,7 @@ set(headers ArrayManagerExecutionSerial.h DeviceAdapterAlgorithmSerial.h DeviceAdapterTagSerial.h + VirtualObjectTransferSerial.h ) vtkm_declare_headers(${headers}) diff --git a/vtkm/cont/serial/internal/VirtualObjectTransferSerial.h b/vtkm/cont/serial/internal/VirtualObjectTransferSerial.h new file mode 100644 index 000000000..fcfda6ac8 --- /dev/null +++ b/vtkm/cont/serial/internal/VirtualObjectTransferSerial.h @@ -0,0 +1,43 @@ +//============================================================================ +// 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_cont_serial_internal_VirtualObjectTransferSerial_h +#define vtk_m_cont_serial_internal_VirtualObjectTransferSerial_h + +#include +#include +#include + + +namespace vtkm { +namespace cont { +namespace internal { + +template +struct VirtualObjectTransfer< + VirtualObject, TargetClass, vtkm::cont::DeviceAdapterTagSerial> : + public VirtualObjectTransferShareWithControl +{ +}; + +} +} +} // vtkm::cont::internal + +#endif // vtk_m_cont_serial_internal_VirtualObjectTransferSerial_h diff --git a/vtkm/cont/serial/testing/CMakeLists.txt b/vtkm/cont/serial/testing/CMakeLists.txt index 824131aeb..e97a35fd0 100644 --- a/vtkm/cont/serial/testing/CMakeLists.txt +++ b/vtkm/cont/serial/testing/CMakeLists.txt @@ -25,5 +25,6 @@ set(unit_tests UnitTestSerialDataSetExplicit.cxx UnitTestSerialDataSetSingleType.cxx UnitTestSerialDeviceAdapter.cxx + UnitTestSerialVirtualObjectCache.cxx ) vtkm_unit_tests(TBB SOURCES ${unit_tests}) diff --git a/vtkm/cont/serial/testing/UnitTestSerialVirtualObjectCache.cxx b/vtkm/cont/serial/testing/UnitTestSerialVirtualObjectCache.cxx new file mode 100644 index 000000000..04e9b51e2 --- /dev/null +++ b/vtkm/cont/serial/testing/UnitTestSerialVirtualObjectCache.cxx @@ -0,0 +1,37 @@ +//============================================================================ +// 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. +//============================================================================ +#include + +namespace { + +void TestVirtualObjectCache() +{ + using DeviceAdapterList = vtkm::ListTagBase; + + vtkm::cont::testing::TestingVirtualObjectCache::Run(); +} + +} // anonymous namespace + + +int UnitTestSerialVirtualObjectCache(int, char *[]) +{ + return vtkm::cont::testing::Testing::Run(TestVirtualObjectCache); +} diff --git a/vtkm/cont/tbb/DeviceAdapterTBB.h b/vtkm/cont/tbb/DeviceAdapterTBB.h index f783e9e9f..43a028e47 100644 --- a/vtkm/cont/tbb/DeviceAdapterTBB.h +++ b/vtkm/cont/tbb/DeviceAdapterTBB.h @@ -25,6 +25,7 @@ #ifdef VTKM_ENABLE_TBB #include #include +#include #endif #endif //vtk_m_cont_tbb_DeviceAdapterTBB_h diff --git a/vtkm/cont/tbb/internal/CMakeLists.txt b/vtkm/cont/tbb/internal/CMakeLists.txt index 23de84c6f..20dd71f27 100644 --- a/vtkm/cont/tbb/internal/CMakeLists.txt +++ b/vtkm/cont/tbb/internal/CMakeLists.txt @@ -23,6 +23,7 @@ set(headers DeviceAdapterAlgorithmTBB.h DeviceAdapterTagTBB.h FunctorsTBB.h + VirtualObjectTransferTBB.h ) if (VTKm_ENABLE_TBB) diff --git a/vtkm/cont/tbb/internal/VirtualObjectTransferTBB.h b/vtkm/cont/tbb/internal/VirtualObjectTransferTBB.h new file mode 100644 index 000000000..63349cde5 --- /dev/null +++ b/vtkm/cont/tbb/internal/VirtualObjectTransferTBB.h @@ -0,0 +1,43 @@ +//============================================================================ +// 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_cont_tbb_internal_VirtualObjectTransferTBB_h +#define vtk_m_cont_tbb_internal_VirtualObjectTransferTBB_h + +#include +#include +#include + + +namespace vtkm { +namespace cont { +namespace internal { + +template +struct VirtualObjectTransfer< + VirtualObject, TargetClass, vtkm::cont::DeviceAdapterTagTBB> : + public VirtualObjectTransferShareWithControl +{ +}; + +} +} +} // vtkm::cont::internal + +#endif // vtk_m_cont_tbb_internal_VirtualObjectTransferTBB_h diff --git a/vtkm/cont/tbb/testing/CMakeLists.txt b/vtkm/cont/tbb/testing/CMakeLists.txt index a94396121..f6d15b058 100644 --- a/vtkm/cont/tbb/testing/CMakeLists.txt +++ b/vtkm/cont/tbb/testing/CMakeLists.txt @@ -25,5 +25,6 @@ set(unit_tests UnitTestTBBDataSetExplicit.cxx UnitTestTBBDataSetSingleType.cxx UnitTestTBBDeviceAdapter.cxx + UnitTestTBBVirtualObjectCache.cxx ) vtkm_unit_tests(TBB SOURCES ${unit_tests}) diff --git a/vtkm/cont/tbb/testing/UnitTestTBBVirtualObjectCache.cxx b/vtkm/cont/tbb/testing/UnitTestTBBVirtualObjectCache.cxx new file mode 100644 index 000000000..b76c1d2b4 --- /dev/null +++ b/vtkm/cont/tbb/testing/UnitTestTBBVirtualObjectCache.cxx @@ -0,0 +1,39 @@ +//============================================================================ +// 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. +//============================================================================ +#include + +namespace { + +void TestVirtualObjectCache() +{ + using DeviceAdapterList = + vtkm::ListTagBase; + + vtkm::cont::testing::TestingVirtualObjectCache::Run(); +} + +} // anonymous namespace + + +int UnitTestTBBVirtualObjectCache(int, char *[]) +{ + return vtkm::cont::testing::Testing::Run(TestVirtualObjectCache); +} diff --git a/vtkm/cont/testing/CMakeLists.txt b/vtkm/cont/testing/CMakeLists.txt index abe2cc8b0..6ea6a1078 100644 --- a/vtkm/cont/testing/CMakeLists.txt +++ b/vtkm/cont/testing/CMakeLists.txt @@ -28,6 +28,7 @@ set(headers TestingDataSetExplicit.h TestingDataSetSingleType.h TestingFancyArrayHandles.h + TestingVirtualObjectCache.h ) vtkm_declare_headers(${headers}) diff --git a/vtkm/cont/testing/TestingDeviceAdapter.h b/vtkm/cont/testing/TestingDeviceAdapter.h index e2bbc89f4..2c3f47061 100644 --- a/vtkm/cont/testing/TestingDeviceAdapter.h +++ b/vtkm/cont/testing/TestingDeviceAdapter.h @@ -38,6 +38,7 @@ #include #include +#include #include @@ -311,6 +312,59 @@ public: vtkm::exec::AtomicArray AArray; }; + class VirtualObjectTransferKernel + { + public: + struct Interface + { + using FooSig = vtkm::Id (const void*); + + template + VTKM_EXEC + void Bind(const T *target) + { + this->Target = target; + this->FooPtr = [](const void* t){ return static_cast(t)->Foo(); }; + } + + VTKM_EXEC + vtkm::Id Foo() const + { + return this->FooPtr(this->Target); + } + + const void *Target; + FooSig *FooPtr; + }; + + struct Concrete + { + vtkm::Id Foo() const + { + return this->Value; + } + + vtkm::Id Value = 0; + }; + + VirtualObjectTransferKernel(const Interface &vo, IdArrayHandle &result) + : Virtual(vo), Result(result.PrepareForInPlace(DeviceAdapterTag())) + { } + + VTKM_EXEC + void operator()(vtkm::Id) const + { + this->Result.Set(0, this->Virtual.Foo()); + } + + VTKM_CONT void SetErrorMessageBuffer( + const vtkm::exec::internal::ErrorMessageBuffer&) + { } + + private: + Interface Virtual; + IdPortalType Result; + }; private: @@ -472,6 +526,41 @@ private: VTKM_TEST_ASSERT(valid_runtime, "runtime detection failed for device"); } + VTKM_CONT + static void TestVirtualObjectTransfer() + { + using VirtualObject = typename VirtualObjectTransferKernel::Interface; + using TargetType = typename VirtualObjectTransferKernel::Concrete; + using Transfer = vtkm::cont::internal::VirtualObjectTransfer< + VirtualObject, TargetType, DeviceAdapterTag>; + + IdArrayHandle result; + result.Allocate(1); + result.GetPortalControl().Set(0, 0); + + TargetType target; + target.Value = 5; + + VirtualObject vo; + void* state = Transfer::Create(vo, &target); + + + std::cout << "-------------------------------------------" << std::endl; + std::cout << "Testing VirtualObjectTransfer" << std::endl; + + Algorithm::Schedule(VirtualObjectTransferKernel(vo, result), 1); + VTKM_TEST_ASSERT(result.GetPortalConstControl().Get(0) == 5, + "Did not get expected result"); + + target.Value = 10; + Transfer::Update(state, &target); + Algorithm::Schedule(VirtualObjectTransferKernel(vo, result), 1); + VTKM_TEST_ASSERT(result.GetPortalConstControl().Get(0) == 10, + "Did not get expected result"); + + Transfer::Cleanup(state); + } + static VTKM_CONT void TestAlgorithmSchedule() { std::cout << "-------------------------------------------" << std::endl; @@ -1892,6 +1981,7 @@ private: TestOutOfMemory(); TestTimer(); TestRuntime(); + TestVirtualObjectTransfer(); TestAlgorithmSchedule(); TestErrorExecution(); diff --git a/vtkm/cont/testing/TestingVirtualObjectCache.h b/vtkm/cont/testing/TestingVirtualObjectCache.h new file mode 100644 index 000000000..2c823d341 --- /dev/null +++ b/vtkm/cont/testing/TestingVirtualObjectCache.h @@ -0,0 +1,220 @@ +//============================================================================ +// 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_cont_testing_TestingVirtualObjectCache_h +#define vtk_m_cont_testing_TestingVirtualObjectCache_h + +#include +#include +#include +#include +#include +#include + + +#define ARRAY_LEN 8 + +namespace vtkm { +namespace cont { +namespace testing { + +namespace virtual_object_detail +{ + +class Transformer +{ +public: + template + VTKM_EXEC void Bind(const T *target) + { + this->Concrete = target; + this->Caller = + [](const void *concrete, vtkm::FloatDefault val) { + return static_cast(concrete)->operator()(val); + }; + } + + VTKM_EXEC + vtkm::FloatDefault operator()(vtkm::FloatDefault val) const + { + return this->Caller(this->Concrete, val); + } + +private: + using Signature = vtkm::FloatDefault (const void*, vtkm::FloatDefault); + + const void *Concrete; + Signature *Caller; +}; + +struct Square +{ + VTKM_EXEC + vtkm::FloatDefault operator()(vtkm::FloatDefault val) const + { + return val * val; + } +}; + +struct Multiply +{ + VTKM_EXEC + vtkm::FloatDefault operator()(vtkm::FloatDefault val) const + { + return val * this->Multiplicand; + } + + vtkm::FloatDefault Multiplicand; +}; + +} // virtual_object_detail + +template +class TestingVirtualObjectCache +{ +private: + using FloatArrayHandle = vtkm::cont::ArrayHandle; + using ArrayTransform = + vtkm::cont::ArrayHandleTransform; + using TransformerCache = + vtkm::cont::VirtualObjectCache; + + + class TestStage1 + { + public: + TestStage1(const FloatArrayHandle &input, TransformerCache &manager) + : Input(&input), Manager(&manager) + { } + + template + void operator()(DeviceAdapter device) const + { + using Algorithm = vtkm::cont::DeviceAdapterAlgorithm; + std::cout << "\tDeviceAdapter: " + << vtkm::cont::DeviceAdapterTraits::GetName() + << std::endl; + + for (int n = 0; n < 2; ++n) + { + ArrayTransform transformed(*this->Input, this->Manager->GetVirtualObject(device)); + FloatArrayHandle output; + Algorithm::Copy(transformed, output); + auto portal = output.GetPortalConstControl(); + for (vtkm::Id i = 0; i < ARRAY_LEN; ++i) + { + VTKM_TEST_ASSERT(portal.Get(i) == FloatDefault(i*i), + "\tIncorrect result"); + } + std::cout << "\tSuccess." << std::endl; + + if (n == 0) + { + std::cout << "\tReleaseResources and test again..." << std::endl; + this->Manager->ReleaseResources(); + } + } + } + + private: + const FloatArrayHandle *Input; + TransformerCache *Manager; + }; + + class TestStage2 + { + public: + TestStage2(const FloatArrayHandle &input, virtual_object_detail::Multiply &mul, + TransformerCache &manager) + : Input(&input), Mul(&mul), Manager(&manager) + { } + + template + void operator()(DeviceAdapter device) const + { + using Algorithm = vtkm::cont::DeviceAdapterAlgorithm; + std::cout << "\tDeviceAdapter: " + << vtkm::cont::DeviceAdapterTraits::GetName() + << std::endl; + + this->Mul->Multiplicand = 2; + this->Manager->SetRefreshFlag(true); + for (int n = 0; n < 2; ++n) + { + ArrayTransform transformed(*this->Input, this->Manager->GetVirtualObject(device)); + FloatArrayHandle output; + Algorithm::Copy(transformed, output); + auto portal = output.GetPortalConstControl(); + for (vtkm::Id i = 0; i < ARRAY_LEN; ++i) + { + VTKM_TEST_ASSERT(portal.Get(i) == FloatDefault(i) * this->Mul->Multiplicand, + "\tIncorrect result"); + } + std::cout << "\tSuccess." << std::endl; + + if (n == 0) + { + std::cout << "\tUpdate and test again..." << std::endl; + this->Mul->Multiplicand = 3; + this->Manager->SetRefreshFlag(true); + } + } + } + + private: + const FloatArrayHandle *Input; + virtual_object_detail::Multiply *Mul; + TransformerCache *Manager; + }; + +public: + static void Run() + { + vtkm::cont::ArrayHandle input; + input.Allocate(ARRAY_LEN); + auto portal = input.GetPortalControl(); + for (vtkm::Id i = 0; i < ARRAY_LEN; ++i) + { + portal.Set(i, vtkm::FloatDefault(i)); + } + + TransformerCache manager; + + std::cout << "Testing with concrete type 1 (Square)..." << std::endl; + virtual_object_detail::Square sqr; + manager.Bind(&sqr, DeviceAdapterList()); + vtkm::ListForEach(TestStage1(input, manager), DeviceAdapterList()); + + std::cout << "Reset..." << std::endl; + manager.Reset(); + + std::cout << "Testing with concrete type 2 (Multiply)..." << std::endl; + virtual_object_detail::Multiply mul; + manager.Bind(&mul, DeviceAdapterList()); + vtkm::ListForEach(TestStage2(input, mul, manager), DeviceAdapterList()); + } +}; + +} +} +} // vtkm::cont::testing + +#endif diff --git a/vtkm/cont/testing/UnitTestDeviceAdapterAlgorithmGeneral.cxx b/vtkm/cont/testing/UnitTestDeviceAdapterAlgorithmGeneral.cxx index 9d14f0bf8..f4c15632d 100644 --- a/vtkm/cont/testing/UnitTestDeviceAdapterAlgorithmGeneral.cxx +++ b/vtkm/cont/testing/UnitTestDeviceAdapterAlgorithmGeneral.cxx @@ -95,10 +95,16 @@ public: : Superclass(storage) { } }; +template +struct VirtualObjectTransfer< + VirtualObject, TargetClass, vtkm::cont::DeviceAdapterTagTestAlgorithmGeneral> : + public VirtualObjectTransferShareWithControl +{ +}; } } -} // namespace vtkm::cont::testing +} // namespace vtkm::cont::internal int UnitTestDeviceAdapterAlgorithmGeneral(int, char *[]) { From 82d02e46efb704e67bcdc9548b6c6e0fd8f87f7d Mon Sep 17 00:00:00 2001 From: Sujin Philip Date: Wed, 29 Mar 2017 12:11:09 -0400 Subject: [PATCH 2/3] Modify ImplicitFunctions to use Virtual Methods --- vtkm/CMakeLists.txt | 1 - vtkm/ImplicitFunctions.h | 450 ------------------ .../benchmarking/BenchmarkFieldAlgorithms.cxx | 239 +++++++++- vtkm/cont/CMakeLists.txt | 3 + .../ImplicitFunction.cxx} | 26 +- vtkm/cont/ImplicitFunction.h | 341 +++++++++++++ vtkm/cont/ImplicitFunction.hxx | 204 ++++++++ vtkm/cont/cuda/testing/CMakeLists.txt | 1 + .../testing/UnitTestCudaImplicitFunction.cu | 37 ++ vtkm/cont/serial/testing/CMakeLists.txt | 1 + .../UnitTestSerialImplicitFunction.cxx | 38 ++ vtkm/cont/tbb/testing/CMakeLists.txt | 1 + .../testing/UnitTestTBBImplicitFunction.cxx | 37 ++ vtkm/cont/testing/CMakeLists.txt | 1 + vtkm/cont/testing/TestingImplicitFunction.h | 187 ++++++++ vtkm/exec/CMakeLists.txt | 1 + vtkm/exec/ImplicitFunction.h | 143 ++++++ vtkm/testing/CMakeLists.txt | 2 - vtkm/testing/TestingImplicitFunctions.h | 213 --------- vtkm/worklet/Clip.h | 35 +- vtkm/worklet/ExtractGeometry.h | 17 +- vtkm/worklet/ExtractPoints.h | 20 +- vtkm/worklet/testing/UnitTestClipping.cxx | 4 +- .../testing/UnitTestExtractGeometry.cxx | 6 +- .../worklet/testing/UnitTestExtractPoints.cxx | 6 +- 25 files changed, 1282 insertions(+), 732 deletions(-) delete mode 100644 vtkm/ImplicitFunctions.h rename vtkm/{testing/UnitTestImplicitFunctions.cxx => cont/ImplicitFunction.cxx} (62%) create mode 100644 vtkm/cont/ImplicitFunction.h create mode 100644 vtkm/cont/ImplicitFunction.hxx create mode 100644 vtkm/cont/cuda/testing/UnitTestCudaImplicitFunction.cu create mode 100644 vtkm/cont/serial/testing/UnitTestSerialImplicitFunction.cxx create mode 100644 vtkm/cont/tbb/testing/UnitTestTBBImplicitFunction.cxx create mode 100644 vtkm/cont/testing/TestingImplicitFunction.h create mode 100644 vtkm/exec/ImplicitFunction.h delete mode 100644 vtkm/testing/TestingImplicitFunctions.h diff --git a/vtkm/CMakeLists.txt b/vtkm/CMakeLists.txt index bd8802158..4934cc538 100644 --- a/vtkm/CMakeLists.txt +++ b/vtkm/CMakeLists.txt @@ -34,7 +34,6 @@ set(headers Bounds.h CellShape.h CellTraits.h - ImplicitFunctions.h ListTag.h Math.h Matrix.h diff --git a/vtkm/ImplicitFunctions.h b/vtkm/ImplicitFunctions.h deleted file mode 100644 index 92bb6db46..000000000 --- a/vtkm/ImplicitFunctions.h +++ /dev/null @@ -1,450 +0,0 @@ -//============================================================================ -// 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_ImplicitFunctions_h -#define vtk_m_ImplicitFunctions_h - -#include -#include -#include -#include - -namespace vtkm { - -/// \brief Implicit function for a plane -class Plane -{ -public: - VTKM_CONT - Plane() - : Origin(FloatDefault(0)), - Normal(FloatDefault(0), FloatDefault(0), FloatDefault(1)) - { } - - VTKM_CONT - explicit Plane(const vtkm::Vec &normal) - : Origin(FloatDefault(0)), - Normal(normal) - { } - - VTKM_CONT - Plane(const vtkm::Vec &origin, - const vtkm::Vec &normal) - : Origin(origin), Normal(normal) - { } - - VTKM_EXEC_CONT - const vtkm::Vec& GetOrigin() const - { - return this->Origin; - } - - VTKM_EXEC_CONT - const vtkm::Vec& GetNormal() const - { - return this->Normal; - } - - VTKM_EXEC_CONT - FloatDefault Value(FloatDefault x, FloatDefault y, FloatDefault z) const - { - return ((x - this->Origin[0]) * this->Normal[0]) + - ((y - this->Origin[1]) * this->Normal[1]) + - ((z - this->Origin[2]) * this->Normal[2]); - } - - VTKM_EXEC_CONT - FloatDefault Value(const vtkm::Vec &x) const - { - return this->Value(x[0], x[1], x[2]); - } - - VTKM_EXEC_CONT - vtkm::Vec Gradient(FloatDefault, FloatDefault, FloatDefault) const - { - return this->Normal; - } - - VTKM_EXEC_CONT - vtkm::Vec Gradient(const vtkm::Vec&) const - { - return this->Normal; - } - -private: - vtkm::Vec Origin; - vtkm::Vec Normal; -}; - - -/// \brief Implicit function for a sphere -class Sphere -{ -public: - VTKM_CONT - Sphere() : Radius(FloatDefault(0.2)), Center(FloatDefault(0)) - { } - - VTKM_CONT - explicit Sphere(FloatDefault radius) : Radius(radius), Center(FloatDefault(0)) - { } - - VTKM_CONT - Sphere(vtkm::Vec center, FloatDefault radius) - : Radius(radius), Center(center) - { } - - VTKM_EXEC_CONT - FloatDefault GetRadius() const - { - return this->Radius; - } - - VTKM_EXEC_CONT - const vtkm::Vec& GetCenter() const - { - return this->Center; - } - - VTKM_EXEC_CONT - FloatDefault Value(FloatDefault x, FloatDefault y, FloatDefault z) const - { - return ((x - this->Center[0]) * (x - this->Center[0]) + - (y - this->Center[1]) * (y - this->Center[1]) + - (z - this->Center[2]) * (z - this->Center[2])) - - (this->Radius * this->Radius); - } - - VTKM_EXEC_CONT - FloatDefault Value(const vtkm::Vec &x) const - { - return this->Value(x[0], x[1], x[2]); - } - - VTKM_EXEC_CONT - vtkm::Vec Gradient(FloatDefault x, FloatDefault y, FloatDefault z) - const - { - return this->Gradient(vtkm::Vec(x, y, z)); - } - - VTKM_EXEC_CONT - vtkm::Vec Gradient(const vtkm::Vec &x) const - { - return FloatDefault(2) * (x - this->Center); - } - -private: - FloatDefault Radius; - vtkm::Vec Center; -}; - -/// \brief Implicit function for a box -class Box -{ -public: - VTKM_CONT - Box() : MinPoint(vtkm::Vec(FloatDefault(0), FloatDefault(0), FloatDefault(0))), - MaxPoint(vtkm::Vec(FloatDefault(1), FloatDefault(1), FloatDefault(1))) - { } - - VTKM_CONT - Box(vtkm::Vec minPoint, vtkm::Vec maxPoint) - : MinPoint(minPoint), MaxPoint(maxPoint) - { } - - VTKM_CONT - Box(FloatDefault xmin, FloatDefault xmax, - FloatDefault ymin, FloatDefault ymax, - FloatDefault zmin, FloatDefault zmax) - { - MinPoint[0] = xmin; MaxPoint[0] = xmax; - MinPoint[1] = ymin; MaxPoint[1] = ymax; - MinPoint[2] = zmin; MaxPoint[2] = zmax; - } - - VTKM_EXEC_CONT - const vtkm::Vec& GetMinPoint() const - { - return this->MinPoint; - } - - VTKM_EXEC_CONT - const vtkm::Vec& GetMaxPoint() const - { - return this->MaxPoint; - } - - VTKM_EXEC_CONT - FloatDefault Value(const vtkm::Vec &x) const - { - FloatDefault minDistance = vtkm::NegativeInfinity32(); - FloatDefault diff, t, dist; - FloatDefault distance = FloatDefault(0.0); - vtkm::IdComponent inside = 1; - - for (vtkm::IdComponent d = 0; d < 3; d++) - { - diff = this->MaxPoint[d] - this->MinPoint[d]; - if (diff != FloatDefault(0.0)) - { - t = (x[d] - this->MinPoint[d]) / diff; - // Outside before the box - if (t < FloatDefault(0.0)) - { - inside = 0; - dist = this->MinPoint[d] - x[d]; - } - // Outside after the box - else if (t > FloatDefault(1.0)) - { - inside = 0; - dist = x[d] - this->MaxPoint[d]; - } - else - { - // Inside the box in lower half - if (t <= FloatDefault(0.5)) - { - dist = MinPoint[d] - x[d]; - } - // Inside the box in upper half - else - { - dist = x[d] - MaxPoint[d]; - } - if (dist > minDistance) - { - minDistance = dist; - } - } - } - else - { - dist = vtkm::Abs(x[d] - MinPoint[d]); - if (dist > FloatDefault(0.0)) - { - inside = 0; - } - } - if (dist > FloatDefault(0.0)) - { - distance += dist*dist; - } - } - - distance = vtkm::Sqrt(distance); - if (inside) - { - return minDistance; - } - else - { - return distance; - } - } - - VTKM_EXEC_CONT - FloatDefault Value(FloatDefault x, FloatDefault y, FloatDefault z) const - { - return this->Value(vtkm::Vec(x, y, z)); - } - - VTKM_EXEC_CONT - vtkm::Vec Gradient(const vtkm::Vec &x) const - { - vtkm::IdComponent minAxis = 0; - FloatDefault dist = 0.0; - FloatDefault minDist = vtkm::Infinity32(); - vtkm::Vec location; - vtkm::Vec normal; - vtkm::Vec inside(FloatDefault(0), FloatDefault(0), FloatDefault(0)); - vtkm::Vec outside(FloatDefault(0), FloatDefault(0), FloatDefault(0)); - vtkm::Vec center((this->MaxPoint[0] + this->MinPoint[0]) * FloatDefault(0.5), - (this->MaxPoint[1] + this->MinPoint[1]) * FloatDefault(0.5), - (this->MaxPoint[2] + this->MinPoint[2]) * FloatDefault(0.5)); - - // Compute the location of the point with respect to the box - // Point will lie in one of 27 separate regions around or within the box - // Gradient vector is computed differently in each of the regions. - for (vtkm::IdComponent d = 0; d < 3; d++) - { - if (x[d] < this->MinPoint[d]) - { - // Outside the box low end - location[d] = 0; - outside[d] = -1.0; - } - else if (x[d] > this->MaxPoint[d]) - { - // Outside the box high end - location[d] = 2; - outside[d] = 1.0; - } - else - { - location[d] = 1; - if (x[d] <= center[d]) - { - // Inside the box low end - dist = x[d] - this->MinPoint[d]; - inside[d] = -1.0; - } - else - { - // Inside the box high end - dist = this->MaxPoint[d] - x[d]; - inside[d] = 1.0; - } - if (dist < minDist) // dist is negative - { - minDist = dist; - minAxis = d; - } - } - } - - vtkm::Id indx = location[0] + 3*location[1] + 9*location[2]; - switch (indx) - { - // verts - gradient points away from center point - case 0: case 2: case 6: case 8: case 18: case 20: case 24: case 26: - for (vtkm::IdComponent d = 0; d < 3; d++) - { - normal[d] = x[d] - center[d]; - } - vtkm::Normalize(normal); - break; - - // edges - gradient points out from axis of cube - case 1: case 3: case 5: case 7: - case 9: case 11: case 15: case 17: - case 19: case 21: case 23: case 25: - for (vtkm::IdComponent d = 0; d < 3; d++) - { - if (outside[d] != 0.0) - { - normal[d] = x[d] - center[d]; - } - else - { - normal[d] = 0.0; - } - } - vtkm::Normalize(normal); - break; - - // faces - gradient points perpendicular to face - case 4: case 10: case 12: case 14: case 16: case 22: - for (vtkm::IdComponent d = 0; d < 3; d++) - { - normal[d] = outside[d]; - } - break; - - // interior - gradient is perpendicular to closest face - case 13: - normal[0] = normal[1] = normal[2] = 0.0; - normal[minAxis] = inside[minAxis]; - break; - default: - VTKM_ASSERT(false); - break; - } - return normal; - } - - VTKM_EXEC_CONT - vtkm::Vec Gradient(FloatDefault x, FloatDefault y, FloatDefault z) - const - { - return this->Gradient(vtkm::Vec(x, y, z)); - } - -private: - vtkm::Vec MinPoint; - vtkm::Vec MaxPoint; -}; - -/// \brief A function object that evaluates the contained implicit function -template -class ImplicitFunctionValue -{ -public: - VTKM_CONT - ImplicitFunctionValue() - : Function() - { } - - VTKM_CONT - explicit ImplicitFunctionValue(const ImplicitFunction &func) - : Function(func) - { } - - VTKM_EXEC_CONT - FloatDefault operator()(const vtkm::Vec x) const - { - return this->Function.Value(x); - } - - VTKM_EXEC_CONT - FloatDefault operator()(FloatDefault x, FloatDefault y, FloatDefault z) const - { - return this->Function.Value(x, y, z); - } - -private: - ImplicitFunction Function; -}; - -/// \brief A function object that computes the gradient of the contained implicit -/// function and the specified point. -template -class ImplicitFunctionGradient -{ -public: - VTKM_CONT - ImplicitFunctionGradient() - : Function() - { } - - VTKM_CONT - explicit ImplicitFunctionGradient(const ImplicitFunction &func) - : Function(func) - { } - - VTKM_EXEC_CONT - FloatDefault operator()(const vtkm::Vec x) const - { - return this->Function.Gradient(x); - } - - VTKM_EXEC_CONT - FloatDefault operator()(FloatDefault x, FloatDefault y, FloatDefault z) const - { - return this->Function.Gradient(x, y, z); - } - -private: - ImplicitFunction Function; -}; - -} // namespace vtkm - -#endif // vtk_m_ImplicitFunctions_h diff --git a/vtkm/benchmarking/BenchmarkFieldAlgorithms.cxx b/vtkm/benchmarking/BenchmarkFieldAlgorithms.cxx index 7421512be..414e1e47f 100644 --- a/vtkm/benchmarking/BenchmarkFieldAlgorithms.cxx +++ b/vtkm/benchmarking/BenchmarkFieldAlgorithms.cxx @@ -25,6 +25,7 @@ #include #include #include +#include #include #include @@ -48,12 +49,14 @@ const static std::string DIVIDER(40, '-'); enum BenchmarkName { BLACK_SCHOLES = 1, MATH = 1 << 1, - FUSED_MATH = 1 << 3, - INTERPOLATE_FIELD = 1 << 4, + FUSED_MATH = 1 << 2, + INTERPOLATE_FIELD = 1 << 3, + IMPLICIT_FUNCTION = 1 << 4, ALL = BLACK_SCHOLES | MATH | FUSED_MATH | - INTERPOLATE_FIELD + INTERPOLATE_FIELD | + IMPLICIT_FUNCTION }; template @@ -268,6 +271,51 @@ public: } }; +template +class EvaluateImplicitFunction : public vtkm::worklet::WorkletMapField +{ +public: + typedef void ControlSignature(FieldIn, FieldOut); + typedef void ExecutionSignature(_1, _2); + + EvaluateImplicitFunction(const ImplicitFunction &function) + : Function(function) + { } + + template + VTKM_EXEC + void operator()(const VecType &point, ScalarType &val) const + { + val = this->Function.Value(point); + } + +private: + ImplicitFunction Function; +}; + +template +class Evaluate2ImplicitFunctions : public vtkm::worklet::WorkletMapField +{ +public: + typedef void ControlSignature(FieldIn, FieldOut); + typedef void ExecutionSignature(_1, _2); + + Evaluate2ImplicitFunctions(const T1 &f1, const T2 &f2) + : Function1(f1), Function2(f2) + { } + + template + VTKM_EXEC + void operator()(const VecType &point, ScalarType &val) const + { + val = this->Function1.Value(point) + this->Function2.Value(point); + } + +private: + T1 Function1; + T2 Function2; +}; + struct ValueTypes : vtkm::ListTagBase{}; struct InterpValueTypes : vtkm::ListTagBase> Points; + vtkm::cont::ArrayHandle Result; + vtkm::cont::Sphere Sphere1, Sphere2; + }; + + static ImplicitFunctionBenchData MakeImplicitFunctionBenchData() + { + vtkm::Id count = ARRAY_SIZE; + vtkm::FloatDefault bounds[6] = { -2.0f, 2.0f, -2.0f, 2.0f, -2.0f, 2.0f }; + + ImplicitFunctionBenchData data; + data.Points.Allocate(count); + data.Result.Allocate(count); + + std::default_random_engine rangen; + std::uniform_real_distribution distx(bounds[0], bounds[1]); + std::uniform_real_distribution disty(bounds[2], bounds[3]); + std::uniform_real_distribution distz(bounds[4], bounds[5]); + + auto portal = data.Points.GetPortalControl(); + for (vtkm::Id i = 0; i < count; ++i) + { + portal.Set(i, vtkm::make_Vec(distx(rangen), disty(rangen), distz(rangen))); + } + + data.Sphere1 = vtkm::cont::Sphere({0.22f, 0.33f, 0.44f}, 0.55f); + data.Sphere2 = vtkm::cont::Sphere({0.22f, 0.33f, 0.11f}, 0.77f); + + return data; + } + + template + struct BenchImplicitFunction + { + BenchImplicitFunction() + : Internal(MakeImplicitFunctionBenchData()) + { } + + VTKM_CONT + vtkm::Float64 operator()() + { + using EvalWorklet = EvaluateImplicitFunction; + using EvalDispatcher = vtkm::worklet::DispatcherMapField; + + EvalWorklet eval(Internal.Sphere1); + + vtkm::cont::Timer timer; + EvalDispatcher(eval).Invoke(this->Internal.Points, this->Internal.Result); + return timer.GetElapsedTime(); + } + + VTKM_CONT + std::string Description() const + { + std::stringstream description; + description << "Implicit Function (vtkm::cont::Sphere) on " + << Internal.Points.GetNumberOfValues() + << " points"; + return description.str(); + } + + ImplicitFunctionBenchData Internal; + }; + + template + struct BenchDynamicImplicitFunction + { + BenchDynamicImplicitFunction() + : Internal(MakeImplicitFunctionBenchData()) + { } + + VTKM_CONT + vtkm::Float64 operator()() + { + using EvalWorklet = EvaluateImplicitFunction; + using EvalDispatcher = vtkm::worklet::DispatcherMapField; + + EvalWorklet eval(Internal.Sphere1.PrepareForExecution(DeviceAdapterTag())); + + vtkm::cont::Timer timer; + EvalDispatcher(eval).Invoke(this->Internal.Points, this->Internal.Result); + return timer.GetElapsedTime(); + } + + VTKM_CONT + std::string Description() const + { + std::stringstream description; + description << "Implicit Function (DynamicImplicitFunction) on " + << Internal.Points.GetNumberOfValues() + << " points"; + return description.str(); + } + + ImplicitFunctionBenchData Internal; + }; + + template + struct Bench2ImplicitFunctions + { + Bench2ImplicitFunctions() + : Internal(MakeImplicitFunctionBenchData()) + { } + + VTKM_CONT + vtkm::Float64 operator()() + { + using EvalWorklet = Evaluate2ImplicitFunctions; + using EvalDispatcher = vtkm::worklet::DispatcherMapField; + + EvalWorklet eval(Internal.Sphere1, Internal.Sphere2); + + vtkm::cont::Timer timer; + EvalDispatcher(eval).Invoke(this->Internal.Points, this->Internal.Result); + return timer.GetElapsedTime(); + } + + VTKM_CONT + std::string Description() const + { + std::stringstream description; + description << "Implicit Function 2x(vtkm::cont::Sphere) on " + << Internal.Points.GetNumberOfValues() + << " points"; + return description.str(); + } + + ImplicitFunctionBenchData Internal; + }; + + template + struct Bench2DynamicImplicitFunctions + { + Bench2DynamicImplicitFunctions() + : Internal(MakeImplicitFunctionBenchData()) + { } + + VTKM_CONT + vtkm::Float64 operator()() + { + using EvalWorklet = Evaluate2ImplicitFunctions; + using EvalDispatcher = vtkm::worklet::DispatcherMapField; + + EvalWorklet eval(Internal.Sphere1.PrepareForExecution(DeviceAdapterTag()), + Internal.Sphere2.PrepareForExecution(DeviceAdapterTag())); + + vtkm::cont::Timer timer; + EvalDispatcher(eval).Invoke(this->Internal.Points, this->Internal.Result); + return timer.GetElapsedTime(); + } + + VTKM_CONT + std::string Description() const + { + std::stringstream description; + description << "Implicit Function 2x(DynamicImplicitFunction) on " + << Internal.Points.GetNumberOfValues() + << " points"; + return description.str(); + } + + ImplicitFunctionBenchData Internal; + }; + + VTKM_MAKE_BENCHMARK(ImplicitFunction, BenchImplicitFunction); + VTKM_MAKE_BENCHMARK(ImplicitFunctionDynamic, BenchDynamicImplicitFunction); + VTKM_MAKE_BENCHMARK(ImplicitFunction2, Bench2ImplicitFunctions); + VTKM_MAKE_BENCHMARK(ImplicitFunctionDynamic2, Bench2DynamicImplicitFunctions); public: @@ -687,6 +906,16 @@ public: VTKM_RUN_BENCHMARK(EdgeInterpDynamic, InterpValueTypes()); } + if (benchmarks & IMPLICIT_FUNCTION) { + using FloatDefaultType = vtkm::ListTagBase; + + std::cout << "\nBenchmarking Implicit Function\n"; + VTKM_RUN_BENCHMARK(ImplicitFunction, FloatDefaultType()); + VTKM_RUN_BENCHMARK(ImplicitFunctionDynamic, FloatDefaultType()); + VTKM_RUN_BENCHMARK(ImplicitFunction2, FloatDefaultType()); + VTKM_RUN_BENCHMARK(ImplicitFunctionDynamic2, FloatDefaultType()); + } + return 0; } }; @@ -723,6 +952,10 @@ int main(int argc, char *argv[]) { benchmarks |= vtkm::benchmarking::INTERPOLATE_FIELD; } + else if (arg == "implicit_function") + { + benchmarks |= vtkm::benchmarking::IMPLICIT_FUNCTION; + } else { std::cout << "Unrecognized benchmark: " << argv[i] << std::endl; diff --git a/vtkm/cont/CMakeLists.txt b/vtkm/cont/CMakeLists.txt index 1decadf56..fcb269af0 100644 --- a/vtkm/cont/CMakeLists.txt +++ b/vtkm/cont/CMakeLists.txt @@ -64,6 +64,7 @@ set(headers ErrorExecution.h ErrorInternal.h Field.h + ImplicitFunction.h RuntimeDeviceInformation.h RuntimeDeviceTracker.h Storage.h @@ -78,6 +79,7 @@ set(headers set(header_impls ArrayHandle.hxx CellSetStructured.hxx + ImplicitFunction.hxx StorageBasic.hxx ) @@ -88,6 +90,7 @@ set(sources CoordinateSystem.cxx DynamicArrayHandle.cxx Field.cxx + ImplicitFunction.cxx internal/SimplePolymorphicContainer.cxx StorageBasic.cxx ) diff --git a/vtkm/testing/UnitTestImplicitFunctions.cxx b/vtkm/cont/ImplicitFunction.cxx similarity index 62% rename from vtkm/testing/UnitTestImplicitFunctions.cxx rename to vtkm/cont/ImplicitFunction.cxx index 4e7b8a3c6..8bba1a17c 100644 --- a/vtkm/testing/UnitTestImplicitFunctions.cxx +++ b/vtkm/cont/ImplicitFunction.cxx @@ -1,30 +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 2015 Sandia Corporation. -// Copyright 2015 UT-Battelle, LLC. -// Copyright 2015 Los Alamos National Security. +// 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. -// -//============================================================================= +//============================================================================ +#include -#include -#include +namespace vtkm { +namespace cont { + +ImplicitFunction::~ImplicitFunction() = default; -int UnitTestImplicitFunctions(int, char *[]) -{ - return vtkm::cont::testing::Testing::Run( - vtkm::testing::TestingImplicitFunctions::Run); } +} // vtkm::cont diff --git a/vtkm/cont/ImplicitFunction.h b/vtkm/cont/ImplicitFunction.h new file mode 100644 index 000000000..a2bb33376 --- /dev/null +++ b/vtkm/cont/ImplicitFunction.h @@ -0,0 +1,341 @@ +//============================================================================ +// 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_cont_ImplicitFunction_h +#define vtk_m_cont_ImplicitFunction_h + +#include +#include + +#include + + +namespace vtkm { +namespace cont { + +class VTKM_CONT_EXPORT ImplicitFunction +{ +public: + virtual ~ImplicitFunction(); + + template + vtkm::exec::ImplicitFunction PrepareForExecution(DeviceAdapter device) const + { + if (!this->Cache->GetValid()) + { + this->SetDefaultDevices(); + } + return this->Cache->GetVirtualObject(device); + } + + void Modified() + { + this->Cache->SetRefreshFlag(true); + } + +protected: + using CacheType = vtkm::cont::VirtualObjectCache; + + ImplicitFunction() : Cache(new CacheType) + { + } + + ImplicitFunction(ImplicitFunction &&other) + : Cache(std::move(other.Cache)) + { + } + + ImplicitFunction& operator=(ImplicitFunction &&other) + { + if (this != &other) + { + this->Cache = std::move(other.Cache); + } + return *this; + } + + virtual void SetDefaultDevices() const = 0; + + std::unique_ptr Cache; +}; + + +template +class VTKM_ALWAYS_EXPORT ImplicitFunctionImpl : public ImplicitFunction +{ +public: + template + void ResetDevices(DeviceAdapterList devices) + { + this->Cache->Bind(static_cast(this), devices); + } + +protected: + ImplicitFunctionImpl() = default; + ImplicitFunctionImpl(const ImplicitFunctionImpl &) : ImplicitFunction() + { + } + + // Cannot default due to a bug in VS2013 + ImplicitFunctionImpl(ImplicitFunctionImpl &&other) + : ImplicitFunction(std::move(other)) + { + } + + ImplicitFunctionImpl& operator=(const ImplicitFunctionImpl &) + { + return *this; + } + + // Cannot default due to a bug in VS2013 + ImplicitFunctionImpl& operator=(ImplicitFunctionImpl &&other) + { + ImplicitFunction::operator=(std::move(other)); + return *this; + } + + void SetDefaultDevices() const override + { + this->Cache->Bind(static_cast(this)); + } +}; + + +//============================================================================ +// ImplicitFunctions: + +//============================================================================ +/// \brief Implicit function for a box +class VTKM_ALWAYS_EXPORT Box : public ImplicitFunctionImpl +{ +public: + Box() : MinPoint(vtkm::Vec(FloatDefault(0), FloatDefault(0), FloatDefault(0))), + MaxPoint(vtkm::Vec(FloatDefault(1), FloatDefault(1), FloatDefault(1))) + { } + + Box(vtkm::Vec minPoint, vtkm::Vec maxPoint) + : MinPoint(minPoint), MaxPoint(maxPoint) + { } + + Box(FloatDefault xmin, FloatDefault xmax, + FloatDefault ymin, FloatDefault ymax, + FloatDefault zmin, FloatDefault zmax) + { + MinPoint[0] = xmin; MaxPoint[0] = xmax; + MinPoint[1] = ymin; MaxPoint[1] = ymax; + MinPoint[2] = zmin; MaxPoint[2] = zmax; + } + + void SetMinPoint(const vtkm::Vec &point) + { + this->MinPoint = point; + this->Modified(); + } + + void SetMaxPoint(const vtkm::Vec &point) + { + this->MaxPoint = point; + this->Modified(); + } + + const vtkm::Vec& GetMinPoint() const + { + return this->MinPoint; + } + + const vtkm::Vec& GetMaxPoint() const + { + return this->MaxPoint; + } + + VTKM_EXEC_CONT + FloatDefault Value(const vtkm::Vec &x) const; + + VTKM_EXEC_CONT + FloatDefault Value(FloatDefault x, FloatDefault y, FloatDefault z) const + { + return this->Value(vtkm::Vec(x, y, z)); + } + + VTKM_EXEC_CONT + vtkm::Vec Gradient(const vtkm::Vec &x) const; + + VTKM_EXEC_CONT + vtkm::Vec Gradient(FloatDefault x, FloatDefault y, FloatDefault z) + const + { + return this->Gradient(vtkm::Vec(x, y, z)); + } + +private: + vtkm::Vec MinPoint; + vtkm::Vec MaxPoint; +}; + + +//============================================================================ +/// \brief Implicit function for a plane +class VTKM_ALWAYS_EXPORT Plane : public ImplicitFunctionImpl +{ +public: + Plane() + : Origin(FloatDefault(0)), + Normal(FloatDefault(0), FloatDefault(0), FloatDefault(1)) + { } + + explicit Plane(const vtkm::Vec &normal) + : Origin(FloatDefault(0)), + Normal(normal) + { } + + Plane(const vtkm::Vec &origin, + const vtkm::Vec &normal) + : Origin(origin), Normal(normal) + { } + + void SetOrigin(const vtkm::Vec &origin) + { + this->Origin = origin; + this->Modified(); + } + + void SetNormal(const vtkm::Vec &normal) + { + this->Normal = normal; + this->Modified(); + } + + const vtkm::Vec& GetOrigin() const + { + return this->Origin; + } + + const vtkm::Vec& GetNormal() const + { + return this->Normal; + } + + VTKM_EXEC_CONT + FloatDefault Value(FloatDefault x, FloatDefault y, FloatDefault z) const + { + return ((x - this->Origin[0]) * this->Normal[0]) + + ((y - this->Origin[1]) * this->Normal[1]) + + ((z - this->Origin[2]) * this->Normal[2]); + } + + VTKM_EXEC_CONT + FloatDefault Value(const vtkm::Vec &x) const + { + return this->Value(x[0], x[1], x[2]); + } + + VTKM_EXEC_CONT + vtkm::Vec Gradient(FloatDefault, FloatDefault, FloatDefault) const + { + return this->Normal; + } + + VTKM_EXEC_CONT + vtkm::Vec Gradient(const vtkm::Vec&) const + { + return this->Normal; + } + +private: + vtkm::Vec Origin; + vtkm::Vec Normal; +}; + + +//============================================================================ +/// \brief Implicit function for a sphere +class VTKM_ALWAYS_EXPORT Sphere : public ImplicitFunctionImpl +{ +public: + Sphere() : Radius(FloatDefault(0.2)), Center(FloatDefault(0)) + { } + + explicit Sphere(FloatDefault radius) : Radius(radius), Center(FloatDefault(0)) + { } + + Sphere(vtkm::Vec center, FloatDefault radius) + : Radius(radius), Center(center) + { } + + void SetRadius(FloatDefault radius) + { + this->Radius = radius; + this->Modified(); + } + + void GetCenter(const vtkm::Vec ¢er) + { + this->Center = center; + this->Modified(); + } + + FloatDefault GetRadius() const + { + return this->Radius; + } + + const vtkm::Vec& GetCenter() const + { + return this->Center; + } + + VTKM_EXEC_CONT + FloatDefault Value(FloatDefault x, FloatDefault y, FloatDefault z) const + { + return ((x - this->Center[0]) * (x - this->Center[0]) + + (y - this->Center[1]) * (y - this->Center[1]) + + (z - this->Center[2]) * (z - this->Center[2])) - + (this->Radius * this->Radius); + } + + VTKM_EXEC_CONT + FloatDefault Value(const vtkm::Vec &x) const + { + return this->Value(x[0], x[1], x[2]); + } + + VTKM_EXEC_CONT + vtkm::Vec Gradient(FloatDefault x, FloatDefault y, FloatDefault z) + const + { + return this->Gradient(vtkm::Vec(x, y, z)); + } + + VTKM_EXEC_CONT + vtkm::Vec Gradient(const vtkm::Vec &x) const + { + return FloatDefault(2) * (x - this->Center); + } + +private: + FloatDefault Radius; + vtkm::Vec Center; +}; + +} +} // vtkm::cont + +#include + +#endif // vtk_m_cont_ImplicitFunction_h diff --git a/vtkm/cont/ImplicitFunction.hxx b/vtkm/cont/ImplicitFunction.hxx new file mode 100644 index 000000000..ffe053393 --- /dev/null +++ b/vtkm/cont/ImplicitFunction.hxx @@ -0,0 +1,204 @@ +//============================================================================ +// 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. +//============================================================================ +#include + +namespace vtkm { +namespace cont { + +//============================================================================ +VTKM_EXEC_CONT +inline FloatDefault +Box::Value(const vtkm::Vec &x) const +{ + FloatDefault minDistance = vtkm::NegativeInfinity32(); + FloatDefault diff, t, dist; + FloatDefault distance = FloatDefault(0.0); + vtkm::IdComponent inside = 1; + + for (vtkm::IdComponent d = 0; d < 3; d++) + { + diff = this->MaxPoint[d] - this->MinPoint[d]; + if (diff != FloatDefault(0.0)) + { + t = (x[d] - this->MinPoint[d]) / diff; + // Outside before the box + if (t < FloatDefault(0.0)) + { + inside = 0; + dist = this->MinPoint[d] - x[d]; + } + // Outside after the box + else if (t > FloatDefault(1.0)) + { + inside = 0; + dist = x[d] - this->MaxPoint[d]; + } + else + { + // Inside the box in lower half + if (t <= FloatDefault(0.5)) + { + dist = MinPoint[d] - x[d]; + } + // Inside the box in upper half + else + { + dist = x[d] - MaxPoint[d]; + } + if (dist > minDistance) + { + minDistance = dist; + } + } + } + else + { + dist = vtkm::Abs(x[d] - MinPoint[d]); + if (dist > FloatDefault(0.0)) + { + inside = 0; + } + } + if (dist > FloatDefault(0.0)) + { + distance += dist*dist; + } + } + + distance = vtkm::Sqrt(distance); + if (inside) + { + return minDistance; + } + else + { + return distance; + } +} + +//============================================================================ +VTKM_EXEC_CONT +inline vtkm::Vec +Box::Gradient(const vtkm::Vec &x) const +{ + vtkm::IdComponent minAxis = 0; + FloatDefault dist = 0.0; + FloatDefault minDist = vtkm::Infinity32(); + vtkm::Vec location; + vtkm::Vec normal; + vtkm::Vec inside(FloatDefault(0), FloatDefault(0), FloatDefault(0)); + vtkm::Vec outside(FloatDefault(0), FloatDefault(0), FloatDefault(0)); + vtkm::Vec center((this->MaxPoint[0] + this->MinPoint[0]) * FloatDefault(0.5), + (this->MaxPoint[1] + this->MinPoint[1]) * FloatDefault(0.5), + (this->MaxPoint[2] + this->MinPoint[2]) * FloatDefault(0.5)); + + // Compute the location of the point with respect to the box + // Point will lie in one of 27 separate regions around or within the box + // Gradient vector is computed differently in each of the regions. + for (vtkm::IdComponent d = 0; d < 3; d++) + { + if (x[d] < this->MinPoint[d]) + { + // Outside the box low end + location[d] = 0; + outside[d] = -1.0; + } + else if (x[d] > this->MaxPoint[d]) + { + // Outside the box high end + location[d] = 2; + outside[d] = 1.0; + } + else + { + location[d] = 1; + if (x[d] <= center[d]) + { + // Inside the box low end + dist = x[d] - this->MinPoint[d]; + inside[d] = -1.0; + } + else + { + // Inside the box high end + dist = this->MaxPoint[d] - x[d]; + inside[d] = 1.0; + } + if (dist < minDist) // dist is negative + { + minDist = dist; + minAxis = d; + } + } + } + + vtkm::Id indx = location[0] + 3*location[1] + 9*location[2]; + switch (indx) + { + // verts - gradient points away from center point + case 0: case 2: case 6: case 8: case 18: case 20: case 24: case 26: + for (vtkm::IdComponent d = 0; d < 3; d++) + { + normal[d] = x[d] - center[d]; + } + vtkm::Normalize(normal); + break; + + // edges - gradient points out from axis of cube + case 1: case 3: case 5: case 7: + case 9: case 11: case 15: case 17: + case 19: case 21: case 23: case 25: + for (vtkm::IdComponent d = 0; d < 3; d++) + { + if (outside[d] != 0.0) + { + normal[d] = x[d] - center[d]; + } + else + { + normal[d] = 0.0; + } + } + vtkm::Normalize(normal); + break; + + // faces - gradient points perpendicular to face + case 4: case 10: case 12: case 14: case 16: case 22: + for (vtkm::IdComponent d = 0; d < 3; d++) + { + normal[d] = outside[d]; + } + break; + + // interior - gradient is perpendicular to closest face + case 13: + normal[0] = normal[1] = normal[2] = 0.0; + normal[minAxis] = inside[minAxis]; + break; + default: + VTKM_ASSERT(false); + break; + } + return normal; +} + + +} +} // vtkm::cont diff --git a/vtkm/cont/cuda/testing/CMakeLists.txt b/vtkm/cont/cuda/testing/CMakeLists.txt index b937fb3c9..22ea32bbb 100644 --- a/vtkm/cont/cuda/testing/CMakeLists.txt +++ b/vtkm/cont/cuda/testing/CMakeLists.txt @@ -25,6 +25,7 @@ set(unit_tests UnitTestCudaDataSetExplicit.cu UnitTestCudaDataSetSingleType.cu UnitTestCudaDeviceAdapter.cu + UnitTestCudaImplicitFunction.cu UnitTestCudaMath.cu UnitTestCudaVirtualObjectCache.cu ) diff --git a/vtkm/cont/cuda/testing/UnitTestCudaImplicitFunction.cu b/vtkm/cont/cuda/testing/UnitTestCudaImplicitFunction.cu new file mode 100644 index 000000000..7f5112184 --- /dev/null +++ b/vtkm/cont/cuda/testing/UnitTestCudaImplicitFunction.cu @@ -0,0 +1,37 @@ +//============================================================================ +// 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. +//============================================================================ + +#include + +namespace { + +void TestImplicitFunctions() +{ + vtkm::cont::testing::TestingImplicitFunction testing; + testing.Run(vtkm::cont::DeviceAdapterTagCuda()); +} + +} // anonymous namespace + + +int UnitTestCudaImplicitFunction(int, char *[]) +{ + return vtkm::cont::testing::Testing::Run(TestImplicitFunctions); +} diff --git a/vtkm/cont/serial/testing/CMakeLists.txt b/vtkm/cont/serial/testing/CMakeLists.txt index e97a35fd0..4e17f939e 100644 --- a/vtkm/cont/serial/testing/CMakeLists.txt +++ b/vtkm/cont/serial/testing/CMakeLists.txt @@ -25,6 +25,7 @@ set(unit_tests UnitTestSerialDataSetExplicit.cxx UnitTestSerialDataSetSingleType.cxx UnitTestSerialDeviceAdapter.cxx + UnitTestSerialImplicitFunction.cxx UnitTestSerialVirtualObjectCache.cxx ) vtkm_unit_tests(TBB SOURCES ${unit_tests}) diff --git a/vtkm/cont/serial/testing/UnitTestSerialImplicitFunction.cxx b/vtkm/cont/serial/testing/UnitTestSerialImplicitFunction.cxx new file mode 100644 index 000000000..367a2f414 --- /dev/null +++ b/vtkm/cont/serial/testing/UnitTestSerialImplicitFunction.cxx @@ -0,0 +1,38 @@ +//============================================================================ +// 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. +//============================================================================ + +#include + +namespace { + +void TestImplicitFunctions() +{ + vtkm::cont::testing::TestingImplicitFunction testing; + testing.Run(vtkm::cont::DeviceAdapterTagSerial()); +} + +} // anonymous namespace + + +int UnitTestSerialImplicitFunction(int, char *[]) +{ + return vtkm::cont::testing::Testing::Run(TestImplicitFunctions); +} + diff --git a/vtkm/cont/tbb/testing/CMakeLists.txt b/vtkm/cont/tbb/testing/CMakeLists.txt index f6d15b058..b69e76211 100644 --- a/vtkm/cont/tbb/testing/CMakeLists.txt +++ b/vtkm/cont/tbb/testing/CMakeLists.txt @@ -25,6 +25,7 @@ set(unit_tests UnitTestTBBDataSetExplicit.cxx UnitTestTBBDataSetSingleType.cxx UnitTestTBBDeviceAdapter.cxx + UnitTestTBBImplicitFunction.cxx UnitTestTBBVirtualObjectCache.cxx ) vtkm_unit_tests(TBB SOURCES ${unit_tests}) diff --git a/vtkm/cont/tbb/testing/UnitTestTBBImplicitFunction.cxx b/vtkm/cont/tbb/testing/UnitTestTBBImplicitFunction.cxx new file mode 100644 index 000000000..b6bd19b72 --- /dev/null +++ b/vtkm/cont/tbb/testing/UnitTestTBBImplicitFunction.cxx @@ -0,0 +1,37 @@ +//============================================================================ +// 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. +//============================================================================ + +#include + +namespace { + +void TestImplicitFunctions() +{ + vtkm::cont::testing::TestingImplicitFunction testing; + testing.Run(vtkm::cont::DeviceAdapterTagTBB()); +} + +} // anonymous namespace + + +int UnitTestTBBImplicitFunction(int, char *[]) +{ + return vtkm::cont::testing::Testing::Run(TestImplicitFunctions); +} diff --git a/vtkm/cont/testing/CMakeLists.txt b/vtkm/cont/testing/CMakeLists.txt index 6ea6a1078..8d3a8ab9e 100644 --- a/vtkm/cont/testing/CMakeLists.txt +++ b/vtkm/cont/testing/CMakeLists.txt @@ -28,6 +28,7 @@ set(headers TestingDataSetExplicit.h TestingDataSetSingleType.h TestingFancyArrayHandles.h + TestingImplicitFunction.h TestingVirtualObjectCache.h ) diff --git a/vtkm/cont/testing/TestingImplicitFunction.h b/vtkm/cont/testing/TestingImplicitFunction.h new file mode 100644 index 000000000..d366e4b8a --- /dev/null +++ b/vtkm/cont/testing/TestingImplicitFunction.h @@ -0,0 +1,187 @@ +//============================================================================ +// Copyright (c) Kitware, Inc. +// All rights reserved. +// See LICENSE.txt for details. +// This software is distributed WITHOUT ANY WARRANTY; without even +// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR +// PURPOSE. See the above copyright notice for more information. +// +// 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_testing_TestingImplicitFunction_h +#define vtk_m_cont_testing_TestingImplicitFunction_h + +#include +#include +#include +#include +#include +#include + +#include + +#include +#include + +#include + + +namespace vtkm { +namespace cont { +namespace testing { + +namespace implicit_function_detail { + +class EvaluateImplicitFunction : public vtkm::worklet::WorkletMapField +{ +public: + typedef void ControlSignature(FieldIn, FieldOut); + typedef void ExecutionSignature(_1, _2); + + EvaluateImplicitFunction(const vtkm::exec::ImplicitFunction &function) + : Function(function) + { } + + template + VTKM_EXEC + void operator()(const VecType &point, ScalarType &val) const + { + val = this->Function.Value(point); + } + +private: + vtkm::exec::ImplicitFunction Function; +}; + +template +void EvaluateOnCoordinates(vtkm::cont::CoordinateSystem points, + const vtkm::cont::ImplicitFunction &function, + vtkm::cont::ArrayHandle &values, + DeviceAdapter device) +{ + typedef vtkm::worklet::DispatcherMapField + EvalDispatcher; + + EvaluateImplicitFunction eval(function.PrepareForExecution(device)); + EvalDispatcher(eval).Invoke(points, values); +} + +template +bool TestArrayEqual(const vtkm::cont::ArrayHandle &result, + const std::array &expected) +{ + if (result.GetNumberOfValues() != N) + { + return false; + } + + vtkm::cont::ArrayHandle::PortalConstControl portal = + result.GetPortalConstControl(); + for (std::size_t i = 0; i < N; ++i) + { + if (!test_equal(portal.Get(static_cast(i)), expected[i])) + { + return false; + } + } + return true; +} + +} // anonymous namespace + +class TestingImplicitFunction +{ +public: + TestingImplicitFunction() + : Input(vtkm::cont::testing::MakeTestDataSet().Make3DExplicitDataSet2()) + { + } + + template + void Run(DeviceAdapter device) + { + this->TestSphere(device); + this->TestPlane(device); + this->TestBox(device); + } + +private: + template + void TestSphere(DeviceAdapter device) + { + std::cout << "Testing vtkm::cont::Sphere on " + << vtkm::cont::DeviceAdapterTraits::GetName() + << "\n"; + + vtkm::cont::Sphere sphere({0.0f, 0.0f, 0.0f}, 1.0f); + vtkm::cont::ArrayHandle values; + implicit_function_detail::EvaluateOnCoordinates( + this->Input.GetCoordinateSystem(0), sphere, values, device); + + std::array expected = + { {-1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 2.0f, 1.0f} }; + VTKM_TEST_ASSERT(implicit_function_detail::TestArrayEqual(values, expected), + "Result does not match expected values"); + } + + template + void TestPlane(DeviceAdapter device) + { + std::cout << "Testing vtkm::cont::Plane on " + << vtkm::cont::DeviceAdapterTraits::GetName() + << "\n"; + + vtkm::cont::Plane plane({0.5f, 0.5f, 0.5f}, {1.0f, 0.0f, 1.0f}); + vtkm::cont::ArrayHandle values; + + implicit_function_detail::EvaluateOnCoordinates( + this->Input.GetCoordinateSystem(0), plane, values, device); + std::array expected1 = + { {-1.0f, 0.0f, 1.0f, 0.0f, -1.0f, 0.0f, 1.0f, 0.0f} }; + VTKM_TEST_ASSERT(implicit_function_detail::TestArrayEqual(values, expected1), + "Result does not match expected values"); + + plane.SetNormal({-1.0f, 0.0f, -1.0f}); + implicit_function_detail::EvaluateOnCoordinates( + this->Input.GetCoordinateSystem(0), plane, values, device); + std::array expected2 = + { {1.0f, 0.0f, -1.0f, 0.0f, 1.0f, 0.0f, -1.0f, 0.0f} }; + VTKM_TEST_ASSERT(implicit_function_detail::TestArrayEqual(values, expected2), + "Result does not match expected values"); + } + + template + void TestBox(DeviceAdapter device) + { + std::cout << "Testing vtkm::cont::Box on " + << vtkm::cont::DeviceAdapterTraits::GetName() + << "\n"; + + vtkm::cont::Box box({0.0f, -0.5f, -0.5f}, {1.5f, 1.5f, 0.5f}); + vtkm::cont::ArrayHandle values; + implicit_function_detail::EvaluateOnCoordinates( + this->Input.GetCoordinateSystem(0), box, values, device); + + std::array expected = + { {0.0f, -0.5f, 0.5f, 0.5f, 0.0f, -0.5f, 0.5f, 0.5f} }; + VTKM_TEST_ASSERT(implicit_function_detail::TestArrayEqual(values, expected), + "Result does not match expected values"); + } + + vtkm::cont::DataSet Input; +}; + +} +} +} // vtmk::cont::testing + +#endif //vtk_m_cont_testing_TestingImplicitFunction_h diff --git a/vtkm/exec/CMakeLists.txt b/vtkm/exec/CMakeLists.txt index c5d9908c0..8324f88f7 100644 --- a/vtkm/exec/CMakeLists.txt +++ b/vtkm/exec/CMakeLists.txt @@ -30,6 +30,7 @@ set(headers ExecutionObjectBase.h ExecutionWholeArray.h FunctorBase.h + ImplicitFunction.h Jacobian.h ParametricCoordinates.h ) diff --git a/vtkm/exec/ImplicitFunction.h b/vtkm/exec/ImplicitFunction.h new file mode 100644 index 000000000..58693f6dd --- /dev/null +++ b/vtkm/exec/ImplicitFunction.h @@ -0,0 +1,143 @@ +//============================================================================ +// 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_ImplicitFunction_h +#define vtk_m_exec_ImplicitFunction_h + +#include + + +namespace vtkm { +namespace exec { + +class VTKM_ALWAYS_EXPORT ImplicitFunction +{ +public: + ImplicitFunction() + : Function(nullptr), ValueCaller(nullptr), GradientCaller(nullptr) + { } + + 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 &x) const + { + return this->ValueCaller(this->Function, x[0], x[1], x[2]); + } + + VTKM_EXEC + vtkm::Vec Gradient(FloatDefault x, FloatDefault y, + FloatDefault z) const + { + return this->GradientCaller(this->Function, x, y, z); + } + + VTKM_EXEC + vtkm::Vec Gradient(const vtkm::Vec &x) const + { + return this->GradientCaller(this->Function, x[0], x[1], x[2]); + } + + template + VTKM_EXEC + void Bind(const T *function) + { + this->Function = function; + this->ValueCaller = + [](const void *t, FloatDefault x, FloatDefault y, FloatDefault z) { + return static_cast(t)->Value(x, y, z); + }; + this->GradientCaller = + [](const void *t, FloatDefault x, FloatDefault y, FloatDefault z) { + return static_cast(t)->Gradient(x, y, z); + }; + } + +private: + using ValueCallerSig = + FloatDefault(const void*, FloatDefault, FloatDefault, FloatDefault); + using GradientCallerSig = + vtkm::Vec(const void*, FloatDefault, FloatDefault, FloatDefault); + + const void *Function; + ValueCallerSig *ValueCaller; + GradientCallerSig *GradientCaller; +}; + +/// \brief A function object that evaluates the contained implicit function +class VTKM_ALWAYS_EXPORT ImplicitFunctionValue +{ +public: + ImplicitFunctionValue() = default; + + explicit ImplicitFunctionValue(const vtkm::exec::ImplicitFunction &func) + : Function(func) + { } + + VTKM_EXEC + FloatDefault operator()(const vtkm::Vec x) const + { + return this->Function.Value(x); + } + + VTKM_EXEC + FloatDefault operator()(FloatDefault x, FloatDefault y, FloatDefault z) const + { + return this->Function.Value(x, y, z); + } + +private: + vtkm::exec::ImplicitFunction Function; +}; + +/// \brief A function object that computes the gradient of the contained implicit +/// function and the specified point. +class VTKM_ALWAYS_EXPORT ImplicitFunctionGradient +{ +public: + ImplicitFunctionGradient() = default; + + explicit ImplicitFunctionGradient(const vtkm::exec::ImplicitFunction &func) + : Function(func) + { } + + VTKM_EXEC + vtkm::Vec operator()(const vtkm::Vec x) const + { + return this->Function.Gradient(x); + } + + VTKM_EXEC + vtkm::Vec operator()(FloatDefault x, FloatDefault y, FloatDefault z) const + { + return this->Function.Gradient(x, y, z); + } + +private: + vtkm::exec::ImplicitFunction Function; +}; + +} +} // vtkm::exec + +#endif // vtk_m_exec_ImplicitFunction_h diff --git a/vtkm/testing/CMakeLists.txt b/vtkm/testing/CMakeLists.txt index d55764fe2..8a5734d03 100644 --- a/vtkm/testing/CMakeLists.txt +++ b/vtkm/testing/CMakeLists.txt @@ -20,7 +20,6 @@ set(headers Testing.h - TestingImplicitFunctions.h TestingMath.h OptionParser.h VecTraitsTests.h @@ -34,7 +33,6 @@ set(unit_tests UnitTestBounds.cxx UnitTestCellShape.cxx UnitTestExceptions.cxx - UnitTestImplicitFunctions.cxx UnitTestListTag.cxx UnitTestMath.cxx UnitTestMatrix.cxx diff --git a/vtkm/testing/TestingImplicitFunctions.h b/vtkm/testing/TestingImplicitFunctions.h deleted file mode 100644 index f33badaac..000000000 --- a/vtkm/testing/TestingImplicitFunctions.h +++ /dev/null @@ -1,213 +0,0 @@ -//============================================================================ -// 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_testing_TestingImplicitFunctions_h -#define vtk_m_testing_TestingImplicitFunctions_h - -#include -#include - -#include -#include -#include - -#include -#include - -#include - -namespace vtkm { -namespace testing { - -namespace { - -inline vtkm::FloatDefault GetRandomValue() -{ - static std::mt19937 gen; - std::uniform_real_distribution dist(-7.0, 7.0); - return dist(gen); -} - -} - -template -struct TestingImplicitFunctions -{ - typedef vtkm::Vec FVec3; - - template - class EvaluateImplicitFunction : public vtkm::worklet::WorkletMapField - { - public: - typedef void ControlSignature(FieldIn, FieldOut); - typedef void ExecutionSignature(_1, _2); - - EvaluateImplicitFunction(const ImplicitFunction &function) - : Function(function) - { } - - template - VTKM_EXEC - void operator()(const VecType &point, ScalarType &val) const - { - val = this->Function.Value(point); - } - - private: - ImplicitFunction Function; - }; - - - static void TestBoxValue() - { - typedef EvaluateImplicitFunction EvalWorklet; - typedef vtkm::worklet::DispatcherMapField - EvalDispatcher; - - FVec3 half(vtkm::Abs(GetRandomValue()), vtkm::Abs(GetRandomValue()), vtkm::Abs(GetRandomValue())); - FVec3 minPoint(GetRandomValue(), GetRandomValue(), GetRandomValue()); - FVec3 maxPoint(minPoint[0] + 2.f * half[0], minPoint[1] + 2.f * half[1], minPoint[2] + 2.f * half[2]); - - vtkm::Box box(minPoint, maxPoint); - - FVec3 data[8] = { minPoint, - maxPoint, - minPoint + FVec3(FloatDefault(0), half[1], half[2]), - minPoint + FVec3(half[0], FloatDefault(0), half[2]), - minPoint + FVec3(half[0], half[1], FloatDefault(0)), - minPoint + half, - minPoint - half, - maxPoint + half - }; - vtkm::cont::ArrayHandle points = vtkm::cont::make_ArrayHandle(data, 8); - - EvalWorklet eval(box); - vtkm::cont::ArrayHandle values; - EvalDispatcher(eval).Invoke(points, values); - - vtkm::cont::ArrayHandle::PortalConstControl portal = - values.GetPortalConstControl(); - - bool success = test_equal(portal.Get(0), FloatDefault(0.0) ) && - test_equal(portal.Get(1), FloatDefault(0.0) ) && - test_equal(portal.Get(2), FloatDefault(0.0) ) && - test_equal(portal.Get(3), FloatDefault(0.0) ) && - test_equal(portal.Get(4), FloatDefault(0.0) ) && - (portal.Get(5) < 0.0) && - (portal.Get(6) > 0.0) && - (portal.Get(7) > 0.0); - - VTKM_TEST_ASSERT(success, "Box: did not get expected results."); - } - - static void TestSphereValue() - { - typedef EvaluateImplicitFunction EvalWorklet; - typedef vtkm::worklet::DispatcherMapField - EvalDispatcher; - - FVec3 center(GetRandomValue(), GetRandomValue(), GetRandomValue()); - FloatDefault radius = vtkm::Abs(GetRandomValue()); - vtkm::Sphere sphere(center, radius); - - FloatDefault r = radius; - FVec3 cube[8] = { center + FVec3(0, 0, 0), - center + FVec3(r, 0, 0), - center + FVec3(0, r, 0), - center + FVec3(r, r, 0), - center + FVec3(0, 0, r), - center + FVec3(r, 0, r), - center + FVec3(0, r, r), - center + FVec3(r, r, r) - }; - vtkm::cont::ArrayHandle points = vtkm::cont::make_ArrayHandle(cube, 8); - - EvalWorklet eval(sphere); - vtkm::cont::ArrayHandle values; - EvalDispatcher(eval).Invoke(points, values); - - vtkm::cont::ArrayHandle::PortalConstControl portal = - values.GetPortalConstControl(); - - bool success = (portal.Get(0) == -(r*r)) && - test_equal(portal.Get(1), FloatDefault(0.0) ) && - test_equal(portal.Get(2), FloatDefault(0.0) ) && - test_equal(portal.Get(4), FloatDefault(0.0) ) && - (portal.Get(3) > 0.0) && - (portal.Get(5) > 0.0) && - (portal.Get(6) > 0.0) && - (portal.Get(7) > 0.0); - - VTKM_TEST_ASSERT(success, "Sphere: did not get expected results."); - } - - static void TestPlaneValue() - { - typedef EvaluateImplicitFunction EvalWorklet; - typedef vtkm::worklet::DispatcherMapField - EvalDispatcher; - - FVec3 origin(GetRandomValue(), GetRandomValue(), GetRandomValue()); - FVec3 normal(GetRandomValue(), GetRandomValue(), GetRandomValue()); - vtkm::Plane plane(origin, normal); - - FloatDefault t[] = { -2.0, -1.0, 0.0, 1.0, 2.0 }; - vtkm::cont::ArrayHandle points; - points.Allocate(5); - for (int i = 0; i < 5; ++i) - { - points.GetPortalControl().Set(i, origin + (t[i] * normal)); - } - - EvalWorklet eval(plane); - vtkm::cont::ArrayHandle values; - EvalDispatcher(eval).Invoke(points, values); - - vtkm::cont::ArrayHandle::PortalConstControl portal = - values.GetPortalConstControl(); - bool success = (portal.Get(0) < 0.0) && - (portal.Get(1) < 0.0) && - test_equal(portal.Get(2), FloatDefault(0.0) ) && - (portal.Get(3) > 0.0) && - (portal.Get(4) > 0.0); - - VTKM_TEST_ASSERT(success, "Plane: did not get expected results."); - } - - static void Run() - { - for (int i = 0; i < 50; ++i) - { - TestSphereValue(); - } - for (int i = 0; i < 50; ++i) - { - TestPlaneValue(); - } - for (int i = 0; i < 50; ++i) - { - TestBoxValue(); - } - } -}; - -} -} // namespace vtkm::testing - -#endif // vtk_m_testing_TestingImplicitFunctions_h diff --git a/vtkm/worklet/Clip.h b/vtkm/worklet/Clip.h index b813be1c0..ba80c0916 100644 --- a/vtkm/worklet/Clip.h +++ b/vtkm/worklet/Clip.h @@ -29,8 +29,8 @@ #include #include #include +#include #include -#include #include #include @@ -529,13 +529,13 @@ public: return output; } - template + template class ClipWithImplicitFunction { public: VTKM_CONT ClipWithImplicitFunction(Clip *clipper, const vtkm::cont::DynamicCellSet &cellSet, - ImplicitFunction function, vtkm::cont::CellSetExplicit<> *result) + const vtkm::exec::ImplicitFunction &function, vtkm::cont::CellSetExplicit<> *result) : Clipper(clipper), CellSet(&cellSet), Function(function), Result(result) { } @@ -546,40 +546,37 @@ public: vtkm::cont::ArrayHandleTransform< vtkm::FloatDefault, ArrayHandleType, - vtkm::ImplicitFunctionValue > clipScalars(handle, - this->Function); + vtkm::exec::ImplicitFunctionValue> clipScalars(handle, this->Function); - *this->Result = this->Clipper->Run(*this->CellSet, - clipScalars, - 0.0, - DeviceAdapter()); + // Clip at locations where the implicit function evaluates to 0 + *this->Result = this->Clipper->Run(*this->CellSet, + clipScalars, + 0.0, + DeviceAdapter()); } private: Clip *Clipper; const vtkm::cont::DynamicCellSet *CellSet; - const vtkm::ImplicitFunctionValue Function; + const vtkm::exec::ImplicitFunctionValue Function; vtkm::cont::CellSetExplicit<> *Result; }; template vtkm::cont::CellSetExplicit<> Run(const vtkm::cont::DynamicCellSetBase &cellSet, - const ImplicitFunction &clipFunction, + const vtkm::cont::ImplicitFunction &clipFunction, const vtkm::cont::CoordinateSystem &coords, DeviceAdapter device) { - (void) device; vtkm::cont::CellSetExplicit<> output; - ClipWithImplicitFunction clip(this, - cellSet, - clipFunction, - &output); - CastAndCall(coords, clip); + ClipWithImplicitFunction clip(this, + cellSet, + clipFunction.PrepareForExecution(device), + &output); + CastAndCall(coords, clip); return output; } diff --git a/vtkm/worklet/ExtractGeometry.h b/vtkm/worklet/ExtractGeometry.h index f89dd77ce..621040bec 100644 --- a/vtkm/worklet/ExtractGeometry.h +++ b/vtkm/worklet/ExtractGeometry.h @@ -28,7 +28,7 @@ #include #include #include -#include +#include namespace vtkm { namespace worklet { @@ -40,7 +40,6 @@ public: //////////////////////////////////////////////////////////////////////////////////// // Worklet to identify cells within volume of interest - template class ExtractCellsByVOI : public vtkm::worklet::WorkletMapPointToCell { public: @@ -53,7 +52,7 @@ public: ExtractCellsByVOI() : Function() {} VTKM_CONT - explicit ExtractCellsByVOI(const ImplicitFunction &function) + explicit ExtractCellsByVOI(const vtkm::exec::ImplicitFunction &function) : Function(function) {} template @@ -76,7 +75,7 @@ public: } private: - ImplicitFunction Function; + vtkm::exec::ImplicitFunction Function; }; //////////////////////////////////////////////////////////////////////////////////// @@ -99,22 +98,20 @@ public: //////////////////////////////////////////////////////////////////////////////////// // Extract cells by implicit function permutes input data template vtkm::cont::CellSetPermutation Run( const CellSetType &cellSet, const vtkm::cont::CoordinateSystem &coordinates, - const ImplicitFunction &implicitFunction, - DeviceAdapter) + const vtkm::cont::ImplicitFunction &implicitFunction, + DeviceAdapter device) { typedef vtkm::cont::CellSetPermutation OutputType; // Worklet output will be a boolean passFlag array - typedef ExtractCellsByVOI ExtractCellsWorklet; vtkm::cont::ArrayHandle passFlags; - ExtractCellsWorklet worklet(implicitFunction); - DispatcherMapTopology dispatcher(worklet); + ExtractCellsByVOI worklet(implicitFunction.PrepareForExecution(device)); + DispatcherMapTopology dispatcher(worklet); dispatcher.Invoke(cellSet, coordinates, passFlags); diff --git a/vtkm/worklet/ExtractPoints.h b/vtkm/worklet/ExtractPoints.h index e63953a48..665f3b949 100644 --- a/vtkm/worklet/ExtractPoints.h +++ b/vtkm/worklet/ExtractPoints.h @@ -26,7 +26,7 @@ #include #include #include -#include +#include namespace vtkm { namespace worklet { @@ -38,7 +38,6 @@ public: //////////////////////////////////////////////////////////////////////////////////// // Worklet to identify points within volume of interest - template class ExtractPointsByVOI : public vtkm::worklet::WorkletMapCellToPoint { public: @@ -48,7 +47,7 @@ public: typedef _3 ExecutionSignature(_2); VTKM_CONT - ExtractPointsByVOI(const ImplicitFunction &function) : + ExtractPointsByVOI(const vtkm::exec::ImplicitFunction &function) : Function(function) {} VTKM_EXEC @@ -62,7 +61,7 @@ public: } private: - ImplicitFunction Function; + vtkm::exec::ImplicitFunction Function; }; //////////////////////////////////////////////////////////////////////////////////// @@ -90,23 +89,20 @@ public: //////////////////////////////////////////////////////////////////////////////////// // Extract points by implicit function template vtkm::cont::CellSetSingleType<> Run( const CellSetType &cellSet, const vtkm::cont::CoordinateSystem &coordinates, - const ImplicitFunction &implicitFunction, - DeviceAdapter) + const vtkm::cont::ImplicitFunction &implicitFunction, + DeviceAdapter device) { typedef typename vtkm::cont::DeviceAdapterAlgorithm DeviceAlgorithm; + // Worklet output will be a boolean passFlag array vtkm::cont::ArrayHandle passFlags; - // Worklet output will be a boolean passFlag array - typedef ExtractPointsByVOI ExtractPointsWorklet; - - ExtractPointsWorklet worklet(implicitFunction); - DispatcherMapTopology dispatcher(worklet); + ExtractPointsByVOI worklet(implicitFunction.PrepareForExecution(device)); + DispatcherMapTopology dispatcher(worklet); dispatcher.Invoke(cellSet, coordinates, passFlags); vtkm::cont::ArrayHandleCounting indices = diff --git a/vtkm/worklet/testing/UnitTestClipping.cxx b/vtkm/worklet/testing/UnitTestClipping.cxx index 78e1aacb7..d3caefe36 100644 --- a/vtkm/worklet/testing/UnitTestClipping.cxx +++ b/vtkm/worklet/testing/UnitTestClipping.cxx @@ -30,8 +30,8 @@ #include #include #include +#include #include -#include #include @@ -230,7 +230,7 @@ void TestClippingWithImplicitFunction() { vtkm::Vec center(1, 1, 0); vtkm::FloatDefault radius(0.5); - vtkm::Sphere sphere(center, radius); + vtkm::cont::Sphere sphere(center, radius); vtkm::cont::DataSet ds = MakeTestDatasetStructured(); diff --git a/vtkm/worklet/testing/UnitTestExtractGeometry.cxx b/vtkm/worklet/testing/UnitTestExtractGeometry.cxx index c4009eed9..d08c25ffc 100644 --- a/vtkm/worklet/testing/UnitTestExtractGeometry.cxx +++ b/vtkm/worklet/testing/UnitTestExtractGeometry.cxx @@ -95,7 +95,7 @@ public: // Implicit function vtkm::Vec minPoint(0.5f, 0.0f, 0.0f); vtkm::Vec maxPoint(2.0f, 2.0f, 2.0f); - vtkm::Box box(minPoint, maxPoint); + vtkm::cont::Box box(minPoint, maxPoint); // Output data set with cell set containing extracted cells and all points vtkm::worklet::ExtractGeometry extractGeometry; @@ -215,7 +215,7 @@ public: // Implicit function vtkm::Vec minPoint(1.0f, 1.0f, 1.0f); vtkm::Vec maxPoint(3.0f, 3.0f, 3.0f); - vtkm::Box box(minPoint, maxPoint); + vtkm::cont::Box box(minPoint, maxPoint); // Output data set with cell set containing extracted points vtkm::worklet::ExtractGeometry extractGeometry; @@ -255,7 +255,7 @@ public: // Implicit function vtkm::Vec center(2.f, 2.f, 2.f); vtkm::FloatDefault radius(1.8f); - vtkm::Sphere sphere(center, radius); + vtkm::cont::Sphere sphere(center, radius); // Output data set with cell set containing extracted cells vtkm::worklet::ExtractGeometry extractGeometry; diff --git a/vtkm/worklet/testing/UnitTestExtractPoints.cxx b/vtkm/worklet/testing/UnitTestExtractPoints.cxx index 01cb25189..b5e8cd68f 100644 --- a/vtkm/worklet/testing/UnitTestExtractPoints.cxx +++ b/vtkm/worklet/testing/UnitTestExtractPoints.cxx @@ -78,7 +78,7 @@ public: // Implicit function vtkm::Vec minPoint(1.f, 1.f, 1.f); vtkm::Vec maxPoint(3.f, 3.f, 3.f); - vtkm::Box box(minPoint, maxPoint); + vtkm::cont::Box box(minPoint, maxPoint); // Output dataset contains input coordinate system and point data vtkm::cont::DataSet outDataSet; @@ -108,7 +108,7 @@ public: // Implicit function vtkm::Vec center(2.f, 2.f, 2.f); vtkm::FloatDefault radius(1.8f); - vtkm::Sphere sphere(center, radius); + vtkm::cont::Sphere sphere(center, radius); // Output dataset contains input coordinate system and point data vtkm::cont::DataSet outDataSet; @@ -138,7 +138,7 @@ public: // Implicit function vtkm::Vec minPoint(0.f, 0.f, 0.f); vtkm::Vec maxPoint(1.f, 1.f, 1.f); - vtkm::Box box(minPoint, maxPoint); + vtkm::cont::Box box(minPoint, maxPoint); // Output dataset contains input coordinate system and point data vtkm::cont::DataSet outDataSet; From 4049b5b25eb7834f8e46e8bc123cb7c62e7f35e7 Mon Sep 17 00:00:00 2001 From: Sujin Philip Date: Tue, 18 Apr 2017 11:01:17 -0400 Subject: [PATCH 3/3] Add ClipWithImplicitFunction Filter Rename current Clip filter to ClipWithField. --- vtkm/filter/CMakeLists.txt | 6 +- vtkm/filter/{Clip.h => ClipWithField.h} | 14 +-- vtkm/filter/{Clip.hxx => ClipWithField.hxx} | 26 ++-- vtkm/filter/ClipWithImplicitFunction.h | 76 ++++++++++++ vtkm/filter/ClipWithImplicitFunction.hxx | 104 ++++++++++++++++ vtkm/filter/testing/CMakeLists.txt | 3 +- ...er.cxx => UnitTestClipWithFieldFilter.cxx} | 6 +- ...UnitTestClipWithImplicitFunctionFilter.cxx | 112 ++++++++++++++++++ vtkm/worklet/Clip.h | 18 +-- 9 files changed, 324 insertions(+), 41 deletions(-) rename vtkm/filter/{Clip.h => ClipWithField.h} (89%) rename vtkm/filter/{Clip.hxx => ClipWithField.hxx} (80%) create mode 100644 vtkm/filter/ClipWithImplicitFunction.h create mode 100644 vtkm/filter/ClipWithImplicitFunction.hxx rename vtkm/filter/testing/{UnitTestClipFilter.cxx => UnitTestClipWithFieldFilter.cxx} (96%) create mode 100644 vtkm/filter/testing/UnitTestClipWithImplicitFunctionFilter.cxx diff --git a/vtkm/filter/CMakeLists.txt b/vtkm/filter/CMakeLists.txt index c719e9041..53b605774 100644 --- a/vtkm/filter/CMakeLists.txt +++ b/vtkm/filter/CMakeLists.txt @@ -21,7 +21,8 @@ set(headers CellAverage.h CleanGrid.h - Clip.h + ClipWithField.h + ClipWithImplicitFunction.h ContourTreeUniform.h ExternalFaces.h FieldMetadata.h @@ -53,7 +54,8 @@ set(headers set(header_template_sources CellAverage.hxx CleanGrid.hxx - Clip.hxx + ClipWithField.h + ClipWithImplicitFunction.hxx ContourTreeUniform.hxx ExternalFaces.hxx FilterCell.hxx diff --git a/vtkm/filter/Clip.h b/vtkm/filter/ClipWithField.h similarity index 89% rename from vtkm/filter/Clip.h rename to vtkm/filter/ClipWithField.h index 5f52431cd..c520f10f4 100644 --- a/vtkm/filter/Clip.h +++ b/vtkm/filter/ClipWithField.h @@ -18,8 +18,8 @@ // this software. //============================================================================ -#ifndef vtk_m_filter_Clip_h -#define vtk_m_filter_Clip_h +#ifndef vtk_m_filter_ClipWithField_h +#define vtk_m_filter_ClipWithField_h #include #include @@ -27,11 +27,11 @@ namespace vtkm { namespace filter { -class Clip : public vtkm::filter::FilterDataSetWithField +class ClipWithField : public vtkm::filter::FilterDataSetWithField { public: VTKM_CONT - Clip(); + ClipWithField(); VTKM_CONT void SetClipValue(vtkm::Float64 value){ this->ClipValue = value; } @@ -64,7 +64,7 @@ private: }; template<> -class FilterTraits +class FilterTraits { //currently the Clip filter only works on scalar data. public: typedef TypeListTagScalarAll InputFieldTypeList; @@ -75,6 +75,6 @@ public: } // namespace vtkm::filter -#include +#include -#endif // vtk_m_filter_Clip_h +#endif // vtk_m_filter_ClipWithField_h diff --git a/vtkm/filter/Clip.hxx b/vtkm/filter/ClipWithField.hxx similarity index 80% rename from vtkm/filter/Clip.hxx rename to vtkm/filter/ClipWithField.hxx index 656880780..aa046273c 100644 --- a/vtkm/filter/Clip.hxx +++ b/vtkm/filter/ClipWithField.hxx @@ -31,8 +31,8 @@ namespace filter { //----------------------------------------------------------------------------- inline VTKM_CONT -Clip::Clip(): - vtkm::filter::FilterDataSetWithField(), +ClipWithField::ClipWithField(): + vtkm::filter::FilterDataSetWithField(), ClipValue(0), Worklet() { @@ -45,11 +45,12 @@ template inline VTKM_CONT -vtkm::filter::ResultDataSet Clip::DoExecute(const vtkm::cont::DataSet& input, - const vtkm::cont::ArrayHandle& field, - const vtkm::filter::FieldMetadata& fieldMeta, - const vtkm::filter::PolicyBase& policy, - const DeviceAdapter& device) +vtkm::filter::ResultDataSet ClipWithField::DoExecute( + const vtkm::cont::DataSet& input, + const vtkm::cont::ArrayHandle& field, + const vtkm::filter::FieldMetadata& fieldMeta, + const vtkm::filter::PolicyBase& policy, + const DeviceAdapter& device) { if(fieldMeta.IsPointField() == false) { @@ -95,11 +96,12 @@ template inline VTKM_CONT -bool Clip::DoMapField(vtkm::filter::ResultDataSet& result, - const vtkm::cont::ArrayHandle& input, - const vtkm::filter::FieldMetadata& fieldMeta, - const vtkm::filter::PolicyBase&, - const DeviceAdapter& device) +bool ClipWithField::DoMapField( + vtkm::filter::ResultDataSet& result, + const vtkm::cont::ArrayHandle& input, + const vtkm::filter::FieldMetadata& fieldMeta, + const vtkm::filter::PolicyBase&, + const DeviceAdapter& device) { if(fieldMeta.IsPointField() == false) { diff --git a/vtkm/filter/ClipWithImplicitFunction.h b/vtkm/filter/ClipWithImplicitFunction.h new file mode 100644 index 000000000..d25442d2b --- /dev/null +++ b/vtkm/filter/ClipWithImplicitFunction.h @@ -0,0 +1,76 @@ +//============================================================================ +// 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_filter_ClipWithImplicitFunction_h +#define vtk_m_filter_ClipWithImplicitFunction_h + +#include +#include +#include + +#include + + +namespace vtkm { +namespace filter { + +class ClipWithImplicitFunction : public vtkm::filter::FilterDataSet +{ +public: + template + void SetImplicitFunction(const std::shared_ptr &func, + const vtkm::filter::PolicyBase& policy); + + template + void SetImplicitFunction(const std::shared_ptr &func) + { + this->Function = func; + } + + std::shared_ptr GetImplicitFunction() const + { + return this->Function; + } + + template + vtkm::filter::ResultDataSet DoExecute(const vtkm::cont::DataSet& input, + const vtkm::filter::PolicyBase& policy, + const DeviceAdapter& tag); + + //Map a new field onto the resulting dataset after running the filter. + //This call is only valid after Execute has been called. + template + bool DoMapField(vtkm::filter::ResultDataSet& result, + const vtkm::cont::ArrayHandle& input, + const vtkm::filter::FieldMetadata& fieldMeta, + const vtkm::filter::PolicyBase& policy, + const DeviceAdapter& tag); + +private: + std::shared_ptr Function; + vtkm::worklet::Clip Worklet; +}; + +} +} // namespace vtkm::filter + +#include + +#endif // vtk_m_filter_ClipWithImplicitFunction_h diff --git a/vtkm/filter/ClipWithImplicitFunction.hxx b/vtkm/filter/ClipWithImplicitFunction.hxx new file mode 100644 index 000000000..1226b602e --- /dev/null +++ b/vtkm/filter/ClipWithImplicitFunction.hxx @@ -0,0 +1,104 @@ +//============================================================================ +// 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. +//============================================================================ + +#include +#include +#include + +#include + +namespace vtkm { +namespace filter { + +//----------------------------------------------------------------------------- +template +inline +void ClipWithImplicitFunction::SetImplicitFunction( + const std::shared_ptr &func, + const vtkm::filter::PolicyBase&) +{ + func->ResetDevice(DerivedPolicy::DeviceAdapterList); + this->Function = func; +} + +//----------------------------------------------------------------------------- +template +inline vtkm::filter::ResultDataSet ClipWithImplicitFunction::DoExecute( + const vtkm::cont::DataSet& input, + const vtkm::filter::PolicyBase& policy, + const DeviceAdapter& device) +{ + //get the cells and coordinates of the dataset + const vtkm::cont::DynamicCellSet& cells = + input.GetCellSet(this->GetActiveCellSetIndex()); + + const vtkm::cont::CoordinateSystem& inputCoords = + input.GetCoordinateSystem(this->GetActiveCoordinateSystemIndex()); + + + vtkm::cont::CellSetExplicit<> outputCellSet = + this->Worklet.Run( vtkm::filter::ApplyPolicy(cells, policy), + *this->Function, + inputCoords, + device + ); + + // compute output coordinates + vtkm::cont::CoordinateSystem outputCoords; + outputCoords.SetData(this->Worklet.ProcessField(inputCoords, device)); + + //create the output data + vtkm::cont::DataSet output; + output.AddCellSet( outputCellSet ); + output.AddCoordinateSystem( outputCoords ); + + vtkm::filter::ResultDataSet result(output); + return result; +} + +//----------------------------------------------------------------------------- +template +inline bool ClipWithImplicitFunction::DoMapField( + vtkm::filter::ResultDataSet& result, + const vtkm::cont::ArrayHandle& input, + const vtkm::filter::FieldMetadata& fieldMeta, + const vtkm::filter::PolicyBase&, + const DeviceAdapter& device) +{ + if(fieldMeta.IsPointField() == false) + { + //not a point field, we can't map it + return false; + } + + vtkm::cont::DynamicArrayHandle output = + this->Worklet.ProcessField( input, device); + + //use the same meta data as the input so we get the same field name, etc. + result.GetDataSet().AddField( fieldMeta.AsField(output) ); + return true; +} + +} +} diff --git a/vtkm/filter/testing/CMakeLists.txt b/vtkm/filter/testing/CMakeLists.txt index bd154a1b3..bf9371878 100644 --- a/vtkm/filter/testing/CMakeLists.txt +++ b/vtkm/filter/testing/CMakeLists.txt @@ -21,7 +21,8 @@ set(unit_tests UnitTestCellAverageFilter.cxx UnitTestCleanGrid.cxx - UnitTestClipFilter.cxx + UnitTestClipWithFieldFilter.cxx + UnitTestClipWithImplicitFunctionFilter.cxx UnitTestContourTreeUniformFilter.cxx UnitTestExternalFacesFilter.cxx UnitTestFieldMetadata.cxx diff --git a/vtkm/filter/testing/UnitTestClipFilter.cxx b/vtkm/filter/testing/UnitTestClipWithFieldFilter.cxx similarity index 96% rename from vtkm/filter/testing/UnitTestClipFilter.cxx rename to vtkm/filter/testing/UnitTestClipWithFieldFilter.cxx index e94fc7407..06cb024f6 100644 --- a/vtkm/filter/testing/UnitTestClipFilter.cxx +++ b/vtkm/filter/testing/UnitTestClipWithFieldFilter.cxx @@ -18,7 +18,7 @@ // this software. //============================================================================ -#include +#include #include #include @@ -67,7 +67,7 @@ void TestClipExplicit() vtkm::cont::DataSet ds = MakeTestDatasetExplicit(); vtkm::filter::ResultDataSet result; - vtkm::filter::Clip clip; + vtkm::filter::ClipWithField clip; clip.SetClipValue(0.5); result = clip.Execute( ds, std::string("scalars")); @@ -110,7 +110,7 @@ void TestClip() } -int UnitTestClipFilter(int, char *[]) +int UnitTestClipWithFieldFilter(int, char *[]) { return vtkm::cont::testing::Testing::Run(TestClip); } diff --git a/vtkm/filter/testing/UnitTestClipWithImplicitFunctionFilter.cxx b/vtkm/filter/testing/UnitTestClipWithImplicitFunctionFilter.cxx new file mode 100644 index 000000000..4a4bec559 --- /dev/null +++ b/vtkm/filter/testing/UnitTestClipWithImplicitFunctionFilter.cxx @@ -0,0 +1,112 @@ +//============================================================================ +// 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. +//============================================================================ + +#include + +#include +#include +#include + +namespace { + +typedef vtkm::Vec Coord3D; + + +vtkm::cont::DataSet MakeTestDatasetStructured() +{ + static const vtkm::Id xdim = 3, ydim = 3; + static const vtkm::Id2 dim(xdim, ydim); + static const vtkm::Id numVerts = xdim * ydim; + + vtkm::Float32 scalars[numVerts]; + for (vtkm::Id i = 0; i < numVerts; ++i) + { + scalars[i] = 1.0f; + } + scalars[4] = 0.0f; + + vtkm::cont::DataSet ds; + vtkm::cont::DataSetBuilderUniform builder; + ds = builder.Create(dim); + + vtkm::cont::DataSetFieldAdd fieldAdder; + fieldAdder.AddPointField(ds, "scalars", scalars, numVerts); + + return ds; +} + +void TestClipStructured() +{ + std::cout << "Testing ClipWithImplicitFunction Filter on Structured data" + << std::endl; + + vtkm::cont::DataSet ds = MakeTestDatasetStructured(); + + vtkm::Vec center(1, 1, 0); + vtkm::FloatDefault radius(0.5); + auto sphere = std::make_shared(center, radius); + + vtkm::filter::ResultDataSet result; + vtkm::filter::ClipWithImplicitFunction clip; + clip.SetImplicitFunction(sphere); + + result = clip.Execute(ds); + clip.MapFieldOntoOutput(result, ds.GetField("scalars")); + + + const vtkm::cont::DataSet& outputData = result.GetDataSet(); + VTKM_TEST_ASSERT(outputData.GetNumberOfCellSets() == 1, + "Wrong number of cellsets in the output dataset"); + VTKM_TEST_ASSERT(outputData.GetNumberOfCoordinateSystems() == 1, + "Wrong number of coordinate systems in the output dataset"); + VTKM_TEST_ASSERT(outputData.GetNumberOfFields() == 1, + "Wrong number of fields in the output dataset"); + VTKM_TEST_ASSERT(outputData.GetCellSet().GetNumberOfCells() == 12, + "Wrong number of cells in the output dataset"); + + vtkm::cont::DynamicArrayHandle temp = outputData.GetField("scalars").GetData(); + vtkm::cont::ArrayHandle resultArrayHandle; + temp.CopyTo(resultArrayHandle); + + VTKM_TEST_ASSERT(resultArrayHandle.GetNumberOfValues() == 13, + "Wrong number of points in the output dataset"); + + vtkm::Float32 expected[13] = { 1, 1, 1, 1, 0, 1, 1, 1, 1, 0.25, 0.25, 0.25, + 0.25 }; + for (int i = 0; i < 13; ++i) + { + VTKM_TEST_ASSERT( + test_equal(resultArrayHandle.GetPortalConstControl().Get(i), expected[i]), + "Wrong result for ClipWithImplicitFunction fliter on sturctured quads data"); + } +} + +void TestClip() +{ + //todo: add more clip tests + TestClipStructured(); +} + +} // anonymous namespace + +int UnitTestClipWithImplicitFunctionFilter(int, char *[]) +{ + return vtkm::cont::testing::Testing::Run(TestClip); +} diff --git a/vtkm/worklet/Clip.h b/vtkm/worklet/Clip.h index ba80c0916..d2cc3020a 100644 --- a/vtkm/worklet/Clip.h +++ b/vtkm/worklet/Clip.h @@ -177,22 +177,6 @@ struct EdgeInterpolation }; -VTKM_EXEC_CONT -std::ostream& operator<<(std::ostream &out, const ClipStats &val) -{ - return out << std::endl << "{ Cells: " << val.NumberOfCells - << ", Indices: " << val.NumberOfIndices - << ", NewPoints: " << val.NumberOfNewPoints << " }"; -} - -VTKM_EXEC_CONT -std::ostream& operator<<(std::ostream &out, const EdgeInterpolation &val) -{ - return out << std::endl << "{ Vertex1: " << val.Vertex1 - << ", Vertex2: " << val.Vertex2 - << ", Weight: " << val.Weight << " }"; -} - class Clip { public: @@ -543,6 +527,8 @@ public: VTKM_CONT void operator()(const ArrayHandleType &handle) const { + // Evaluate the implicit function on the input coordinates using + // ArrayHandleTransform vtkm::cont::ArrayHandleTransform< vtkm::FloatDefault, ArrayHandleType,