                                                                            make_uint3(ctx->grid_size.x - 1,
                                                                                       ctx->grid_size.y - 1,
                                                                                       ctx->grid_size.z - 1));
+#endif
 
       CHECKED_CUDA(cudaEventRecord(ctx->baseEvent, stream));
       CHECKED_CUDA(cudaStreamWaitEvent(ctx->indexStream, ctx->baseEvent, 0));

      //        case it might be better to just merge the two launches and use e.g. block.y to tell which
      //        pyramid to reduce.
      bool sb = true;
      for (unsigned i = 1; i < ctx->levels - 3; i++) {
        const auto blocks = (ctx->level_sizes[i] + 31) / 32;
        reduce1 << <blocks, 5 * 32, 0, stream >> > (ctx->index_pyramid + ctx->level_offsets[i],    // Each block will write 32 uvec4's into this
                                                    ctx->index_sidebands[sb ? 1 : 0],              // Each block will write 32 uint32's into this

    if (ctx->indexed) {
      // 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.
#if 1
      unsigned i0 = 1;
      ::reduceBaseIndexedMultiWarp<float><<<ctx->chunk_total, 32*6, 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,
                                                                                 size_t(field_size.x),
                                                                                 size_t(field_size.x) * field_size.y,
                                                                                 threshold,
                                                                                 ctx->chunks,
                                                                                 make_uint3(ctx->grid_size.x - 1,
                                                                                            ctx->grid_size.y - 1,
                                                                                            ctx->grid_size.z - 1));
#elif 1
      unsigned i0 = 1;
      ::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
