     }
   }
 
+
+  template<uint32_t L>
+  void runExtract(uint32_t* out_d, const std::vector<uint32_t>& offsets, const uint32_t* hp_d, uint32_t N, cudaStream_t stream)
+  {
+    OffsetBlob offsetBlob;
+    std::memcpy(&offsetBlob.data0, offsets.data(), sizeof(uint32_t)*std::min(uint32_t(12), L));
+
+    int minGridSize = 0;
+    int blockSize = 0;
+    auto rv = cudaOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize, extract<L>);
+    assert(rv == cudaSuccess);
+
+    auto blocks = std::min(minGridSize, int((N + blockSize - 1) / blockSize));
+    ::extract<L><<<blocks, blockSize, 0, stream>>>(out_d, hp_d, offsetBlob);
+  }
+
   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
  }

}



#define HP5_WARP_COUNT 5
#define HP5_WARP_SIZE 32

namespace {

