  */
 
 #include <nvbio/sufsort/sufsort_priv.h>
+#include <nvbio/basic/primitives.h>
 #ifdef _OPENMP
 #include <omp.h>
 #endif

    const uint64*           keys,
          uint8*            flags)
{
    const uint32 blockdim = 128;
    const uint32 n_blocks = util::divide_ri( n, blockdim );

    build_head_flags_kernel<<<n_blocks,blockdim>>>(
        n,
        keys,
        flags );
}

uint32 extract_radix_16(

    const uint32*           keys,
          uint8*            flags)
{
    const uint32 n_quads = util::divide_ri( n, 4u );

    // use a for_each to automate support for older compute capabilities with limited grid sizes
    nvbio::for_each<device_tag>(
        n_quads,
        thrust::make_counting_iterator<uint32>(0u),
        build_head_flags_functor<uint32>( n, keys, flags ) );
}

// build a set of head flags looking at adjacent keys
