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

    offsets[levels.size() + 3] = o; // Final size
  }


  template<uint32_t L>
  std::array<uint32_t, L> arrayFromVector(const std::vector<uint32_t> src)
  {
    std::array<uint32_t, L> rv;
    for (uint32_t i = 0; i < L; i++) rv[i] = src[i];
    return rv;
  }

  template<>
  std::array<uint32_t, 0> arrayFromVector<0>(const std::vector<uint32_t> src)
  {
    std::array<uint32_t, 0> rv = {};
    return rv;
  }


    return key;
  }

  __device__
  inline uint32_t processDataElement(uint32_t& offset, uint32_t key, const uint4 element)
  {
    if (element.x <= key) {
      key -= element.x;
      offset++;
      if (element.y <= key) {
        key -= element.y;
        offset++;
        if (element.z <= key) {
          key -= element.z;
          offset++;
          if (element.w <= key) {
            key -= element.w;
            offset++;
          }
        }
      }
    }
    return key;
  }

  template<uint32_t Q>
  struct OffsetBlob
  {
    uint32_t data[Q];
  };

  template<uint32_t L, uint32_t Q>
  __global__
  __launch_bounds__(128)
  void extract(uint32_t* __restrict__ out_d,
               const uint4* __restrict__ hp_d,
               const OffsetBlob<Q> offsets)
  {
    const uint32_t index = 128 * blockIdx.x + threadIdx.x;
    if (hp_d[0].x <= index) return;
