     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) {
+    for (uint32_t k = blockDim.x * blockIdx.x; k < N; k += gridDim.x * blockDim.x) {
       const uint32_t index = k + threadIdx.x;
       if (index < N) {
         uint32_t offset = 0;

    }
  }

  void scratchLayout(std::vector<uint32_t>& levels, std::vector<uint32_t>& offsets, uint32_t N)
  {
    if (N == 0) return;


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

  


}

