Add Virtual Methods Framework

vtkm/cont/CMakeLists.txt

@@ -72,6 +72,7 @@ set(headers
   StorageListTag.h
   Timer.h
   TryExecute.h
+  VirtualObjectCache.h
   )
 
 set(header_impls

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 <vtkm/cont/DeviceAdapterListTag.h>
+#include <vtkm/cont/ErrorBadType.h>
+#include <vtkm/cont/ErrorBadValue.h>
+#include <vtkm/cont/internal/DeviceAdapterTag.h>
+#include <vtkm/cont/internal/VirtualObjectTransfer.h>
+
+#include <array>
+#include <type_traits>
+
+
+#define VTKM_MAX_DEVICE_ADAPTER_ID 8
+
+namespace vtkm {
+namespace cont {
+
+/// \brief Implements VTK-m's execution side <em> Virtual Methods </em>
+/// 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 <tt>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 <tt>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:
+///    <tt>
+///       template<typename TargetClass>
+///       VTKM_EXEC void Bind(const TargetClass *deviceTarget);
+///    </tt>
+///    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 <tt>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<typename VirtualObject>
+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<typename TargetClass,
+           typename DeviceAdapterList = VTKM_DEFAULT_DEVICE_ADAPTER_LIST_TAG>
+  void Bind(const TargetClass *target,
+            DeviceAdapterList devices = DeviceAdapterList())
+  {
+    this->Reset();
+
+    this->Target = target;
+    vtkm::ListForEach(
+      CreateTransferInterface<TargetClass>(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<typename DeviceAdapter>
+  VirtualObject GetVirtualObject(DeviceAdapter)
+  {
+    using DeviceInfo = vtkm::cont::DeviceAdapterTraits<DeviceAdapter>;
+
+    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<std::size_t>(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<std::size_t>(this->CurrentDevice)];
+  }
+
+  template<typename TargetClass>
+  class CreateTransferInterface
+  {
+  private:
+    template<typename DeviceAdapter>
+    using EnableIfValid =
+      std::enable_if<vtkm::cont::DeviceAdapterTraits<DeviceAdapter>::Valid>;
+
+    template<typename DeviceAdapter>
+    using EnableIfInvalid =
+      std::enable_if<!vtkm::cont::DeviceAdapterTraits<DeviceAdapter>::Valid>;
+
+  public:
+    CreateTransferInterface(TransferInterface *transfers)
+      : Transfers(transfers)
+    { }
+
+    // Use SFINAE to create entries for valid device adapters only
+    template<typename DeviceAdapter>
+    typename EnableIfValid<DeviceAdapter>::type operator()(DeviceAdapter) const
+    {
+      using DeviceInfo = vtkm::cont::DeviceAdapterTraits<DeviceAdapter>;
+
+      if (DeviceInfo::GetId() >= 0 &&
+          DeviceInfo::GetId() < VTKM_MAX_DEVICE_ADAPTER_ID)
+      {
+        using TransferImpl =
+          internal::VirtualObjectTransfer<VirtualObject, TargetClass, DeviceAdapter>;
+
+        std::size_t id = static_cast<std::size_t>(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 DeviceAdapter>
+    typename EnableIfInvalid<DeviceAdapter>::type operator()(DeviceAdapter) const
+    { }
+
+  private:
+    TransferInterface *Transfers;
+  };
+
+  const void *Target;
+  std::array<TransferInterface, VTKM_MAX_DEVICE_ADAPTER_ID> 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

vtkm/cont/cuda/DeviceAdapterCuda.h

@@ -30,6 +30,7 @@
 
 #include <vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.h>
 #include <vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h>
+#include <vtkm/cont/cuda/internal/VirtualObjectTransferCuda.h>
 #endif
 
 #endif //vtk_m_cont_cuda_DeviceAdapterCuda_h

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})

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 <vtkm/cont/cuda/ErrorCuda.h>
+#include <vtkm/cont/cuda/internal/DeviceAdapterTagCuda.h>
+#include <vtkm/cont/internal/VirtualObjectTransfer.h>
+
+
+namespace vtkm {
+namespace cont {
+namespace internal {
+
+namespace detail {
+
+template<typename VirtualObject, typename TargetClass>
+__global__
+void CreateKernel(VirtualObject *object, const TargetClass *target)
+{
+  object->Bind(target);
+}
+
+} // detail
+
+template<typename VirtualObject, typename TargetClass>
+struct VirtualObjectTransfer<VirtualObject, TargetClass,
+                             vtkm::cont::DeviceAdapterTagCuda>
+{
+  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

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})

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 <vtkm/cont/testing/TestingVirtualObjectCache.h>
+
+namespace {
+
+void TestVirtualObjectCache()
+{
+  using DeviceAdapterList =
+    vtkm::ListTagBase<vtkm::cont::DeviceAdapterTagSerial,
+                      vtkm::cont::DeviceAdapterTagCuda>;
+
+  vtkm::cont::testing::TestingVirtualObjectCache<DeviceAdapterList>::Run();
+}
+
+} // anonymous namespace
+
+
+int UnitTestCudaVirtualObjectCache(int, char *[])
+{
+  return vtkm::cont::testing::Testing::Run(TestVirtualObjectCache);
+}

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})

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 <typename VirtualObject, typename TargetClass, typename DeviceAdapter>
+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

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<typename VirtualObject, typename TargetClass>
+struct VirtualObjectTransferShareWithControl
+{
+  static void* Create(VirtualObject &object, const void *target)
+  {
+    static int dummyState = 0;
+    object.Bind(static_cast<const TargetClass*>(target));
+    return &dummyState;
+  }
+
+  static void Update(void*, const void*)
+  {
+  }
+
+  static void Cleanup(void*)
+  {
+  }
+};
+
+}
+}
+} // vtkm::cont::internal
+
+#endif // vtk_m_cont_internal_VirtualObjectTransferShareWithControl_h

vtkm/cont/serial/DeviceAdapterSerial.h

@@ -23,5 +23,6 @@
 #include <vtkm/cont/serial/internal/DeviceAdapterTagSerial.h>
 #include <vtkm/cont/serial/internal/ArrayManagerExecutionSerial.h>
 #include <vtkm/cont/serial/internal/DeviceAdapterAlgorithmSerial.h>
+#include <vtkm/cont/serial/internal/VirtualObjectTransferSerial.h>
 
 #endif //vtk_m_cont_serial_DeviceAdapterSerial_h

vtkm/cont/serial/internal/CMakeLists.txt

@@ -22,6 +22,7 @@ set(headers
   ArrayManagerExecutionSerial.h
   DeviceAdapterAlgorithmSerial.h
   DeviceAdapterTagSerial.h
+  VirtualObjectTransferSerial.h
   )
 
 vtkm_declare_headers(${headers})

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 <vtkm/cont/serial/internal/DeviceAdapterTagSerial.h>
+#include <vtkm/cont/internal/VirtualObjectTransfer.h>
+#include <vtkm/cont/internal/VirtualObjectTransferShareWithControl.h>
+
+
+namespace vtkm {
+namespace cont {
+namespace internal {
+
+template<typename VirtualObject, typename TargetClass>
+struct VirtualObjectTransfer<
+  VirtualObject, TargetClass, vtkm::cont::DeviceAdapterTagSerial> :
+  public VirtualObjectTransferShareWithControl<VirtualObject, TargetClass>
+{
+};
+
+}
+}
+} // vtkm::cont::internal
+
+#endif // vtk_m_cont_serial_internal_VirtualObjectTransferSerial_h

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})

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 <vtkm/cont/testing/TestingVirtualObjectCache.h>
+
+namespace {
+
+void TestVirtualObjectCache()
+{
+  using DeviceAdapterList = vtkm::ListTagBase<vtkm::cont::DeviceAdapterTagSerial>;
+
+  vtkm::cont::testing::TestingVirtualObjectCache<DeviceAdapterList>::Run();
+}
+
+} // anonymous namespace
+
+
+int UnitTestSerialVirtualObjectCache(int, char *[])
+{
+  return vtkm::cont::testing::Testing::Run(TestVirtualObjectCache);
+}

vtkm/cont/tbb/DeviceAdapterTBB.h

@@ -25,6 +25,7 @@
 #ifdef VTKM_ENABLE_TBB
 #include <vtkm/cont/tbb/internal/ArrayManagerExecutionTBB.h>
 #include <vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h>
+#include <vtkm/cont/tbb/internal/VirtualObjectTransferTBB.h>
 #endif
 
 #endif //vtk_m_cont_tbb_DeviceAdapterTBB_h

vtkm/cont/tbb/internal/CMakeLists.txt

@@ -23,6 +23,7 @@ set(headers
   DeviceAdapterAlgorithmTBB.h
   DeviceAdapterTagTBB.h
   FunctorsTBB.h
+  VirtualObjectTransferTBB.h
   )
 
 if (VTKm_ENABLE_TBB)

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 <vtkm/cont/tbb/internal/DeviceAdapterTagTBB.h>
+#include <vtkm/cont/internal/VirtualObjectTransfer.h>
+#include <vtkm/cont/internal/VirtualObjectTransferShareWithControl.h>
+
+
+namespace vtkm {
+namespace cont {
+namespace internal {
+
+template<typename VirtualObject, typename TargetClass>
+struct VirtualObjectTransfer<
+  VirtualObject, TargetClass, vtkm::cont::DeviceAdapterTagTBB> :
+  public VirtualObjectTransferShareWithControl<VirtualObject, TargetClass>
+{
+};
+
+}
+}
+} // vtkm::cont::internal
+
+#endif // vtk_m_cont_tbb_internal_VirtualObjectTransferTBB_h

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})

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 <vtkm/cont/testing/TestingVirtualObjectCache.h>
+
+namespace {
+
+void TestVirtualObjectCache()
+{
+  using DeviceAdapterList =
+    vtkm::ListTagBase<vtkm::cont::DeviceAdapterTagSerial,
+                      vtkm::cont::DeviceAdapterTagTBB>;
+
+  vtkm::cont::testing::TestingVirtualObjectCache<DeviceAdapterList>::Run();
+}
+
+} // anonymous namespace
+
+
+int UnitTestTBBVirtualObjectCache(int, char *[])
+{
+  return vtkm::cont::testing::Testing::Run(TestVirtualObjectCache);
+}

