 
         io::ReadDataCUDA reads_data_cuda( *reads_data );
 
+        // create a host-side read batch
+        io::ReadDataView<const uint32*,const uint32*,const char*,const char*> host_reads_view( *reads_data );
+
         // create a device-side read batch
-        io::ReadDataView<const uint32*,const uint32*,const char*,const char*> reads_view( reads_data_cuda );
+        io::ReadDataView<const uint32*,const uint32*,const char*,const char*> reads_view_cuda( reads_data_cuda );
 
         thrust::device_vector<uint32> counter(1);
         counter[0] = 0;

    // constructor
    //
    // \param count     pointer to the global counter
    NVBIO_FORCEINLINE NVBIO_DEVICE
    CountDelegate(uint32* count) : m_count( count ) {}

    // main functor operator
    //
    NVBIO_FORCEINLINE NVBIO_DEVICE
    void operator() (const uint2 range) const { atomicAdd( m_count, range.y + 1u - range.x ); }

private:
    uint32* m_count;    // global counter
};

//
// k-mer counting kernel

            cudaEventRecord( start, 0 );

            count_kernel<<<n_blocks,blockdim>>>(
                reads_view_cuda,
                fmindex_cuda,
                20u,
                0u,
