     template<typename REDOP, int N, typename T, bool EXCLUSIVE>
     __global__ void
     __launch_bounds__(THREADS_PER_BLOCK,MIN_BLOCKS_PER_SM)
-    reduction_kernel(const Realm::AffineAccessor<typename REDOP::LHS,N,T> dst,
-                     const Realm::AffineAccessor<typename REDOP::RHS,N,T> src,
-                     const DimOrder<N> order,
-                     const Realm::Point<N,T> lo,
-                     const Realm::Point<N,T> extents,
-                     const size_t max_offset)
+    fold_kernel(const Realm::AffineAccessor<typename REDOP::RHS,N,T> dst,
+                const Realm::AffineAccessor<typename REDOP::RHS,N,T> src,
+                const DimOrder<N> order,
+                const Realm::Point<N,T> lo,
+                const Realm::Point<N,T> extents,
+                const size_t max_offset)
+    {
+      size_t offset = blockIdx.x * blockDim.x + threadIdx.x;
+      if (offset >= max_offset)
+        return;
+      Point<N,T> point = lo;
+      for (int i = 0; i < N; i++)
+      {
+        const int index = order.index[i];
+        point[index] += (offset / extents[index]);
+        offset = offset % extents[index];
+      }
+      REDOP::template fold<EXCLUSIVE>(dst[point], src[point]);
+    }
+
+    template<typename REDOP, int N, typename T, bool EXCLUSIVE>
+    __global__ void
+    __launch_bounds__(THREADS_PER_BLOCK,MIN_BLOCKS_PER_SM)
+    apply_kernel(const Realm::AffineAccessor<typename REDOP::LHS,N,T> dst,
+                 const Realm::AffineAccessor<typename REDOP::RHS,N,T> src,
+                 const DimOrder<N> order,
+                 const Realm::Point<N,T> lo,
+                 const Realm::Point<N,T> extents,
+                 const size_t max_offset)
     {
       size_t offset = blockIdx.x * blockDim.x + threadIdx.x;
       if (offset >= max_offset)

          *((const Realm::RegionInstance*)buffer);
        buffer += sizeof(src);
        const bool exclusive = *((const bool*)buffer);
        
        const Realm::AffineAccessor<typename REDOP::LHS,N,T> 
          src_accessor(src, field_id, space.bounds);
        const Realm::AffineAccessor<typename REDOP::RHS,N,T> 
          dst_accessor(dst, field_id, space.bounds);

        // Compute the order of dimensions to walk based on sorting the strides
        // for the source accessor, we'll optimistically assume the two 
        // instances are laid out the same way
        std::map<ptrdiff_t,int> strides;
        for (int i = 0; i < N; i++)
        {

        offset = offset % extents[index];
      }
      REDOP::template apply<EXCLUSIVE>(dst[point], src[point]);
    } 

    template<typename REDOP>
    struct ReductionRunner {
