diff --git a/src/libs/ascent/CMakeLists.txt b/src/libs/ascent/CMakeLists.txt
index 9066e3e4b..e1e1fa0fd 100644
--- a/src/libs/ascent/CMakeLists.txt
+++ b/src/libs/ascent/CMakeLists.txt
@@ -264,7 +264,6 @@ if(VTKM_FOUND)
     set(ascent_vtkh_dep_headers
         runtimes/ascent_data_object.hpp
         runtimes/ascent_vtkh_data_adapter.hpp
-        runtimes/ascent_vtkh_device_utils.hpp
         runtimes/ascent_vtkh_collection.hpp
         runtimes/flow_filters/ascent_runtime_vtkh_filters.hpp
         runtimes/flow_filters/ascent_runtime_vtkh_utils.hpp
diff --git a/src/libs/ascent/runtimes/ascent_vtkh_data_adapter.cpp b/src/libs/ascent/runtimes/ascent_vtkh_data_adapter.cpp
index e1926199c..f7fde18f7 100644
--- a/src/libs/ascent/runtimes/ascent_vtkh_data_adapter.cpp
+++ b/src/libs/ascent/runtimes/ascent_vtkh_data_adapter.cpp
@@ -11,7 +11,6 @@
 ///
 //-----------------------------------------------------------------------------
 #include "ascent_vtkh_data_adapter.hpp"
-#include "ascent_vtkh_device_utils.hpp"
 
 // standard lib includes
 #include <iostream>
@@ -33,8 +32,10 @@
 #include <vtkm/cont/DataSet.h>
 #include <vtkm/cont/Algorithm.h>
 #include <vtkm/cont/ArrayCopy.h>
+#include <vtkm/cont/ArrayHandle.h>
 #include <vtkm/cont/ArrayHandleExtractComponent.h>
 #include <vtkm/cont/CoordinateSystem.h>
+#include <vtkm/worklet/DispatcherMapField.h>
 #include <vtkh/DataSet.hpp>
 
 // other ascent includes
@@ -1674,14 +1675,15 @@ std::cerr << "AF 11" << std::endl;
             Node n_tmp;
 std::cerr << "AF 12" << std::endl;
             n_tmp.set_external(DataType::float64(num_vals),ptr);
-           const uint64_t *input = n_vals.value();
-           double output2[num_vals];
+           const unsigned long long *input = n_vals.value();
+           vtkm::cont::ArrayHandle<unsigned long long> input_arr = vtkm::cont::make_ArrayHandle(vtkm_arr,input_arr);
 std::cerr<<  " before calling cast" <<std::endl;
-           castUint64ToFloat64(input, output2, num_vals);
+            vtkm::worklet::DispatcherMapField<vtkh::VTKMTypeCast>().Invoke(input_arr,vtkm_arr);
+          // VTKHDeviceAdapter::castUint64ToFloat64(input, output2, num_vals);
 std::cerr<<  " after calling cast" <<std::endl;
 
 std::cerr << "AF 13" << std::endl;
-            n_vals.to_float64_array(n_tmp);
+//            n_vals.to_float64_array(n_tmp);
 
 std::cerr << "AF 14" << std::endl;
             // add field to dataset
diff --git a/src/libs/ascent/runtimes/ascent_vtkh_device_utils.cpp b/src/libs/ascent/runtimes/ascent_vtkh_device_utils.cpp
deleted file mode 100644
index df129c881..000000000
--- a/src/libs/ascent/runtimes/ascent_vtkh_device_utils.cpp
+++ /dev/null
@@ -1,68 +0,0 @@
-//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~//
-// Copyright (c) Lawrence Livermore National Security, LLC and other Ascent
-// Project developers. See top-level LICENSE AND COPYRIGHT files for dates and
-// other details. No copyright assignment is required to contribute to Ascent.
-//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~//
-
-
-//-----------------------------------------------------------------------------
-///
-/// file: ascent_vtkh_device_utils.cpp
-///
-//-----------------------------------------------------------------------------
-#include "ascent_vtkh_device_utils.hpp"
-#include "ascent/runtimes/expressions/ascent_array.hpp"
-
-// standard lib includes
-#include <iostream>
-#include <string.h>
-#include <limits.h>
-#include <cstdlib>
-#include <sstream>
-#include <type_traits>
-
-
-// VTKm includes
-#define VTKM_USE_DOUBLE_PRECISION
-
-
-using namespace std;
-using namespace conduit;
-
-//-----------------------------------------------------------------------------
-// -- begin ascent:: --
-//-----------------------------------------------------------------------------
-namespace ascent
-{
-
-//-----------------------------------------------------------------------------
-// -- begin detail:: --
-//-----------------------------------------------------------------------------
-namespace detail
-{
-
-};
-//-----------------------------------------------------------------------------
-// -- end detail:: --
-//-----------------------------------------------------------------------------
-
-//-----------------------------------------------------------------------------
-// VTKHDataAdapter device methods
-//-----------------------------------------------------------------------------
-
-// Definition of the cast function with __host__ __device__
-ASCENT_EXEC void 
-VTKHDeviceAdapter::castUint64ToFloat64(const uint64_t* input, double* output, int size) {
-    Array<double> int_to_double(input, size);
-    output = int_to_double.get_ptr(memory_space);
-
-    // init device array
-    ascent::forall<for_policy>(0, size, [=] ASCENT_LAMBDA(index_t i)
-    {
-        output[i] = static_cast<double>(input[i]);
-    });
-    ASCENT_DEVICE_ERROR_CHECK();
-};
-//-----------------------------------------------------------------------------
-// -- end ascent:: --
-//-----------------------------------------------------------------------------
diff --git a/src/libs/ascent/runtimes/ascent_vtkh_device_utils.hpp b/src/libs/ascent/runtimes/ascent_vtkh_device_utils.hpp
deleted file mode 100644
index 1785f0c50..000000000
--- a/src/libs/ascent/runtimes/ascent_vtkh_device_utils.hpp
+++ /dev/null
@@ -1,185 +0,0 @@
-//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~//
-// Copyright (c) Lawrence Livermore National Security, LLC and other Ascent
-// Project developers. See top-level LICENSE AND COPYRIGHT files for dates and
-// other details. No copyright assignment is required to contribute to Ascent.
-//~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~//
-
-#ifndef ASCENT_VTKH_DEVICE_UTILS_HPP
-#define ASCENT_VTKH_DEVICE_UTILS_HPP
-
-//-----------------------------------------------------------------------------
-///
-/// file: ascent_vtkh_device_utils.hpp
-///
-//-----------------------------------------------------------------------------
-#if defined(ASCENT_RAJA_ENABLED)
-#include <RAJA/RAJA.hpp>
-#endif
-
-#include <vtkm/cont/DataSet.h>
-
-#include <ascent_exports.h>
-#include <conduit.hpp>
-
-//-----------------------------------------------------------------------------
-// -- begin ascent:: --
-//-----------------------------------------------------------------------------
-namespace ascent
-{
-
-using index_t = conduit::index_t;
-
-struct EmptyPolicy
-{};
-
-#if defined(ASCENT_CUDA_ENABLED)
-//---------------------------------------------------------------------------//
-// CUDA decorators
-//---------------------------------------------------------------------------//
-#define ASCENT_EXEC inline __host__ __device__
-// Note: there is a performance hit for doing both host and device
-// the cuda compiler calls this on then host as a std::function call for each i
-// in the for loop, and that basically works out to a virtual function
-// call. Thus for small loops, the know overhead is about 3x
-#define ASCENT_LAMBDA __device__ __host__
-#if defined(ASCENT_RAJA_ENABLED)
-#define BLOCK_SIZE 128
-using for_policy = RAJA::cuda_exec<BLOCK_SIZE>;
-using reduce_policy = RAJA::cuda_reduce;
-using atomic_policy = RAJA::cuda_atomic;
-#else
-using for_policy    = EmptyPolicy;
-using reduce_policy = EmptyPolicy;
-using atomic_policy = EmptyPolicy;
-#endif
-
-#elif defined(ASCENT_HIP_ENABLED) // && ?
-//---------------------------------------------------------------------------//
-// HIP decorators
-//---------------------------------------------------------------------------//
-#define ASCENT_EXEC inline __host__ __device__
-#define ASCENT_LAMBDA __device__ __host__
-#if defined(ASCENT_RAJA_ENABLED)
-#define BLOCK_SIZE 256
-using for_policy = RAJA::hip_exec<BLOCK_SIZE>;
-using reduce_policy = RAJA::hip_reduce;
-using atomic_policy = RAJA::hip_atomic;
-#else
-using for_policy    = EmptyPolicy;
-using reduce_policy = EmptyPolicy;
-using atomic_policy = EmptyPolicy;
-#endif
-
-#else
-//---------------------------------------------------------------------------//
-// Non-device decorators
-//---------------------------------------------------------------------------//
-#define ASCENT_EXEC   inline
-#define ASCENT_LAMBDA
-#if defined(ASCENT_RAJA_ENABLED)
-using for_policy = RAJA::seq_exec;
-using reduce_policy = RAJA::seq_reduce;
-using atomic_policy = RAJA::seq_atomic;
-#else
-using for_policy    = EmptyPolicy;
-#endif
-#endif
-
-
-
-//---------------------------------------------------------------------------//
-// Device Error Checks
-//---------------------------------------------------------------------------//
-#if defined(ASCENT_CUDA_ENABLED)
-//---------------------------------------------------------------------------//
-// cuda error check
-//---------------------------------------------------------------------------//
-inline void cuda_error_check(const char *file, const int line )
-{
-  cudaError err = cudaGetLastError();
-  if ( cudaSuccess != err )
-  {
-    std::cerr<<"CUDA error reported at: "<<file<<":"<<line;
-    std::cerr<<" : "<<cudaGetErrorString(err)<<"\n";
-    //exit( -1 );
-  }
-}
-#define ASCENT_DEVICE_ERROR_CHECK() cuda_error_check(__FILE__,__LINE__);
-
-#elif defined(ASCENT_HIP_ENABLED)
-//---------------------------------------------------------------------------//
-// hip error check
-//---------------------------------------------------------------------------//
-inline void hip_error_check(const char *file, const int line )
-{
-  hipError_t err = hipGetLastError();
-  if ( hipSuccess != err )
-  {
-    std::cerr<<"HIP error reported at: "<<file<<":"<<line;
-    std::cerr<<" : "<<hipGetErrorName(err)<<"\n";
-    //exit( -1 );
-  }
-}
-#define ASCENT_DEVICE_ERROR_CHECK() hip_error_check(__FILE__,__LINE__);
-#else
-//---------------------------------------------------------------------------//
-// non-device error check (no op)
-//---------------------------------------------------------------------------//
-#define ASCENT_DEVICE_ERROR_CHECK()
-#endif
-
-//---------------------------------------------------------------------------//
-// forall
-//---------------------------------------------------------------------------//
-#if defined(ASCENT_RAJA_ENABLED)
-template <typename ExecPolicy, typename Kernel>
-inline void forall(const index_t& begin,
-                   const index_t& end,
-                   Kernel&& kernel) noexcept
-{
-    RAJA::forall<ExecPolicy>(RAJA::RangeSegment(begin, end),
-                             std::forward<Kernel>(kernel));
-}
-#else
-template <typename ExecPolicy, typename Kernel>
-inline void forall(const index_t& begin,
-                   const index_t& end,
-                   Kernel&& kernel) noexcept
-{
-    for(index_t i = begin; i < end; ++i)
-    {
-        kernel(i);
-    };
-}
-#endif
-
-//-----------------------------------------------------------------------------
-// -- start VTKHDeviceAdapter --
-//-----------------------------------------------------------------------------
-class ASCENT_API VTKHDeviceAdapter
-{
-
-public: 
-  // Definition of the cast function with __host__ __device__
-  template <typename T, typename S> ASCENT_EXEC static void castUint64ToFloat64(const T* input, S* output, int size)
-  {
-      forall<for_policy>(0, size, [=] ASCENT_LAMBDA(index_t i)
-      {
-          output[i] = static_cast<S>(input[i]);
-      });
-      ASCENT_DEVICE_ERROR_CHECK();
-  }
-
-//-----------------------------------------------------------------------------
-// -- end VTKHDeviceAdapter --
-//-----------------------------------------------------------------------------
-};
-
-//-----------------------------------------------------------------------------
-// -- end ascent:: --
-//-----------------------------------------------------------------------------
-};
-//-----------------------------------------------------------------------------
-// -- end header --
-//-----------------------------------------------------------------------------
-#endif
diff --git a/src/libs/vtkh/utils/vtkm_array_utils.hpp b/src/libs/vtkh/utils/vtkm_array_utils.hpp
index dae3d6226..0d3f04ab0 100644
--- a/src/libs/vtkh/utils/vtkm_array_utils.hpp
+++ b/src/libs/vtkh/utils/vtkm_array_utils.hpp
@@ -12,5 +12,19 @@ GetVTKMPointer(vtkm::cont::ArrayHandle<T> &handle)
   return handle.WritePortal().GetArray();
 }
 
+struct VTKmTypeCast : public vtkm::worklet::WorkletMapField
+{
+    using ControlSignature = void(FieldIn input, FieldOut output);
+    using ExecutionSignature = void(_1, _2);
+
+    // Use VTKM_EXEC for the operator() function to make it run on both host and device
+    template<typename InType, typename OutType>
+    VTKM_EXEC void operator()(const InType& input, OutType& output) const
+    {
+        // Cast input to the output type and assign it
+        output = static_cast<OutType>(input);
+    }
+};
+
 }//namespace vtkh
 #endif