vtkm/cont/testing/CMakeLists.txt

@@ -28,6 +28,7 @@ set(headers
   TestingDataSetExplicit.h
   TestingDataSetSingleType.h
   TestingFancyArrayHandles.h
+  TestingVirtualObjectCache.h
   )
 
 vtkm_declare_headers(${headers})

vtkm/cont/testing/TestingDeviceAdapter.h

@@ -38,6 +38,7 @@
 #include <vtkm/cont/DeviceAdapterAlgorithm.h>
 
 #include <vtkm/cont/internal/DeviceAdapterError.h>
+#include <vtkm/cont/internal/VirtualObjectTransfer.h>
 
 #include <vtkm/cont/testing/Testing.h>
 
@@ -311,6 +312,59 @@ public:
     vtkm::exec::AtomicArray<T,DeviceAdapterTag> AArray;
   };
 
+  class VirtualObjectTransferKernel
+  {
+  public:
+    struct Interface
+    {
+      using FooSig = vtkm::Id (const void*);
+
+      template<typename T>
+      VTKM_EXEC
+      void Bind(const T *target)
+      {
+        this->Target = target;
+        this->FooPtr = [](const void* t){ return static_cast<const T*>(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();

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 <vtkm/cont/ArrayHandle.h>
+#include <vtkm/cont/ArrayHandleTransform.h>
+#include <vtkm/cont/DeviceAdapterAlgorithm.h>
+#include <vtkm/cont/testing/Testing.h>
+#include <vtkm/cont/VirtualObjectCache.h>
+#include <vtkm/Types.h>
+
+
+#define ARRAY_LEN 8
+
+namespace vtkm {
+namespace cont {
+namespace testing {
+
+namespace virtual_object_detail
+{
+
+class Transformer
+{
+public:
+  template<typename T>
+  VTKM_EXEC void Bind(const T *target)
+  {
+    this->Concrete = target;
+    this->Caller =
+      [](const void *concrete, vtkm::FloatDefault val) {
+        return static_cast<const T*>(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<typename DeviceAdapterList>
+class TestingVirtualObjectCache
+{
+private:
+  using FloatArrayHandle = vtkm::cont::ArrayHandle<vtkm::FloatDefault>;
+  using ArrayTransform =
+      vtkm::cont::ArrayHandleTransform<vtkm::FloatDefault,
+                                       FloatArrayHandle,
+                                       virtual_object_detail::Transformer>;
+  using TransformerCache =
+    vtkm::cont::VirtualObjectCache<virtual_object_detail::Transformer>;
+
+
+  class TestStage1
+  {
+  public:
+    TestStage1(const FloatArrayHandle &input, TransformerCache &manager)
+      : Input(&input), Manager(&manager)
+    { }
+
+    template<typename DeviceAdapter>
+    void operator()(DeviceAdapter device) const
+    {
+      using Algorithm = vtkm::cont::DeviceAdapterAlgorithm<DeviceAdapter>;
+      std::cout << "\tDeviceAdapter: "
+                << vtkm::cont::DeviceAdapterTraits<DeviceAdapter>::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<typename DeviceAdapter>
+    void operator()(DeviceAdapter device) const
+    {
+      using Algorithm = vtkm::cont::DeviceAdapterAlgorithm<DeviceAdapter>;
+      std::cout << "\tDeviceAdapter: "
+                << vtkm::cont::DeviceAdapterTraits<DeviceAdapter>::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<vtkm::FloatDefault> 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

vtkm/cont/testing/UnitTestDeviceAdapterAlgorithmGeneral.cxx

@@ -95,10 +95,16 @@ public:
     : Superclass(storage) {  }
 };
 
+template<typename VirtualObject, typename TargetClass>
+struct VirtualObjectTransfer<
+  VirtualObject, TargetClass, vtkm::cont::DeviceAdapterTagTestAlgorithmGeneral> :
+  public VirtualObjectTransferShareWithControl<VirtualObject, TargetClass>
+{
+};
 
 }
 }
-} // namespace vtkm::cont::testing
+} // namespace vtkm::cont::internal
 
 int UnitTestDeviceAdapterAlgorithmGeneral(int, char *[])
 {