 #define HP5_EXTRACT_BLOCKS (22*16)
 
 namespace {
+
+
+  __global__ __launch_bounds__(128) void reduceBase(uint32_t* __restrict__ hp_d,
+                                                    uint32_t* __restrict__ sb_d,
+                                                    const uint32_t* __restrict__ src,
+                                                    uint32_t N)
+  {
+    const uint32_t offsetIn = blockDim.x * blockIdx.x + threadIdx.x;
+    const uint32_t lane = threadIdx.x %  HP5_WARP_SIZE;
+    const uint32_t value = offsetIn < N ? src[offsetIn] : 0;
+    const uint32_t warpMask = __ballot(value != 0);
+    if (lane == 0) {
+      const uint32_t offsetOut = offsetIn / HP5_WARP_SIZE;
+      hp_d[offsetOut] = warpMask;
+      sb_d[offsetOut] = __popc(warpMask);
+    }
+  }
 
   // Reads 160 values, outputs HP level of 128 values, and 32 sideband values.
 

        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;
      }
    }
  }

  __launch_bounds__(128)
  void extract(uint32_t* __restrict__ out_d,
               const uint4* __restrict__ hp_d,
               const OffsetBlob offsetBlob)
  {

    __shared__ uint32_t offsets[12];
