         //std::fprintf(stderr, "sketch ptr for index %zu is %p + %zu (%p)\n", index, (void *)dmem_, (void *)(static_cast<uint8_t *>(dmem_) + entrysize_ * index));
         return static_cast<void *>(static_cast<uint8_t *>(dmem_) + entrysize_ * index);
     }
-    virtual void process_hlls() { // Default, unpacked 8-bit HLLs
+    void process_hlls() {
+        switch(storage_) {
+        case DEFAULT_6_BIT:
+            process([dmem=dmem_](void *drmem, size_t row_index, size_t round_nrows, int tpb, int nblocks){
+                 // 6 bit kernel
+            });
+            break;
+        case PACKED_4_BIT:
+            process([dmem=dmem_](void *drmem, size_t row_index, size_t round_nrows, int tpb, int nblocks){
+                 // 4 bit kernel
+            });
+            break;
+        case WIDE_8_BIT:
+            process([dmem=dmem_](void *drmem, size_t row_index, size_t round_nrows, int tpb, int nblocks){
+                 // 8 bit kernel
+            });
+            break;
+        default: __builtin_unreachable();
+        }
+    }
+    void process_bbmh() {
+        process([dmem=dmem_](void *drmem, size_t row_index, size_t round_nrows, int tpb, int nblocks){
+            // BBit kernel
+        });
+    }
+    template<typename F>
+    template<typename F>
+    void process(const F &f) { // Default, unpacked 8-bit HLLs
         if(entrysize_ < 1024) throw std::runtime_error("entrysize must be at least 1024. (4 for intrinsics, 256 for threads per block");
         assert(entrysize_ % 256 == 0);
         size_t rind    = 0; // Row index
         constexpr int tpb = 256;
         int nblocks = nelem_ * numrows_ * entrysize_ / 256;
         std::thread copy_and_flush;
-        bool copied = true;
         cudaError_t ce;
         for(size_t tranche = 0, ntranches = (nelem_ - 1 + numrows_) / numrows_; tranche < ntranches; ++tranche) {
             size_t round_nrows = std::min(numrows_, nelem_ - numrows_);
+            f(drmem_ + (nelem_ * elemsz() * rind) /* local destination */,
+              rind, round_nrows, tpb, nblocks);
 #if 0
             original_hll_compare<<<nblocks, tpb, 64 * sizeof(uint32_t)>>>(
                 dmem_,

INLINE uint64_t nchoose2(T x) {
    return (uint64_t(x) * uint64_t(x - 1)) / 2;
}


template<typename T>

#endif
            cudaDeviceSynchronize();
            if(copy_and_flush.joinable()) copy_and_flush.join();
            // At this point, the CPU buffer is available for loading.
            // TODO: double the memory on the device and compute the next portion
            // while transferring
            size_t nb = nelem_ * elemsz() * round_nrows;
            std::fprintf(stderr, "drmem %p. nb: %zu. round nr: %zu\n", (void *)drmem_, nb, round_nrows);
            std::fprintf(stderr, "cmem %p\n", (void *)cmem_);
