 
   template<class FieldType, bool indexed>
   __device__ void reduceBaseKernel(uint8_t* __restrict__ index_case_ptr,
-                                   uint4* __restrict__ out_index_level0_ptr,
-                                   uint4* __restrict__ out_vertex_level0_ptr,
+                                   uchar4* __restrict__ out_index_level0_ptr,
+                                   uchar4* __restrict__ out_vertex_level0_ptr,
                                    uint32_t* __restrict__ out_index_sideband_ptr,
                                    uint32_t* __restrict__ out_vertex_sideband_ptr,
                                    const uint8_t* __restrict__   index_count,

  ctx->level_offsets[0] = 4 + 32; // First, level offsets (16 uint32's = 4 uvec4's), then pyramid apex (32 uvec4's).
  for (unsigned l = 1; l < ctx->levels - 3; l++) {
    ctx->level_sizes[l] = (ctx->level_sizes[l - 1] + 4) / 5;
    ctx->level_offsets[l] = ctx->level_offsets[l - 1] + ctx->level_sizes[l - 1];
  }
  ctx->total_size = ctx->level_offsets[ctx->levels - 4] + ctx->level_sizes[ctx->levels - 4];


      // FIXME: Try to add extra reduction passes on the tail of this kernel (2 or 3 levels in one go
      //        have worked well before) and see if that reduces memory pressure.
      ::reduceBaseIndexed<float> << <ctx->chunk_total, 32, 0, stream >> > (ctx->index_cases_d,                   // block writes 5*5*32=800 uint8's
                                                                           (uchar4*)(ctx->vertex_pyramid + ctx->level_offsets[0]), // block writes 5*32 uvec4's
                                                                           ctx->vertex_sidebands[0],                    // block writes 5*32 uint32's
                                                                           (uchar4*)(ctx->index_pyramid + ctx->level_offsets[0]), // block writes 5*32 uvec4's
                                                                           ctx->index_sidebands[0],                    // block writes 5*32 uint32's
                                                                           ctx->tables->index_count,
                                                                           field_d,
