     {
         block_scan.template ConsumeTile<false, false>(block_offset, prefix_op);
         block_offset += AgentScanT::TILE_ITEMS;
+    }
+
+    // Process the remaining partial tile (if any).
+    if (block_offset < num_counts)
+    {
+        block_scan.template ConsumeTile<false, true>(block_offset, prefix_op,
+                                                     num_counts - block_offset);
     }
 }
 

        const ValueT    *d_values_in,
        ValueT          *d_values_out,
        OffsetT         *d_spine,
        int             spine_length,
        int             &current_bit,
        PassConfigT     &pass_config)
    {

                _CubLog("Invoking upsweep_kernel<<<%d, %d, 0, %lld>>>(), %d items per thread, %d SM occupancy, current bit %d, bit_grain %d\n",
                pass_config.even_share.grid_size, pass_config.upsweep_config.block_threads, (long long) stream,
                pass_config.upsweep_config.items_per_thread, pass_config.upsweep_config.sm_occupancy, current_bit, pass_bits);

            // Spine length written by the upsweep kernel in the current pass.
            int pass_spine_length = pass_config.even_share.grid_size * pass_config.radix_digits;

            // Invoke upsweep_kernel with same grid size as downsweep_kernel
            thrust::cuda_cub::launcher::triple_chevron(
