                                                          levels[0],
                                                          in_d,
                                                          N);
-
-    //::reduce1<true><<<levels[0], 160, 0, stream>>>(reinterpret_cast<uint4*>(scratch_d) + offsets[0],
-    //                                               scratch_d + 4 * offsets[L + 1],
-    //                                               in_d,
-    //                                               N);
     for (size_t i = 1; i < levels.size(); i++) {
       ::reduce1<<<(levels[i] + 31)/32, 160, 0, stream>>>(reinterpret_cast<uint4*>(scratch_d + offsets[i]),
                                                          scratch_d + offsets[L + 1 + ((i + 0) & 1)],

  auto L = levels.size();
  if (L == 0) {
    assert(false);
/*    reduceApex<true, true><<<1, 128, 0, stream>>>(reinterpret_cast<uint4*>(scratch_d),
                                                  sum_d,
                                                  in_d,
                                                  N);*/
  }
  else {
    ::reduceBase<<<(levels[0] + 3)/4, 4*32, 0, stream>>>(scratch_d + offsets[0],

        uint32_t offset = 0;
        uint32_t key = index;
        offset = processHistoElement(key, 5 * offset, T);
        offset = processHistoElement(key, 5 * offset, *(const uint4*)(hp_d + 8 + 4 * offset));
        offset = processHistoElement(key, 5 * offset, *(const uint4*)(hp_d + 28 + 4 * offset));
        for (uint32_t i = L; 1 < i; i--) {
          offset = processDataElement(key, 5 * offset, *(const uint4*)(hp_d + offsets[i - 1] + 4 * offset));
        }
        out_d[index] = processMaskElement(key, 32 * offset, hp_d[offsets[0] + offset]);
      }
