diff --git a/vtkm/cont/ArrayHandle.h b/vtkm/cont/ArrayHandle.h
index 33558b904..345669d7b 100644
--- a/vtkm/cont/ArrayHandle.h
+++ b/vtkm/cont/ArrayHandle.h
@@ -302,24 +302,21 @@ public:
   {
     VTKM_IS_DEVICE_ADAPTER_TAG(DeviceAdapterTag);
 
-    if (this->Internals->ExecutionArrayValid)
-    {
-      // Nothing to do, data already loaded.
-    }
-    else if (this->Internals->ControlArrayValid)
-    {
-      this->PrepareForDevice(DeviceAdapterTag());
-      this->Internals->ExecutionArray->LoadDataForInput(
-        this->Internals->ControlArray);
-      this->Internals->ExecutionArrayValid = true;
-    }
-    else
+    if (!this->Internals->ControlArrayValid
+        && !this->Internals->ExecutionArrayValid)
     {
       throw vtkm::cont::ErrorControlBadValue(
         "ArrayHandle has no data when PrepareForInput called.");
     }
-    return this->Internals->ExecutionArray->GetPortalConstExecution(
-             DeviceAdapterTag());
+
+    this->PrepareForDevice(DeviceAdapterTag());
+    typename ExecutionTypes<DeviceAdapterTag>::PortalConst portal =
+        this->Internals->ExecutionArray->PrepareForInput(
+          !this->Internals->ExecutionArrayValid, DeviceAdapterTag());
+
+    this->Internals->ExecutionArrayValid = true;
+
+    return portal;
   }
 
   /// Prepares (allocates) this array to be used as an output from an operation
@@ -342,8 +339,9 @@ public:
     this->Internals->ControlArrayValid = false;
 
     this->PrepareForDevice(DeviceAdapterTag());
-    this->Internals->ExecutionArray->AllocateArrayForOutput(
-      this->Internals->ControlArray, numberOfValues);
+    typename ExecutionTypes<DeviceAdapterTag>::Portal portal =
+        this->Internals->ExecutionArray->PrepareForOutput(numberOfValues,
+                                                          DeviceAdapterTag());
 
     // We are assuming that the calling code will fill the array using the
     // iterators we are returning, so go ahead and mark the execution array as
@@ -356,7 +354,7 @@ public:
     // assumption anyway.)
     this->Internals->ExecutionArrayValid = true;
 
-    return this->Internals->ExecutionArray->GetPortalExecution(DeviceAdapterTag());
+    return portal;
   }
 
   /// Prepares this array to be used in an in-place operation (both as input
@@ -372,32 +370,26 @@ public:
   {
     VTKM_IS_DEVICE_ADAPTER_TAG(DeviceAdapterTag);
 
-    // This code is similar to PrepareForInput except that we have to give a
-    // writable portal instead of the const portal to the execution array
-    // manager so that the data can (potentially) be written to.
-    if (this->Internals->ExecutionArrayValid)
-    {
-      // Nothing to do, data already loaded.
-    }
-    else if (this->Internals->ControlArrayValid)
-    {
-      this->PrepareForDevice(DeviceAdapterTag());
-      this->Internals->ExecutionArray->LoadDataForInPlace(
-        this->Internals->ControlArray);
-      this->Internals->ExecutionArrayValid = true;
-    }
-    else
+    if (!this->Internals->ControlArrayValid
+        && !this->Internals->ExecutionArrayValid)
     {
       throw vtkm::cont::ErrorControlBadValue(
         "ArrayHandle has no data when PrepareForInput called.");
     }
 
+    this->PrepareForDevice(DeviceAdapterTag());
+    typename ExecutionTypes<DeviceAdapterTag>::Portal portal =
+        this->Internals->ExecutionArray->PrepareForInPlace(
+          !this->Internals->ExecutionArrayValid, DeviceAdapterTag());
+
+    this->Internals->ExecutionArrayValid = true;
+
     // Invalidate any control arrays since their data will become invalid when
     // the execution data is overwritten. Don't actually release the control
     // array. It may be shared as the execution array.
     this->Internals->ControlArrayValid = false;
 
-    return this->Internals->ExecutionArray->GetPortalExecution(DeviceAdapterTag());
+    return portal;
   }
 
 // private:
@@ -451,14 +443,14 @@ public:
       }
 
     VTKM_ASSERT_CONT(this->Internals->ExecutionArray == NULL);
-    VTKM_ASSERT_CONT(this->Internals->ExecutionArrayValid == false);
+    VTKM_ASSERT_CONT(!this->Internals->ExecutionArrayValid);
     // Need to change some state that does not change the logical state from
     // an external point of view.
     InternalStruct *internals
         = const_cast<InternalStruct*>(this->Internals.get());
     internals->ExecutionArray.reset(
           new vtkm::cont::internal::ArrayHandleExecutionManager<
-            T, StorageTag, DeviceAdapterTag>);
+            T, StorageTag, DeviceAdapterTag>(internals->ControlArray));
   }
 
   /// Synchronizes the control array with the execution array. If either the
diff --git a/vtkm/cont/ArrayHandleCompositeVector.h b/vtkm/cont/ArrayHandleCompositeVector.h
index 9fb18b9ed..d61251499 100644
--- a/vtkm/cont/ArrayHandleCompositeVector.h
+++ b/vtkm/cont/ArrayHandleCompositeVector.h
@@ -408,38 +408,44 @@ public:
   typedef ArrayPortalCompositeVector<SignatureWithPortals> PortalConstExecution;
 
   VTKM_CONT_EXPORT
-  ArrayTransfer() : StorageValid(false) {  }
+  ArrayTransfer(StorageType &storage) : Storage(storage) {  }
 
   VTKM_CONT_EXPORT
   vtkm::Id GetNumberOfValues() const {
-    VTKM_ASSERT_CONT(this->StorageValid);
     return this->Storage.GetNumberOfValues();
   }
 
   VTKM_CONT_EXPORT
-  void LoadDataForInput(const StorageType &controlArray)
+  PortalConstExecution PrepareForInput(bool vtkmNotUsed(updateData)) const
   {
-    this->Storage = controlArray;
-    this->StorageValid = true;
+    return
+        PortalConstExecution(
+          this->Storage.GetArrays().StaticTransformCont(
+            detail::CompositeVectorArrayToPortalExec<DeviceAdapterTag>()),
+          this->Storage.GetSourceComponents());
   }
 
   VTKM_CONT_EXPORT
-  void LoadDataForInPlace(StorageType &vtkmNotUsed(controlArray))
+  PortalExecution PrepareForInPlace(bool vtkmNotUsed(updateData))
   {
+    // It may be the case a composite vector could be used for in place
+    // operations, but this is not implemented currently.
     throw vtkm::cont::ErrorControlBadValue(
           "Composite vector arrays cannot be used for output or in place.");
   }
 
   VTKM_CONT_EXPORT
-  void AllocateArrayForOutput(StorageType &vtkmNotUsed(controlArray),
-                              vtkm::Id vtkmNotUsed(numberOfValues))
+  PortalExecution PrepareForOutput(vtkm::Id vtkmNotUsed(numberOfValues))
   {
+    // It may be the case a composite vector could be used for output if you
+    // want the delegate arrays to be resized, but this is not implemented
+    // currently.
     throw vtkm::cont::ErrorControlBadValue(
           "Composite vector arrays cannot be used for output.");
   }
 
   VTKM_CONT_EXPORT
-  void RetrieveOutputData(StorageType &vtkmNotUsed(controlArray)) const
+  void RetrieveOutputData(StorageType &vtkmNotUsed(storage)) const
   {
     throw vtkm::cont::ErrorControlBadValue(
           "Composite vector arrays cannot be used for output.");
@@ -452,32 +458,13 @@ public:
           "Composite vector arrays cannot be resized.");
   }
 
-  VTKM_CONT_EXPORT
-  PortalExecution GetPortalExecution()
-  {
-    throw vtkm::cont::ErrorControlBadValue(
-          "Composite vector arrays are read-only. (Get the const portal.)");
-  }
-
-  VTKM_CONT_EXPORT
-  PortalConstExecution GetPortalConstExecution() const
-  {
-    VTKM_ASSERT_CONT(this->StorageValid);
-    return
-        PortalConstExecution(
-          this->Storage.GetArrays().StaticTransformCont(
-            detail::CompositeVectorArrayToPortalExec<DeviceAdapterTag>()),
-          this->Storage.GetSourceComponents());
-  }
-
   VTKM_CONT_EXPORT
   void ReleaseResources() {
     this->Storage.ReleaseResources();
   }
 
 private:
