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

    return key;
  }


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

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

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

