                                                 ctx->sum_d,
                                                 ctx->index_sidebands[sb ? 0 : 1],
                                                 ctx->level_sizes[ctx->levels - 4]);
+      // FIXME: Try to combine these two kernels into one, as they are independent.
       reduceApex << <1, 4 * 32, 0, stream >> > (ctx->vertex_pyramid + 8,
                                                 ctx->sum_d + 1,
                                                 ctx->vertex_sidebands[sb ? 0 : 1],

  {
     // FIXME: no point in up-traversal if cell doesn't change.
     // FIXME: If extract 3 in a go, organize traversals to reduce up-traversal.


    uint32_t ix = blockDim.x * blockIdx.x + threadIdx.x;

  if (buildPyramid) {
    // Indexed pyramid buildup
    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.
      ::reduceBaseIndexed<float> << <ctx->chunk_total, 32, 0, stream >> > (ctx->index_cases_d,                   // block writes 5*5*32=800 uint8's
                                                                           ctx->vertex_pyramid + ctx->level_offsets[0], // block writes 5*32 uvec4's
                                                                           ctx->vertex_sidebands[0],                    // block writes 5*32 uint32's
