                          + cell.z * field_slice_stride
                          + cell.y * field_row_stride
                          + cell.x;
-    uint6 prev = fetch(ptr, field_end_d, field_row_stride, threshold, thread);
+    uint2 prev = fetch(ptr, field_end_d, field_row_stride, threshold, thread);
     ptr += field_slice_stride;
 
     bool xmask = cell.x < cells.x;

      unsigned sum = 0;
      bool zmask = cell.z < cells.z;
      if (zmask) {
        uint6 next = fetch(ptr, field_end_d, field_row_stride, threshold, thread);
        ptr += field_slice_stride;
        if (xmask && thread < 31) {


          uint5 cases = mergeY(mergeZ(prev, next));
          prev = next;

          uint5 counts{
            cell.y + 0u < cells.y ? index_count[cases.e0] : 0u,

    uint32_t e4;
  };


  template<typename T>
  __device__ uint2 fetch(const T* ptr, const T* end, const size_t stride, const float t, unsigned thread)
  {
    auto qtr = ptr;
    unsigned t0 = ((ptr < end) && (ptr[0] < t) ? 1 : 0);
    ptr += stride;
    t0 |= ((ptr < end) && (ptr[0] < t) ? (1<<8) : 0);
    ptr += stride;
    t0 |= ((ptr < end) && (ptr[0] < t) ? (1<<16) : 0);
    ptr += stride;
    t0 |= ((ptr < end) && (ptr[0] < t) ? (1<<24) : 0);
    ptr += stride;

    unsigned t1 = (t0 >> 24) | ((ptr < end) && (ptr[0] < t) ? (1<<8) : 0);
    ptr += stride;
    t1 |= ((ptr < end) && (ptr[0] < t) ? (1<<16) : 0);
    ptr += stride;

    unsigned r0 = t0;// (t0 << 0) | (t1 << 8) | (t2 << 16) | (t3 << 24);
    unsigned r1 = t1;// (t3 << 0) | (t4 << 8) | (t5 << 16);

    return make_uint2( r0, r1 );
  }

  __device__ uint2 mergeZ(const uint2 z0, const uint2 z1)
  {
    return make_uint2((z1.x << 4) | (z0.x),
                      (z1.y << 4) | (z0.y));
  }

  
  template<class FieldType>
  __global__ __launch_bounds__(32) void reduce_base(uint8_t* __restrict__           index_cases_d,
                                                    uint4*   __restrict__           out_level0_d,
