           *((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);
+        buffer += sizeof(exclusive);
+        const bool fold = *((const bool*)buffer);
+
         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
+            src_accessor(src, 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, if we're wrong
+        // it will still be correct, just slow
         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 {

    template<typename REDOP, int N, typename T, bool EXCLUSIVE>
    __global__ void
    __launch_bounds__(THREADS_PER_BLOCK,MIN_BLOCKS_PER_SM)
    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)
