                                     ctx->grid_size.y - 1,
                                     ctx->grid_size.z - 1);
   // Indexed pyramid buildup
-  unsigned next_level = ~0u;
+  unsigned nl = ~0u;
   auto build_mode = ctx->build_mode;
   if (build_mode == BaseLevelBuildMode::TripleLevelSingleWarpChunk && ctx->levels == 5) {
     // There are only two levels below the apex


  }



  // Reads 160 values, outputs HP level of 128 values, and 32 sideband values.
  __global__  __launch_bounds__(5 * 32) void reduce1(uint4* __restrict__          hp1_d,  //< Each block will write 32 uvec4's into this
                                                     uint32_t* __restrict__       sb1_d,  //< Each block will write 32 values into this.
                                                     const uint32_t               n1,     //< Number of uvec4's in hp1_d
                                                     const uint32_t* __restrict__ sb0_d,  //< Each block will read 5*32=160 values from here
                                                     const uint32_t               n0)     //< Number of elements in sb0_d
  {
    const uint32_t offset0 = 5 * 32 * blockIdx.x + threadIdx.x;


  case BaseLevelBuildMode::TripleLevelSingleWarpChunk: {
    uint32_t bn = (ctx->level_sizes[2] + 31) / 32;
    uint32_t tn = 32 * 5;
    nl = 3;
    if (ctx->indexed) {
      reduceBaseTriple<float, true><<<bn, tn, 0, stream>>>(ctx->index_cases_d,
                                                            (uchar4*)(ctx->vertex_pyramid + ctx->level_offsets[0]),
