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

    uint32_t N = hp_d[0];
    uint4 T = *(const uint4*)(hp_d + 4);

    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;


#define HP5_WARP_COUNT 5
#define HP5_WARP_SIZE 32

namespace {