-  bool StorageValid;
-  StorageType Storage;
+  StorageType &Storage;
 };
 
 } // namespace internal
diff --git a/vtkm/cont/StorageImplicit.h b/vtkm/cont/StorageImplicit.h
index 231cf0f50..fab14ef38 100644
--- a/vtkm/cont/StorageImplicit.h
+++ b/vtkm/cont/StorageImplicit.h
@@ -122,32 +122,29 @@ public:
   typedef PortalConstControl PortalConstExecution;
 
   VTKM_CONT_EXPORT
-  ArrayTransfer() : PortalValid(false) {  }
+  ArrayTransfer(StorageType &storage) : Storage(storage) {  }
 
   VTKM_CONT_EXPORT
   vtkm::Id GetNumberOfValues() const
   {
-    VTKM_ASSERT_CONT(this->PortalValid);
-    return this->Portal.GetNumberOfValues();
+    return this->Storage.GetNumberOfValues();
   }
 
   VTKM_CONT_EXPORT
-  void LoadDataForInput(const StorageType& controlArray)
+  PortalConstExecution PrepareForInput(bool vtkmNotUsed(updateData))
   {
-    this->Portal = controlArray.GetPortalConst();
-    this->PortalValid = true;
+    return this->Storage.GetPortalConst();
   }
 
   VTKM_CONT_EXPORT
-  void LoadDataForInPlace(StorageType &vtkmNotUsed(controlArray))
+  PortalExecution PrepareForInPlace(bool vtkmNotUsed(updateData))
   {
     throw vtkm::cont::ErrorControlBadValue(
           "Implicit arrays cannot be used for output or in place.");
   }
 
   VTKM_CONT_EXPORT
-  void AllocateArrayForOutput(StorageType &vtkmNotUsed(controlArray),
-                              vtkm::Id vtkmNotUsed(numberOfValues))
+  PortalExecution PrepareForOutput(vtkm::Id vtkmNotUsed(numberOfValues))
   {
     throw vtkm::cont::ErrorControlBadValue(
           "Implicit arrays cannot be used for output.");
@@ -165,25 +162,11 @@ public:
     throw vtkm::cont::ErrorControlBadValue("Implicit arrays cannot be resized.");
   }
 
-  VTKM_CONT_EXPORT
-  PortalExecution GetPortalExecution()
-  {
-    throw vtkm::cont::ErrorControlBadValue(
-          "Implicit arrays are read-only.  (Get the const portal.)");
-  }
-  VTKM_CONT_EXPORT
-  PortalConstExecution GetPortalConstExecution() const
-  {
-    VTKM_ASSERT_CONT(this->PortalValid);
-    return this->Portal;
-  }
-
   VTKM_CONT_EXPORT
   void ReleaseResources() {  }
 
 private:
-  PortalConstExecution Portal;
-  bool PortalValid;
+  StorageType &Storage;
 };
 
 } // namespace internal
diff --git a/vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.h b/vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.h
index 589334962..283e723de 100644
--- a/vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.h
+++ b/vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.h
@@ -47,11 +47,16 @@ public:
   typedef typename Superclass::PortalConstType PortalConstType;
   typedef typename Superclass::StorageType StorageType;
 
-  VTKM_CONT_EXPORT void LoadDataForInput(const StorageType &arrayPortal)
+  VTKM_CONT_EXPORT
+  ArrayManagerExecution(StorageType &storage)
+    : Superclass(storage) {  }
+
+  VTKM_CONT_EXPORT
+  PortalConstType PrepareForInput(bool updateData)
   {
     try
       {
-      this->Superclass::LoadDataForInput(arrayPortal);
+      return this->Superclass::PrepareForInput(updateData);
       }
     catch (vtkm::cont::ErrorControlOutOfMemory error)
       {
@@ -65,14 +70,31 @@ public:
       }
   }
 
