 
   auto L = levels.size();
   if (L == 0) {
-    reduceApex<true, true><<<1, 128, 0, stream>>>(reinterpret_cast<uint4*>(scratch_d),
+    assert(false);
+/*    reduceApex<true, true><<<1, 128, 0, stream>>>(reinterpret_cast<uint4*>(scratch_d),
                                                   sum_d,
                                                   in_d,
-                                                  N);
+                                                  N);*/
   }
   else {
-    ::reduce1<true><<<levels[0], 160, 0, stream>>>(reinterpret_cast<uint4*>(scratch_d) + offsets[0],
-                                                   scratch_d + 4 * offsets[L + 1],
-                                                   in_d,
-                                                   N);
+    ::reduceBase<<<levels[0], 128, 0, stream>>>(scratch_d + 4 * offsets[0],
+                                                scratch_d + 4 * offsets[L + 1],
+                                                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<false><<<levels[i], 160, 0, stream>>>(reinterpret_cast<uint4*>(scratch_d) + offsets[i],
                                                       scratch_d + 4*offsets[L + 1 + ((i + 0) & 1)],

        offset = processHistoElement(key, 5 * offset, T);
        offset = processHistoElement(key, 5 * offset, hp_d[2 + offset]);
        offset = processHistoElement(key, 5 * offset, hp_d[7 + offset]);
        for (uint32_t i = L; 0 < i; i--) {
          offset = processDataElement(key, 5 * offset, hp_d[offsets[i - 1] + offset]);
        }
        out_d[index] = offset;
      }
    }
  }

    uint32_t value[L];
  };*/

  __device__ __forceinline__ uint32_t processHistoElement(uint32_t& key, uint32_t offset, const uint4 element)
  {
    assert(element.x <= element.y);
    assert(element.y <= element.z);
