             BLOCK_THREADS,
             ITEMS_PER_THREAD,
             ValueT,
-            ChainedPolicyT::ActivePolicy::DownsweepPolicy::RADIX_BITS,
-            ChainedPolicyT::ActivePolicy::DownsweepPolicy::MEMOIZE_OUTER_SCAN,
-            ChainedPolicyT::ActivePolicy::DownsweepPolicy::INNER_SCAN_ALGORITHM>
+            ChainedPolicyT::ActivePolicy::SingleTilePolicy::RADIX_BITS,
+            ChainedPolicyT::ActivePolicy::SingleTilePolicy::MEMOIZE_OUTER_SCAN,
+            ChainedPolicyT::ActivePolicy::SingleTilePolicy::INNER_SCAN_ALGORITHM>
         BlockRadixSortT;
 
     // BlockLoad type (keys)

 */
template <
    typename                ChainedPolicyT,                 ///< Chained tuning policy
    bool                    IS_DESCENDING,                     ///< Whether or not the sorted-order is high-to-low
    typename                KeyT,                           ///< Key type
    typename                ValueT,                         ///< Value type
    typename                OffsetT>                        ///< Signed integer type for global offsets
__launch_bounds__ (int(ChainedPolicyT::ActivePolicy::DownsweepPolicy::BLOCK_THREADS), 1)
__global__ void DeviceRadixSortSingleTileKernel(
    KeyT                    *d_keys_in,                     ///< [in] Input keys buffer
    KeyT                    *d_keys_out,                    ///< [in] Output keys buffer

    // Constants
    enum
    {
        BLOCK_THREADS           = ChainedPolicyT::ActivePolicy::SingleTilePolicy::BLOCK_THREADS,
        ITEMS_PER_THREAD        = ChainedPolicyT::ActivePolicy::SingleTilePolicy::ITEMS_PER_THREAD,
        KEYS_ONLY               = Equals<ValueT, NullType>::VALUE,
    };