-  VTKM_CONT_EXPORT void AllocateArrayForOutput(
-      vtkm::cont::internal::Storage<ValueType,StorageTag>
-        &container,
-      vtkm::Id numberOfValues)
+  VTKM_CONT_EXPORT
+  PortalType PrepareForInPlace(bool updateData)
   {
     try
       {
-      this->Superclass::AllocateArrayForOutput(container, numberOfValues);
+      return this->Superclass::PrepareForInPlace(updateData);
+      }
+    catch (vtkm::cont::ErrorControlOutOfMemory error)
+      {
+      // Thrust does not seem to be clearing the CUDA error, so do it here.
+      cudaError_t cudaError = cudaPeekAtLastError();
+      if (cudaError == cudaErrorMemoryAllocation)
+        {
+        cudaGetLastError();
+        }
+      throw error;
+      }
+  }
+
+  VTKM_CONT_EXPORT
+  PortalType PrepareForOutput(vtkm::Id numberOfValues)
+  {
+    try
+      {
+      return this->Superclass::PrepareForOutput(numberOfValues);
       }
     catch (vtkm::cont::ErrorControlOutOfMemory error)
       {
diff --git a/vtkm/cont/cuda/internal/ArrayManagerExecutionThrustDevice.h b/vtkm/cont/cuda/internal/ArrayManagerExecutionThrustDevice.h
index dd3ae0341..6080eefa8 100644
--- a/vtkm/cont/cuda/internal/ArrayManagerExecutionThrustDevice.h
+++ b/vtkm/cont/cuda/internal/ArrayManagerExecutionThrustDevice.h
@@ -90,12 +90,14 @@ public:
   typedef vtkm::exec::cuda::internal::ArrayPortalFromThrust< T > PortalType;
   typedef vtkm::exec::cuda::internal::ConstArrayPortalFromThrust< const T > PortalConstType;
 
-  VTKM_CONT_EXPORT ArrayManagerExecutionThrustDevice():
-    Array()
+  VTKM_CONT_EXPORT
+  ArrayManagerExecutionThrustDevice(StorageType &storage)
+    : Storage(storage), Array()
   {
 
   }
 
+  VTKM_CONT_EXPORT
   ~ArrayManagerExecutionThrustDevice()
   {
     this->ReleaseResources();
@@ -103,62 +105,83 @@ public:
 
   /// Returns the size of the array.
   ///
-  VTKM_CONT_EXPORT vtkm::Id GetNumberOfValues() const {
+  VTKM_CONT_EXPORT
+  vtkm::Id GetNumberOfValues() const {
     return this->Array.size();
   }
 
   /// Allocates the appropriate size of the array and copies the given data
   /// into the array.
   ///
-  VTKM_CONT_EXPORT void LoadDataForInput(const StorageType &storage)
+  VTKM_CONT_EXPORT
+  PortalConstType PrepareForInput(bool updateData)
   {
-    //don't bind to the texture yet, as we could have allocate the array
-    //on a previous call with AllocateArrayForOutput and now are directly
-    //calling get portal const
-    try
-      {
-      this->Array.assign(
-            vtkm::cont::ArrayPortalToIteratorBegin(storage.GetPortalConst()),
-            vtkm::cont::ArrayPortalToIteratorEnd(storage.GetPortalConst()));
-      }
-    catch (std::bad_alloc error)
-      {
-      throw vtkm::cont::ErrorControlOutOfMemory(error.what());
-      }
+    if (updateData)
+    {
+      this->CopyToExecution();
+    }
+    else // !updateData
+    {
+      // The data in this->Array should already be valid. We are not messing
+      // with texture yet, so we don't have to worry about being in the wrong
+      // memory space.
+    }
+
+    return PortalConstType(this->Array.data(),
+                           this->Array.data() + this->Array.size());
   }
 
   /// Allocates the appropriate size of the array and copies the given data
   /// into the array.
   ///
-  template<class PortalControl>
-  VTKM_CONT_EXPORT void LoadDataForInPlace(PortalControl arrayPortal)
+  VTKM_CONT_EXPORT
+  PortalType PrepareForInPlace(bool updateData)
   {
-    this->LoadDataForInput(arrayPortal);
+    if (updateData)
+    {
+      this->CopyToExecution();
+    }
+    else // !updateData
+    {
+      // The data in this->Array should already be valid. We are not messing
+      // with texture yet, so we don't have to worry about being in the wrong
+      // memory space.
+    }
+
+    return PortalType(this->Array.data(),
+                      this->Array.data() + this->Array.size());
   }
 
   /// Allocates the array to the given size.
   ///
-  VTKM_CONT_EXPORT void AllocateArrayForOutput(
-      StorageType &vtkmNotUsed(container),
-      vtkm::Id numberOfValues)
+  VTKM_CONT_EXPORT
+  PortalType PrepareForOutput(vtkm::Id numberOfValues)
   {
-    try
-      {
+    if (numberOfValues > this->GetNumberOfValues())
+    {
       // Resize to 0 first so that you don't have to copy data when resizing
       // to a larger size.
       this->Array.clear();
+    }
+
+    try
+      {
       this->Array.resize(numberOfValues);
       }
     catch (std::bad_alloc error)
       {
       throw vtkm::cont::ErrorControlOutOfMemory(error.what());
       }
+
+    return PortalType(this->Array.data(),
+                      this->Array.data() + this->Array.size());
   }
 
   /// Allocates enough space in \c storage and copies the data in the
   /// device vector into it.
   ///
-  VTKM_CONT_EXPORT void RetrieveOutputData(StorageType &storage) const
+  VTKM_CONT_EXPORT
+  void RetrieveOutputData(StorageType &storage) const
   {
     storage.Allocate(this->Array.size());
     ::thrust::copy(this->Array.data(),
@@ -177,18 +200,6 @@ public:
     this->Array.resize(numberOfValues);
   }
 
-  VTKM_CONT_EXPORT PortalType GetPortal()
-  {
-    return PortalType( this->Array.data(),
-                       this->Array.data() + this->Array.size());
-  }
-
-  VTKM_CONT_EXPORT PortalConstType GetPortalConst() const
-  {
-    return PortalConstType( this->Array.data(),
-                            this->Array.data() + this->Array.size());
-  }
-
 
   /// Frees all memory.
   ///
@@ -205,8 +216,28 @@ private:
   void operator=(
       ArrayManagerExecutionThrustDevice<T, StorageTag> &);
 
+  StorageType &Storage;
+
   ::thrust::system::cuda::vector<ValueType,
                                  UninitializedAllocator<ValueType> > Array;
+
+  VTKM_CONT_EXPORT
+  void CopyToExecution()
+  {
+    //don't bind to the texture yet, as we could have allocate the array
+    //on a previous call with AllocateArrayForOutput and now are directly
+    //calling get portal const
+    try
+    {
+      this->Array.assign(
+            vtkm::cont::ArrayPortalToIteratorBegin(this->Storage.GetPortalConst()),
+            vtkm::cont::ArrayPortalToIteratorEnd(this->Storage.GetPortalConst()));
+    }
+    catch (std::bad_alloc error)
+    {
+      throw vtkm::cont::ErrorControlOutOfMemory(error.what());
+    }
+  }
 };
 
 
diff --git a/vtkm/cont/cuda/testing/UnitTestDeviceAdapterCuda.cu b/vtkm/cont/cuda/testing/UnitTestDeviceAdapterCuda.cu
index 66c038b02..3d6a8e34c 100644
--- a/vtkm/cont/cuda/testing/UnitTestDeviceAdapterCuda.cu
+++ b/vtkm/cont/cuda/testing/UnitTestDeviceAdapterCuda.cu
@@ -26,38 +26,6 @@
 #include <vtkm/cont/testing/TestingDeviceAdapter.h>
 #include <vtkm/cont/cuda/internal/testing/Testing.h>
 
-namespace vtkm {
-namespace cont {
-namespace testing {
-
-template<>
-struct CopyInto<vtkm::cont::DeviceAdapterTagCuda>
-{
-  template<typename T, typename StorageTagType>
-  VTKM_CONT_EXPORT
-  void operator()( vtkm::cont::internal::ArrayManagerExecution<
-                    T,
-                    StorageTagType,
-                    vtkm::cont::DeviceAdapterTagCuda>& manager,
-                 T* start)
-  {
-    typedef vtkm::cont::internal::Storage< T, StorageTagType > StorageType;
-    StorageType outputArray;
-    std::cout << "now calling RetrieveOutputData: " << std::endl;
-    manager.RetrieveOutputData( outputArray );
-
-    vtkm::cont::ArrayPortalToIterators<
-                typename StorageType::PortalConstType>
-      iterators(outputArray.GetPortalConst());
-     std::copy(iterators.GetBegin(), iterators.GetEnd(), start);
-  }
-};
-
-
-}
-}
-}
-
 int UnitTestDeviceAdapterCuda(int, char *[])
 {
   int result =  vtkm::cont::testing::TestingDeviceAdapter
diff --git a/vtkm/cont/internal/ArrayHandleExecutionManager.h b/vtkm/cont/internal/ArrayHandleExecutionManager.h
index 8a4ef73d0..e146cf158 100644
--- a/vtkm/cont/internal/ArrayHandleExecutionManager.h
+++ b/vtkm/cont/internal/ArrayHandleExecutionManager.h
@@ -65,39 +65,74 @@ public:
   /// Returns the number of values stored in the array.  Results are undefined
   /// if data has not been loaded or allocated.
   ///
-  virtual vtkm::Id GetNumberOfValues() const = 0;
+  VTKM_CONT_EXPORT
+  vtkm::Id GetNumberOfValues() const {
+    return this->GetNumberOfValuesImpl();
+  }
 
-  /// Allocates a large enough array in the execution environment and copies
-  /// the given data to that array. The allocated array can later be accessed
-  /// via the GetPortalConstExecution method. If control and execution share
-  /// arrays, then this method may save the iterators to be returned in the \c
-  /// GetPortalConst methods.
+  /// Prepares the data for use as input in the execution environment. If the
+  /// flag \c updateData is true, then data is transferred to the execution
+  /// environment. Otherwise, this transfer should be skipped.
   ///
-  virtual void LoadDataForInput(const StorageType &storage) = 0;
+  /// Returns a constant array portal valid in the execution environment.
+  ///
+  template<typename DeviceAdapter>
+  VTKM_CONT_EXPORT
+  typename ExecutionTypes<DeviceAdapter>::PortalConst
+  PrepareForInput(bool updateData, DeviceAdapter) {
+    this->VerifyDeviceAdapter(DeviceAdapter());
 
-  /// Allocates a large enough array in the execution environment and copies
-  /// the given data to that array. The allocated array can later be accessed
-  /// via the GetPortalExection method. If control and execution share arrays,
-  /// then this method may save the iterators of the storage to be returned
-  /// in the \c GetPortal* methods.
+    typename ExecutionTypes<DeviceAdapter>::PortalConst portal;
+    this->PrepareForInputImpl(updateData, &portal);
+    return portal;
+  }
+
+  /// Prepares the data for use as both input and output in the execution
+  /// environment. If the flag \c updateData is true, then data is transferred
+  /// to the execution environment. Otherwise, this transfer should be skipped.
   ///
-  virtual void LoadDataForInPlace(StorageType &storage) = 0;
+  /// Returns a read-write array portal valid in the execution environment.
+  ///
+  template<typename DeviceAdapter>
+  VTKM_CONT_EXPORT
+  typename ExecutionTypes<DeviceAdapter>::Portal
+  PrepareForInPlace(bool updateData, DeviceAdapter) {
+    this->VerifyDeviceAdapter(DeviceAdapter());
+
+    typename ExecutionTypes<DeviceAdapter>::Portal portal;
+    this->PrepareForInPlaceImpl(updateData, &portal);
+    return portal;
+  }
 
   /// Allocates an array in the execution environment of the specified size. If
   /// control and execution share arrays, then this class can allocate data
-  /// using the given Storage and remember its iterators so that it can be used
-  /// directly in the execution environment.
+  /// using the given Storage it can be used directly in the execution
+  /// environment.
   ///
-  virtual void AllocateArrayForOutput(StorageType &storage,
-                                      vtkm::Id numberOfValues) = 0;
+  /// Returns a writable array portal valid in the execution environment.
+  ///
+  template<typename DeviceAdapter>
+  VTKM_CONT_EXPORT
+  typename ExecutionTypes<DeviceAdapter>::Portal
+  PrepareForOutput(vtkm::Id numberOfValues, DeviceAdapter) {
+    this->VerifyDeviceAdapter(DeviceAdapter());
 
-  /// Allocates data in the given Storage and copies data held
-  /// in the execution environment (managed by this class) into the control
-  /// array. If control and execution share arrays, this can be no operation.
-  /// This method should only be called after AllocateArrayForOutput is
+    typename ExecutionTypes<DeviceAdapter>::Portal portal;
+    this->PrepareForOutputImpl(numberOfValues, &portal);
+    return portal;
+  }
+
+  /// Allocates data in the given Storage and copies data held in the execution
+  /// environment (managed by this class) into the storage object. The
+  /// reference to the storage given is the same as that passed to the
+  /// constructor. If control and execution share arrays, this can be no
+  /// operation. This method should only be called after PrepareForOutput is
   /// called.
   ///
-  virtual void RetrieveOutputData(StorageType &storage) const = 0;
+  VTKM_CONT_EXPORT
+  void RetrieveOutputData(StorageType &storage) const {
+    this->RetrieveOutputDataImpl(storage);
+  }
 
   /// \brief Reduces the size of the array without changing its values.
   ///
@@ -108,43 +143,18 @@ public:
   /// (returned from GetNumberOfValues). That is, this method can only be used
   /// to shorten the array, not lengthen.
   ///
-  virtual void Shrink(vtkm::Id numberOfValues) = 0;
-
-  /// Returns an array portal that can be used in the execution environment.
-  /// This portal was defined in either LoadDataForInput or
-  /// AllocateArrayForOutput. If control and environment share memory space,
-  /// this class may return the iterator from the \c controlArray.
-  ///
-  template<typename DeviceAdapter>
   VTKM_CONT_EXPORT
-  typename ExecutionTypes<DeviceAdapter>::Portal
-  GetPortalExecution(DeviceAdapter device)
-  {
-    this->VerifyDeviceAdapter(device);
-
-    typename ExecutionTypes<DeviceAdapter>::Portal portal;
-    this->GetPortalExecutionImpl(&portal);
-    return portal;
-  }
-
-  /// Const version of GetPortal.
-  ///
-  template<typename DeviceAdapter>
-  VTKM_CONT_EXPORT
-  typename ExecutionTypes<DeviceAdapter>::PortalConst
-  GetPortalConstExecution(DeviceAdapter device) const
-  {
-    this->VerifyDeviceAdapter(device);
-
-    typename ExecutionTypes<DeviceAdapter>::PortalConst portal;
-    this->GetPortalConstExecutionImpl(&portal);
-    return portal;
+  void Shrink(vtkm::Id numberOfValues) {
+    this->ShrinkImpl(numberOfValues);
   }
 
   /// Frees any resources (i.e. memory) allocated for the exeuction
   /// environment, if any.
   ///
-  virtual void ReleaseResources() = 0;
+  VTKM_CONT_EXPORT
+  void ReleaseResources() {
+    this->ReleaseResourcesImpl();
+  }
 
   template<typename DeviceAdapter>
   VTKM_CONT_EXPORT
@@ -155,10 +165,22 @@ public:
   }
 
 protected:
-  virtual void GetPortalExecutionImpl(void *portalExecution) = 0;
+  virtual vtkm::Id GetNumberOfValuesImpl() const = 0;
 
-  virtual void GetPortalConstExecutionImpl(
-    void *portalConstExecution) const = 0;
+  virtual void PrepareForInputImpl(bool updateData,
+                                   void *portalExecutionVoid) = 0;
+
+  virtual void PrepareForInPlaceImpl(bool updateData,
+                                     void *portalExecutionVoid) = 0;
+
+  virtual void PrepareForOutputImpl(vtkm::Id numberOfValues,
+                                    void *portalExecution) = 0;
+
+  virtual void RetrieveOutputDataImpl(StorageType &storage) const = 0;
+
+  virtual void ShrinkImpl(Id numberOfValues) = 0;
+
+  virtual void ReleaseResourcesImpl() = 0;
 
   virtual bool IsDeviceAdapterImpl(
     const vtkm::cont::internal::DeviceAdapterId &id) const = 0;
@@ -197,65 +219,59 @@ public:
   typedef typename ArrayTransferType::PortalControl PortalControl;
   typedef typename ArrayTransferType::PortalConstControl PortalConstControl;
 
+  typedef typename ArrayTransferType::PortalExecution PortalExecution;
+  typedef typename ArrayTransferType::PortalConstExecution PortalConstExecution;
+
   VTKM_CONT_EXPORT
-  vtkm::Id GetNumberOfValues() const
+  ArrayHandleExecutionManager(StorageType &storage)
+    : Transfer(storage) {  }
+
+protected:
+  VTKM_CONT_EXPORT
+  vtkm::Id GetNumberOfValuesImpl() const
   {
     return this->Transfer.GetNumberOfValues();
   }
 
   VTKM_CONT_EXPORT
-  void LoadDataForInput(const StorageType &storage)
+  void PrepareForInputImpl(bool updateData, void *portalExecutionVoid)
   {
-    this->Transfer.LoadDataForInput(storage);
+    PortalConstExecution portal = this->Transfer.PrepareForInput(updateData);
+    *reinterpret_cast<PortalConstExecution *>(portalExecutionVoid) = portal;
   }
 
   VTKM_CONT_EXPORT
-  void LoadDataForInPlace(StorageType &storage)
+  void PrepareForInPlaceImpl(bool updateData, void *portalExecutionVoid)
   {
-    this->Transfer.LoadDataForInPlace(storage);
+    PortalExecution portal = this->Transfer.PrepareForInPlace(updateData);
+    *reinterpret_cast<PortalExecution *>(portalExecutionVoid) = portal;
   }
 
   VTKM_CONT_EXPORT
-  void AllocateArrayForOutput(StorageType &storage, Id numberOfValues)
+  void PrepareForOutputImpl(vtkm::Id numberOfValues, void *portalExecutionVoid)
   {
-    this->Transfer.AllocateArrayForOutput(storage, numberOfValues);
+    PortalExecution portal = this->Transfer.PrepareForOutput(numberOfValues);
+    *reinterpret_cast<PortalExecution *>(portalExecutionVoid) = portal;
   }
 
   VTKM_CONT_EXPORT
-  void RetrieveOutputData(StorageType &storage) const
+  void RetrieveOutputDataImpl(StorageType &storage) const
   {
     this->Transfer.RetrieveOutputData(storage);
   }
 
   VTKM_CONT_EXPORT
-  void Shrink(Id numberOfValues)
+  void ShrinkImpl(Id numberOfValues)
   {
     this->Transfer.Shrink(numberOfValues);
   }
 
   VTKM_CONT_EXPORT
-  void ReleaseResources()
+  void ReleaseResourcesImpl()
   {
     this->Transfer.ReleaseResources();
   }
 
-protected:
-  VTKM_CONT_EXPORT
-  void GetPortalExecutionImpl(void *portalExecutionVoid)
-  {
-    typedef typename ArrayTransferType::PortalExecution PortalType;
-    PortalType portalExecution = this->Transfer.GetPortalExecution();
-    *reinterpret_cast<PortalType *>(portalExecutionVoid) = portalExecution;
-  }
-
-  VTKM_CONT_EXPORT
-  void GetPortalConstExecutionImpl(void *portalExecutionVoid) const
-  {
-    typedef typename ArrayTransferType::PortalConstExecution PortalType;
-    PortalType portalExecution = this->Transfer.GetPortalConstExecution();
-    *reinterpret_cast<PortalType *>(portalExecutionVoid) = portalExecution;
-  }
-
   VTKM_CONT_EXPORT
   bool IsDeviceAdapterImpl(const DeviceAdapterId &id) const
   {
diff --git a/vtkm/cont/internal/ArrayManagerExecution.h b/vtkm/cont/internal/ArrayManagerExecution.h
index f34bc3ced..62524943d 100644
--- a/vtkm/cont/internal/ArrayManagerExecution.h
+++ b/vtkm/cont/internal/ArrayManagerExecution.h
@@ -70,42 +70,59 @@ public:
   typedef vtkm::exec::internal::ArrayPortalFromIterators<const ValueType*>
       PortalConstType;
 
+  /// All ArrayManagerExecution classes must have a constructor that takes a
+  /// storage reference. The reference may be saved (and will remain valid
+  /// throughout the life of the ArrayManagerExecution). Copying storage
+  /// objects should be avoided (copy references or pointers only). The
+  /// reference can also, of course, be ignored.
+  ///
+  VTKM_CONT_EXPORT
+  ArrayManagerExecution(vtkm::cont::internal::Storage<T,StorageTag> &storage);
+
   /// Returns the number of values stored in the array.  Results are undefined
   /// if data has not been loaded or allocated.
   ///
-  VTKM_CONT_EXPORT vtkm::Id GetNumberOfValues() const;
+  VTKM_CONT_EXPORT
+  vtkm::Id GetNumberOfValues() const;
 
-  /// Allocates a large enough array in the execution environment and copies
-  /// the data from the storage to that array. The allocated array can later be
-  /// accessed via the GetPortalConst method. If control and execution share
-  /// arrays, then this method may save the iterators to be returned in the \c
-  /// GetPortalConst method.
+  /// Prepares the data for use as input in the execution environment. If the
+  /// flag \c updateData is true, then data is transferred to the execution
+  /// environment. Otherwise, this transfer should be skipped.
   ///
-  VTKM_CONT_EXPORT void LoadDataForInput(const StorageType &storage);
-
-  /// Allocates a large enough array in the execution environment and copies
-  /// the data from the storage to that array. The allocated array can later be
-  /// accessed via the GetPortal method. If control and execution share arrays,
-  /// then this method may save the iterators of the storage to be returned in
-  /// the \c GetPortal* methods.
+  /// Returns a constant array portal valid in the execution environment.
   ///
-  VTKM_CONT_EXPORT void LoadDataForInPlace(StorageType &storage);
+  VTKM_CONT_EXPORT
+  PortalConstExecution PrepareForInput(bool updateData);
 
-  /// Allocates an array in the execution environment of the specified size.
-  /// If control and execution share arrays, then this class can allocate
-  /// data using the given Storage object and remember its iterators
-  /// so that it can be used directly in the execution environment.
+  /// Prepares the data for use as both input and output in the execution
+  /// environment. If the flag \c updateData is true, then data is transferred
+  /// to the execution environment. Otherwise, this transfer should be skipped.
   ///
-  VTKM_CONT_EXPORT void AllocateArrayForOutput(StorageType &storage,
-                                               vtkm::Id numberOfValues);
+  /// Returns a read-write array portal valid in the execution environment.
+  ///
+  VTKM_CONT_EXPORT
+  PortalExecution LoadDataForInPlace(bool updateData);
 
-  /// Allocates data in the given Storage and copies data held
-  /// in the execution environment (managed by this class) into the control
-  /// array. If control and execution share arrays, this can be no operation.
-  /// This method should only be called after AllocateArrayForOutput is
+  /// Allocates an array in the execution environment of the specified size. If
+  /// control and execution share arrays, then this class can allocate data
+  /// using the given Storage it can be used directly in the execution
+  /// environment.
+  ///
+  /// Returns a writable array portal valid in the execution environment.
+  ///
+  VTKM_CONT_EXPORT
+  PortalExecution PrepareForOutput(vtkm::Id numberOfValues);
+
+  /// Allocates data in the given Storage and copies data held in the execution
+  /// environment (managed by this class) into the storage object. The
+  /// reference to the storage given is the same as that passed to the
+  /// constructor. If control and execution share arrays, this can be no
+  /// operation. This method should only be called after PrepareForOutput is
   /// called.
   ///
-  VTKM_CONT_EXPORT void RetrieveOutputData(StorageType &storage) const;
+  VTKM_CONT_EXPORT
+  void RetrieveOutputData(
+      vtkm::cont::internal::Storage<T,StorageTag>  &storage) const;
 
   /// \brief Reduces the size of the array without changing its values.
   ///
@@ -116,23 +133,14 @@ public:
   /// (returned from GetNumberOfValues). That is, this method can only be used
   /// to shorten the array, not lengthen.
   ///
-  VTKM_CONT_EXPORT void Shrink(vtkm::Id numberOfValues);
-
-  /// Returns an array portal that can be used in the execution environment.
-  /// This portal was defined in either LoadDataForInput or
-  /// AllocateArrayForOutput. If control and environment share memory space,
-  /// this class may return the iterator from the \c controlArray.
-  ///
-  VTKM_CONT_EXPORT PortalType GetPortal();
-
-  /// Const version of GetPortal.
-  ///
-  VTKM_CONT_EXPORT PortalConstType GetPortalConst() const;
+  VTKM_CONT_EXPORT
+  void Shrink(vtkm::Id numberOfValues);
 
   /// Frees any resources (i.e. memory) allocated for the exeuction
   /// environment, if any.
   ///
-  VTKM_CONT_EXPORT void ReleaseResources();
+  VTKM_CONT_EXPORT
+  void ReleaseResources();
 };
 #else // VTKM_DOXGEN_ONLY
 ;
diff --git a/vtkm/cont/internal/ArrayManagerExecutionSerial.h b/vtkm/cont/internal/ArrayManagerExecutionSerial.h
index 67d8fdcda..5f48363e8 100644
--- a/vtkm/cont/internal/ArrayManagerExecutionSerial.h
+++ b/vtkm/cont/internal/ArrayManagerExecutionSerial.h
@@ -39,6 +39,10 @@ public:
   typedef typename Superclass::ValueType ValueType;
   typedef typename Superclass::PortalType PortalType;
   typedef typename Superclass::PortalConstType PortalConstType;
+
+  VTKM_CONT_EXPORT
+  ArrayManagerExecution(typename Superclass::StorageType &storage)
+    : Superclass(storage) {  }
 };
 
 }
diff --git a/vtkm/cont/internal/ArrayManagerExecutionShareWithControl.h b/vtkm/cont/internal/ArrayManagerExecutionShareWithControl.h
index 7124636b8..71e94063d 100644
--- a/vtkm/cont/internal/ArrayManagerExecutionShareWithControl.h
+++ b/vtkm/cont/internal/ArrayManagerExecutionShareWithControl.h
@@ -25,8 +25,6 @@
 #include <vtkm/cont/Assert.h>
 #include <vtkm/cont/Storage.h>
 
-#include <vtkm/cont/internal/ArrayPortalShrink.h>
-
 #include <algorithm>
 
 namespace vtkm {
@@ -44,101 +42,69 @@ class ArrayManagerExecutionShareWithControl
 public:
   typedef T ValueType;
   typedef vtkm::cont::internal::Storage<ValueType, StorageTag> StorageType;
-  typedef vtkm::cont::internal::ArrayPortalShrink<
-      typename StorageType::PortalType> PortalType;
-  typedef vtkm::cont::internal::ArrayPortalShrink<
-      typename StorageType::PortalConstType> PortalConstType;
+  typedef typename StorageType::PortalType PortalType;
+  typedef typename StorageType::PortalConstType PortalConstType;
 
-  VTKM_CONT_EXPORT ArrayManagerExecutionShareWithControl()
-    : PortalValid(false), ConstPortalValid(false) { }
+  VTKM_CONT_EXPORT
+  ArrayManagerExecutionShareWithControl(StorageType &storage)
+    : Storage(storage) { }
 
-  /// Returns the size of the saved portal.
+  /// Returns the size of the storage.
   ///
-  VTKM_CONT_EXPORT vtkm::Id GetNumberOfValues() const
+  VTKM_CONT_EXPORT
+  vtkm::Id GetNumberOfValues() const
   {
-    VTKM_ASSERT_CONT(this->ConstPortalValid);
-    return this->ConstPortal.GetNumberOfValues();
+    return this->Storage.GetNumberOfValues();
   }
 
-  /// Saves the given iterators to be returned later.
+  /// Returns the constant portal from the storage.
   ///
-  VTKM_CONT_EXPORT void LoadDataForInput(const StorageType& storage)
+  VTKM_CONT_EXPORT
+  PortalConstType PrepareForInput(bool vtkmNotUsed(uploadData)) const
   {
-    this->ConstPortal = storage.GetPortalConst();
-    this->ConstPortalValid = true;
-
-    // Non-const versions not defined.
-    this->PortalValid = false;
+    return this->Storage.GetPortalConst();
   }
 
-  /// Saves the given iterators to be returned later.
+  /// Returns the read-write portal from the storage.
   ///
-  VTKM_CONT_EXPORT void LoadDataForInPlace(StorageType &storage)
+  VTKM_CONT_EXPORT
+  PortalType PrepareForInPlace(bool vtkmNotUsed(uploadData))
   {
-    // This only works if there is a valid cast from non-const to const
-    // iterator.
-    this->LoadDataForInput(storage);
-
-    this->Portal = storage.GetPortal();
-    this->PortalValid = true;
+    return this->Storage.GetPortal();
   }
 
-  /// Actually just allocates memory in the given \p controlArray.
+  /// Allocates data in the storage and return the portal to that.
   ///
-  VTKM_CONT_EXPORT void AllocateArrayForOutput(StorageType &controlArray,
-                                               vtkm::Id numberOfValues)
+  VTKM_CONT_EXPORT
+  PortalType PrepareForOutput(vtkm::Id numberOfValues)
   {
-    controlArray.Allocate(numberOfValues);
-
-    this->Portal = controlArray.GetPortal();
-    this->PortalValid = true;
-
-    this->ConstPortal = controlArray.GetPortalConst();
-    this->ConstPortalValid = true;
+    this->Storage.Allocate(numberOfValues);
+    return this->Storage.GetPortal();
   }
 
   /// This method is a no-op (except for a few checks). Any data written to
   /// this class's portals should already be written to the given \c
   /// controlArray (under correct operation).
   ///
-  VTKM_CONT_EXPORT void RetrieveOutputData(StorageType &controlArray) const
+  VTKM_CONT_EXPORT
+  void RetrieveOutputData(StorageType &storage) const
   {
-    VTKM_ASSERT_CONT(this->ConstPortalValid);
-    controlArray.Shrink(this->ConstPortal.GetNumberOfValues());
+    (void)storage;
+    VTKM_ASSERT_CONT(&storage == &this->Storage);
   }
 
-  /// Adjusts saved end iterators to resize array.
+  /// Shrinks the storage.
   ///
-  VTKM_CONT_EXPORT void Shrink(vtkm::Id numberOfValues)
+  VTKM_CONT_EXPORT
+  void Shrink(vtkm::Id numberOfValues)
   {
-    VTKM_ASSERT_CONT(this->ConstPortalValid);
-    this->ConstPortal.Shrink(numberOfValues);
-
-    if (this->PortalValid)
-    {
-      this->Portal.Shrink(numberOfValues);
-    }
-  }
-
-  /// Returns the portal previously saved from a \c Storage.
-  ///
-  VTKM_CONT_EXPORT PortalType GetPortal()
-  {
-    VTKM_ASSERT_CONT(this->PortalValid);
-    return this->Portal;
-  }
-
-  /// Const version of GetPortal.
-  ///
-  VTKM_CONT_EXPORT PortalConstType GetPortalConst() const
-  {
-    VTKM_ASSERT_CONT(this->ConstPortalValid);
-    return this->ConstPortal;
+    this->Storage.Shrink(numberOfValues);
   }
 
   /// A no-op.
   ///
-  VTKM_CONT_EXPORT void ReleaseResources() { }
+  VTKM_CONT_EXPORT
+  void ReleaseResources() { }
 
 private:
   // Not implemented.
@@ -147,11 +113,7 @@ private:
   void operator=(
       ArrayManagerExecutionShareWithControl<T, StorageTag> &);
 
-  PortalType Portal;
-  bool PortalValid;
-
-  PortalConstType ConstPortal;
-  bool ConstPortalValid;
+  StorageType &Storage;
 };
 
 }
diff --git a/vtkm/cont/internal/ArrayTransfer.h b/vtkm/cont/internal/ArrayTransfer.h
index ac02820a3..384674c79 100644
--- a/vtkm/cont/internal/ArrayTransfer.h
+++ b/vtkm/cont/internal/ArrayTransfer.h
@@ -66,54 +66,65 @@ public:
   typedef typename ArrayManagerType::PortalType PortalExecution;
   typedef typename ArrayManagerType::PortalConstType PortalConstExecution;
 
+  VTKM_CONT_EXPORT
+  ArrayTransfer(StorageType &storage) : ArrayManager(storage) {  }
 
   /// Returns the number of values stored in the array.  Results are undefined
   /// if data has not been loaded or allocated.
   ///
-  VTKM_CONT_EXPORT vtkm::Id GetNumberOfValues() const
+  VTKM_CONT_EXPORT
+  vtkm::Id GetNumberOfValues() const
   {
     return this->ArrayManager.GetNumberOfValues();
   }
 
-  /// Allocates a large enough array in the execution environment and copies
-  /// the given data to that array. The allocated array can later be accessed
-  /// via the GetPortalConstExecution method. If control and execution share
-  /// arrays, then this method may save the iterators to be returned in the \c
-  /// GetPortalConst methods.
+  /// Prepares the data for use as input in the execution environment. If the
+  /// flag \c updateData is true, then data is transferred to the execution
+  /// environment. Otherwise, this transfer is (or may be) skipped.
   ///
-  VTKM_CONT_EXPORT void LoadDataForInput(const StorageType &storage)
+  /// Returns a constant array portal valid in the execution environment.
+  ///
+  VTKM_CONT_EXPORT
+  PortalConstExecution PrepareForInput(bool updateData)
   {
-    this->ArrayManager.LoadDataForInput(storage);
+    return this->ArrayManager.PrepareForInput(updateData);
   }
 
-  /// Allocates a large enough array in the execution environment and copies
-  /// the given data to that array. The allocated array can later be accessed
-  /// via the GetPortalExection method. If control and execution share arrays,
-  /// then this method may save the iterators of the storage to be returned
-  /// in the \c GetPortal* methods.
+  /// Prepares the data for use as both input and output in the execution
+  /// environment. If the flag \c updateData is true, then data is transferred
+  /// to the execution environment. Otherwise, this transfer is (or may be)
+  /// skipped.
   ///
-  VTKM_CONT_EXPORT void LoadDataForInPlace(StorageType &storage)
+  /// Returns a read-write array portal valid in the execution environment.
+  ///
+  VTKM_CONT_EXPORT
+  PortalExecution PrepareForInPlace(bool updateData)
   {
-    this->ArrayManager.LoadDataForInPlace(storage);
+    return this->ArrayManager.PrepareForInPlace(updateData);
   }
 
   /// Allocates an array in the execution environment of the specified size. If
   /// control and execution share arrays, then this class can allocate data
-  /// using the given Storage and remember its iterators so that it can be used
-  /// directly in the execution environment.
+  /// using the given Storage it can be used directly in the execution
+  /// environment.
   ///
-  VTKM_CONT_EXPORT void AllocateArrayForOutput(StorageType &storage,
-                                               vtkm::Id numberOfValues)
+  /// Returns a writable array portal valid in the execution environment.
+  ///
+  VTKM_CONT_EXPORT
+  PortalExecution PrepareForOutput(vtkm::Id numberOfValues)
   {
-    this->ArrayManager.AllocateArrayForOutput(storage, numberOfValues);
+    return this->ArrayManager.PrepareForOutput(numberOfValues);
   }
 
   /// Allocates data in the given Storage and copies data held in the execution
-  /// environment (managed by this class) into the control array. If control
-  /// and execution share arrays, this can be no operation. This method should
-  /// only be called after AllocateArrayForOutput is called.
+  /// environment (managed by this class) into the storage object. The
+  /// reference to the storage given is the same as that passed to the
+  /// constructor. If control and execution share arrays, this can be no
+  /// operation. This method should only be called after PrepareForOutput is
+  /// called.
   ///
-  VTKM_CONT_EXPORT void RetrieveOutputData(StorageType &storage) const
+  VTKM_CONT_EXPORT
+  void RetrieveOutputData(StorageType &storage) const
   {
     this->ArrayManager.RetrieveOutputData(storage);
   }
@@ -127,32 +138,17 @@ public:
   /// (returned from GetNumberOfValues). That is, this method can only be used
   /// to shorten the array, not lengthen.
   ///
-  VTKM_CONT_EXPORT void Shrink(vtkm::Id numberOfValues)
+  VTKM_CONT_EXPORT
+  void Shrink(vtkm::Id numberOfValues)
   {
     this->ArrayManager.Shrink(numberOfValues);
   }
 
-  /// Returns an array portal that can be used in the execution environment.
-  /// This portal was defined in either LoadDataForInput or
-  /// AllocateArrayForOutput. If control and environment share memory space,
-  /// this class may return the iterator from the \c controlArray.
-  ///
-  VTKM_CONT_EXPORT PortalExecution GetPortalExecution()
-  {
-    return this->ArrayManager.GetPortal();
-  }
-
-  /// Const version of GetPortal.
-  ///
-  VTKM_CONT_EXPORT PortalConstExecution GetPortalConstExecution() const
-  {
-    return this->ArrayManager.GetPortalConst();
-  }
-
   /// Frees any resources (i.e. memory) allocated for the exeuction
   /// environment, if any.
   ///
-  VTKM_CONT_EXPORT void ReleaseResources()
+  VTKM_CONT_EXPORT
+  void ReleaseResources()
   {
     this->ArrayManager.ReleaseResources();
   }
diff --git a/vtkm/cont/internal/testing/UnitTestArrayManagerExecutionShareWithControl.cxx b/vtkm/cont/internal/testing/UnitTestArrayManagerExecutionShareWithControl.cxx
index 5bcf7fab3..3e9bb68df 100644
--- a/vtkm/cont/internal/testing/UnitTestArrayManagerExecutionShareWithControl.cxx
+++ b/vtkm/cont/internal/testing/UnitTestArrayManagerExecutionShareWithControl.cxx
@@ -66,26 +66,25 @@ struct TemplatedTests
 
   bool CheckManager(ArrayManagerType &manager, const ValueType &value)
   {
-    return CheckPortal(manager.GetPortalConst(), value);
+    return CheckPortal(manager.PrepareForInput(false), value);
   }
 
   void InputData()
   {
     const ValueType INPUT_VALUE(45);
 
-    StorageType controlArray;
-    controlArray.Allocate(ARRAY_SIZE);
-    SetStorage(controlArray, INPUT_VALUE);
+    StorageType storage;
+    storage.Allocate(ARRAY_SIZE);
+    SetStorage(storage, INPUT_VALUE);
 
-    ArrayManagerType executionArray;
-    executionArray.LoadDataForInput(controlArray);
+    ArrayManagerType executionArray(storage);
 
     // Although the ArrayManagerExecutionShareWithControl class wraps the
     // control array portal in a different array portal, it should still
     // give the same iterator (to avoid any unnecessary indirection).
     VTKM_TEST_ASSERT(
-          vtkm::cont::ArrayPortalToIteratorBegin(controlArray.GetPortalConst()) ==
-          vtkm::cont::ArrayPortalToIteratorBegin(executionArray.GetPortalConst()),
+          vtkm::cont::ArrayPortalToIteratorBegin(storage.GetPortalConst()) ==
+          vtkm::cont::ArrayPortalToIteratorBegin(executionArray.PrepareForInput(true)),
           "Execution array manager not holding control array iterators.");
 
     VTKM_TEST_ASSERT(CheckManager(executionArray, INPUT_VALUE),
@@ -96,23 +95,22 @@ struct TemplatedTests
   {
     const ValueType INPUT_VALUE(50);
 
-    StorageType controlArray;
-    controlArray.Allocate(ARRAY_SIZE);
-    SetStorage(controlArray, INPUT_VALUE);
+    StorageType storage;
+    storage.Allocate(ARRAY_SIZE);
+    SetStorage(storage, INPUT_VALUE);
 
-    ArrayManagerType executionArray;
-    executionArray.LoadDataForInPlace(controlArray);
+    ArrayManagerType executionArray(storage);
 
     // Although the ArrayManagerExecutionShareWithControl class wraps the
     // control array portal in a different array portal, it should still
     // give the same iterator (to avoid any unnecessary indirection).
     VTKM_TEST_ASSERT(
-          vtkm::cont::ArrayPortalToIteratorBegin(controlArray.GetPortal()) ==
-          vtkm::cont::ArrayPortalToIteratorBegin(executionArray.GetPortal()),
+          vtkm::cont::ArrayPortalToIteratorBegin(storage.GetPortal()) ==
+          vtkm::cont::ArrayPortalToIteratorBegin(executionArray.PrepareForInPlace(true)),
           "Execution array manager not holding control array iterators.");
     VTKM_TEST_ASSERT(
-          vtkm::cont::ArrayPortalToIteratorBegin(controlArray.GetPortalConst()) ==
-          vtkm::cont::ArrayPortalToIteratorBegin(executionArray.GetPortalConst()),
+          vtkm::cont::ArrayPortalToIteratorBegin(storage.GetPortalConst()) ==
+          vtkm::cont::ArrayPortalToIteratorBegin(executionArray.PrepareForInput(false)),
           "Execution array manager not holding control array iterators.");
 
     VTKM_TEST_ASSERT(CheckManager(executionArray, INPUT_VALUE),
@@ -123,21 +121,20 @@ struct TemplatedTests
   {
     const ValueType OUTPUT_VALUE(12);
 
-    StorageType controlArray;
+    StorageType storage;
 
-    ArrayManagerType executionArray;
-    executionArray.AllocateArrayForOutput(controlArray, ARRAY_SIZE);
+    ArrayManagerType executionArray(storage);
 
     vtkm::cont::ArrayPortalToIterators<typename ArrayManagerType::PortalType>
-        iterators(executionArray.GetPortal());
+        iterators(executionArray.PrepareForOutput(ARRAY_SIZE));
     std::fill(iterators.GetBegin(), iterators.GetEnd(), OUTPUT_VALUE);
 
     VTKM_TEST_ASSERT(CheckManager(executionArray, OUTPUT_VALUE),
                      "Did not get correct array back.");
 
-    executionArray.RetrieveOutputData(controlArray);
+    executionArray.RetrieveOutputData(storage);
 
-    VTKM_TEST_ASSERT(CheckStorage(controlArray, OUTPUT_VALUE),
+    VTKM_TEST_ASSERT(CheckStorage(storage, OUTPUT_VALUE),
                      "Did not get the right value in the storage.");
   }
 
diff --git a/vtkm/cont/testing/TestingDeviceAdapter.h b/vtkm/cont/testing/TestingDeviceAdapter.h
index b85ef45f4..13d0f6ffc 100644
--- a/vtkm/cont/testing/TestingDeviceAdapter.h
+++ b/vtkm/cont/testing/TestingDeviceAdapter.h
@@ -103,29 +103,6 @@ struct SortGreater
 };
 }
 
-//in the namespace vtkm::cont::testing so device adapters
-//that don't use shared memory space can specialize this
-template<typename DeviceTagType>
-struct CopyInto
-{
-  template<typename T, typename StorageTagType>
-  VTKM_CONT_EXPORT
-  void operator()( vtkm::cont::internal::ArrayManagerExecution<
-                    T,
-                    StorageTagType,
-                    DeviceTagType>& manager,
-                 T* start)
-  {
-    typedef vtkm::cont::internal::ArrayManagerExecution< T,
-            StorageTagType, DeviceTagType> ArrayManagerExecution;
-
-    vtkm::cont::ArrayPortalToIterators<
-        typename ArrayManagerExecution::PortalConstType>
-      iterators(manager.GetPortalConst());
-     std::copy(iterators.GetBegin(), iterators.GetEnd(), start);
-  }
-};
-
 
 
 #define ERROR_MESSAGE "Got an error."
@@ -380,43 +357,57 @@ private:
     std::cout << "Testing ArrayManagerExecution" << std::endl;
 
     typedef vtkm::cont::internal::ArrayManagerExecution<
-        vtkm::FloatDefault,StorageTagBasic,DeviceAdapterTag>
+        vtkm::Id,StorageTagBasic,DeviceAdapterTag>
         ArrayManagerExecution;
 
-    typedef vtkm::cont::internal::Storage<vtkm::FloatDefault,StorageTagBasic>
-        StorageType;
+    typedef vtkm::cont::internal::Storage<vtkm::Id,StorageTagBasic> StorageType;
 
     // Create original input array.
-    vtkm::FloatDefault inputArray[ARRAY_SIZE*2];
-    for (vtkm::Id index = 0; index < ARRAY_SIZE*2; index++)
-    {
-      inputArray[index] = vtkm::FloatDefault(index);
-    }
+    StorageType storage;
+    storage.Allocate(ARRAY_SIZE*2);
 
-    StorageType storage(inputArray, ARRAY_SIZE*2);
+    StorageType::PortalType portal = storage.GetPortal();
+    VTKM_TEST_ASSERT(portal.GetNumberOfValues() == ARRAY_SIZE*2,
+                     "Storage portal has unexpected size.");
 
-    ArrayManagerExecution inputManager;
-    inputManager.LoadDataForInput(storage);
-
-    // Change size.
-    inputManager.Shrink(ARRAY_SIZE);
-
-    // Copy array back. The issue is we need to know if we are accessing
-    // an array manger that shares memory with the control side. If so
-    // it doesn't support RetrieveOutputData.
-    // the naive way is to use ArrayPortalToIteratorBegin but that fails
-    // since it only works with portals from arrayhandles, as the
-    // arrayhandle does all the syncing.
-    //The solution is to a class that the cuda device adapter can specialize
-    //that handles copying back into control space
-    vtkm::FloatDefault outputArray[ARRAY_SIZE];
-    CopyInto<DeviceAdapterTag>()(inputManager, outputArray);
-
-    // Check array.
     for (vtkm::Id index = 0; index < ARRAY_SIZE; index++)
     {
-      VTKM_TEST_ASSERT(outputArray[index] == index,
-                       "Did not get correct values from array.");
+      portal.Set(index, TestValue(index, vtkm::Id()));
+    }
+
+    ArrayManagerExecution manager(storage);
+
+    // Do an operation just so we know the values are placed in the execution
+    // environment and they change. We are only calling on half the array
+    // because we are about to shrink.
+    Algorithm::Schedule(AddArrayKernel(manager.PrepareForInPlace(true)),
+                        ARRAY_SIZE);
+
+    // Change size.
+    manager.Shrink(ARRAY_SIZE);
+
+    VTKM_TEST_ASSERT(manager.GetNumberOfValues() == ARRAY_SIZE,
+                     "Shrink did not set size of array manager correctly.");
+
+    // Get the array back and check its values. We have to get it back into
+    // the same storage since some ArrayManagerExecution classes will expect
+    // that.
+    manager.RetrieveOutputData(storage);
+
+    VTKM_TEST_ASSERT(storage.GetNumberOfValues() == ARRAY_SIZE,
+                     "Storage has wrong number of values after execution "
+                     "array shrink.");
+
+    // Check array.
+    StorageType::PortalConstType checkPortal = storage.GetPortalConst();
+    VTKM_TEST_ASSERT(checkPortal.GetNumberOfValues() == ARRAY_SIZE,
+                     "Storage portal wrong size.");
+
+    for (vtkm::Id index = 0; index < ARRAY_SIZE; index++)
+    {
+      VTKM_TEST_ASSERT(
+            checkPortal.Get(index) == TestValue(index, vtkm::Id()) + index,
+            "Did not get correct values from array.");
     }
   }
 
@@ -485,14 +476,13 @@ private:
     {
       std::cout << "Allocating execution array" << std::endl;
       IdStorage storage;
-      IdArrayManagerExecution manager;
-      manager.AllocateArrayForOutput(storage, 1);
+      IdArrayManagerExecution manager(storage);
 
       std::cout << "Running clear." << std::endl;
-      Algorithm::Schedule(ClearArrayKernel(manager.GetPortal()), 1);
+      Algorithm::Schedule(ClearArrayKernel(manager.PrepareForOutput(1)), 1);
 
       std::cout << "Running add." << std::endl;
-      Algorithm::Schedule(AddArrayKernel(manager.GetPortal()), 1);
+      Algorithm::Schedule(AddArrayKernel(manager.PrepareForInPlace(false)), 1);
 
       std::cout << "Checking results." << std::endl;
       manager.RetrieveOutputData(storage);
@@ -511,14 +501,15 @@ private:
     {
       std::cout << "Allocating execution array" << std::endl;
       IdStorage storage;
-      IdArrayManagerExecution manager;
-      manager.AllocateArrayForOutput(storage, ARRAY_SIZE);
+      IdArrayManagerExecution manager(storage);
 
       std::cout << "Running clear." << std::endl;
-      Algorithm::Schedule(ClearArrayKernel(manager.GetPortal()), ARRAY_SIZE);
+      Algorithm::Schedule(ClearArrayKernel(manager.PrepareForOutput(ARRAY_SIZE)),
+                          ARRAY_SIZE);
 
       std::cout << "Running add." << std::endl;
-      Algorithm::Schedule(AddArrayKernel(manager.GetPortal()), ARRAY_SIZE);
+      Algorithm::Schedule(AddArrayKernel(manager.PrepareForInPlace(false)),
+                          ARRAY_SIZE);
 
       std::cout << "Checking results." << std::endl;
       manager.RetrieveOutputData(storage);
@@ -538,17 +529,19 @@ private:
     {
       std::cout << "Allocating execution array" << std::endl;
       IdStorage storage;
-      IdArrayManagerExecution manager;
+      IdArrayManagerExecution manager(storage);
       vtkm::Id DIM_SIZE = vtkm::Id(std::pow(ARRAY_SIZE, 1/3.0f));
-      manager.AllocateArrayForOutput(storage,
-                                     DIM_SIZE * DIM_SIZE * DIM_SIZE);
       vtkm::Id3 maxRange(DIM_SIZE);
 
       std::cout << "Running clear." << std::endl;
-      Algorithm::Schedule(ClearArrayKernel(manager.GetPortal()), maxRange);
+      Algorithm::Schedule(
+            ClearArrayKernel(manager.PrepareForOutput(
+                               DIM_SIZE * DIM_SIZE * DIM_SIZE)),
+            maxRange);
 
       std::cout << "Running add." << std::endl;
-      Algorithm::Schedule(AddArrayKernel(manager.GetPortal()), maxRange);
+      Algorithm::Schedule(AddArrayKernel(manager.PrepareForInPlace(false)),
+                          maxRange);
 
       std::cout << "Checking results." << std::endl;
       manager.RetrieveOutputData(storage);
diff --git a/vtkm/cont/testing/UnitTestDeviceAdapterAlgorithmGeneral.cxx b/vtkm/cont/testing/UnitTestDeviceAdapterAlgorithmGeneral.cxx
index db2267ea7..74644b87d 100644
--- a/vtkm/cont/testing/UnitTestDeviceAdapterAlgorithmGeneral.cxx
+++ b/vtkm/cont/testing/UnitTestDeviceAdapterAlgorithmGeneral.cxx
@@ -90,6 +90,9 @@ public:
   typedef typename Superclass::ValueType ValueType;
   typedef typename Superclass::PortalType PortalType;
   typedef typename Superclass::PortalConstType PortalConstType;
+
+  ArrayManagerExecution(vtkm::cont::internal::Storage<T, StorageTag> &storage)
+    : Superclass(storage) {  }
 };
 
 
diff --git a/vtkm/exec/CMakeLists.txt b/vtkm/exec/CMakeLists.txt
index b0ad1586e..f86d67bb7 100644
--- a/vtkm/exec/CMakeLists.txt
+++ b/vtkm/exec/CMakeLists.txt
@@ -30,5 +30,9 @@ add_subdirectory(arg)
 vtkm_declare_headers(${impl_headers} ${headers})
 
 
+#-----------------------------------------------------------------------------
+add_subdirectory(cuda)
+
+
 #-----------------------------------------------------------------------------
 # add_subdirectory(testing)
diff --git a/vtkm/exec/cuda/CMakeLists.txt b/vtkm/exec/cuda/CMakeLists.txt
new file mode 100644
index 000000000..61cb609a7
--- /dev/null
+++ b/vtkm/exec/cuda/CMakeLists.txt
@@ -0,0 +1,31 @@
+##============================================================================
+##  Copyright (c) Kitware, Inc.
+##  All rights reserved.
+##  See LICENSE.txt for details.
+##  This software is distributed WITHOUT ANY WARRANTY; without even
+##  the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
+##  PURPOSE.  See the above copyright notice for more information.
+##
+##  Copyright 2014 Sandia Corporation.
+##  Copyright 2014 UT-Battelle, LLC.
+##  Copyright 2014. Los Alamos National Security
+##
+##  Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
+##  the U.S. Government retains certain rights in this software.
+##
+##  Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
+##  Laboratory (LANL), the U.S. Government retains certain rights in
+##  this software.
+##============================================================================
+
+set(headers
+  )
+
+#-----------------------------------------------------------------------------
+add_subdirectory(internal)
+
+# vtkm_declare_headers(${impl_headers} ${headers})
+
+
+#-----------------------------------------------------------------------------
+# add_subdirectory(testing)
diff --git a/vtkm/exec/cuda/internal/ArrayPortalFromThrust.h b/vtkm/exec/cuda/internal/ArrayPortalFromThrust.h
index 9c3d9ffc5..196e0a678 100644
--- a/vtkm/exec/cuda/internal/ArrayPortalFromThrust.h
+++ b/vtkm/exec/cuda/internal/ArrayPortalFromThrust.h
@@ -43,6 +43,7 @@
 #endif // gcc version >= 4.6
 #endif // gcc && !CUDA
 
+#include <boost/type_traits/integral_constant.hpp>
 #include <boost/utility/enable_if.hpp>
 
 namespace vtkm {
@@ -163,7 +164,7 @@ public:
   VTKM_CONT_EXPORT
   IteratorType GetIteratorEnd() const { return this->EndIterator.get(); }
 
-private:
+//private:
   PointerType BeginIterator;
   PointerType EndIterator;
 
diff --git a/vtkm/exec/cuda/internal/CMakeLists.txt b/vtkm/exec/cuda/internal/CMakeLists.txt
new file mode 100644
index 000000000..dc90f88b7
--- /dev/null
+++ b/vtkm/exec/cuda/internal/CMakeLists.txt
@@ -0,0 +1,34 @@
+##============================================================================
+##  Copyright (c) Kitware, Inc.
+##  All rights reserved.
+##  See LICENSE.txt for details.
+##  This software is distributed WITHOUT ANY WARRANTY; without even
+##  the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
+##  PURPOSE.  See the above copyright notice for more information.
+##
+##  Copyright 2014 Sandia Corporation.
+##  Copyright 2014 UT-Battelle, LLC.
+##  Copyright 2014. Los Alamos National Security
+##
+##  Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
+##  the U.S. Government retains certain rights in this software.
+##
+##  Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
+##  Laboratory (LANL), the U.S. Government retains certain rights in
+##  this software.
+##============================================================================
+
+set(headers
+  ArrayPortalFromThrust.h
+  )
+
+vtkm_disable_troublesome_thrust_warnings()
+
+#-----------------------------------------------------------------------------
+CUDA_INCLUDE_DIRECTORIES(${Boost_INCLUDE_DIRS})
+
+vtkm_declare_headers(CUDA ${headers})
+
+
+#-----------------------------------------------------------------------------
+# add_subdirectory(testing)