 
     size_t i = 0;
     if (1 < L) {
-      ::reduceBase2<<<(levels[1] + 31) / 32, 8 * 32, 0, stream>>>(scratch_d + offsets[1],
+      ::reduceBase2b<<<(levels[1] + 31) / 32, 8 * 32, 0, stream>>>(scratch_d + offsets[1],
                                                                   scratch_d + offsets[L + 1 + (sb ? 1 : 0)],
                                                                   levels[1],
                                                                   scratch_d + offsets[0],

    }
  }

  __global__ __launch_bounds__(256) void reduceBase2(uint32_t* __restrict__ hp2_d,
                                                     uint32_t* __restrict__ sb2_d,
                                                     const uint32_t n2,

#if 1
      // Prefetch data for next row into L2
      if (i + 1 < 20) {
        asm volatile("prefetch.global.L1 [%0];"::PTR_REG(sb0_d + offset0 + 256));
      }
#endif

