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

    const uint8*            flags,
          uint32*           comp_flags)
{
    const uint32 blockdim = 128;
    const uint32 n_words  = util::divide_ri( n, 32u );
    const uint32 n_blocks = util::divide_ri( n_words, blockdim );

    pack_flags_kernel<<<n_blocks,blockdim>>>(
        n,
        flags,
        comp_flags );
}

// build a set of head flags looking at adjacent keys
//
__global__ void build_head_flags_kernel(
    const uint32            n,
    const uint32*           keys,
          uint8*            flags)
{
    const uint32 thread_id = (threadIdx.x + blockIdx.x * blockDim.x);
    const uint32 idx = 4 * thread_id;
    if (idx >= n)
        return;

    // load the previous key
    const uint32 key_p = idx ? keys[idx-1] : 0xFFFFFFFF;

    // load the next 4 keys
    const uint4  key  = ((const uint4*)keys)[thread_id];
    const uchar4 flag = ((const uchar4*)flags)[thread_id];

    // and write the corresponding 4 flags
    ((uchar4*)flags)[thread_id] = make_uchar4(
        (key.x != key_p) ? 1u : flag.x,
        (key.y != key.x) ? 1u : flag.y,
        (key.z != key.y) ? 1u : flag.z,
        (key.w != key.z) ? 1u : flag.w );
}

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

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

uint32 extract_radix_16(
