   }
 
   bool sb = true;
-  for (unsigned i = next_level; 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
-                                            ctx->level_sizes[i],                           // Number of uvec4's in level i
-                                            ctx->index_sidebands[sb ? 0 : 1],              // Input, each block will read 5*32=160 values from here
-                                            ctx->level_sizes[i - 1]);                      // Number of sideband elements from level i-1
+
+  // Do double reductions as long as it goes...
+  for (; nl + 1 < ctx->levels - 3; nl += 2) {
+    const auto blocks = (ctx->level_sizes[nl + 1] + 31) / 32;
+    reduceDouble<<<blocks, 5 * 32, 0, stream>>>(ctx->index_pyramid + ctx->level_offsets[nl],
+                                                ctx->index_pyramid + ctx->level_offsets[nl + 1],
+                                                ctx->index_sidebands[sb ? 1 : 0],
+                                                ctx->level_sizes[nl],
+                                                ctx->level_sizes[nl + 1],
+                                                ctx->index_sidebands[sb ? 0 : 1],
+                                                ctx->level_sizes[nl - 1]);
+    if (ctx->indexed) {
+      reduceDouble<<<blocks, 5 * 32, 0, ctx->indexStream>>>(ctx->vertex_pyramid + ctx->level_offsets[nl],
+                                                            ctx->vertex_pyramid + ctx->level_offsets[nl + 1],
+                                                            ctx->vertex_sidebands[sb ? 1 : 0],
+                                                            ctx->level_sizes[nl],
+                                                            ctx->level_sizes[nl + 1],
+                                                            ctx->vertex_sidebands[sb ? 0 : 1],
+                                                            ctx->level_sizes[nl - 1]);
+    }
+    sb = !sb;
+  }
+
+  // Do a single reduction if one remaining.
+  for (; nl < ctx->levels - 3; nl++) {
+    const auto blocks = (ctx->level_sizes[nl] + 31) / 32;
+    reduceSingle<<<blocks, 5 * 32, 0, stream>>>(ctx->index_pyramid + ctx->level_offsets[nl],
+                                                ctx->index_sidebands[sb ? 1 : 0],
+                                                ctx->level_sizes[nl],
+                                                ctx->index_sidebands[sb ? 0 : 1],
+                                                ctx->level_sizes[nl - 1]);
     if(ctx->indexed) {
-      reduce1<<<blocks, 5 * 32, 0, ctx->indexStream>>>(ctx->vertex_pyramid + ctx->level_offsets[i], // Each block will write 32 uvec4's into this
-                                                        ctx->vertex_sidebands[sb ? 1 : 0],           // Each block will write 32 uint32's into this
-                                                        ctx->level_sizes[i],                         // Number of uvec4's in level i
-                                                        ctx->vertex_sidebands[sb ? 0 : 1],           // Input, each block will read 5*32=160 values from here
-                                                        ctx->level_sizes[i - 1]);                    // Number of sideband elements from level i-1
+      reduceSingle<<<blocks, 5 * 32, 0, ctx->indexStream>>>(ctx->vertex_pyramid + ctx->level_offsets[nl],
+                                                            ctx->vertex_sidebands[sb ? 1 : 0],
+                                                            ctx->level_sizes[nl],
+                                                            ctx->vertex_sidebands[sb ? 0 : 1],
+                                                            ctx->level_sizes[nl - 1]);
     }
     sb = !sb;
   }


  }



  // 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::DoubleLevelMultiWarpChunk: {
    uint32_t bn = ctx->chunk_total;
    uint32_t tn = 32 * 6;
    nl = 2;
    if (ctx->indexed) {
      reduceBaseDoubleMultiWarp<float, true><<<bn, tn, 0, stream>>>(ctx->index_cases_d,
                                                                    (uchar4*)(ctx->vertex_pyramid + ctx->level_offsets[0]),
