     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>>>(
+    // use a for_each to automate support older compute capabilities with limited grid sizes
+    nvbio::for_each<device_tag>(
         n,
-        keys,
-        flags );
+        thrust::make_counting_iterator<uint32>(0u),
+        build_head_flags_functor<uint64>( n, keys, flags ) );
 }
 
 uint32 extract_radix_16(

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

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

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

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

    // load the next key
    const uint64 key = keys[thread_id];
    const uint8 flag = flags[thread_id];

    // and write the corresponding flag out
    flags[thread_id] = (key != key_p) ? 1u : flag;
}

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

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

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

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