     if (10 < L) offsets[10] = offsetBlob.dataA;
     if (11 < L) offsets[11] = offsetBlob.dataB;
 
-    // Traverse apex.
-    uint32_t offset = 0;
-    uint32_t key = index;
-    offset = processHistoElement(key, 5 * offset, hp_d[1]);
-    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 N = hp_d[0].x;
+    uint4 T = hp_d[1];
+
+    for (uint32_t k = HP5_EXTRACT_BLOCKSIZE * blockIdx.x; k < N; k += HP5_EXTRACT_BLOCKS * HP5_EXTRACT_BLOCKSIZE) {
+      const uint32_t index = k + threadIdx.x;
+      if (index < N) {
+        uint32_t offset = 0;
+        uint32_t key = index;
+        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;
+      }
+    }
   }
 
   void scratchLayout(std::vector<uint32_t>& levels, std::vector<uint32_t>& offsets, uint32_t N)

               const uint4* __restrict__ hp_d,
               OffsetBlob offsetBlob)
  {
    const uint32_t index = 128 * blockIdx.x + threadIdx.x;
    if (hp_d[0].x <= index) return;

    __shared__ uint32_t offsets[12];
    if (0 < L) offsets[0] = offsetBlob.data0;


#define HP5_WARP_COUNT 5
#define HP5_WARP_SIZE 32

#define HP5_EXTRACT_BLOCKSIZE 128
#define HP5_EXTRACT_BLOCKS (22*16)

namespace {

