     return key;
   }
 
-
-  template<uint32_t L>
+  __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 std::array<uint32_t,L> offsets)
+               const OffsetBlob<Q> offsets)
   {
     const uint32_t index = 128 * blockIdx.x + threadIdx.x;
     if (hp_d[0].x <= index) return;


  switch (L)
  {
  case 0: ::extract<0><<<(N + 127) / 128, 127, 0, stream>>>(out_d, reinterpret_cast<uint4*>(scratch_d), arrayFromVector<0>(offsets)); break;
  default:
    abort();
  }

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

