                                                                     scratch_d + offsets[L + 1 + (sb ? 0 : 1)],
                                                                     levels[i - 1]);
       sb = !sb;
-
     }
 
     for (; i < L; i++) {

  }


  template<bool mask_input, bool write_sum>
  __global__
  __launch_bounds__(128)

      sb = !sb;
    }

    for (; i + 2 < L; i += 3) {
      ::reduce3 << <(levels[i + 2] + 31) / 32, 160, 0, stream >> > (reinterpret_cast<uint4*>(scratch_d + offsets[i + 2]),
                                                                    scratch_d + offsets[L + 1 + (sb ? 1 : 0)],
                                                                    levels[i + 2],
                                                                    reinterpret_cast<uint4*>(scratch_d + offsets[i+1]),
                                                                    levels[i + 1],
                                                                    reinterpret_cast<uint4*>(scratch_d + offsets[i]),
                                                                    levels[i],
                                                                    scratch_d + offsets[L + 1 + (sb ? 0 : 1)],
                                                                    levels[i - 1]);
      sb = !sb;
    }

    for (; i + 1 < L; i += 2) {
      ::reduce2 << <(levels[i + 1] + 31) / 32, 160, 0, stream >> > (reinterpret_cast<uint4*>(scratch_d + offsets[i + 1]),
                                                                    scratch_d + offsets[L + 1 + (sb ? 1 : 0)],
