From 6330f93b1cdf4903e5f3eb8611fafe8f5cb42dd1 Mon Sep 17 00:00:00 2001
From: dumerrill <dumerrill@nvidia.com>
Date: Fri, 20 Feb 2015 18:18:50 -0500
Subject: [PATCH] Update radix sort policies

Former-commit-id: aca50f0a12128a6f8b2f15d83fcb26eed4a781ca
---
 .cproject                                     |  10 +-
 .../agent_histogram.cuh}                      |  30 +--
 .../agent_radix_sort_downsweep.cuh}           |  49 ++--
 .../agent_radix_sort_upsweep.cuh}             |  26 +--
 .../agent_reduce.cuh}                         |  32 +--
 .../agent_reduce_by_key.cuh}                  |  44 ++--
 .../agent_rle.cuh}                            |  32 +--
 .../agent_scan.cuh}                           |  44 ++--
 .../agent_select_if.cuh}                      |  40 ++--
 .../agent_spmv.cuh}                           |  26 +--
 .../single_pass_scan_operators.cuh}           |   0
 cub/block/block_load.cuh                      |  27 +--
 cub/block/block_store.cuh                     | 113 +++++++--
 .../block_unzip_channels_sweep.cuh            | 217 ------------------
 cub/device/device_histogram.cuh               |  10 +-
 cub/device/device_partition.cuh               |   6 +-
 cub/device/device_radix_sort.cuh              |  18 +-
 cub/device/device_reduce.cuh                  |  30 +--
 cub/device/device_run_length_encode.cuh       |   6 +-
 cub/device/device_scan.cuh                    |  10 +-
 cub/device/device_select.cuh                  |   8 +-
 cub/device/device_spmv.cuh                    |   2 +-
 ...am_dispatch.cuh => dispatch_histogram.cuh} |  34 +--
 ...t_dispatch.cuh => dispatch_radix_sort.cuh} | 204 ++++++++++------
 ...educe_dispatch.cuh => dispatch_reduce.cuh} |  50 ++--
 ...ispatch.cuh => dispatch_reduce_by_key.cuh} |  44 ++--
 ...vice_rle_dispatch.cuh => dispatch_rle.cuh} |  42 ++--
 ...ce_scan_dispatch.cuh => dispatch_scan.cuh} |  56 ++---
 ...ct_dispatch.cuh => dispatch_select_if.cuh} |  44 ++--
 ...ce_spmv_dispatch.cuh => dispatch_spmv.cuh} |  28 +--
 test/test_device_radix_sort.cu                |   3 +
 31 files changed, 579 insertions(+), 706 deletions(-)
 rename cub/{block_sweep/block_histogram_sweep.cuh => agent/agent_histogram.cuh} (95%)
 rename cub/{block_sweep/block_radix_sort_downsweep.cuh => agent/agent_radix_sort_downsweep.cuh} (93%)
 rename cub/{block_sweep/block_radix_sort_upsweep.cuh => agent/agent_radix_sort_upsweep.cuh} (94%)
 rename cub/{block_sweep/block_reduce_sweep.cuh => agent/agent_reduce.cuh} (91%)
 rename cub/{block_sweep/block_reduce_by_key_sweep.cuh => agent/agent_reduce_by_key.cuh} (94%)
 rename cub/{block_sweep/block_rle_sweep.cuh => agent/agent_rle.cuh} (96%)
 rename cub/{block_sweep/block_scan_sweep.cuh => agent/agent_scan.cuh} (89%)
 rename cub/{block_sweep/block_select_sweep.cuh => agent/agent_select_if.cuh} (94%)
 rename cub/{block_sweep/block_spmv_sweep.cuh => agent/agent_spmv.cuh} (93%)
 rename cub/{block_sweep/block_scan_prefix_operators.cuh => agent/single_pass_scan_operators.cuh} (100%)
 delete mode 100644 cub/block_sweep/block_unzip_channels_sweep.cuh
 rename cub/device/dispatch/{device_histogram_dispatch.cuh => dispatch_histogram.cuh} (98%)
 rename cub/device/dispatch/{device_radix_sort_dispatch.cuh => dispatch_radix_sort.cuh} (77%)
 rename cub/device/dispatch/{device_reduce_dispatch.cuh => dispatch_reduce.cuh} (96%)
 rename cub/device/dispatch/{device_reduce_by_key_dispatch.cuh => dispatch_reduce_by_key.cuh} (92%)
 rename cub/device/dispatch/{device_rle_dispatch.cuh => dispatch_rle.cuh} (92%)
 rename cub/device/dispatch/{device_scan_dispatch.cuh => dispatch_scan.cuh} (91%)
 rename cub/device/dispatch/{device_select_dispatch.cuh => dispatch_select_if.cuh} (92%)
 rename cub/device/dispatch/{device_spmv_dispatch.cuh => dispatch_spmv.cuh} (98%)

diff --git a/.cproject b/.cproject
index 2efaabc2e4..5ec455b7c4 100644
--- a/.cproject
+++ b/.cproject
@@ -20,16 +20,16 @@
 					<folderInfo id="cdt.managedbuild.toolchain.gnu.cygwin.base.1260156311.1722659113" name="/" resourcePath="">
 						<toolChain id="cdt.managedbuild.toolchain.gnu.cygwin.base.1285519857" name="Cygwin GCC" superClass="cdt.managedbuild.toolchain.gnu.cygwin.base">
 							<targetPlatform archList="all" binaryParser="org.eclipse.cdt.core.Cygwin_PE" id="cdt.managedbuild.target.gnu.platform.cygwin.base.365256368" name="Debug Platform" osList="win32" superClass="cdt.managedbuild.target.gnu.platform.cygwin.base"/>
-							<builder buildPath="${workspace_loc:/PrivateCub}/Default" id="cdt.managedbuild.target.gnu.builder.cygwin.base.446293184" name="Gnu Make Builder.Default" superClass="cdt.managedbuild.target.gnu.builder.cygwin.base"/>
+							<builder buildPath="${workspace_loc:/PrivateCub}/Default" id="cdt.managedbuild.target.gnu.builder.cygwin.base.446293184" keepEnvironmentInBuildfile="false" name="Gnu Make Builder" superClass="cdt.managedbuild.target.gnu.builder.cygwin.base"/>
 							<tool id="cdt.managedbuild.tool.gnu.assembler.cygwin.base.243775805" name="GCC Assembler" superClass="cdt.managedbuild.tool.gnu.assembler.cygwin.base">
 								<inputType id="cdt.managedbuild.tool.gnu.assembler.input.1808657577" superClass="cdt.managedbuild.tool.gnu.assembler.input"/>
 							</tool>
 							<tool id="cdt.managedbuild.tool.gnu.archiver.cygwin.base.880994059" name="GCC Archiver" superClass="cdt.managedbuild.tool.gnu.archiver.cygwin.base"/>
 							<tool id="cdt.managedbuild.tool.gnu.cpp.compiler.cygwin.base.586941236" name="Cygwin C++ Compiler" superClass="cdt.managedbuild.tool.gnu.cpp.compiler.cygwin.base">
-								<option id="gnu.cpp.compiler.option.include.paths.1901125882" superClass="gnu.cpp.compiler.option.include.paths" valueType="includePath">
+								<option id="gnu.cpp.compiler.option.include.paths.1901125882" name="Include paths (-I)" superClass="gnu.cpp.compiler.option.include.paths" valueType="includePath">
 									<listOptionValue builtIn="false" value="&quot;${CUDA_PATH}/include&quot;"/>
 								</option>
-								<option id="gnu.cpp.compiler.option.preprocessor.def.690074796" superClass="gnu.cpp.compiler.option.preprocessor.def" valueType="definedSymbols">
+								<option id="gnu.cpp.compiler.option.preprocessor.def.690074796" name="Defined symbols (-D)" superClass="gnu.cpp.compiler.option.preprocessor.def" valueType="definedSymbols">
 									<listOptionValue builtIn="false" value="__device__"/>
 									<listOptionValue builtIn="false" value="__global__"/>
 									<listOptionValue builtIn="false" value="__shared__"/>
@@ -44,10 +44,10 @@
 								<inputType id="cdt.managedbuild.tool.gnu.cpp.compiler.input.cygwin.1654082299" superClass="cdt.managedbuild.tool.gnu.cpp.compiler.input.cygwin"/>
 							</tool>
 							<tool id="cdt.managedbuild.tool.gnu.c.compiler.cygwin.base.1149397878" name="Cygwin C Compiler" superClass="cdt.managedbuild.tool.gnu.c.compiler.cygwin.base">
-								<option id="gnu.c.compiler.option.include.paths.18887661" superClass="gnu.c.compiler.option.include.paths" valueType="includePath">
+								<option id="gnu.c.compiler.option.include.paths.18887661" name="Include paths (-I)" superClass="gnu.c.compiler.option.include.paths" valueType="includePath">
 									<listOptionValue builtIn="false" value="&quot;${CUDA_PATH}/include&quot;"/>
 								</option>
-								<option id="gnu.c.compiler.option.preprocessor.def.symbols.788454999" superClass="gnu.c.compiler.option.preprocessor.def.symbols" valueType="definedSymbols">
+								<option id="gnu.c.compiler.option.preprocessor.def.symbols.788454999" name="Defined symbols (-D)" superClass="gnu.c.compiler.option.preprocessor.def.symbols" valueType="definedSymbols">
 									<listOptionValue builtIn="false" value="__device__"/>
 									<listOptionValue builtIn="false" value="__global__"/>
 									<listOptionValue builtIn="false" value="__shared__"/>
diff --git a/cub/block_sweep/block_histogram_sweep.cuh b/cub/agent/agent_histogram.cuh
similarity index 95%
rename from cub/block_sweep/block_histogram_sweep.cuh
rename to cub/agent/agent_histogram.cuh
index 219b865a1c..e07412bd0a 100644
--- a/cub/block_sweep/block_histogram_sweep.cuh
+++ b/cub/agent/agent_histogram.cuh
@@ -28,7 +28,7 @@
 
 /**
  * \file
- * cub::BlockHistogramSweep implements a stateful abstraction of CUDA thread blocks for participating in device-wide histogram across a range of tiles.
+ * cub::AgentHistogram implements a stateful abstraction of CUDA thread blocks for participating in device-wide histogram across a range of tiles.
  */
 
 #pragma once
@@ -64,7 +64,7 @@ enum BlockHistogramMemoryPreference
 
 
 /**
- * Parameterizable tuning policy type for BlockHistogramSweep
+ * Parameterizable tuning policy type for AgentHistogram
  */
 template <
     int                             _BLOCK_THREADS,                 ///< Threads per thread block
@@ -74,7 +74,7 @@ template <
     bool                            _RLE_COMPRESS,                  ///< Whether to perform localized RLE to compress samples before histogramming
     BlockHistogramMemoryPreference  _MEM_PREFERENCE,                ///< Whether to prefer privatized shared-memory bins (versus privatized global-memory bins)
     bool                            _WORK_STEALING>                 ///< Whether to dequeue tiles from a global work queue
-struct BlockHistogramSweepPolicy
+struct AgentHistogramPolicy
 {
     enum
     {
@@ -95,10 +95,10 @@ struct BlockHistogramSweepPolicy
  ******************************************************************************/
 
 /**
- * \brief BlockHistogramSweep implements a stateful abstraction of CUDA thread blocks for participating in device-wide histogram across a range of tiles.
+ * \brief AgentHistogram implements a stateful abstraction of CUDA thread blocks for participating in device-wide histogram across a range of tiles.
  */
 template <
-    typename    BlockHistogramSweepPolicyT,     ///< Parameterized BlockHistogramSweepPolicy tuning policy type
+    typename    AgentHistogramPolicyT,     ///< Parameterized AgentHistogramPolicy tuning policy type
     int         PRIVATIZED_SMEM_BINS,           ///< Number of privatized shared-memory histogram bins of any channel.  Zero indicates privatized counters to be maintained in global memory.
     int         NUM_CHANNELS,                   ///< Number of channels interleaved in the input data.  Supports up to four channels.
     int         NUM_ACTIVE_CHANNELS,            ///< Number of channels actively being histogrammed
@@ -108,7 +108,7 @@ template <
     typename    OutputDecodeOpT,                ///< The transform operator type for determining output bin-ids from privatized counter indices, one for each channel
     typename    OffsetT,                        ///< Signed integer type for global offsets
     int         PTX_ARCH = CUB_PTX_ARCH>        ///< PTX compute capability
-struct BlockHistogramSweep
+struct AgentHistogram
 {
     //---------------------------------------------------------------------
     // Types and constants
@@ -123,25 +123,25 @@ struct BlockHistogramSweep
     /// Constants
     enum
     {
-        BLOCK_THREADS           = BlockHistogramSweepPolicyT::BLOCK_THREADS,
+        BLOCK_THREADS           = AgentHistogramPolicyT::BLOCK_THREADS,
 
-        PIXELS_PER_THREAD       = BlockHistogramSweepPolicyT::PIXELS_PER_THREAD,
+        PIXELS_PER_THREAD       = AgentHistogramPolicyT::PIXELS_PER_THREAD,
         TILE_PIXELS             = PIXELS_PER_THREAD * BLOCK_THREADS,
 
         SAMPLES_PER_THREAD      = PIXELS_PER_THREAD * NUM_CHANNELS,
         TILE_SAMPLES            = SAMPLES_PER_THREAD * BLOCK_THREADS,
 
-        RLE_COMPRESS            = BlockHistogramSweepPolicyT::RLE_COMPRESS,
+        RLE_COMPRESS            = AgentHistogramPolicyT::RLE_COMPRESS,
 
         MEM_PREFERENCE          = (PRIVATIZED_SMEM_BINS > 0) ?
-                                        BlockHistogramSweepPolicyT::MEM_PREFERENCE :
+                                        AgentHistogramPolicyT::MEM_PREFERENCE :
                                         GMEM,
 
-        WORK_STEALING           = BlockHistogramSweepPolicyT::WORK_STEALING,
+        WORK_STEALING           = AgentHistogramPolicyT::WORK_STEALING,
     };
 
     /// Cache load modifier for reading input elements
-    static const CacheLoadModifier LOAD_MODIFIER = BlockHistogramSweepPolicyT::LOAD_MODIFIER;
+    static const CacheLoadModifier LOAD_MODIFIER = AgentHistogramPolicyT::LOAD_MODIFIER;
 
 
     /// Input iterator wrapper type (for applying cache modifier)
@@ -159,7 +159,7 @@ struct BlockHistogramSweep
             WrappedSampleIteratorT,
             BLOCK_THREADS,
             SAMPLES_PER_THREAD,
-            BlockHistogramSweepPolicyT::LOAD_ALGORITHM>
+            AgentHistogramPolicyT::LOAD_ALGORITHM>
         BlockLoadSampleT;
 
     /// Parameterized BlockLoad type for pixels
@@ -167,7 +167,7 @@ struct BlockHistogramSweep
             WrappedPixelIteratorT,
             BLOCK_THREADS,
             PIXELS_PER_THREAD,
-            BlockHistogramSweepPolicyT::LOAD_ALGORITHM>
+            AgentHistogramPolicyT::LOAD_ALGORITHM>
         BlockLoadPixelT;
 
 
@@ -621,7 +621,7 @@ struct BlockHistogramSweep
     /**
      * Constructor
      */
-    __device__ __forceinline__ BlockHistogramSweep(
+    __device__ __forceinline__ AgentHistogram(
         TempStorage         &temp_storage,                                      ///< Reference to temp_storage
         SampleIteratorT     d_samples,                                          ///< Input data to reduce
         int                 (&num_output_bins)[NUM_ACTIVE_CHANNELS],            ///< The number bins per final output histogram
diff --git a/cub/block_sweep/block_radix_sort_downsweep.cuh b/cub/agent/agent_radix_sort_downsweep.cuh
similarity index 93%
rename from cub/block_sweep/block_radix_sort_downsweep.cuh
rename to cub/agent/agent_radix_sort_downsweep.cuh
index 24b4e42c94..316ca955ed 100644
--- a/cub/block_sweep/block_radix_sort_downsweep.cuh
+++ b/cub/agent/agent_radix_sort_downsweep.cuh
@@ -28,7 +28,7 @@
 
 /**
  * \file
- * BlockRadixSortDownsweep implements a stateful abstraction of CUDA thread blocks for participating in device-wide radix sort downsweep across a range of tiles.
+ * AgentRadixSortDownsweep implements a stateful abstraction of CUDA thread blocks for participating in device-wide radix sort downsweep across a range of tiles.
  */
 
 
@@ -65,26 +65,24 @@ enum RadixSortScatterAlgorithm
 
 
 /**
- * Parameterizable tuning policy type for BlockRadixSortDownsweep
+ * Parameterizable tuning policy type for AgentRadixSortDownsweep
  */
 template <
     int                         _BLOCK_THREADS,             ///< Threads per thread block
     int                         _ITEMS_PER_THREAD,          ///< Items per thread (per tile of input)
     BlockLoadAlgorithm          _LOAD_ALGORITHM,            ///< The BlockLoad algorithm to use
     CacheLoadModifier           _LOAD_MODIFIER,             ///< Cache load modifier for reading keys (and values)
-    bool                        _EXCHANGE_TIME_SLICING,     ///< Whether or not to time-slice key/value exchanges through shared memory to lower shared memory pressure
     bool                        _MEMOIZE_OUTER_SCAN,        ///< Whether or not to buffer outer raking scan partials to incur fewer shared memory reads at the expense of higher register pressure.  See BlockScanAlgorithm::BLOCK_SCAN_RAKING_MEMOIZE for more details.
     BlockScanAlgorithm          _INNER_SCAN_ALGORITHM,      ///< The BlockScan algorithm algorithm to use
     RadixSortScatterAlgorithm   _SCATTER_ALGORITHM,         ///< The scattering strategy to use
     cudaSharedMemConfig         _SMEM_CONFIG,               ///< Shared memory bank mode
     int                         _RADIX_BITS>                ///< The number of radix bits, i.e., log2(bins)
-struct BlockRadixSortDownsweepPolicy
+struct AgentRadixSortDownsweepPolicy
 {
     enum
     {
         BLOCK_THREADS           = _BLOCK_THREADS,           ///< Threads per thread block
         ITEMS_PER_THREAD        = _ITEMS_PER_THREAD,        ///< Items per thread (per tile of input)
-        EXCHANGE_TIME_SLICING   = _EXCHANGE_TIME_SLICING,   ///< Whether or not to time-slice key/value exchanges through shared memory to lower shared memory pressure
         RADIX_BITS              = _RADIX_BITS,              ///< The number of radix bits, i.e., log2(bins)
         MEMOIZE_OUTER_SCAN      = _MEMOIZE_OUTER_SCAN,      ///< Whether or not to buffer outer raking scan partials to incur fewer shared memory reads at the expense of higher register pressure.  See BlockScanAlgorithm::BLOCK_SCAN_RAKING_MEMOIZE for more details.
     };
@@ -102,15 +100,15 @@ struct BlockRadixSortDownsweepPolicy
  ******************************************************************************/
 
 /**
- * \brief BlockRadixSortDownsweep implements a stateful abstraction of CUDA thread blocks for participating in device-wide radix sort downsweep across a range of tiles.
+ * \brief AgentRadixSortDownsweep implements a stateful abstraction of CUDA thread blocks for participating in device-wide radix sort downsweep across a range of tiles.
  */
 template <
-    typename BlockRadixSortDownsweepPolicy,             ///< Parameterized BlockRadixSortDownsweepPolicy tuning policy type
+    typename AgentRadixSortDownsweepPolicy,             ///< Parameterized AgentRadixSortDownsweepPolicy tuning policy type
     bool     DESCENDING,                                ///< Whether or not the sorted-order is high-to-low
     typename KeyT,                                       ///< KeyT type
     typename ValueT,                                     ///< ValueT type
     typename OffsetT>                                   ///< Signed integer type for global offsets
-struct BlockRadixSortDownsweep
+struct AgentRadixSortDownsweep
 {
     //---------------------------------------------------------------------
     // Type definitions and constants
@@ -122,19 +120,18 @@ struct BlockRadixSortDownsweep
     static const UnsignedBits MIN_KEY = Traits<KeyT>::MIN_KEY;
     static const UnsignedBits MAX_KEY = Traits<KeyT>::MAX_KEY;
 
-    static const BlockLoadAlgorithm         LOAD_ALGORITHM          = BlockRadixSortDownsweepPolicy::LOAD_ALGORITHM;
-    static const CacheLoadModifier          LOAD_MODIFIER           = BlockRadixSortDownsweepPolicy::LOAD_MODIFIER;
-    static const BlockScanAlgorithm         INNER_SCAN_ALGORITHM    = BlockRadixSortDownsweepPolicy::INNER_SCAN_ALGORITHM;
-    static const RadixSortScatterAlgorithm  SCATTER_ALGORITHM       = BlockRadixSortDownsweepPolicy::SCATTER_ALGORITHM;
-    static const cudaSharedMemConfig        SMEM_CONFIG             = BlockRadixSortDownsweepPolicy::SMEM_CONFIG;
+    static const BlockLoadAlgorithm         LOAD_ALGORITHM          = AgentRadixSortDownsweepPolicy::LOAD_ALGORITHM;
+    static const CacheLoadModifier          LOAD_MODIFIER           = AgentRadixSortDownsweepPolicy::LOAD_MODIFIER;
+    static const BlockScanAlgorithm         INNER_SCAN_ALGORITHM    = AgentRadixSortDownsweepPolicy::INNER_SCAN_ALGORITHM;
+    static const RadixSortScatterAlgorithm  SCATTER_ALGORITHM       = AgentRadixSortDownsweepPolicy::SCATTER_ALGORITHM;
+    static const cudaSharedMemConfig        SMEM_CONFIG             = AgentRadixSortDownsweepPolicy::SMEM_CONFIG;
 
     enum
     {
-        BLOCK_THREADS           = BlockRadixSortDownsweepPolicy::BLOCK_THREADS,
-        ITEMS_PER_THREAD        = BlockRadixSortDownsweepPolicy::ITEMS_PER_THREAD,
-        EXCHANGE_TIME_SLICING   = BlockRadixSortDownsweepPolicy::EXCHANGE_TIME_SLICING,
-        RADIX_BITS              = BlockRadixSortDownsweepPolicy::RADIX_BITS,
-        MEMOIZE_OUTER_SCAN      = BlockRadixSortDownsweepPolicy::MEMOIZE_OUTER_SCAN,
+        BLOCK_THREADS           = AgentRadixSortDownsweepPolicy::BLOCK_THREADS,
+        ITEMS_PER_THREAD        = AgentRadixSortDownsweepPolicy::ITEMS_PER_THREAD,
+        RADIX_BITS              = AgentRadixSortDownsweepPolicy::RADIX_BITS,
+        MEMOIZE_OUTER_SCAN      = AgentRadixSortDownsweepPolicy::MEMOIZE_OUTER_SCAN,
         TILE_ITEMS              = BLOCK_THREADS * ITEMS_PER_THREAD,
 
         RADIX_DIGITS            = 1 << RADIX_BITS,
@@ -174,30 +171,26 @@ struct BlockRadixSortDownsweep
         KeysItr,
         BLOCK_THREADS,
         ITEMS_PER_THREAD,
-        LOAD_ALGORITHM,
-        EXCHANGE_TIME_SLICING> BlockLoadKeys;
+        LOAD_ALGORITHM> BlockLoadKeys;
 
     // BlockLoad type (values)
     typedef BlockLoad<
         ValuesItr,
         BLOCK_THREADS,
         ITEMS_PER_THREAD,
-        LOAD_ALGORITHM,
-        EXCHANGE_TIME_SLICING> BlockLoadValues;
+        LOAD_ALGORITHM> BlockLoadValues;
 
     // BlockExchange type (keys)
     typedef BlockExchange<
         UnsignedBits,
         BLOCK_THREADS,
-        ITEMS_PER_THREAD,
-        EXCHANGE_TIME_SLICING> BlockExchangeKeys;
+        ITEMS_PER_THREAD> BlockExchangeKeys;
 
     // BlockExchange type (values)
     typedef BlockExchange<
         ValueT,
         BLOCK_THREADS,
-        ITEMS_PER_THREAD,
-        EXCHANGE_TIME_SLICING> BlockExchangeValues;
+        ITEMS_PER_THREAD> BlockExchangeValues;
 
 
     /**
@@ -636,7 +629,7 @@ struct BlockRadixSortDownsweep
     /**
      * Constructor
      */
-    __device__ __forceinline__ BlockRadixSortDownsweep(
+    __device__ __forceinline__ AgentRadixSortDownsweep(
         TempStorage &temp_storage,
         OffsetT      bin_offset,
         KeyT        *d_keys_in,
@@ -661,7 +654,7 @@ struct BlockRadixSortDownsweep
     /**
      * Constructor
      */
-    __device__ __forceinline__ BlockRadixSortDownsweep(
+    __device__ __forceinline__ AgentRadixSortDownsweep(
         TempStorage &temp_storage,
         OffsetT     num_items,
         OffsetT     *d_spine,
diff --git a/cub/block_sweep/block_radix_sort_upsweep.cuh b/cub/agent/agent_radix_sort_upsweep.cuh
similarity index 94%
rename from cub/block_sweep/block_radix_sort_upsweep.cuh
rename to cub/agent/agent_radix_sort_upsweep.cuh
index 00e41fd500..88f951840e 100644
--- a/cub/block_sweep/block_radix_sort_upsweep.cuh
+++ b/cub/agent/agent_radix_sort_upsweep.cuh
@@ -28,7 +28,7 @@
 
 /**
  * \file
- * BlockRadixSortUpsweep implements a stateful abstraction of CUDA thread blocks for participating in device-wide radix sort upsweep across a range of tiles.
+ * AgentRadixSortUpsweep implements a stateful abstraction of CUDA thread blocks for participating in device-wide radix sort upsweep across a range of tiles.
  */
 
 #pragma once
@@ -51,14 +51,14 @@ namespace cub {
  ******************************************************************************/
 
 /**
- * Parameterizable tuning policy type for BlockRadixSortUpsweep
+ * Parameterizable tuning policy type for AgentRadixSortUpsweep
  */
 template <
     int                 _BLOCK_THREADS,     ///< Threads per thread block
     int                 _ITEMS_PER_THREAD,  ///< Items per thread (per tile of input)
     CacheLoadModifier   _LOAD_MODIFIER,     ///< Cache load modifier for reading keys
     int                 _RADIX_BITS>        ///< The number of radix bits, i.e., log2(bins)
-struct BlockRadixSortUpsweepPolicy
+struct AgentRadixSortUpsweepPolicy
 {
     enum
     {
@@ -76,13 +76,13 @@ struct BlockRadixSortUpsweepPolicy
  ******************************************************************************/
 
 /**
- * \brief BlockRadixSortUpsweep implements a stateful abstraction of CUDA thread blocks for participating in device-wide radix sort upsweep across a range of tiles.
+ * \brief AgentRadixSortUpsweep implements a stateful abstraction of CUDA thread blocks for participating in device-wide radix sort upsweep across a range of tiles.
  */
 template <
-    typename BlockRadixSortUpsweepPolicy,   ///< Parameterized BlockRadixSortUpsweepPolicy tuning policy type
+    typename AgentRadixSortUpsweepPolicy,   ///< Parameterized AgentRadixSortUpsweepPolicy tuning policy type
     typename KeyT,                          ///< KeyT type
     typename OffsetT>                       ///< Signed integer type for global offsets
-struct BlockRadixSortUpsweep
+struct AgentRadixSortUpsweep
 {
 
     //---------------------------------------------------------------------
@@ -97,13 +97,13 @@ struct BlockRadixSortUpsweep
     // Integer type for packing DigitCounters into columns of shared memory banks
     typedef unsigned int PackedCounter;
 
-    static const CacheLoadModifier LOAD_MODIFIER = BlockRadixSortUpsweepPolicy::LOAD_MODIFIER;
+    static const CacheLoadModifier LOAD_MODIFIER = AgentRadixSortUpsweepPolicy::LOAD_MODIFIER;
 
     enum
     {
-        RADIX_BITS              = BlockRadixSortUpsweepPolicy::RADIX_BITS,
-        BLOCK_THREADS           = BlockRadixSortUpsweepPolicy::BLOCK_THREADS,
-        KEYS_PER_THREAD         = BlockRadixSortUpsweepPolicy::ITEMS_PER_THREAD,
+        RADIX_BITS              = AgentRadixSortUpsweepPolicy::RADIX_BITS,
+        BLOCK_THREADS           = AgentRadixSortUpsweepPolicy::BLOCK_THREADS,
+        KEYS_PER_THREAD         = AgentRadixSortUpsweepPolicy::ITEMS_PER_THREAD,
 
         RADIX_DIGITS            = 1 << RADIX_BITS,
 
@@ -186,7 +186,7 @@ struct BlockRadixSortUpsweep
     {
         // BucketKeys
         static __device__ __forceinline__ void BucketKeys(
-            BlockRadixSortUpsweep     &cta,
+            AgentRadixSortUpsweep     &cta,
             UnsignedBits                    keys[KEYS_PER_THREAD])
         {
             cta.Bucket(keys[COUNT]);
@@ -201,7 +201,7 @@ struct BlockRadixSortUpsweep
     struct Iterate<MAX, MAX>
     {
         // BucketKeys
-        static __device__ __forceinline__ void BucketKeys(BlockRadixSortUpsweep &cta, UnsignedBits keys[KEYS_PER_THREAD]) {}
+        static __device__ __forceinline__ void BucketKeys(AgentRadixSortUpsweep &cta, UnsignedBits keys[KEYS_PER_THREAD]) {}
     };
 
 
@@ -373,7 +373,7 @@ struct BlockRadixSortUpsweep
     /**
      * Constructor
      */
-    __device__ __forceinline__ BlockRadixSortUpsweep(
+    __device__ __forceinline__ AgentRadixSortUpsweep(
         TempStorage &temp_storage,
         KeyT        *d_keys_in,
         int         current_bit,
diff --git a/cub/block_sweep/block_reduce_sweep.cuh b/cub/agent/agent_reduce.cuh
similarity index 91%
rename from cub/block_sweep/block_reduce_sweep.cuh
rename to cub/agent/agent_reduce.cuh
index 21c6cfb16a..27735c0353 100644
--- a/cub/block_sweep/block_reduce_sweep.cuh
+++ b/cub/agent/agent_reduce.cuh
@@ -28,7 +28,7 @@
 
 /**
  * \file
- * cub::BlockReduceSweep implements a stateful abstraction of CUDA thread blocks for participating in device-wide reduction across a range of tiles.
+ * cub::AgentReduce implements a stateful abstraction of CUDA thread blocks for participating in device-wide reduction across a range of tiles.
  */
 
 #pragma once
@@ -57,7 +57,7 @@ namespace cub {
  ******************************************************************************/
 
 /**
- * Parameterizable tuning policy type for BlockReduceSweep
+ * Parameterizable tuning policy type for AgentReduce
  */
 template <
     int                     _BLOCK_THREADS,         ///< Threads per thread block
@@ -66,7 +66,7 @@ template <
     BlockReduceAlgorithm    _BLOCK_ALGORITHM,       ///< Cooperative block-wide reduction algorithm to use
     CacheLoadModifier       _LOAD_MODIFIER,         ///< Cache load modifier for reading input elements
     GridMappingStrategy     _GRID_MAPPING>          ///< How to map tiles of input onto thread blocks
-struct BlockReduceSweepPolicy
+struct AgentReducePolicy
 {
     enum
     {
@@ -87,18 +87,18 @@ struct BlockReduceSweepPolicy
  ******************************************************************************/
 
 /**
- * \brief BlockReduceSweep implements a stateful abstraction of CUDA thread blocks for participating in device-wide reduction across a range of tiles.
+ * \brief AgentReduce implements a stateful abstraction of CUDA thread blocks for participating in device-wide reduction across a range of tiles.
  *
  * Each thread reduces only the values it loads. If \p FIRST_TILE, this
  * partial reduction is stored into \p thread_aggregate.  Otherwise it is
  * accumulated into \p thread_aggregate.
  */
 template <
-    typename BlockReduceSweepPolicy,        ///< Parameterized BlockReduceSweepPolicy tuning policy type
+    typename AgentReducePolicy,        ///< Parameterized AgentReducePolicy tuning policy type
     typename InputIteratorT,                ///< Random-access iterator type for input
     typename OffsetT,                       ///< Signed integer type for global offsets
     typename ReductionOp>                   ///< Binary reduction operator type having member <tt>T operator()(const T &a, const T &b)</tt>
-struct BlockReduceSweep
+struct AgentReduce
 {
 
     //---------------------------------------------------------------------
@@ -109,20 +109,20 @@ struct BlockReduceSweep
     typedef typename std::iterator_traits<InputIteratorT>::value_type T;
 
     /// Vector type of T for data movement
-    typedef typename CubVector<T, BlockReduceSweepPolicy::VECTOR_LOAD_LENGTH>::Type VectorT;
+    typedef typename CubVector<T, AgentReducePolicy::VECTOR_LOAD_LENGTH>::Type VectorT;
 
     /// Input iterator wrapper type (for applying cache modifier)
     typedef typename If<IsPointer<InputIteratorT>::VALUE,
-            CacheModifiedInputIterator<BlockReduceSweepPolicy::LOAD_MODIFIER, T, OffsetT>,  // Wrap the native input pointer with CacheModifiedInputIterator
+            CacheModifiedInputIterator<AgentReducePolicy::LOAD_MODIFIER, T, OffsetT>,  // Wrap the native input pointer with CacheModifiedInputIterator
             InputIteratorT>::Type                                                            // Directly use the supplied input iterator type
         WrappedInputIteratorT;
 
     /// Constants
     enum
     {
-        BLOCK_THREADS       = BlockReduceSweepPolicy::BLOCK_THREADS,
-        ITEMS_PER_THREAD    = BlockReduceSweepPolicy::ITEMS_PER_THREAD,
-        VECTOR_LOAD_LENGTH  = CUB_MIN(ITEMS_PER_THREAD, BlockReduceSweepPolicy::VECTOR_LOAD_LENGTH),
+        BLOCK_THREADS       = AgentReducePolicy::BLOCK_THREADS,
+        ITEMS_PER_THREAD    = AgentReducePolicy::ITEMS_PER_THREAD,
+        VECTOR_LOAD_LENGTH  = CUB_MIN(ITEMS_PER_THREAD, AgentReducePolicy::VECTOR_LOAD_LENGTH),
         TILE_ITEMS          = BLOCK_THREADS * ITEMS_PER_THREAD,
 
         // Can vectorize according to the policy if the input iterator is a native pointer to a primitive type
@@ -130,11 +130,11 @@ struct BlockReduceSweep
 
     };
 
-    static const CacheLoadModifier    LOAD_MODIFIER   = BlockReduceSweepPolicy::LOAD_MODIFIER;
-    static const BlockReduceAlgorithm BLOCK_ALGORITHM = BlockReduceSweepPolicy::BLOCK_ALGORITHM;
+    static const CacheLoadModifier    LOAD_MODIFIER   = AgentReducePolicy::LOAD_MODIFIER;
+    static const BlockReduceAlgorithm BLOCK_ALGORITHM = AgentReducePolicy::BLOCK_ALGORITHM;
 
     /// Parameterized BlockReduce primitive
-    typedef BlockReduce<T, BLOCK_THREADS, BlockReduceSweepPolicy::BLOCK_ALGORITHM> BlockReduceT;
+    typedef BlockReduce<T, BLOCK_THREADS, AgentReducePolicy::BLOCK_ALGORITHM> BlockReduceT;
 
     /// Shared memory type required by this thread block
     typedef typename BlockReduceT::TempStorage _TempStorage;
@@ -187,7 +187,7 @@ struct BlockReduceSweep
     /**
      * Constructor
      */
-    __device__ __forceinline__ BlockReduceSweep(
+    __device__ __forceinline__ AgentReduce(
         TempStorage&            temp_storage,       ///< Reference to temp_storage
         InputIteratorT          d_in,               ///< Input data to reduce
         ReductionOp             reduction_op)       ///< Binary reduction operator
@@ -243,7 +243,7 @@ struct BlockReduceSweep
 
             // Vector Input iterator wrapper type (for applying cache modifier)
             T *d_in_unqualified = const_cast<T*>(d_in) + block_offset + (threadIdx.x * VECTOR_LOAD_LENGTH);
-            CacheModifiedInputIterator<BlockReduceSweepPolicy::LOAD_MODIFIER, VectorT, OffsetT> d_vec_in(
+            CacheModifiedInputIterator<AgentReducePolicy::LOAD_MODIFIER, VectorT, OffsetT> d_vec_in(
                 reinterpret_cast<VectorT*>(d_in_unqualified));
 
             #pragma unroll
diff --git a/cub/block_sweep/block_reduce_by_key_sweep.cuh b/cub/agent/agent_reduce_by_key.cuh
similarity index 94%
rename from cub/block_sweep/block_reduce_by_key_sweep.cuh
rename to cub/agent/agent_reduce_by_key.cuh
index b4c71bb81b..e701fff9c0 100644
--- a/cub/block_sweep/block_reduce_by_key_sweep.cuh
+++ b/cub/agent/agent_reduce_by_key.cuh
@@ -28,14 +28,14 @@
 
 /**
  * \file
- * cub::BlockReduceSweepByKey implements a stateful abstraction of CUDA thread blocks for participating in device-wide reduce-value-by-key.
+ * cub::AgentReduceByKey implements a stateful abstraction of CUDA thread blocks for participating in device-wide reduce-value-by-key.
  */
 
 #pragma once
 
 #include <iterator>
 
-#include "block_scan_prefix_operators.cuh"
+#include "single_pass_scan_operators.cuh"
 #include "../block/block_load.cuh"
 #include "../block/block_store.cuh"
 #include "../block/block_scan.cuh"
@@ -58,7 +58,7 @@ namespace cub {
  ******************************************************************************/
 
 /**
- * Parameterizable tuning policy type for BlockReduceSweepByKey
+ * Parameterizable tuning policy type for AgentReduceByKey
  */
 template <
     int                         _BLOCK_THREADS,                 ///< Threads per thread block
@@ -67,7 +67,7 @@ template <
     CacheLoadModifier           _LOAD_MODIFIER,                 ///< Cache load modifier for reading input elements
     bool                        _TWO_PHASE_SCATTER,             ///< Whether or not to coalesce output values in shared memory before scattering them to global
     BlockScanAlgorithm          _SCAN_ALGORITHM>                ///< The BlockScan algorithm to use
-struct BlockReduceSweepByKeyPolicy
+struct AgentReduceByKeyPolicy
 {
     enum
     {
@@ -87,10 +87,10 @@ struct BlockReduceSweepByKeyPolicy
  ******************************************************************************/
 
 /**
- * \brief BlockReduceSweepByKey implements a stateful abstraction of CUDA thread blocks for participating in device-wide reduce-value-by-key across a range of tiles
+ * \brief AgentReduceByKey implements a stateful abstraction of CUDA thread blocks for participating in device-wide reduce-value-by-key across a range of tiles
  */
 template <
-    typename    BlockReduceSweepByKeyPolicy,    ///< Parameterized BlockReduceSweepByKeyPolicy tuning policy type
+    typename    AgentReduceByKeyPolicy,    ///< Parameterized AgentReduceByKeyPolicy tuning policy type
     typename    KeysInputIteratorT,             ///< Random-access input iterator type for keys
     typename    UniqueOutputIteratorT,          ///< Random-access output iterator type for keys
     typename    ValuesInputIteratorT,           ///< Random-access input iterator type for values
@@ -98,7 +98,7 @@ template <
     typename    EqualityOp,                     ///< KeyT equality operator type
     typename    ReductionOp,                    ///< ValueT reduction operator type
     typename    OffsetT>                        ///< Signed integer type for global offsets
-struct BlockReduceSweepByKey
+struct AgentReduceByKey
 {
     //---------------------------------------------------------------------
     // Types and constants
@@ -119,17 +119,17 @@ struct BlockReduceSweepByKey
     // Constants
     enum
     {
-        BLOCK_THREADS       = BlockReduceSweepByKeyPolicy::BLOCK_THREADS,
+        BLOCK_THREADS       = AgentReduceByKeyPolicy::BLOCK_THREADS,
         WARPS               = BLOCK_THREADS / CUB_PTX_WARP_THREADS,
-        ITEMS_PER_THREAD    = BlockReduceSweepByKeyPolicy::ITEMS_PER_THREAD,
-        TWO_PHASE_SCATTER   = (BlockReduceSweepByKeyPolicy::TWO_PHASE_SCATTER) && (ITEMS_PER_THREAD > 1),
+        ITEMS_PER_THREAD    = AgentReduceByKeyPolicy::ITEMS_PER_THREAD,
+        TWO_PHASE_SCATTER   = (AgentReduceByKeyPolicy::TWO_PHASE_SCATTER) && (ITEMS_PER_THREAD > 1),
         TILE_ITEMS          = BLOCK_THREADS * ITEMS_PER_THREAD,
 
         // Whether or not the scan operation has a zero-valued identity value (true if we're performing addition on a primitive type)
         HAS_IDENTITY_ZERO       = (Equals<ReductionOp, cub::Sum>::VALUE) && (Traits<ValueT>::PRIMITIVE),
 
         // Whether or not to sync after loading data
-        SYNC_AFTER_LOAD         = (BlockReduceSweepByKeyPolicy::LOAD_ALGORITHM != BLOCK_LOAD_DIRECT),
+        SYNC_AFTER_LOAD         = (AgentReduceByKeyPolicy::LOAD_ALGORITHM != BLOCK_LOAD_DIRECT),
 
         // Whether or not this is run-length-encoding with a constant iterator as values
         IS_RUN_LENGTH_ENCODE    = (Equals<ValuesInputIteratorT, ConstantInputIterator<ValueT, size_t> >::VALUE) || (Equals<ValuesInputIteratorT, ConstantInputIterator<ValueT, int> >::VALUE) || (Equals<ValuesInputIteratorT, ConstantInputIterator<ValueT, unsigned int> >::VALUE),
@@ -138,13 +138,13 @@ struct BlockReduceSweepByKey
 
     // Cache-modified Input iterator wrapper type (for applying cache modifier) for keys
     typedef typename If<IsPointer<KeysInputIteratorT>::VALUE,
-            CacheModifiedInputIterator<BlockReduceSweepByKeyPolicy::LOAD_MODIFIER, KeyT, OffsetT>,   // Wrap the native input pointer with CacheModifiedValuesInputIterator
+            CacheModifiedInputIterator<AgentReduceByKeyPolicy::LOAD_MODIFIER, KeyT, OffsetT>,   // Wrap the native input pointer with CacheModifiedValuesInputIterator
             KeysInputIteratorT>::Type                                                                 // Directly use the supplied input iterator type
         WrappedKeysInputIteratorT;
 
     // Cache-modified Input iterator wrapper type (for applying cache modifier) for values
     typedef typename If<IsPointer<ValuesInputIteratorT>::VALUE,
-            CacheModifiedInputIterator<BlockReduceSweepByKeyPolicy::LOAD_MODIFIER, ValueT, OffsetT>,  // Wrap the native input pointer with CacheModifiedValuesInputIterator
+            CacheModifiedInputIterator<AgentReduceByKeyPolicy::LOAD_MODIFIER, ValueT, OffsetT>,  // Wrap the native input pointer with CacheModifiedValuesInputIterator
             ValuesInputIteratorT>::Type                                                                // Directly use the supplied input iterator type
         WrappedValuesInputIteratorT;
 
@@ -154,19 +154,19 @@ struct BlockReduceSweepByKey
     // Parameterized BlockLoad type for keys
     typedef BlockLoad<
             WrappedKeysInputIteratorT,
-            BlockReduceSweepByKeyPolicy::BLOCK_THREADS,
-            BlockReduceSweepByKeyPolicy::ITEMS_PER_THREAD,
-            BlockReduceSweepByKeyPolicy::LOAD_ALGORITHM>
+            AgentReduceByKeyPolicy::BLOCK_THREADS,
+            AgentReduceByKeyPolicy::ITEMS_PER_THREAD,
+            AgentReduceByKeyPolicy::LOAD_ALGORITHM>
         BlockLoadKeys;
 
     // Parameterized BlockLoad type for values
     typedef BlockLoad<
             WrappedValuesInputIteratorT,
-            BlockReduceSweepByKeyPolicy::BLOCK_THREADS,
-            BlockReduceSweepByKeyPolicy::ITEMS_PER_THREAD,
+            AgentReduceByKeyPolicy::BLOCK_THREADS,
+            AgentReduceByKeyPolicy::ITEMS_PER_THREAD,
             (IS_RUN_LENGTH_ENCODE) ?
                 BLOCK_LOAD_DIRECT :
-                (BlockLoadAlgorithm) BlockReduceSweepByKeyPolicy::LOAD_ALGORITHM>
+                (BlockLoadAlgorithm) AgentReduceByKeyPolicy::LOAD_ALGORITHM>
         BlockLoadValues;
 
     // Parameterized BlockExchange type for locally compacting items as part of a two-phase scatter
@@ -189,8 +189,8 @@ struct BlockReduceSweepByKey
     // Parameterized BlockScan type
     typedef BlockScan<
             ReductionOffsetPair,
-            BlockReduceSweepByKeyPolicy::BLOCK_THREADS,
-            BlockReduceSweepByKeyPolicy::SCAN_ALGORITHM>
+            AgentReduceByKeyPolicy::BLOCK_THREADS,
+            AgentReduceByKeyPolicy::SCAN_ALGORITHM>
         BlockScanAllocations;
 
     // Callback type for obtaining tile prefix during block scan
@@ -256,7 +256,7 @@ struct BlockReduceSweepByKey
 
     // Constructor
     __device__ __forceinline__
-    BlockReduceSweepByKey(
+    AgentReduceByKey(
         TempStorage                 &temp_storage,      ///< Reference to temp_storage
         KeysInputIteratorT          d_keys_in,          ///< Input keys
         UniqueOutputIteratorT       d_unique_out,       ///< Unique output keys
diff --git a/cub/block_sweep/block_rle_sweep.cuh b/cub/agent/agent_rle.cuh
similarity index 96%
rename from cub/block_sweep/block_rle_sweep.cuh
rename to cub/agent/agent_rle.cuh
index ecebc1495a..cf06eaaa89 100644
--- a/cub/block_sweep/block_rle_sweep.cuh
+++ b/cub/agent/agent_rle.cuh
@@ -28,14 +28,14 @@
 
 /**
  * \file
- * cub::BlockRleSweep implements a stateful abstraction of CUDA thread blocks for participating in device-wide run-length-encode.
+ * cub::AgentRle implements a stateful abstraction of CUDA thread blocks for participating in device-wide run-length-encode.
  */
 
 #pragma once
 
 #include <iterator>
 
-#include "block_scan_prefix_operators.cuh"
+#include "single_pass_scan_operators.cuh"
 #include "../block/block_load.cuh"
 #include "../block/block_store.cuh"
 #include "../block/block_scan.cuh"
@@ -58,7 +58,7 @@ namespace cub {
  ******************************************************************************/
 
 /**
- * Parameterizable tuning policy type for BlockRleSweep
+ * Parameterizable tuning policy type for AgentRle
  */
 template <
     int                         _BLOCK_THREADS,                 ///< Threads per thread block
@@ -67,7 +67,7 @@ template <
     CacheLoadModifier           _LOAD_MODIFIER,                 ///< Cache load modifier for reading input elements
     bool                        _STORE_WARP_TIME_SLICING,       ///< Whether or not only one warp's worth of shared memory should be allocated and time-sliced among block-warps during any store-related data transpositions (versus each warp having its own storage)
     BlockScanAlgorithm          _SCAN_ALGORITHM>                ///< The BlockScan algorithm to use
-struct BlockRleSweepPolicy
+struct AgentRlePolicy
 {
     enum
     {
@@ -90,16 +90,16 @@ struct BlockRleSweepPolicy
  ******************************************************************************/
 
 /**
- * \brief BlockRleSweep implements a stateful abstraction of CUDA thread blocks for participating in device-wide run-length-encode across a range of tiles
+ * \brief AgentRle implements a stateful abstraction of CUDA thread blocks for participating in device-wide run-length-encode across a range of tiles
  */
 template <
-    typename    BlockRleSweepPolicy,      ///< Parameterized BlockRleSweepPolicy tuning policy type
+    typename    AgentRlePolicy,      ///< Parameterized AgentRlePolicy tuning policy type
     typename    InputIteratorT,           ///< Random-access input iterator type for data
     typename    OffsetsOutputIteratorT,   ///< Random-access output iterator type for offset values
     typename    LengthsOutputIteratorT,   ///< Random-access output iterator type for length values
     typename    EqualityOp,               ///< T equality operator type
     typename    OffsetT>                  ///< Signed integer type for global offsets
-struct BlockRleSweep
+struct AgentRle
 {
     //---------------------------------------------------------------------
     // Types and constants
@@ -121,17 +121,17 @@ struct BlockRleSweep
     enum
     {
         WARP_THREADS            = CUB_WARP_THREADS(PTX_ARCH),
-        BLOCK_THREADS           = BlockRleSweepPolicy::BLOCK_THREADS,
-        ITEMS_PER_THREAD        = BlockRleSweepPolicy::ITEMS_PER_THREAD,
+        BLOCK_THREADS           = AgentRlePolicy::BLOCK_THREADS,
+        ITEMS_PER_THREAD        = AgentRlePolicy::ITEMS_PER_THREAD,
         WARP_ITEMS              = WARP_THREADS * ITEMS_PER_THREAD,
         TILE_ITEMS              = BLOCK_THREADS * ITEMS_PER_THREAD,
         WARPS                   = (BLOCK_THREADS + WARP_THREADS - 1) / WARP_THREADS,
 
         /// Whether or not to sync after loading data
-        SYNC_AFTER_LOAD         = (BlockRleSweepPolicy::LOAD_ALGORITHM != BLOCK_LOAD_DIRECT),
+        SYNC_AFTER_LOAD         = (AgentRlePolicy::LOAD_ALGORITHM != BLOCK_LOAD_DIRECT),
 
         /// Whether or not only one warp's worth of shared memory should be allocated and time-sliced among block-warps during any store-related data transpositions (versus each warp having its own storage)
-        STORE_WARP_TIME_SLICING = BlockRleSweepPolicy::STORE_WARP_TIME_SLICING,
+        STORE_WARP_TIME_SLICING = AgentRlePolicy::STORE_WARP_TIME_SLICING,
         ACTIVE_EXCHANGE_WARPS   = (STORE_WARP_TIME_SLICING) ? 1 : WARPS,
     };
 
@@ -168,16 +168,16 @@ struct BlockRleSweep
 
     // Cache-modified Input iterator wrapper type (for applying cache modifier) for data
     typedef typename If<IsPointer<InputIteratorT>::VALUE,
-            CacheModifiedInputIterator<BlockRleSweepPolicy::LOAD_MODIFIER, T, OffsetT>,      // Wrap the native input pointer with CacheModifiedVLengthnputIterator
+            CacheModifiedInputIterator<AgentRlePolicy::LOAD_MODIFIER, T, OffsetT>,      // Wrap the native input pointer with CacheModifiedVLengthnputIterator
             InputIteratorT>::Type                                                            // Directly use the supplied input iterator type
         WrappedInputIteratorT;
 
     // Parameterized BlockLoad type for data
     typedef BlockLoad<
             WrappedInputIteratorT,
-            BlockRleSweepPolicy::BLOCK_THREADS,
-            BlockRleSweepPolicy::ITEMS_PER_THREAD,
-            BlockRleSweepPolicy::LOAD_ALGORITHM>
+            AgentRlePolicy::BLOCK_THREADS,
+            AgentRlePolicy::ITEMS_PER_THREAD,
+            AgentRlePolicy::LOAD_ALGORITHM>
         BlockLoadT;
 
     // Parameterized BlockDiscontinuity type for data
@@ -260,7 +260,7 @@ struct BlockRleSweep
 
     // Constructor
     __device__ __forceinline__
-    BlockRleSweep(
+    AgentRle(
         TempStorage                 &temp_storage,      ///< [in] Reference to temp_storage
         InputIteratorT              d_in,               ///< [in] Pointer to input sequence of data items
         OffsetsOutputIteratorT      d_offsets_out,      ///< [out] Pointer to output sequence of run offsets
diff --git a/cub/block_sweep/block_scan_sweep.cuh b/cub/agent/agent_scan.cuh
similarity index 89%
rename from cub/block_sweep/block_scan_sweep.cuh
rename to cub/agent/agent_scan.cuh
index 7532730a25..ba07574547 100644
--- a/cub/block_sweep/block_scan_sweep.cuh
+++ b/cub/agent/agent_scan.cuh
@@ -28,14 +28,14 @@
 
 /**
  * \file
- * cub::BlockScanSweep implements a stateful abstraction of CUDA thread blocks for participating in device-wide prefix scan across a range of tiles.
+ * cub::AgentScan implements a stateful abstraction of CUDA thread blocks for participating in device-wide prefix scan across a range of tiles.
  */
 
 #pragma once
 
 #include <iterator>
 
-#include "block_scan_prefix_operators.cuh"
+#include "single_pass_scan_operators.cuh"
 #include "../block/block_load.cuh"
 #include "../block/block_store.cuh"
 #include "../block/block_scan.cuh"
@@ -55,25 +55,21 @@ namespace cub {
  ******************************************************************************/
 
 /**
- * Parameterizable tuning policy type for BlockScanSweep
+ * Parameterizable tuning policy type for AgentScan
  */
 template <
     int                         _BLOCK_THREADS,                 ///< Threads per thread block
     int                         _ITEMS_PER_THREAD,              ///< Items per thread (per tile of input)
     BlockLoadAlgorithm          _LOAD_ALGORITHM,                ///< The BlockLoad algorithm to use
-    bool                        _LOAD_WARP_TIME_SLICING,        ///< Whether or not only one warp's worth of shared memory should be allocated and time-sliced among block-warps during any load-related data transpositions (versus each warp having its own storage)
     CacheLoadModifier           _LOAD_MODIFIER,                 ///< Cache load modifier for reading input elements
     BlockStoreAlgorithm         _STORE_ALGORITHM,               ///< The BlockStore algorithm to use
-    bool                        _STORE_WARP_TIME_SLICING,       ///< Whether or not only one warp's worth of shared memory should be allocated and time-sliced among block-warps during any store-related data transpositions (versus each warp having its own storage)
     BlockScanAlgorithm          _SCAN_ALGORITHM>                ///< The BlockScan algorithm to use
-struct BlockScanSweepPolicy
+struct AgentScanPolicy
 {
     enum
     {
         BLOCK_THREADS           = _BLOCK_THREADS,               ///< Threads per thread block
         ITEMS_PER_THREAD        = _ITEMS_PER_THREAD,            ///< Items per thread (per tile of input)
-        LOAD_WARP_TIME_SLICING  = _LOAD_WARP_TIME_SLICING,      ///< Whether or not only one warp's worth of shared memory should be allocated and time-sliced among block-warps during any load-related data transpositions (versus each warp having its own storage)
-        STORE_WARP_TIME_SLICING = _STORE_WARP_TIME_SLICING,     ///< Whether or not only one warp's worth of shared memory should be allocated and time-sliced among block-warps during any store-related data transpositions (versus each warp having its own storage)
     };
 
     static const BlockLoadAlgorithm     LOAD_ALGORITHM          = _LOAD_ALGORITHM;          ///< The BlockLoad algorithm to use
@@ -90,16 +86,16 @@ struct BlockScanSweepPolicy
  ******************************************************************************/
 
 /**
- * \brief BlockScanSweep implements a stateful abstraction of CUDA thread blocks for participating in device-wide prefix scan across a range of tiles.
+ * \brief AgentScan implements a stateful abstraction of CUDA thread blocks for participating in device-wide prefix scan across a range of tiles.
  */
 template <
-    typename BlockScanSweepPolicy,      ///< Parameterized BlockScanSweepPolicy tuning policy type
+    typename AgentScanPolicy,      ///< Parameterized AgentScanPolicy tuning policy type
     typename InputIteratorT,            ///< Random-access input iterator type
     typename OutputIteratorT,           ///< Random-access output iterator type
     typename ScanOp,                    ///< Scan functor type
     typename Identity,                  ///< Identity element type (cub::NullType for inclusive scan)
     typename OffsetT>                   ///< Signed integer type for global offsets
-struct BlockScanSweep
+struct AgentScan
 {
     //---------------------------------------------------------------------
     // Types and constants
@@ -113,7 +109,7 @@ struct BlockScanSweep
 
     // Input iterator wrapper type (for applying cache modifier)
     typedef typename If<IsPointer<InputIteratorT>::VALUE,
-            CacheModifiedInputIterator<BlockScanSweepPolicy::LOAD_MODIFIER, T, OffsetT>,    // Wrap the native input pointer with CacheModifiedInputIterator
+            CacheModifiedInputIterator<AgentScanPolicy::LOAD_MODIFIER, T, OffsetT>,    // Wrap the native input pointer with CacheModifiedInputIterator
             InputIteratorT>::Type                                                            // Directly use the supplied input iterator type
         WrappedInputIteratorT;
 
@@ -121,36 +117,36 @@ struct BlockScanSweep
     enum
     {
         INCLUSIVE           = Equals<Identity, NullType>::VALUE,            // Inclusive scan if no identity type is provided
-        BLOCK_THREADS       = BlockScanSweepPolicy::BLOCK_THREADS,
-        ITEMS_PER_THREAD    = BlockScanSweepPolicy::ITEMS_PER_THREAD,
+        BLOCK_THREADS       = AgentScanPolicy::BLOCK_THREADS,
+        ITEMS_PER_THREAD    = AgentScanPolicy::ITEMS_PER_THREAD,
         TILE_ITEMS          = BLOCK_THREADS * ITEMS_PER_THREAD,
 
         // Whether or not to sync after loading data
-        SYNC_AFTER_LOAD     = (BlockScanSweepPolicy::LOAD_ALGORITHM != BLOCK_LOAD_DIRECT),
+        SYNC_AFTER_LOAD     = (AgentScanPolicy::LOAD_ALGORITHM != BLOCK_LOAD_DIRECT),
 
     };
 
     // Parameterized BlockLoad type
     typedef BlockLoad<
             WrappedInputIteratorT,
-            BlockScanSweepPolicy::BLOCK_THREADS,
-            BlockScanSweepPolicy::ITEMS_PER_THREAD,
-            BlockScanSweepPolicy::LOAD_ALGORITHM>
+            AgentScanPolicy::BLOCK_THREADS,
+            AgentScanPolicy::ITEMS_PER_THREAD,
+            AgentScanPolicy::LOAD_ALGORITHM>
         BlockLoadT;
 
     // Parameterized BlockStore type
     typedef BlockStore<
             OutputIteratorT,
-            BlockScanSweepPolicy::BLOCK_THREADS,
-            BlockScanSweepPolicy::ITEMS_PER_THREAD,
-            BlockScanSweepPolicy::STORE_ALGORITHM>
+            AgentScanPolicy::BLOCK_THREADS,
+            AgentScanPolicy::ITEMS_PER_THREAD,
+            AgentScanPolicy::STORE_ALGORITHM>
         BlockStoreT;
 
     // Parameterized BlockScan type
     typedef BlockScan<
             T,
-            BlockScanSweepPolicy::BLOCK_THREADS,
-            BlockScanSweepPolicy::SCAN_ALGORITHM>
+            AgentScanPolicy::BLOCK_THREADS,
+            AgentScanPolicy::SCAN_ALGORITHM>
         BlockScanT;
 
     // Callback type for obtaining tile prefix during block scan
@@ -293,7 +289,7 @@ struct BlockScanSweep
 
     // Constructor
     __device__ __forceinline__
-    BlockScanSweep(
+    AgentScan(
         TempStorage                 &temp_storage,      ///< Reference to temp_storage
         InputIteratorT              d_in,               ///< Input data
         OutputIteratorT             d_out,              ///< Output data
diff --git a/cub/block_sweep/block_select_sweep.cuh b/cub/agent/agent_select_if.cuh
similarity index 94%
rename from cub/block_sweep/block_select_sweep.cuh
rename to cub/agent/agent_select_if.cuh
index 540ac69f8a..8d5d3e7887 100644
--- a/cub/block_sweep/block_select_sweep.cuh
+++ b/cub/agent/agent_select_if.cuh
@@ -28,14 +28,14 @@
 
 /**
  * \file
- * cub::BlockSelectSweep implements a stateful abstraction of CUDA thread blocks for participating in device-wide select.
+ * cub::AgentSelectIf implements a stateful abstraction of CUDA thread blocks for participating in device-wide select.
  */
 
 #pragma once
 
 #include <iterator>
 
-#include "block_scan_prefix_operators.cuh"
+#include "single_pass_scan_operators.cuh"
 #include "../block/block_load.cuh"
 #include "../block/block_store.cuh"
 #include "../block/block_scan.cuh"
@@ -57,7 +57,7 @@ namespace cub {
  ******************************************************************************/
 
 /**
- * Parameterizable tuning policy type for BlockSelectSweep
+ * Parameterizable tuning policy type for AgentSelectIf
  */
 template <
     int                         _BLOCK_THREADS,                 ///< Threads per thread block
@@ -66,7 +66,7 @@ template <
     CacheLoadModifier           _LOAD_MODIFIER,                 ///< Cache load modifier for reading input elements
     bool                        _STORE_WARP_TIME_SLICING,       ///< Whether or not only one warp's worth of shared memory should be allocated and time-sliced among block-warps during any store-related data transpositions (versus each warp having its own storage)
     BlockScanAlgorithm          _SCAN_ALGORITHM>                ///< The BlockScan algorithm to use
-struct BlockSelectSweepPolicy
+struct AgentSelectIfPolicy
 {
     enum
     {
@@ -88,14 +88,14 @@ struct BlockSelectSweepPolicy
  ******************************************************************************/
 
 /**
- * \brief BlockSelectSweep implements a stateful abstraction of CUDA thread blocks for participating in device-wide selection across a range of tiles
+ * \brief AgentSelectIf implements a stateful abstraction of CUDA thread blocks for participating in device-wide selection across a range of tiles
  *
  * Performs functor-based selection if SelectOp functor type != NullType
  * Otherwise performs flag-based selection if FlagsInputIterator's value type != NullType
  * Otherwise performs discontinuity selection (keep unique)
  */
 template <
-    typename    BlockSelectSweepPolicy,         ///< Parameterized BlockSelectSweepPolicy tuning policy type
+    typename    AgentSelectIfPolicy,         ///< Parameterized AgentSelectIfPolicy tuning policy type
     typename    InputIteratorT,                 ///< Random-access input iterator type for selection items
     typename    FlagsInputIteratorT,                  ///< Random-access input iterator type for selections (NullType* if a selection functor or discontinuity flagging is to be used for selection)
     typename    SelectedOutputIteratorT,                ///< Random-access input iterator type for selected items
@@ -103,7 +103,7 @@ template <
     typename    EqualityOp,                     ///< Equality operator type (NullType if selection functor or selections is to be used for selection)
     typename    OffsetT,                        ///< Signed integer type for global offsets
     bool        KEEP_REJECTS>                   ///< Whether or not we push rejected items to the back of the output
-struct BlockSelectSweep
+struct AgentSelectIf
 {
     //---------------------------------------------------------------------
     // Types and constants
@@ -125,7 +125,7 @@ struct BlockSelectSweep
         USE_SELECT_FLAGS,
         USE_DISCONTINUITY,
 
-        BLOCK_THREADS           = BlockSelectSweepPolicy::BLOCK_THREADS,
+        BLOCK_THREADS           = AgentSelectIfPolicy::BLOCK_THREADS,
 
         /// Number of warp threads
         WARP_THREADS            = CUB_WARP_THREADS(PTX_ARCH),
@@ -133,14 +133,14 @@ struct BlockSelectSweep
         /// Number of active warps
         WARPS                   = (BLOCK_THREADS + WARP_THREADS - 1) / WARP_THREADS,
 
-        ITEMS_PER_THREAD        = BlockSelectSweepPolicy::ITEMS_PER_THREAD,
+        ITEMS_PER_THREAD        = AgentSelectIfPolicy::ITEMS_PER_THREAD,
         TILE_ITEMS              = BLOCK_THREADS * ITEMS_PER_THREAD,
 
         /// Whether or not to sync after loading data
-        SYNC_AFTER_LOAD         = (BlockSelectSweepPolicy::LOAD_ALGORITHM != BLOCK_LOAD_DIRECT),
+        SYNC_AFTER_LOAD         = (AgentSelectIfPolicy::LOAD_ALGORITHM != BLOCK_LOAD_DIRECT),
 
         /// Whether or not only one warp's worth of shared memory should be allocated and time-sliced among block-warps during any store-related data transpositions (versus each warp having its own storage)
-        STORE_WARP_TIME_SLICING = BlockSelectSweepPolicy::STORE_WARP_TIME_SLICING,
+        STORE_WARP_TIME_SLICING = AgentSelectIfPolicy::STORE_WARP_TIME_SLICING,
         ACTIVE_EXCHANGE_WARPS   = (STORE_WARP_TIME_SLICING) ? 1 : WARPS,
 
         SELECT_METHOD           = (!Equals<SelectOp, NullType>::VALUE) ?
@@ -152,30 +152,30 @@ struct BlockSelectSweep
 
     // Input iterator wrapper type (for applying cache modifier)
     typedef typename If<IsPointer<InputIteratorT>::VALUE,
-            CacheModifiedInputIterator<BlockSelectSweepPolicy::LOAD_MODIFIER, T, OffsetT>,      // Wrap the native input pointer with CacheModifiedInputIterator
+            CacheModifiedInputIterator<AgentSelectIfPolicy::LOAD_MODIFIER, T, OffsetT>,      // Wrap the native input pointer with CacheModifiedInputIterator
             InputIteratorT>::Type                                                                // Directly use the supplied input iterator type
         WrappedInputIteratorT;
 
     // FlagT iterator wrapper type
     typedef typename If<IsPointer<FlagsInputIteratorT>::VALUE,
-            CacheModifiedInputIterator<BlockSelectSweepPolicy::LOAD_MODIFIER, FlagT, OffsetT>,   // Wrap the native input pointer with CacheModifiedInputIterator
+            CacheModifiedInputIterator<AgentSelectIfPolicy::LOAD_MODIFIER, FlagT, OffsetT>,   // Wrap the native input pointer with CacheModifiedInputIterator
             FlagsInputIteratorT>::Type                                                                 // Directly use the supplied input iterator type
         WrappedFlagsInputIteratorT;
 
     // Parameterized BlockLoad type for input items
     typedef BlockLoad<
             WrappedInputIteratorT,
-            BlockSelectSweepPolicy::BLOCK_THREADS,
-            BlockSelectSweepPolicy::ITEMS_PER_THREAD,
-            BlockSelectSweepPolicy::LOAD_ALGORITHM>
+            AgentSelectIfPolicy::BLOCK_THREADS,
+            AgentSelectIfPolicy::ITEMS_PER_THREAD,
+            AgentSelectIfPolicy::LOAD_ALGORITHM>
         BlockLoadT;
 
     // Parameterized BlockLoad type for flags
     typedef BlockLoad<
             WrappedFlagsInputIteratorT,
-            BlockSelectSweepPolicy::BLOCK_THREADS,
-            BlockSelectSweepPolicy::ITEMS_PER_THREAD,
-            BlockSelectSweepPolicy::LOAD_ALGORITHM>
+            AgentSelectIfPolicy::BLOCK_THREADS,
+            AgentSelectIfPolicy::ITEMS_PER_THREAD,
+            AgentSelectIfPolicy::LOAD_ALGORITHM>
         BlockLoadFlags;
 
     // Parameterized BlockDiscontinuity type for input items
@@ -249,7 +249,7 @@ struct BlockSelectSweep
 
     // Constructor
     __device__ __forceinline__
-    BlockSelectSweep(
+    AgentSelectIf(
         TempStorage                 &temp_storage,      ///< Reference to temp_storage
         InputIteratorT              d_in,               ///< Input data
         FlagsInputIteratorT         d_flags,            ///< Input flags
diff --git a/cub/block_sweep/block_spmv_sweep.cuh b/cub/agent/agent_spmv.cuh
similarity index 93%
rename from cub/block_sweep/block_spmv_sweep.cuh
rename to cub/agent/agent_spmv.cuh
index e4503ff82d..ce903ccf0d 100644
--- a/cub/block_sweep/block_spmv_sweep.cuh
+++ b/cub/agent/agent_spmv.cuh
@@ -28,7 +28,7 @@
 
 /**
  * \file
- * cub::BlockSpmvSweep implements a stateful abstraction of CUDA thread blocks for participating in device-wide SpMV.
+ * cub::AgentSpmv implements a stateful abstraction of CUDA thread blocks for participating in device-wide SpMV.
  */
 
 #pragma once
@@ -53,7 +53,7 @@ namespace cub {
  ******************************************************************************/
 
 /**
- * Parameterizable tuning policy type for BlockSpmvSweep
+ * Parameterizable tuning policy type for AgentSpmv
  */
 template <
     int                             _BLOCK_THREADS,                         ///< Threads per thread block
@@ -62,7 +62,7 @@ template <
     CacheLoadModifier               _MATRIX_COLUMN_INDICES_LOAD_MODIFIER,   ///< Cache load modifier for reading CSR column-indices
     CacheLoadModifier               _MATRIX_VALUES_LOAD_MODIFIER,           ///< Cache load modifier for reading CSR values
     CacheLoadModifier               _VECTOR_VALUES_LOAD_MODIFIER>           ///< Cache load modifier for reading vector values
-struct BlockSpmvSweepPolicy
+struct AgentSpmvPolicy
 {
     enum
     {
@@ -82,15 +82,15 @@ struct BlockSpmvSweepPolicy
  ******************************************************************************/
 
 /**
- * \brief BlockSpmvSweep implements a stateful abstraction of CUDA thread blocks for participating in device-wide SpMV.
+ * \brief AgentSpmv implements a stateful abstraction of CUDA thread blocks for participating in device-wide SpMV.
  */
 template <
-    typename    BlockSpmvSweepPolicyT,      ///< Parameterized BlockSpmvSweepPolicy tuning policy type
+    typename    AgentSpmvPolicyT,      ///< Parameterized AgentSpmvPolicy tuning policy type
     typename    VertexT,                    ///< Integer type for vertex identifiers
     typename    ValueT,                     ///< Matrix and vector value type
     typename    OffsetT,                    ///< Signed integer type for sequence offsets
     int         PTX_ARCH = CUB_PTX_ARCH>    ///< PTX compute capability
-struct BlockSpmvSweep
+struct AgentSpmv
 {
     //---------------------------------------------------------------------
     // Types and constants
@@ -99,33 +99,33 @@ struct BlockSpmvSweep
     /// Constants
     enum
     {
-        BLOCK_THREADS           = BlockSpmvSweepPolicyT::BLOCK_THREADS,
-        ITEMS_PER_THREAD        = BlockSpmvSweepPolicyT::ITEMS_PER_THREAD,
+        BLOCK_THREADS           = AgentSpmvPolicyT::BLOCK_THREADS,
+        ITEMS_PER_THREAD        = AgentSpmvPolicyT::ITEMS_PER_THREAD,
         TILE_ITEMS              = BLOCK_THREADS * ITEMS_PER_THREAD,
     };
 
     /// Input iterator wrapper types (for applying cache modifiers)
 
     typedef CacheModifiedInputIterator<
-            BlockSpmvSweepPolicyT::MATRIX_ROW_OFFSETS_LOAD_MODIFIER,
+            AgentSpmvPolicyT::MATRIX_ROW_OFFSETS_LOAD_MODIFIER,
             OffsetT,
             OffsetT>
         MatrixValueIteratorT;
 
     typedef CacheModifiedInputIterator<
-            BlockSpmvSweepPolicyT::MATRIX_COLUMN_INDICES_LOAD_MODIFIER,
+            AgentSpmvPolicyT::MATRIX_COLUMN_INDICES_LOAD_MODIFIER,
             VertexT,
             OffsetT>
         MatrixRowOffsetsIteratorT;
 
     typedef CacheModifiedInputIterator<
-            BlockSpmvSweepPolicyT::MATRIX_VALUES_LOAD_MODIFIER,
+            AgentSpmvPolicyT::MATRIX_VALUES_LOAD_MODIFIER,
             ValueT,
             OffsetT>
         MatrixColumnIndicesIteratorT;
 
     typedef CacheModifiedInputIterator<
-            BlockSpmvSweepPolicyT::VECTOR_VALUES_LOAD_MODIFIER,
+            AgentSpmvPolicyT::VECTOR_VALUES_LOAD_MODIFIER,
             ValueT,
             OffsetT>
         VectorValueIteratorT;
@@ -177,7 +177,7 @@ struct BlockSpmvSweep
     /**
      * Constructor
      */
-    __device__ __forceinline__ BlockSpmvSweep(
+    __device__ __forceinline__ AgentSpmv(
         TempStorage &temp_storage,              ///< Reference to temp_storage
         ValueT*     d_matrix_values,            ///< [in] Pointer to the array of \p num_nonzeros values of the corresponding nonzero elements of matrix <b>A</b>.
         OffsetT*    d_matrix_row_offsets,       ///< [in] Pointer to the array of \p m + 1 offsets demarcating the start of every row in \p d_matrix_column_indices and \p d_matrix_values (with the final entry being equal to \p num_nonzeros)
diff --git a/cub/block_sweep/block_scan_prefix_operators.cuh b/cub/agent/single_pass_scan_operators.cuh
similarity index 100%
rename from cub/block_sweep/block_scan_prefix_operators.cuh
rename to cub/agent/single_pass_scan_operators.cuh
diff --git a/cub/block/block_load.cuh b/cub/block/block_load.cuh
index c829b68a06..b7beb52fa5 100644
--- a/cub/block/block_load.cuh
+++ b/cub/block/block_load.cuh
@@ -470,8 +470,7 @@ enum BlockLoadAlgorithm
      * \par Overview
      *
      * A [<em>blocked arrangement</em>](index.html#sec5sec3) of data is read
-     * directly from memory.  The thread block reads items in a parallel "raking" fashion: thread<sub><em>i</em></sub>
-     * reads the <em>i</em><sup>th</sup> segment of consecutive elements.
+     * directly from memory.
      *
      * \par Performance Considerations
      * - The utilization of memory transactions (coalescing) decreases as the
@@ -482,12 +481,10 @@ enum BlockLoadAlgorithm
     /**
      * \par Overview
      *
-     * A [<em>blocked arrangement</em>](index.html#sec5sec3) of data is read directly
+     * A [<em>blocked arrangement</em>](index.html#sec5sec3) of data is read
      * from memory using CUDA's built-in vectorized loads as a coalescing optimization.
-     * The thread block reads items in a parallel "raking" fashion: thread<sub><em>i</em></sub> uses vector loads to
-     * read the <em>i</em><sup>th</sup> segment of consecutive elements.
-     *
-     * For example, <tt>ld.global.v4.s32</tt> instructions will be generated when \p T = \p int and \p ITEMS_PER_THREAD > 4.
+     * For example, <tt>ld.global.v4.s32</tt> instructions will be generated
+     * when \p T = \p int and \p ITEMS_PER_THREAD % 4 == 0.
      *
      * \par Performance Considerations
      * - The utilization of memory transactions (coalescing) remains high until the the
@@ -505,12 +502,8 @@ enum BlockLoadAlgorithm
      * \par Overview
      *
      * A [<em>striped arrangement</em>](index.html#sec5sec3) of data is read
-     * directly from memory and then is locally transposed into a
-     * [<em>blocked arrangement</em>](index.html#sec5sec3). The thread block
-     * reads items in a parallel "strip-mining" fashion:
-     * thread<sub><em>i</em></sub> reads items having stride \p BLOCK_THREADS
-     * between them. cub::BlockExchange is then used to locally reorder the items
-     * into a [<em>blocked arrangement</em>](index.html#sec5sec3).
+     * efficiently from memory and then locally transposed into a
+     * [<em>blocked arrangement</em>](index.html#sec5sec3).
      *
      * \par Performance Considerations
      * - The utilization of memory transactions (coalescing) remains high regardless
@@ -524,12 +517,8 @@ enum BlockLoadAlgorithm
     /**
      * \par Overview
      *
-     * A [<em>warp-striped arrangement</em>](index.html#sec5sec3) of data is read
-     * directly from memory and then is locally transposed into a
-     * [<em>blocked arrangement</em>](index.html#sec5sec3). Each warp reads its own
-     * contiguous segment in a parallel "strip-mining" fashion: lane<sub><em>i</em></sub>
-     * reads items having stride \p WARP_THREADS between them. cub::BlockExchange
-     * is then used to locally reorder the items into a
+     * A [<em>warp-striped arrangement</em>](index.html#sec5sec3) of data is
+     * read efficiently from memory and then locally transposed into a
      * [<em>blocked arrangement</em>](index.html#sec5sec3).
      *
      * \par Usage Considerations
diff --git a/cub/block/block_store.cuh b/cub/block/block_store.cuh
index 9915d5c847..feab038c8e 100644
--- a/cub/block/block_store.cuh
+++ b/cub/block/block_store.cuh
@@ -345,8 +345,7 @@ enum BlockStoreAlgorithm
      * \par Overview
      *
      * A [<em>blocked arrangement</em>](index.html#sec5sec3) of data is written
-     * directly to memory.  The thread block writes items in a parallel "raking" fashion:
-     * thread<sub><em>i</em></sub> writes the <em>i</em><sup>th</sup> segment of consecutive elements.
+     * directly to memory.
      *
      * \par Performance Considerations
      * - The utilization of memory transactions (coalescing) decreases as the
@@ -359,10 +358,8 @@ enum BlockStoreAlgorithm
      *
      * A [<em>blocked arrangement</em>](index.html#sec5sec3) of data is written directly
      * to memory using CUDA's built-in vectorized stores as a coalescing optimization.
-     * The thread block writes items in a parallel "raking" fashion: thread<sub><em>i</em></sub> uses vector stores to
-     * write the <em>i</em><sup>th</sup> segment of consecutive elements.
-     *
-     * For example, <tt>st.global.v4.s32</tt> instructions will be generated when \p T = \p int and \p ITEMS_PER_THREAD > 4.
+     * For example, <tt>st.global.v4.s32</tt> instructions will be generated
+     * when \p T = \p int and \p ITEMS_PER_THREAD % 4 == 0.
      *
      * \par Performance Considerations
      * - The utilization of memory transactions (coalescing) remains high until the the
@@ -379,13 +376,7 @@ enum BlockStoreAlgorithm
     /**
      * \par Overview
      * A [<em>blocked arrangement</em>](index.html#sec5sec3) is locally
-     * transposed into a [<em>striped arrangement</em>](index.html#sec5sec3)
-     * which is then written to memory.  More specifically, cub::BlockExchange
-     * used to locally reorder the items into a
-     * [<em>striped arrangement</em>](index.html#sec5sec3), after which the
-     * thread block writes items in a parallel "strip-mining" fashion: consecutive
-     * items owned by thread<sub><em>i</em></sub> are written to memory with
-     * stride \p BLOCK_THREADS between them.
+     * transposed and then efficiently written to memory as a [<em>striped arrangement</em>](index.html#sec5sec3).
      *
      * \par Performance Considerations
      * - The utilization of memory transactions (coalescing) remains high regardless
@@ -398,13 +389,11 @@ enum BlockStoreAlgorithm
     /**
      * \par Overview
      * A [<em>blocked arrangement</em>](index.html#sec5sec3) is locally
-     * transposed into a [<em>warp-striped arrangement</em>](index.html#sec5sec3)
-     * which is then written to memory.  More specifically, cub::BlockExchange used
-     * to locally reorder the items into a
-     * [<em>warp-striped arrangement</em>](index.html#sec5sec3), after which
-     * each warp writes its own contiguous segment in a parallel "strip-mining" fashion:
-     * consecutive items owned by lane<sub><em>i</em></sub> are written to memory
-     * with stride \p WARP_THREADS between them.
+     * transposed and then efficiently written to memory as a
+     * [<em>warp-striped arrangement</em>](index.html#sec5sec3)
+     *
+     * \par Usage Considerations
+     * - BLOCK_THREADS must be a multiple of WARP_THREADS
      *
      * \par Performance Considerations
      * - The utilization of memory transactions (coalescing) remains high regardless
@@ -413,6 +402,26 @@ enum BlockStoreAlgorithm
      *   direct cub::BLOCK_STORE_DIRECT and cub::BLOCK_STORE_VECTORIZE alternatives.
      */
     BLOCK_STORE_WARP_TRANSPOSE,
+
+    /**
+     * \par Overview
+     * A [<em>blocked arrangement</em>](index.html#sec5sec3) is locally
+     * transposed and then efficiently written to memory as a
+     * [<em>warp-striped arrangement</em>](index.html#sec5sec3)
+     * To reduce the shared memory requirement, only one warp's worth of shared
+     * memory is provisioned and is subsequently time-sliced among warps.
+     *
+     * \par Usage Considerations
+     * - BLOCK_THREADS must be a multiple of WARP_THREADS
+     *
+     * \par Performance Considerations
+     * - The utilization of memory transactions (coalescing) remains high regardless
+     *   of items written per thread.
+     * - Provisions less shared memory temporary storage, but incurs larger
+     *   latencies than the BLOCK_STORE_WARP_TRANSPOSE alternative.
+     */
+    BLOCK_STORE_WARP_TRANSPOSE_TIMESLICED,
+
 };
 
 
@@ -488,7 +497,6 @@ template <
     int                     BLOCK_DIM_X,
     int                     ITEMS_PER_THREAD,
     BlockStoreAlgorithm     ALGORITHM           = BLOCK_STORE_DIRECT,
-    bool                    WARP_TIME_SLICING   = false,
     int                     BLOCK_DIM_Y         = 1,
     int                     BLOCK_DIM_Z         = 1,
     int                     PTX_ARCH            = CUB_PTX_ARCH>
@@ -613,7 +621,7 @@ private:
     struct StoreInternal<BLOCK_STORE_TRANSPOSE, DUMMY>
     {
         // BlockExchange utility type for keys
-        typedef BlockExchange<T, BLOCK_DIM_X, ITEMS_PER_THREAD, WARP_TIME_SLICING, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH> BlockExchange;
+        typedef BlockExchange<T, BLOCK_DIM_X, ITEMS_PER_THREAD, false, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH> BlockExchange;
 
         /// Shared memory storage layout type
         typedef typename BlockExchange::TempStorage _TempStorage;
@@ -672,7 +680,66 @@ private:
         CUB_STATIC_ASSERT((BLOCK_THREADS % WARP_THREADS == 0), "BLOCK_THREADS must be a multiple of WARP_THREADS");
 
         // BlockExchange utility type for keys
-        typedef BlockExchange<T, BLOCK_DIM_X, ITEMS_PER_THREAD, WARP_TIME_SLICING, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH> BlockExchange;
+        typedef BlockExchange<T, BLOCK_DIM_X, ITEMS_PER_THREAD, false, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH> BlockExchange;
+
+        /// Shared memory storage layout type
+        typedef typename BlockExchange::TempStorage _TempStorage;
+
+        /// Alias wrapper allowing storage to be unioned
+        struct TempStorage : Uninitialized<_TempStorage> {};
+
+        /// Thread reference to shared storage
+        _TempStorage &temp_storage;
+
+        /// Linear thread-id
+        int linear_tid;
+
+        /// Constructor
+        __device__ __forceinline__ StoreInternal(
+            TempStorage &temp_storage,
+            int linear_tid)
+        :
+            temp_storage(temp_storage.Alias()),
+            linear_tid(linear_tid)
+        {}
+
+        /// Store items into a linear segment of memory
+        __device__ __forceinline__ void Store(
+            OutputIteratorT   block_itr,                    ///< [in] The thread block's base output iterator for storing to
+            T                 (&items)[ITEMS_PER_THREAD])   ///< [in] Data to store
+        {
+            BlockExchange(temp_storage).BlockedToWarpStriped(items);
+            StoreDirectWarpStriped(linear_tid, block_itr, items);
+        }
+
+        /// Store items into a linear segment of memory, guarded by range
+        __device__ __forceinline__ void Store(
+            OutputIteratorT   block_itr,                    ///< [in] The thread block's base output iterator for storing to
+            T                 (&items)[ITEMS_PER_THREAD],   ///< [in] Data to store
+            int               valid_items)                  ///< [in] Number of valid items to write
+        {
+            BlockExchange(temp_storage).BlockedToWarpStriped(items);
+            StoreDirectWarpStriped(linear_tid, block_itr, items, valid_items);
+        }
+    };
+
+
+    /**
+     * BLOCK_STORE_WARP_TRANSPOSE_TIMESLICED specialization of store helper
+     */
+    template <int DUMMY>
+    struct StoreInternal<BLOCK_STORE_WARP_TRANSPOSE_TIMESLICED, DUMMY>
+    {
+        enum
+        {
+            WARP_THREADS = CUB_WARP_THREADS(PTX_ARCH)
+        };
+
+        // Assert BLOCK_THREADS must be a multiple of WARP_THREADS
+        CUB_STATIC_ASSERT((BLOCK_THREADS % WARP_THREADS == 0), "BLOCK_THREADS must be a multiple of WARP_THREADS");
+
+        // BlockExchange utility type for keys
+        typedef BlockExchange<T, BLOCK_DIM_X, ITEMS_PER_THREAD, true, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH> BlockExchange;
 
         /// Shared memory storage layout type
         typedef typename BlockExchange::TempStorage _TempStorage;
diff --git a/cub/block_sweep/block_unzip_channels_sweep.cuh b/cub/block_sweep/block_unzip_channels_sweep.cuh
deleted file mode 100644
index 5b5a301367..0000000000
--- a/cub/block_sweep/block_unzip_channels_sweep.cuh
+++ /dev/null
@@ -1,217 +0,0 @@
-/******************************************************************************
- * Copyright (c) 2011, Duane Merrill.  All rights reserved.
- * Copyright (c) 2011-2014, NVIDIA CORPORATION.  All rights reserved.
- *
- * Redistribution and use in source and binary forms, with or without
- * modification, are permitted provided that the following conditions are met:
- *     * Redistributions of source code must retain the above copyright
- *       notice, this list of conditions and the following disclaimer.
- *     * Redistributions in binary form must reproduce the above copyright
- *       notice, this list of conditions and the following disclaimer in the
- *       documentation and/or other materials provided with the distribution.
- *     * Neither the name of the NVIDIA CORPORATION nor the
- *       names of its contributors may be used to endorse or promote products
- *       derived from this software without specific prior written permission.
- *
- * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
- * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
- * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
- * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
- * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
- * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
- * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
- * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
- * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
- * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
- *
- ******************************************************************************/
-
-/**
- * \file
- * cub::BlockUnzipChannelsSweep implements a stateful abstraction of CUDA thread blocks for unzipping "pixels" of interleaved data channels
- */
-
-#pragma once
-
-#include <iterator>
-
-#include "../util_type.cuh"
-#include "../grid/grid_mapping.cuh"
-#include "../grid/grid_even_share.cuh"
-#include "../grid/grid_queue.cuh"
-#include "../util_namespace.cuh"
-
-/// Optional outer namespace(s)
-CUB_NS_PREFIX
-
-/// CUB namespace
-namespace cub {
-
-
-
-
-/******************************************************************************
- * Tuning policy
- ******************************************************************************/
-
-/**
- * Parameterizable tuning policy type for BlockScanSweep
- */
-template <
-    int                     _BLOCK_THREADS,                 ///< Threads per thread block
-    int                     _PIXELS_PER_THREAD>             ///< Pixels per thread (per tile of input)
-struct BlockUnzipChannelsSweepPolicy
-{
-    enum
-    {
-        BLOCK_THREADS       = _BLOCK_THREADS,               ///< Threads per thread block
-        PIXELS_PER_THREAD   = _PIXELS_PER_THREAD,           ///< Pixels per thread (per tile of input)
-    };
-};
-
-
-/******************************************************************************
- * Thread block abstractions
- ******************************************************************************/
-
-
-/**
- * \brief BlockUnzipChannelsSweep implements a stateful abstraction of CUDA thread blocks for unzipping interleaved data channels
- */
-template <
-    typename    BlockUnzipChannelsSweepPolicyT, ///< Parameterized BlockUnzipChannelsSweepPolicy tuning policy type
-    int         NUM_CHANNELS,                   ///< Number of channels interleaved in the input data
-    int         NUM_ACTIVE_CHANNELS,            ///< Number of channels actively being histogrammed
-    typename    InputIteratorT,                 ///< Random-access input iterator type for reading samples.
-    typename    OutputIteratorT,                ///< Random-access output iterator type
-    typename    OffsetT>                        ///< Signed integer type for global offsets
-struct BlockUnzipChannelsSweep
-{
-    //---------------------------------------------------------------------
-    // Types and constants
-    //---------------------------------------------------------------------
-
-    // Data type of input iterator
-    typedef typename std::iterator_traits<InputIteratorT>::value_type SampleT;
-
-    enum
-    {
-        BLOCK_THREADS       = BlockUnzipChannelsSweepPolicyT::TILE_PIXELS,
-        PIXELS_PER_THREAD   = BlockUnzipChannelsSweepPolicyT::PIXELS_PER_THREAD,
-        TILE_PIXELS         = BLOCK_THREADS * PIXELS_PER_THREAD,
-        TILE_SAMPLES        = TILE_PIXELS * NUM_CHANNELS,
-    };
-
-    //---------------------------------------------------------------------
-    // Per-thread fields
-    //---------------------------------------------------------------------
-
-    InputIteratorT          d_in;                           ///< Pointer to input data
-    OutputIteratorT         (&d_out)[NUM_ACTIVE_CHANNELS];  ///< Pointers to output data
-
-    //---------------------------------------------------------------------
-    // Interface
-    //---------------------------------------------------------------------
-
-    /**
-     * Constructor
-     */
-    __device__ __forceinline__ BlockUnzipChannelsSweep(
-        InputIteratorT       d_in,                           ///< Pointer to input data
-        OutputIteratorT      (&d_out)[NUM_ACTIVE_CHANNELS])  ///< Pointers to output data
-    :
-        d_in(d_in), d_out(d_out)
-    {}
-
-
-    /**
-     * Process a single tile of input
-     */
-    template <bool FULL_TILE>
-    __device__ __forceinline__ void ConsumeTile(
-        OffsetT             block_offset_in,                                    ///< The offset the tile to consume
-        OffsetT             block_offset_out,                                   ///< The offset the tile to produce in each channel
-        int                 valid_pixels = TILE_PIXELS)                         ///< The number of valid pixels in the tile
-    {
-        if (FULL_TILE)
-        {
-            // Full tile of samples to read and composite
-            SampleT samples[PIXELS_PER_THREAD][NUM_ACTIVE_CHANNELS];
-
-            #pragma unroll
-            for (int PIXEL = 0; PIXEL < PIXELS_PER_THREAD; PIXEL++)
-            {
-                #pragma unroll
-                for (int CHANNEL = 0; CHANNEL < NUM_ACTIVE_CHANNELS; ++CHANNEL)
-                {
-                    samples[PIXEL][CHANNEL] = d_in[block_offset_in + (PIXEL * BLOCK_THREADS * NUM_CHANNELS) + (threadIdx.x * NUM_CHANNELS) + CHANNEL];
-                }
-            }
-
-            __threadfence_block();
-
-            #pragma unroll
-            for (int PIXEL = 0; PIXEL < PIXELS_PER_THREAD; PIXEL++)
-            {
-                #pragma unroll
-                for (int CHANNEL = 0; CHANNEL < NUM_ACTIVE_CHANNELS; ++CHANNEL)
-                {
-                    d_out[CHANNEL][block_offset_out + (PIXEL * BLOCK_THREADS) + threadIdx.x] = samples[PIXEL][CHANNEL];
-                }
-            }
-
-            __threadfence_block();
-        }
-        else
-        {
-            // Only a partially-full tile of samples to read and composite
-            #pragma unroll
-            for (int PIXEL = 0; PIXEL < PIXELS_PER_THREAD; ++PIXEL)
-            {
-                if (PIXEL * BLOCK_THREADS < valid_pixels - threadIdx.x)
-                {
-                    #pragma unroll
-                    for (int CHANNEL = 0; CHANNEL < NUM_ACTIVE_CHANNELS; ++CHANNEL)
-                    {
-                        SampleT sample = d_in[block_offset_in + (PIXEL * BLOCK_THREADS * NUM_CHANNELS) + (threadIdx.x * NUM_CHANNELS) + CHANNEL];
-                        d_out[CHANNEL][block_offset_out + (PIXEL * BLOCK_THREADS) + threadIdx.x] = sample;
-                    }
-                }
-            }
-        }
-    }
-
-
-    /**
-     * \brief Consume striped tiles
-     */
-    __device__ __forceinline__ void ConsumeStriped(
-        OffsetT             block_offset_in,                ///< [in] Threadblock input begin offset
-        OffsetT             block_end,                      ///< [in] Threadblock input end offset (exclusive)
-        OffsetT             block_offset_out)               ///< [in] Threadblock out begin offset
-    {
-        // Consume subsequent full tiles of input
-        while (block_offset_in + TILE_SAMPLES <= block_end)
-        {
-            ConsumeTile<true>(block_offset_in, block_offset_out);
-            block_offset_in += TILE_SAMPLES;
-            block_offset_out += TILE_PIXELS;
-        }
-
-        // Consume a partially-full tile
-        if (block_offset_in < block_end)
-        {
-            int valid_pixels = (block_end - block_offset_in) / NUM_CHANNELS;
-            ConsumeTile<false>(block_offset_in, block_offset_out, valid_pixels);
-        }
-    }
-
-};
-
-
-
-
-
-}               // CUB namespace
-CUB_NS_POSTFIX  // Optional outer namespace(s)
-
diff --git a/cub/device/device_histogram.cuh b/cub/device/device_histogram.cuh
index d9a38ced86..2c143a8dac 100644
--- a/cub/device/device_histogram.cuh
+++ b/cub/device/device_histogram.cuh
@@ -38,7 +38,7 @@
 #include <iterator>
 #include <limits>
 
-#include "dispatch/device_histogram_dispatch.cuh"
+#include "dispatch/dispatch_histogram.cuh"
 #include "../util_namespace.cuh"
 
 /// Optional outer namespace(s)
@@ -456,13 +456,13 @@ struct DeviceHistogram
             // Down-convert OffsetT data type
 
 
-            return DeviceHistogramDispatch<NUM_CHANNELS, NUM_ACTIVE_CHANNELS, SampleIteratorT, CounterT, LevelT, int>::DispatchEven(
+            return DipatchHistogram<NUM_CHANNELS, NUM_ACTIVE_CHANNELS, SampleIteratorT, CounterT, LevelT, int>::DispatchEven(
                 d_temp_storage, temp_storage_bytes, d_samples, d_histogram, num_levels, lower_level, upper_level,
                 (int) num_row_pixels, (int) num_rows, (int) (row_stride_bytes / sizeof(SampleT)),
                 stream, debug_synchronous, is_byte_sample);
         }
 
-        return DeviceHistogramDispatch<NUM_CHANNELS, NUM_ACTIVE_CHANNELS, SampleIteratorT, CounterT, LevelT, OffsetT>::DispatchEven(
+        return DipatchHistogram<NUM_CHANNELS, NUM_ACTIVE_CHANNELS, SampleIteratorT, CounterT, LevelT, OffsetT>::DispatchEven(
             d_temp_storage, temp_storage_bytes, d_samples, d_histogram, num_levels, lower_level, upper_level,
             num_row_pixels, num_rows, (OffsetT) (row_stride_bytes / sizeof(SampleT)),
             stream, debug_synchronous, is_byte_sample);
@@ -849,13 +849,13 @@ struct DeviceHistogram
         if ((sizeof(OffsetT) > sizeof(int)) && (row_stride_bytes * num_rows < std::numeric_limits<int>::max()))
         {
             // Down-convert OffsetT data type
-            return DeviceHistogramDispatch<NUM_CHANNELS, NUM_ACTIVE_CHANNELS, SampleIteratorT, CounterT, LevelT, int>::DispatchRange(
+            return DipatchHistogram<NUM_CHANNELS, NUM_ACTIVE_CHANNELS, SampleIteratorT, CounterT, LevelT, int>::DispatchRange(
                 d_temp_storage, temp_storage_bytes, d_samples, d_histogram, num_levels, d_levels,
                 (int) num_row_pixels, (int) num_rows, (int) (row_stride_bytes / sizeof(SampleT)),
                 stream, debug_synchronous, is_byte_sample);
         }
 
-        return DeviceHistogramDispatch<NUM_CHANNELS, NUM_ACTIVE_CHANNELS, SampleIteratorT, CounterT, LevelT, OffsetT>::DispatchRange(
+        return DipatchHistogram<NUM_CHANNELS, NUM_ACTIVE_CHANNELS, SampleIteratorT, CounterT, LevelT, OffsetT>::DispatchRange(
             d_temp_storage, temp_storage_bytes, d_samples, d_histogram, num_levels, d_levels,
             num_row_pixels, num_rows, (OffsetT) (row_stride_bytes / sizeof(SampleT)),
             stream, debug_synchronous, is_byte_sample);
diff --git a/cub/device/device_partition.cuh b/cub/device/device_partition.cuh
index cfd1e35ea3..eb97184c96 100644
--- a/cub/device/device_partition.cuh
+++ b/cub/device/device_partition.cuh
@@ -37,7 +37,7 @@
 #include <stdio.h>
 #include <iterator>
 
-#include "dispatch/device_select_dispatch.cuh"
+#include "dispatch/dispatch_select_if.cuh"
 #include "../util_namespace.cuh"
 
 /// Optional outer namespace(s)
@@ -139,7 +139,7 @@ struct DevicePartition
         typedef NullType                SelectOp;       // Selection op (not used)
         typedef NullType                EqualityOp;     // Equality operator (not used)
 
-        return DeviceSelectDispatch<InputIteratorT, FlagIterator, OutputIteratorT, NumSelectedIteratorT, SelectOp, EqualityOp, OffsetT, true>::Dispatch(
+        return DispatchSelectIf<InputIteratorT, FlagIterator, OutputIteratorT, NumSelectedIteratorT, SelectOp, EqualityOp, OffsetT, true>::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_in,
@@ -248,7 +248,7 @@ struct DevicePartition
         typedef NullType*               FlagIterator;   // FlagT iterator type (not used)
         typedef NullType                EqualityOp;     // Equality operator (not used)
 
-        return DeviceSelectDispatch<InputIteratorT, FlagIterator, OutputIteratorT, NumSelectedIteratorT, SelectOp, EqualityOp, OffsetT, true>::Dispatch(
+        return DispatchSelectIf<InputIteratorT, FlagIterator, OutputIteratorT, NumSelectedIteratorT, SelectOp, EqualityOp, OffsetT, true>::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_in,
diff --git a/cub/device/device_radix_sort.cuh b/cub/device/device_radix_sort.cuh
index 9050a41c5d..f8fdc6d327 100644
--- a/cub/device/device_radix_sort.cuh
+++ b/cub/device/device_radix_sort.cuh
@@ -37,7 +37,7 @@
 #include <stdio.h>
 #include <iterator>
 
-#include "dispatch/device_radix_sort_dispatch.cuh"
+#include "dispatch/dispatch_radix_sort.cuh"
 #include "../util_namespace.cuh"
 
 /// Optional outer namespace(s)
@@ -163,7 +163,7 @@ struct DeviceRadixSort
         DoubleBuffer<Key>       d_keys(d_keys_in, d_keys_out);
         DoubleBuffer<Value>     d_values(d_values_in, d_values_out);
 
-        return DeviceRadixSortDispatch<false, true, Key, Value, OffsetT>::Dispatch(
+        return DispatchRadixSort<false, true, Key, Value, OffsetT>::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_keys,
@@ -259,7 +259,7 @@ struct DeviceRadixSort
         // Signed integer type for global offsets
         typedef int OffsetT;
 
-        return DeviceRadixSortDispatch<false, false, Key, Value, OffsetT>::Dispatch(
+        return DispatchRadixSort<false, false, Key, Value, OffsetT>::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_keys,
@@ -344,7 +344,7 @@ struct DeviceRadixSort
         DoubleBuffer<Key>       d_keys(d_keys_in, d_keys_out);
         DoubleBuffer<Value>     d_values(d_values_in, d_values_out);
 
-        return DeviceRadixSortDispatch<true, true, Key, Value, OffsetT>::Dispatch(
+        return DispatchRadixSort<true, true, Key, Value, OffsetT>::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_keys,
@@ -435,7 +435,7 @@ struct DeviceRadixSort
         // Signed integer type for global offsets
         typedef int OffsetT;
 
-        return DeviceRadixSortDispatch<true, false, Key, Value, OffsetT>::Dispatch(
+        return DispatchRadixSort<true, false, Key, Value, OffsetT>::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_keys,
@@ -521,7 +521,7 @@ struct DeviceRadixSort
         DoubleBuffer<Key>       d_keys(d_keys_in, d_keys_out);
         DoubleBuffer<NullType>  d_values;
 
-        return DeviceRadixSortDispatch<false, true, Key, NullType, OffsetT>::Dispatch(
+        return DispatchRadixSort<false, true, Key, NullType, OffsetT>::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_keys,
@@ -608,7 +608,7 @@ struct DeviceRadixSort
         // Null value type
         DoubleBuffer<NullType> d_values;
 
-        return DeviceRadixSortDispatch<false, false, Key, NullType, OffsetT>::Dispatch(
+        return DispatchRadixSort<false, false, Key, NullType, OffsetT>::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_keys,
@@ -684,7 +684,7 @@ struct DeviceRadixSort
         DoubleBuffer<Key>       d_keys(d_keys_in, d_keys_out);
         DoubleBuffer<NullType>  d_values;
 
-        return DeviceRadixSortDispatch<true, false, Key, NullType, OffsetT>::Dispatch(
+        return DispatchRadixSort<true, false, Key, NullType, OffsetT>::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_keys,
@@ -767,7 +767,7 @@ struct DeviceRadixSort
         // Null value type
         DoubleBuffer<NullType> d_values;
 
-        return DeviceRadixSortDispatch<true, false, Key, NullType, OffsetT>::Dispatch(
+        return DispatchRadixSort<true, false, Key, NullType, OffsetT>::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_keys,
diff --git a/cub/device/device_reduce.cuh b/cub/device/device_reduce.cuh
index 5222a6d787..d22b49f3d4 100644
--- a/cub/device/device_reduce.cuh
+++ b/cub/device/device_reduce.cuh
@@ -37,8 +37,8 @@
 #include <stdio.h>
 #include <iterator>
 
-#include "dispatch/device_reduce_dispatch.cuh"
-#include "dispatch/device_reduce_by_key_dispatch.cuh"
+#include "dispatch/dispatch_reduce.cuh"
+#include "dispatch/dispatch_reduce_by_key.cuh"
 #include "../util_namespace.cuh"
 
 /// Optional outer namespace(s)
@@ -153,9 +153,9 @@ struct DeviceReduce
         typedef int OffsetT;
 
         // Dispatch type
-        typedef DeviceReduceDispatch<InputIteratorT, OutputIteratorT, OffsetT, ReductionOp> DeviceReduceDispatch;
+        typedef DispatchReduce<InputIteratorT, OutputIteratorT, OffsetT, ReductionOp> DispatchReduce;
 
-        return DeviceReduceDispatch::Dispatch(
+        return DispatchReduce::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_in,
@@ -229,9 +229,9 @@ struct DeviceReduce
         typedef int OffsetT;
 
         // Dispatch type
-        typedef DeviceReduceDispatch<InputIteratorT, OutputIteratorT, OffsetT, cub::Sum> DeviceReduceDispatch;
+        typedef DispatchReduce<InputIteratorT, OutputIteratorT, OffsetT, cub::Sum> DispatchReduce;
 
-        return DeviceReduceDispatch::Dispatch(
+        return DispatchReduce::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_in,
@@ -301,9 +301,9 @@ struct DeviceReduce
         typedef int OffsetT;
 
         // Dispatch type
-        typedef DeviceReduceDispatch<InputIteratorT, OutputIteratorT, OffsetT, cub::Min> DeviceReduceDispatch;
+        typedef DispatchReduce<InputIteratorT, OutputIteratorT, OffsetT, cub::Min> DispatchReduce;
 
-        return DeviceReduceDispatch::Dispatch(
+        return DispatchReduce::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_in,
@@ -382,9 +382,9 @@ struct DeviceReduce
         ArgIndexInputIteratorT d_argmin_in(d_in, 0);
 
         // Dispatch type
-        typedef DeviceReduceDispatch<ArgIndexInputIteratorT, OutputIteratorT, OffsetT, cub::ArgMin> DeviceReduceDispatch;
+        typedef DispatchReduce<ArgIndexInputIteratorT, OutputIteratorT, OffsetT, cub::ArgMin> DispatchReduce;
 
-        return DeviceReduceDispatch::Dispatch(
+        return DispatchReduce::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_argmin_in,
@@ -454,9 +454,9 @@ struct DeviceReduce
         typedef int OffsetT;
 
         // Dispatch type
-        typedef DeviceReduceDispatch<InputIteratorT, OutputIteratorT, OffsetT, cub::Max> DeviceReduceDispatch;
+        typedef DispatchReduce<InputIteratorT, OutputIteratorT, OffsetT, cub::Max> DispatchReduce;
 
-        return DeviceReduceDispatch::Dispatch(
+        return DispatchReduce::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_in,
@@ -535,9 +535,9 @@ struct DeviceReduce
         ArgIndexInputIteratorT d_argmax_in(d_in, 0);
 
         // Dispatch type
-        typedef DeviceReduceDispatch<ArgIndexInputIteratorT, OutputIteratorT, OffsetT, cub::ArgMax> DeviceReduceDispatch;
+        typedef DispatchReduce<ArgIndexInputIteratorT, OutputIteratorT, OffsetT, cub::ArgMax> DispatchReduce;
 
-        return DeviceReduceDispatch::Dispatch(
+        return DispatchReduce::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_argmax_in,
@@ -657,7 +657,7 @@ struct DeviceReduce
         typedef NullType            SelectOp;       // Selection op (not used)
         typedef Equality            EqualityOp;     // Default == operator
 
-        return DeviceReduceByKeyDispatch<KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOp, ReductionOp, OffsetT>::Dispatch(
+        return DispatchReduceByKey<KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOp, ReductionOp, OffsetT>::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_keys_in,
diff --git a/cub/device/device_run_length_encode.cuh b/cub/device/device_run_length_encode.cuh
index 71b17b704b..5499bb24c1 100644
--- a/cub/device/device_run_length_encode.cuh
+++ b/cub/device/device_run_length_encode.cuh
@@ -37,8 +37,8 @@
 #include <stdio.h>
 #include <iterator>
 
-#include "dispatch/device_rle_dispatch.cuh"
-#include "dispatch/device_reduce_by_key_dispatch.cuh"
+#include "dispatch/dispatch_rle.cuh"
+#include "dispatch/dispatch_reduce_by_key.cuh"
 #include "../util_namespace.cuh"
 
 /// Optional outer namespace(s)
@@ -172,7 +172,7 @@ struct DeviceRunLengthEncode
         Value one_val;
         one_val = 1;
 
-        return DeviceReduceByKeyDispatch<InputIteratorT, UniqueOutputIteratorT, LengthsInputIteratorT, LengthsOutputIteratorT, NumRunsOutputIteratorT, EqualityOp, ReductionOp, OffsetT>::Dispatch(
+        return DispatchReduceByKey<InputIteratorT, UniqueOutputIteratorT, LengthsInputIteratorT, LengthsOutputIteratorT, NumRunsOutputIteratorT, EqualityOp, ReductionOp, OffsetT>::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_in,
diff --git a/cub/device/device_scan.cuh b/cub/device/device_scan.cuh
index 4ffe2fbfd7..22c3a12caf 100644
--- a/cub/device/device_scan.cuh
+++ b/cub/device/device_scan.cuh
@@ -37,7 +37,7 @@
 #include <stdio.h>
 #include <iterator>
 
-#include "dispatch/device_scan_dispatch.cuh"
+#include "dispatch/dispatch_scan.cuh"
 #include "../util_namespace.cuh"
 
 /// Optional outer namespace(s)
@@ -145,7 +145,7 @@ struct DeviceScan
         // Scan data type
         typedef typename std::iterator_traits<InputIteratorT>::value_type T;
 
-        return DeviceScanDispatch<InputIteratorT, OutputIteratorT, Sum, T, OffsetT>::Dispatch(
+        return DispatchScan<InputIteratorT, OutputIteratorT, Sum, T, OffsetT>::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_in,
@@ -232,7 +232,7 @@ struct DeviceScan
         // Signed integer type for global offsets
         typedef int OffsetT;
 
-        return DeviceScanDispatch<InputIteratorT, OutputIteratorT, ScanOp, Identity, OffsetT>::Dispatch(
+        return DispatchScan<InputIteratorT, OutputIteratorT, ScanOp, Identity, OffsetT>::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_in,
@@ -309,7 +309,7 @@ struct DeviceScan
         // Signed integer type for global offsets
         typedef int OffsetT;
 
-        return DeviceScanDispatch<InputIteratorT, OutputIteratorT, Sum, NullType, OffsetT>::Dispatch(
+        return DispatchScan<InputIteratorT, OutputIteratorT, Sum, NullType, OffsetT>::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_in,
@@ -393,7 +393,7 @@ struct DeviceScan
         // Signed integer type for global offsets
         typedef int OffsetT;
 
-        return DeviceScanDispatch<InputIteratorT, OutputIteratorT, ScanOp, NullType, OffsetT>::Dispatch(
+        return DispatchScan<InputIteratorT, OutputIteratorT, ScanOp, NullType, OffsetT>::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_in,
diff --git a/cub/device/device_select.cuh b/cub/device/device_select.cuh
index a8b5430520..c38970dd38 100644
--- a/cub/device/device_select.cuh
+++ b/cub/device/device_select.cuh
@@ -37,7 +37,7 @@
 #include <stdio.h>
 #include <iterator>
 
-#include "dispatch/device_select_dispatch.cuh"
+#include "dispatch/dispatch_select_if.cuh"
 #include "../util_namespace.cuh"
 
 /// Optional outer namespace(s)
@@ -146,7 +146,7 @@ struct DeviceSelect
         typedef NullType                SelectOp;       // Selection op (not used)
         typedef NullType                EqualityOp;     // Equality operator (not used)
 
-        return DeviceSelectDispatch<InputIteratorT, FlagIterator, OutputIteratorT, NumSelectedIteratorT, SelectOp, EqualityOp, OffsetT, false>::Dispatch(
+        return DispatchSelectIf<InputIteratorT, FlagIterator, OutputIteratorT, NumSelectedIteratorT, SelectOp, EqualityOp, OffsetT, false>::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_in,
@@ -253,7 +253,7 @@ struct DeviceSelect
         typedef NullType*               FlagIterator;   // FlagT iterator type (not used)
         typedef NullType                EqualityOp;     // Equality operator (not used)
 
-        return DeviceSelectDispatch<InputIteratorT, FlagIterator, OutputIteratorT, NumSelectedIteratorT, SelectOp, EqualityOp, OffsetT, false>::Dispatch(
+        return DispatchSelectIf<InputIteratorT, FlagIterator, OutputIteratorT, NumSelectedIteratorT, SelectOp, EqualityOp, OffsetT, false>::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_in,
@@ -344,7 +344,7 @@ struct DeviceSelect
         typedef NullType                SelectOp;       // Selection op (not used)
         typedef Equality                EqualityOp;     // Default == operator
 
-        return DeviceSelectDispatch<InputIteratorT, FlagIterator, OutputIteratorT, NumSelectedIteratorT, SelectOp, EqualityOp, OffsetT, false>::Dispatch(
+        return DispatchSelectIf<InputIteratorT, FlagIterator, OutputIteratorT, NumSelectedIteratorT, SelectOp, EqualityOp, OffsetT, false>::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_in,
diff --git a/cub/device/device_spmv.cuh b/cub/device/device_spmv.cuh
index b5597c6df1..d82d922693 100644
--- a/cub/device/device_spmv.cuh
+++ b/cub/device/device_spmv.cuh
@@ -38,7 +38,7 @@
 #include <iterator>
 #include <limits>
 
-#include "dispatch/device_histogram_dispatch.cuh"
+#include "dispatch/dispatch_spmv.cuh"
 #include "../util_namespace.cuh"
 
 /// Optional outer namespace(s)
diff --git a/cub/device/dispatch/device_histogram_dispatch.cuh b/cub/device/dispatch/dispatch_histogram.cuh
similarity index 98%
rename from cub/device/dispatch/device_histogram_dispatch.cuh
rename to cub/device/dispatch/dispatch_histogram.cuh
index 3124cfb240..fdc226f3d1 100644
--- a/cub/device/dispatch/device_histogram_dispatch.cuh
+++ b/cub/device/dispatch/dispatch_histogram.cuh
@@ -38,7 +38,7 @@
 #include <iterator>
 #include <limits>
 
-#include "../../block_sweep/block_histogram_sweep.cuh"
+#include "../../agent/agent_histogram.cuh"
 #include "../../util_debug.cuh"
 #include "../../util_device.cuh"
 #include "../../thread/thread_search.cuh"
@@ -87,7 +87,7 @@ __global__ void DeviceHistogramInitKernel(
  * Histogram privatized sweep kernel entry point (multi-block).  Computes privatized histograms, one per thread block.
  */
 template <
-    typename                                            BlockHistogramSweepPolicyT,     ///< Parameterized BlockHistogramSweepPolicy tuning policy type
+    typename                                            AgentHistogramPolicyT,     ///< Parameterized AgentHistogramPolicy tuning policy type
     int                                                 PRIVATIZED_SMEM_BINS,           ///< Maximum number of histogram bins per channel (e.g., up to 256)
     int                                                 NUM_CHANNELS,                   ///< Number of channels interleaved in the input data (may be greater than the number of channels being actively histogrammed)
     int                                                 NUM_ACTIVE_CHANNELS,            ///< Number of channels actively being histogrammed
@@ -96,7 +96,7 @@ template <
     typename                                            PrivatizedDecodeOpT,            ///< The transform operator type for determining privatized counter indices from samples, one for each channel
     typename                                            OutputDecodeOpT,                ///< The transform operator type for determining output bin-ids from privatized counter indices, one for each channel
     typename                                            OffsetT>                        ///< Signed integer type for global offsets
-__launch_bounds__ (int(BlockHistogramSweepPolicyT::BLOCK_THREADS))
+__launch_bounds__ (int(AgentHistogramPolicyT::BLOCK_THREADS))
 __global__ void DeviceHistogramSweepKernel(
     SampleIteratorT                                         d_samples,                          ///< Input data to reduce
     ArrayWrapper<int, NUM_ACTIVE_CHANNELS>                  num_output_bins_wrapper,            ///< The number bins per final output histogram
@@ -112,8 +112,8 @@ __global__ void DeviceHistogramSweepKernel(
     GridQueue<int>                                          tile_queue)                         ///< Drain queue descriptor for dynamically mapping tile data onto thread blocks
 {
     // Thread block type for compositing input tiles
-    typedef BlockHistogramSweep<
-            BlockHistogramSweepPolicyT,
+    typedef AgentHistogram<
+            AgentHistogramPolicyT,
             PRIVATIZED_SMEM_BINS,
             NUM_CHANNELS,
             NUM_ACTIVE_CHANNELS,
@@ -122,12 +122,12 @@ __global__ void DeviceHistogramSweepKernel(
             PrivatizedDecodeOpT,
             OutputDecodeOpT,
             OffsetT>
-        BlockHistogramSweepT;
+        AgentHistogramT;
 
-    // Shared memory for BlockHistogramSweep
-    __shared__ typename BlockHistogramSweepT::TempStorage temp_storage;
+    // Shared memory for AgentHistogram
+    __shared__ typename AgentHistogramT::TempStorage temp_storage;
 
-    BlockHistogramSweepT block_sweep(
+    AgentHistogramT agent(
         temp_storage,
         d_samples,
         num_output_bins_wrapper.array,
@@ -138,10 +138,10 @@ __global__ void DeviceHistogramSweepKernel(
         privatized_decode_op_wrapper.array);
 
     // Initialize counters
-    block_sweep.InitBinCounters();
+    agent.InitBinCounters();
 
     // Consume input tiles
-    block_sweep.ConsumeTiles(
+    agent.ConsumeTiles(
         num_row_pixels,
         num_rows,
         row_stride_samples,
@@ -149,7 +149,7 @@ __global__ void DeviceHistogramSweepKernel(
         tile_queue);
 
     // Store output to global (if necessary)
-    block_sweep.StoreOutput();
+    agent.StoreOutput();
 }
 
 
@@ -171,7 +171,7 @@ template <
     typename    CounterT,                   ///< Integer type for counting sample occurrences per histogram bin
     typename    LevelT,                     ///< Type for specifying bin level boundaries
     typename    OffsetT>                    ///< Signed integer type for global offsets
-struct DeviceHistogramDispatch
+struct DipatchHistogram
 {
     //---------------------------------------------------------------------
     // Types and constants
@@ -339,7 +339,7 @@ struct DeviceHistogramDispatch
     struct Policy110
     {
         // HistogramSweepPolicy
-        typedef BlockHistogramSweepPolicy<
+        typedef AgentHistogramPolicy<
                 512,
                 2,
                 (NUM_CHANNELS == 1) ? BLOCK_LOAD_VECTORIZE : BLOCK_LOAD_DIRECT,
@@ -354,7 +354,7 @@ struct DeviceHistogramDispatch
     struct Policy200
     {
         // HistogramSweepPolicy
-        typedef BlockHistogramSweepPolicy<
+        typedef AgentHistogramPolicy<
                 (NUM_CHANNELS == 1) ? 256 : 128,
                 (NUM_CHANNELS == 1) ? 8 : 3,
                 (NUM_CHANNELS == 1) ? BLOCK_LOAD_VECTORIZE : BLOCK_LOAD_WARP_TRANSPOSE,
@@ -372,7 +372,7 @@ struct DeviceHistogramDispatch
     struct Policy350
     {
         // HistogramSweepPolicy
-        typedef BlockHistogramSweepPolicy<
+        typedef AgentHistogramPolicy<
                 128,
                 (NUM_CHANNELS == 1) ? 8 : 7,
                 (NUM_CHANNELS == 1) ? BLOCK_LOAD_VECTORIZE : BLOCK_LOAD_DIRECT,
@@ -387,7 +387,7 @@ struct DeviceHistogramDispatch
     struct Policy500
     {
         // HistogramSweepPolicy
-        typedef BlockHistogramSweepPolicy<
+        typedef AgentHistogramPolicy<
                 256,
                 8,
                 (NUM_CHANNELS == 1) ? BLOCK_LOAD_VECTORIZE : BLOCK_LOAD_DIRECT,
diff --git a/cub/device/dispatch/device_radix_sort_dispatch.cuh b/cub/device/dispatch/dispatch_radix_sort.cuh
similarity index 77%
rename from cub/device/dispatch/device_radix_sort_dispatch.cuh
rename to cub/device/dispatch/dispatch_radix_sort.cuh
index 92348ab379..7b227529ae 100644
--- a/cub/device/dispatch/device_radix_sort_dispatch.cuh
+++ b/cub/device/dispatch/dispatch_radix_sort.cuh
@@ -37,9 +37,9 @@
 #include <stdio.h>
 #include <iterator>
 
-#include "../../block_sweep/block_radix_sort_upsweep.cuh"
-#include "../../block_sweep/block_radix_sort_downsweep.cuh"
-#include "../../block_sweep/block_scan_sweep.cuh"
+#include "../../agent/agent_radix_sort_upsweep.cuh"
+#include "../../agent/agent_radix_sort_downsweep.cuh"
+#include "../../agent/agent_scan.cuh"
 #include "../../grid/grid_even_share.cuh"
 #include "../../util_debug.cuh"
 #include "../../util_device.cuh"
@@ -59,11 +59,11 @@ namespace cub {
  * Upsweep pass kernel entry point (multi-block).  Computes privatized digit histograms, one per block.
  */
 template <
-    typename                BlockRadixSortUpsweepPolicy,        ///< Parameterized BlockRadixSortUpsweepPolicy tuning policy type
+    typename                AgentRadixSortUpsweepPolicy,        ///< Parameterized AgentRadixSortUpsweepPolicy tuning policy type
     bool                    DESCENDING,                         ///< Whether or not the sorted-order is high-to-low
     typename                Key,                                ///< Key type
     typename                OffsetT>                            ///< Signed integer type for global offsets
-__launch_bounds__ (int(BlockRadixSortUpsweepPolicy::BLOCK_THREADS))
+__launch_bounds__ (int(AgentRadixSortUpsweepPolicy::BLOCK_THREADS))
 __global__ void DeviceRadixSortUpsweepKernel(
     Key                     *d_keys,                            ///< [in] Input keys buffer
     OffsetT                 *d_spine,                           ///< [out] Privatized (per block) digit histograms (striped, i.e., 0s counts from each block, then 1s counts from each block, etc.)
@@ -72,26 +72,26 @@ __global__ void DeviceRadixSortUpsweepKernel(
     int                     num_bits,                           ///< [in] Number of bits of current radix digit
     GridEvenShare<OffsetT>  even_share)                         ///< [in] Even-share descriptor for mapping an equal number of tiles onto each thread block
 {
-    // Parameterize BlockRadixSortUpsweep type for the current configuration
-    typedef BlockRadixSortUpsweep<BlockRadixSortUpsweepPolicy, Key, OffsetT> BlockRadixSortUpsweepT;          // Primary
+    // Parameterize AgentRadixSortUpsweep type for the current configuration
+    typedef AgentRadixSortUpsweep<AgentRadixSortUpsweepPolicy, Key, OffsetT> AgentRadixSortUpsweepT;          // Primary
 
     // Shared memory storage
-    __shared__ typename BlockRadixSortUpsweepT::TempStorage temp_storage;
+    __shared__ typename AgentRadixSortUpsweepT::TempStorage temp_storage;
 
     // Initialize even-share descriptor for this thread block
     even_share.BlockInit();
 
     OffsetT bin_count;
-    BlockRadixSortUpsweepT(temp_storage, d_keys, current_bit, num_bits).ProcessRegion(
+    AgentRadixSortUpsweepT(temp_storage, d_keys, current_bit, num_bits).ProcessRegion(
         even_share.block_offset,
         even_share.block_end,
         bin_count);
 
     // Write out digit counts (striped)
-    if (threadIdx.x < BlockRadixSortUpsweepT::RADIX_DIGITS)
+    if (threadIdx.x < AgentRadixSortUpsweepT::RADIX_DIGITS)
     {
         int bin_idx = (DESCENDING) ?
-            BlockRadixSortUpsweepT::RADIX_DIGITS - threadIdx.x - 1 :
+            AgentRadixSortUpsweepT::RADIX_DIGITS - threadIdx.x - 1 :
             threadIdx.x;
 
         d_spine[(gridDim.x * bin_idx) + blockIdx.x] = bin_count;
@@ -103,31 +103,31 @@ __global__ void DeviceRadixSortUpsweepKernel(
  * Spine scan kernel entry point (single-block).  Computes an exclusive prefix sum over the privatized digit histograms
  */
 template <
-    typename    BlockScanSweepPolicy,       ///< Parameterizable tuning policy type for cub::BlockScanSweep abstraction
+    typename    AgentScanPolicy,       ///< Parameterizable tuning policy type for cub::AgentScan abstraction
     typename    OffsetT>                    ///< Signed integer type for global offsets
-__launch_bounds__ (int(BlockScanSweepPolicy::BLOCK_THREADS), 1)
+__launch_bounds__ (int(AgentScanPolicy::BLOCK_THREADS), 1)
 __global__ void RadixSortScanBinsKernel(
     OffsetT     *d_spine,                   ///< [in,out] Privatized (per block) digit histograms (striped, i.e., 0s counts from each block, then 1s counts from each block, etc.)
     int         num_counts)                 ///< [in] Total number of bin-counts
 {
-    // Parameterize the BlockScanSweep type for the current configuration
-    typedef BlockScanSweep<BlockScanSweepPolicy, OffsetT*, OffsetT*, cub::Sum, OffsetT, OffsetT> BlockScanSweepT;
+    // Parameterize the AgentScan type for the current configuration
+    typedef AgentScan<AgentScanPolicy, OffsetT*, OffsetT*, cub::Sum, OffsetT, OffsetT> AgentScanT;
 
     // Shared memory storage
-    __shared__ typename BlockScanSweepT::TempStorage temp_storage;
+    __shared__ typename AgentScanT::TempStorage temp_storage;
 
     if (blockIdx.x > 0) return;
 
     // Block scan instance
-    BlockScanSweepT block_scan(temp_storage, d_spine, d_spine, cub::Sum(), OffsetT(0)) ;
+    AgentScanT block_scan(temp_storage, d_spine, d_spine, cub::Sum(), OffsetT(0)) ;
 
     // Process full input tiles
     int block_offset = 0;
     BlockScanRunningPrefixOp<OffsetT, Sum> prefix_op(0, Sum());
-    while (block_offset + BlockScanSweepT::TILE_ITEMS <= num_counts)
+    while (block_offset + AgentScanT::TILE_ITEMS <= num_counts)
     {
         block_scan.ConsumeTile<true, false>(block_offset, prefix_op);
-        block_offset += BlockScanSweepT::TILE_ITEMS;
+        block_offset += AgentScanT::TILE_ITEMS;
     }
 }
 
@@ -136,12 +136,12 @@ __global__ void RadixSortScanBinsKernel(
  * Downsweep pass kernel entry point (multi-block).  Scatters keys (and values) into corresponding bins for the current digit place.
  */
 template <
-    typename                BlockRadixSortDownsweepPolicy,          ///< Parameterizable tuning policy type for cub::BlockRadixSortUpsweep abstraction
+    typename                AgentRadixSortDownsweepPolicy,          ///< Parameterizable tuning policy type for cub::AgentRadixSortUpsweep abstraction
     bool                    DESCENDING,                             ///< Whether or not the sorted-order is high-to-low
     typename                Key,                                    ///< Key type
     typename                Value,                                  ///< Value type
     typename                OffsetT>                                ///< Signed integer type for global offsets
-__launch_bounds__ (int(BlockRadixSortDownsweepPolicy::BLOCK_THREADS))
+__launch_bounds__ (int(AgentRadixSortDownsweepPolicy::BLOCK_THREADS))
 __global__ void DeviceRadixSortDownsweepKernel(
     Key                     *d_keys_in,                             ///< [in] Input keys ping buffer
     Key                     *d_keys_out,                            ///< [in] Output keys pong buffer
@@ -153,17 +153,17 @@ __global__ void DeviceRadixSortDownsweepKernel(
     int                     num_bits,                               ///< [in] Number of bits of current radix digit
     GridEvenShare<OffsetT>  even_share)                             ///< [in] Even-share descriptor for mapping an equal number of tiles onto each thread block
 {
-    // Parameterize BlockRadixSortDownsweep type for the current configuration
-    typedef BlockRadixSortDownsweep<BlockRadixSortDownsweepPolicy, DESCENDING, Key, Value, OffsetT> BlockRadixSortDownsweepT;
+    // Parameterize AgentRadixSortDownsweep type for the current configuration
+    typedef AgentRadixSortDownsweep<AgentRadixSortDownsweepPolicy, DESCENDING, Key, Value, OffsetT> AgentRadixSortDownsweepT;
 
     // Shared memory storage
-    __shared__  typename BlockRadixSortDownsweepT::TempStorage temp_storage;
+    __shared__  typename AgentRadixSortDownsweepT::TempStorage temp_storage;
 
     // Initialize even-share descriptor for this thread block
     even_share.BlockInit();
 
     // Process input tiles
-    BlockRadixSortDownsweepT(temp_storage, num_items, d_spine, d_keys_in, d_keys_out, d_values_in, d_values_out, current_bit, num_bits).ProcessRegion(
+    AgentRadixSortDownsweepT(temp_storage, num_items, d_spine, d_keys_in, d_keys_out, d_values_in, d_values_out, current_bit, num_bits).ProcessRegion(
         even_share.block_offset,
         even_share.block_end);
 }
@@ -183,7 +183,7 @@ template <
     typename Key,           ///< Key type
     typename Value,         ///< Value type
     typename OffsetT>       ///< Signed integer type for global offsets
-struct DeviceRadixSortDispatch
+struct DispatchRadixSort
 {
     /******************************************************************************
      * Constants
@@ -202,6 +202,51 @@ struct DeviceRadixSortDispatch
      * Tuning policies
      ******************************************************************************/
 
+    /// SM52
+    struct Policy520
+    {
+        enum {
+            PRIMARY_RADIX_BITS      = 5,
+            ALT_RADIX_BITS          = PRIMARY_RADIX_BITS - 1,
+        };
+/*
+        // Primary UpsweepPolicy (passes having digit-length RADIX_BITS)
+        typedef AgentRadixSortUpsweepPolicy <256,   CUB_MAX(1, 16 / SCALE_FACTOR_4B), LOAD_DEFAULT, PRIMARY_RADIX_BITS> UpsweepPolicyKeys;
+        typedef AgentRadixSortUpsweepPolicy <128,   CUB_MAX(1, 15 / SCALE_FACTOR_4B), LOAD_LDG, PRIMARY_RADIX_BITS> UpsweepPolicyPairs;
+        typedef typename If<KEYS_ONLY, UpsweepPolicyKeys, UpsweepPolicyPairs>::Type UpsweepPolicy;
+
+        // Alternate UpsweepPolicy (passes having digit-length ALT_RADIX_BITS)
+        typedef AgentRadixSortUpsweepPolicy <256,   CUB_MAX(1, 16 / SCALE_FACTOR_4B), LOAD_DEFAULT, ALT_RADIX_BITS> AltUpsweepPolicyKeys;
+        typedef AgentRadixSortUpsweepPolicy <128,   CUB_MAX(1, 15 / SCALE_FACTOR_4B), LOAD_LDG, ALT_RADIX_BITS> AltUpsweepPolicyPairs;
+        typedef typename If<KEYS_ONLY, AltUpsweepPolicyKeys, AltUpsweepPolicyPairs>::Type AltUpsweepPolicy;
+
+        // ScanPolicy
+        typedef AgentScanPolicy <512, 23, BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, BLOCK_STORE_WARP_TRANSPOSE, BLOCK_SCAN_RAKING_MEMOIZE> ScanPolicy;
+
+        // Primary DownsweepPolicy
+        typedef AgentRadixSortDownsweepPolicy <256, CUB_MAX(1, 16 / SCALE_FACTOR_4B),  BLOCK_LOAD_DIRECT, LOAD_LDG, true, BLOCK_SCAN_RAKING_MEMOIZE, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, PRIMARY_RADIX_BITS> DownsweepPolicyKeys;
+        typedef AgentRadixSortDownsweepPolicy <128, CUB_MAX(1, 15 / SCALE_FACTOR_4B), BLOCK_LOAD_DIRECT, LOAD_LDG, true, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, PRIMARY_RADIX_BITS> DownsweepPolicyPairs;
+
+        // Alternate DownsweepPolicy for ALT_RADIX_BITS-bit passes
+        typedef AgentRadixSortDownsweepPolicy <256, CUB_MAX(1, 16 / SCALE_FACTOR_4B),  BLOCK_LOAD_DIRECT, LOAD_LDG, true, BLOCK_SCAN_RAKING_MEMOIZE, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, ALT_RADIX_BITS> AltDownsweepPolicyKeys;
+        typedef AgentRadixSortDownsweepPolicy <128,  CUB_MAX(1, 15 / SCALE_FACTOR_4B), BLOCK_LOAD_DIRECT, LOAD_LDG, true, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, ALT_RADIX_BITS> AltDownsweepPolicyPairs;
+
+        typedef typename If<KEYS_ONLY, DownsweepPolicyKeys, DownsweepPolicyPairs>::Type DownsweepPolicy;
+        typedef typename If<KEYS_ONLY, AltDownsweepPolicyKeys, AltDownsweepPolicyPairs>::Type AltDownsweepPolicy;
+*/
+
+        // ScanPolicy
+        typedef AgentRadixSortUpsweepPolicy <256,   CUB_MAX(1, 16 / SCALE_FACTOR_4B), LOAD_DEFAULT, PRIMARY_RADIX_BITS> UpsweepPolicy;
+        typedef AgentRadixSortUpsweepPolicy <256,   CUB_MAX(1, 16 / SCALE_FACTOR_4B), LOAD_DEFAULT, ALT_RADIX_BITS> AltUpsweepPolicy;
+
+        typedef AgentScanPolicy <512, 23, BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, BLOCK_STORE_WARP_TRANSPOSE, BLOCK_SCAN_RAKING_MEMOIZE> ScanPolicy;
+
+        typedef AgentRadixSortDownsweepPolicy <256, CUB_MAX(1, 16 / SCALE_FACTOR_4B),  BLOCK_LOAD_DIRECT, LOAD_LDG, true, BLOCK_SCAN_RAKING_MEMOIZE, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, PRIMARY_RADIX_BITS> DownsweepPolicy;
+        typedef AgentRadixSortDownsweepPolicy <256, CUB_MAX(1, 16 / SCALE_FACTOR_4B),  BLOCK_LOAD_DIRECT, LOAD_LDG, true, BLOCK_SCAN_RAKING_MEMOIZE, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, ALT_RADIX_BITS> AltDownsweepPolicy;
+
+    };
+
+
     /// SM35
     struct Policy350
     {
@@ -211,26 +256,26 @@ struct DeviceRadixSortDispatch
         };
 
         // Primary UpsweepPolicy (passes having digit-length RADIX_BITS)
-        typedef BlockRadixSortUpsweepPolicy <64,     CUB_MAX(1, 18 / SCALE_FACTOR_4B), LOAD_LDG, PRIMARY_RADIX_BITS> UpsweepPolicyKeys;
-        typedef BlockRadixSortUpsweepPolicy <128,    CUB_MAX(1, 15 / SCALE_FACTOR_4B), LOAD_LDG, PRIMARY_RADIX_BITS> UpsweepPolicyPairs;
-        typedef typename If<KEYS_ONLY, UpsweepPolicyKeys, UpsweepPolicyPairs>::Type PrimaryUpsweepPolicy;
+        typedef AgentRadixSortUpsweepPolicy <64,     CUB_MAX(1, 18 / SCALE_FACTOR_4B), LOAD_LDG, PRIMARY_RADIX_BITS> UpsweepPolicyKeys;
+        typedef AgentRadixSortUpsweepPolicy <128,    CUB_MAX(1, 15 / SCALE_FACTOR_4B), LOAD_LDG, PRIMARY_RADIX_BITS> UpsweepPolicyPairs;
+        typedef typename If<KEYS_ONLY, UpsweepPolicyKeys, UpsweepPolicyPairs>::Type UpsweepPolicy;
 
         // Alternate UpsweepPolicy (passes having digit-length ALT_RADIX_BITS)
-        typedef BlockRadixSortUpsweepPolicy <64,     CUB_MAX(1, 22 / SCALE_FACTOR_4B), LOAD_LDG, ALT_RADIX_BITS> AltUpsweepPolicyKeys;
-        typedef BlockRadixSortUpsweepPolicy <128,    CUB_MAX(1, 15 / SCALE_FACTOR_4B), LOAD_LDG, ALT_RADIX_BITS> AltUpsweepPolicyPairs;
+        typedef AgentRadixSortUpsweepPolicy <64,     CUB_MAX(1, 22 / SCALE_FACTOR_4B), LOAD_LDG, ALT_RADIX_BITS> AltUpsweepPolicyKeys;
+        typedef AgentRadixSortUpsweepPolicy <128,    CUB_MAX(1, 15 / SCALE_FACTOR_4B), LOAD_LDG, ALT_RADIX_BITS> AltUpsweepPolicyPairs;
         typedef typename If<KEYS_ONLY, AltUpsweepPolicyKeys, AltUpsweepPolicyPairs>::Type AltUpsweepPolicy;
 
         // ScanPolicy
-        typedef BlockScanSweepPolicy <1024, 4, BLOCK_LOAD_VECTORIZE, false, LOAD_DEFAULT, BLOCK_STORE_VECTORIZE, false, BLOCK_SCAN_WARP_SCANS> ScanPolicy;
+        typedef AgentScanPolicy <1024, 4, BLOCK_LOAD_VECTORIZE, LOAD_DEFAULT, BLOCK_STORE_VECTORIZE, BLOCK_SCAN_WARP_SCANS> ScanPolicy;
 
         // Primary DownsweepPolicy
-        typedef BlockRadixSortDownsweepPolicy <64,   CUB_MAX(1, 18 / SCALE_FACTOR_4B), BLOCK_LOAD_DIRECT, LOAD_LDG, false, true, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, PRIMARY_RADIX_BITS> DownsweepPolicyKeys;
-        typedef BlockRadixSortDownsweepPolicy <128,  CUB_MAX(1, 15 / SCALE_FACTOR_4B), BLOCK_LOAD_DIRECT, LOAD_LDG, false, true, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, PRIMARY_RADIX_BITS> DownsweepPolicyPairs;
+        typedef AgentRadixSortDownsweepPolicy <64,   CUB_MAX(1, 18 / SCALE_FACTOR_4B), BLOCK_LOAD_DIRECT, LOAD_LDG, true, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, PRIMARY_RADIX_BITS> DownsweepPolicyKeys;
+        typedef AgentRadixSortDownsweepPolicy <128,  CUB_MAX(1, 15 / SCALE_FACTOR_4B), BLOCK_LOAD_DIRECT, LOAD_LDG, true, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, PRIMARY_RADIX_BITS> DownsweepPolicyPairs;
         typedef typename If<KEYS_ONLY, DownsweepPolicyKeys, DownsweepPolicyPairs>::Type DownsweepPolicy;
 
         // Alternate DownsweepPolicy for ALT_RADIX_BITS-bit passes
-        typedef BlockRadixSortDownsweepPolicy <128,  CUB_MAX(1, 11 / SCALE_FACTOR_4B), BLOCK_LOAD_DIRECT, LOAD_LDG, false, true, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, ALT_RADIX_BITS> AltDownsweepPolicyKeys;
-        typedef BlockRadixSortDownsweepPolicy <128,  CUB_MAX(1, 15 / SCALE_FACTOR_4B), BLOCK_LOAD_DIRECT, LOAD_LDG, false, true, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, ALT_RADIX_BITS> AltDownsweepPolicyPairs;
+        typedef AgentRadixSortDownsweepPolicy <128,  CUB_MAX(1, 11 / SCALE_FACTOR_4B), BLOCK_LOAD_DIRECT, LOAD_LDG, true, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, ALT_RADIX_BITS> AltDownsweepPolicyKeys;
+        typedef AgentRadixSortDownsweepPolicy <128,  CUB_MAX(1, 15 / SCALE_FACTOR_4B), BLOCK_LOAD_DIRECT, LOAD_LDG, true, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, ALT_RADIX_BITS> AltDownsweepPolicyPairs;
         typedef typename If<KEYS_ONLY, AltDownsweepPolicyKeys, AltDownsweepPolicyPairs>::Type AltDownsweepPolicy;
     };
 
@@ -244,26 +289,26 @@ struct DeviceRadixSortDispatch
         };
 
         // UpsweepPolicy
-        typedef BlockRadixSortUpsweepPolicy <256, CUB_MAX(1, 7 / SCALE_FACTOR_4B), LOAD_DEFAULT, PRIMARY_RADIX_BITS> UpsweepPolicyKeys;
-        typedef BlockRadixSortUpsweepPolicy <256, CUB_MAX(1, 5 / SCALE_FACTOR_4B), LOAD_DEFAULT, PRIMARY_RADIX_BITS> UpsweepPolicyPairs;
-        typedef typename If<KEYS_ONLY, UpsweepPolicyKeys, UpsweepPolicyPairs>::Type PrimaryUpsweepPolicy;
+        typedef AgentRadixSortUpsweepPolicy <256, CUB_MAX(1, 7 / SCALE_FACTOR_4B), LOAD_DEFAULT, PRIMARY_RADIX_BITS> UpsweepPolicyKeys;
+        typedef AgentRadixSortUpsweepPolicy <256, CUB_MAX(1, 5 / SCALE_FACTOR_4B), LOAD_DEFAULT, PRIMARY_RADIX_BITS> UpsweepPolicyPairs;
+        typedef typename If<KEYS_ONLY, UpsweepPolicyKeys, UpsweepPolicyPairs>::Type UpsweepPolicy;
 
         // Alternate UpsweepPolicy for ALT_RADIX_BITS-bit passes
-        typedef BlockRadixSortUpsweepPolicy <256, CUB_MAX(1, 7 / SCALE_FACTOR_4B), LOAD_DEFAULT, ALT_RADIX_BITS> AltUpsweepPolicyKeys;
-        typedef BlockRadixSortUpsweepPolicy <256, CUB_MAX(1, 5 / SCALE_FACTOR_4B), LOAD_DEFAULT, ALT_RADIX_BITS> AltUpsweepPolicyPairs;
+        typedef AgentRadixSortUpsweepPolicy <256, CUB_MAX(1, 7 / SCALE_FACTOR_4B), LOAD_DEFAULT, ALT_RADIX_BITS> AltUpsweepPolicyKeys;
+        typedef AgentRadixSortUpsweepPolicy <256, CUB_MAX(1, 5 / SCALE_FACTOR_4B), LOAD_DEFAULT, ALT_RADIX_BITS> AltUpsweepPolicyPairs;
         typedef typename If<KEYS_ONLY, AltUpsweepPolicyKeys, AltUpsweepPolicyPairs>::Type AltUpsweepPolicy;
 
         // ScanPolicy
-        typedef BlockScanSweepPolicy <1024, 4, BLOCK_LOAD_VECTORIZE, false, LOAD_DEFAULT, BLOCK_STORE_VECTORIZE, false, BLOCK_SCAN_RAKING_MEMOIZE> ScanPolicy;
+        typedef AgentScanPolicy <1024, 4, BLOCK_LOAD_VECTORIZE, LOAD_DEFAULT, BLOCK_STORE_VECTORIZE, BLOCK_SCAN_RAKING_MEMOIZE> ScanPolicy;
 
         // DownsweepPolicy
-        typedef BlockRadixSortDownsweepPolicy <128, CUB_MAX(1, 14 / SCALE_FACTOR_4B), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, PRIMARY_RADIX_BITS> DownsweepPolicyKeys;
-        typedef BlockRadixSortDownsweepPolicy <128, CUB_MAX(1, 10 / SCALE_FACTOR_4B), BLOCK_LOAD_TRANSPOSE, LOAD_DEFAULT, false, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, PRIMARY_RADIX_BITS> DownsweepPolicyPairs;
+        typedef AgentRadixSortDownsweepPolicy <128, CUB_MAX(1, 14 / SCALE_FACTOR_4B), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, PRIMARY_RADIX_BITS> DownsweepPolicyKeys;
+        typedef AgentRadixSortDownsweepPolicy <128, CUB_MAX(1, 10 / SCALE_FACTOR_4B), BLOCK_LOAD_TRANSPOSE, LOAD_DEFAULT, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, PRIMARY_RADIX_BITS> DownsweepPolicyPairs;
         typedef typename If<KEYS_ONLY, DownsweepPolicyKeys, DownsweepPolicyPairs>::Type DownsweepPolicy;
 
         // Alternate DownsweepPolicy for ALT_RADIX_BITS-bit passes
-        typedef BlockRadixSortDownsweepPolicy <128, CUB_MAX(1, 14 / SCALE_FACTOR_4B), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, ALT_RADIX_BITS> AltDownsweepPolicyKeys;
-        typedef BlockRadixSortDownsweepPolicy <128, CUB_MAX(1, 10 / SCALE_FACTOR_4B), BLOCK_LOAD_TRANSPOSE, LOAD_DEFAULT, false, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, ALT_RADIX_BITS> AltDownsweepPolicyPairs;
+        typedef AgentRadixSortDownsweepPolicy <128, CUB_MAX(1, 14 / SCALE_FACTOR_4B), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, ALT_RADIX_BITS> AltDownsweepPolicyKeys;
+        typedef AgentRadixSortDownsweepPolicy <128, CUB_MAX(1, 10 / SCALE_FACTOR_4B), BLOCK_LOAD_TRANSPOSE, LOAD_DEFAULT, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, ALT_RADIX_BITS> AltDownsweepPolicyPairs;
         typedef typename If<KEYS_ONLY, AltDownsweepPolicyKeys, AltDownsweepPolicyPairs>::Type AltDownsweepPolicy;
     };
 
@@ -277,26 +322,26 @@ struct DeviceRadixSortDispatch
         };
 
         // Primary UpsweepPolicy (passes having digit-length RADIX_BITS)
-        typedef BlockRadixSortUpsweepPolicy <64, CUB_MAX(1, 18 / SCALE_FACTOR_4B), LOAD_DEFAULT, PRIMARY_RADIX_BITS> UpsweepPolicyKeys;
-        typedef BlockRadixSortUpsweepPolicy <128, CUB_MAX(1, 13 / SCALE_FACTOR_4B), LOAD_DEFAULT, PRIMARY_RADIX_BITS> UpsweepPolicyPairs;
-        typedef typename If<KEYS_ONLY, UpsweepPolicyKeys, UpsweepPolicyPairs>::Type PrimaryUpsweepPolicy;
+        typedef AgentRadixSortUpsweepPolicy <64, CUB_MAX(1, 18 / SCALE_FACTOR_4B), LOAD_DEFAULT, PRIMARY_RADIX_BITS> UpsweepPolicyKeys;
+        typedef AgentRadixSortUpsweepPolicy <128, CUB_MAX(1, 13 / SCALE_FACTOR_4B), LOAD_DEFAULT, PRIMARY_RADIX_BITS> UpsweepPolicyPairs;
+        typedef typename If<KEYS_ONLY, UpsweepPolicyKeys, UpsweepPolicyPairs>::Type UpsweepPolicy;
 
         // Alternate UpsweepPolicy for ALT_RADIX_BITS-bit passes
-        typedef BlockRadixSortUpsweepPolicy <64, CUB_MAX(1, 18 / SCALE_FACTOR_4B), LOAD_DEFAULT, ALT_RADIX_BITS> AltUpsweepPolicyKeys;
-        typedef BlockRadixSortUpsweepPolicy <128, CUB_MAX(1, 13 / SCALE_FACTOR_4B), LOAD_DEFAULT, ALT_RADIX_BITS> AltUpsweepPolicyPairs;
+        typedef AgentRadixSortUpsweepPolicy <64, CUB_MAX(1, 18 / SCALE_FACTOR_4B), LOAD_DEFAULT, ALT_RADIX_BITS> AltUpsweepPolicyKeys;
+        typedef AgentRadixSortUpsweepPolicy <128, CUB_MAX(1, 13 / SCALE_FACTOR_4B), LOAD_DEFAULT, ALT_RADIX_BITS> AltUpsweepPolicyPairs;
         typedef typename If<KEYS_ONLY, AltUpsweepPolicyKeys, AltUpsweepPolicyPairs>::Type AltUpsweepPolicy;
 
         // ScanPolicy
-        typedef BlockScanSweepPolicy <512, 4, BLOCK_LOAD_VECTORIZE, false, LOAD_DEFAULT, BLOCK_STORE_VECTORIZE, false, BLOCK_SCAN_RAKING_MEMOIZE> ScanPolicy;
+        typedef AgentScanPolicy <512, 4, BLOCK_LOAD_VECTORIZE, LOAD_DEFAULT, BLOCK_STORE_VECTORIZE, BLOCK_SCAN_RAKING_MEMOIZE> ScanPolicy;
 
         // DownsweepPolicy
-        typedef BlockRadixSortDownsweepPolicy <64, CUB_MAX(1, 18 / SCALE_FACTOR_4B), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, PRIMARY_RADIX_BITS> DownsweepPolicyKeys;
-        typedef BlockRadixSortDownsweepPolicy <128, CUB_MAX(1, 13 / SCALE_FACTOR_4B), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, PRIMARY_RADIX_BITS> DownsweepPolicyPairs;
+        typedef AgentRadixSortDownsweepPolicy <64, CUB_MAX(1, 18 / SCALE_FACTOR_4B), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, PRIMARY_RADIX_BITS> DownsweepPolicyKeys;
+        typedef AgentRadixSortDownsweepPolicy <128, CUB_MAX(1, 13 / SCALE_FACTOR_4B), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, PRIMARY_RADIX_BITS> DownsweepPolicyPairs;
         typedef typename If<KEYS_ONLY, DownsweepPolicyKeys, DownsweepPolicyPairs>::Type DownsweepPolicy;
 
         // Alternate DownsweepPolicy for ALT_RADIX_BITS-bit passes
-        typedef BlockRadixSortDownsweepPolicy <64, CUB_MAX(1, 18 / SCALE_FACTOR_4B), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, ALT_RADIX_BITS> AltDownsweepPolicyKeys;
-        typedef BlockRadixSortDownsweepPolicy <128, CUB_MAX(1, 13 / SCALE_FACTOR_4B), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, ALT_RADIX_BITS> AltDownsweepPolicyPairs;
+        typedef AgentRadixSortDownsweepPolicy <64, CUB_MAX(1, 18 / SCALE_FACTOR_4B), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, ALT_RADIX_BITS> AltDownsweepPolicyKeys;
+        typedef AgentRadixSortDownsweepPolicy <128, CUB_MAX(1, 13 / SCALE_FACTOR_4B), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, ALT_RADIX_BITS> AltDownsweepPolicyPairs;
         typedef typename If<KEYS_ONLY, AltDownsweepPolicyKeys, AltDownsweepPolicyPairs>::Type AltDownsweepPolicy;
     };
 
@@ -310,26 +355,26 @@ struct DeviceRadixSortDispatch
         };
 
         // UpsweepPolicy
-        typedef BlockRadixSortUpsweepPolicy <128, CUB_MAX(1, 19 / SCALE_FACTOR_4B), LOAD_DEFAULT, PRIMARY_RADIX_BITS> UpsweepPolicyKeys;
-        typedef BlockRadixSortUpsweepPolicy <128, CUB_MAX(1, 19 / SCALE_FACTOR_4B), LOAD_DEFAULT, PRIMARY_RADIX_BITS> UpsweepPolicyPairs;
-        typedef typename If<KEYS_ONLY, UpsweepPolicyKeys, UpsweepPolicyPairs>::Type PrimaryUpsweepPolicy;
+        typedef AgentRadixSortUpsweepPolicy <128, CUB_MAX(1, 19 / SCALE_FACTOR_4B), LOAD_DEFAULT, PRIMARY_RADIX_BITS> UpsweepPolicyKeys;
+        typedef AgentRadixSortUpsweepPolicy <128, CUB_MAX(1, 19 / SCALE_FACTOR_4B), LOAD_DEFAULT, PRIMARY_RADIX_BITS> UpsweepPolicyPairs;
+        typedef typename If<KEYS_ONLY, UpsweepPolicyKeys, UpsweepPolicyPairs>::Type UpsweepPolicy;
 
         // Alternate UpsweepPolicy for ALT_RADIX_BITS-bit passes
-        typedef BlockRadixSortUpsweepPolicy <128, CUB_MAX(1, 15 / SCALE_FACTOR_4B), LOAD_DEFAULT, ALT_RADIX_BITS> AltUpsweepPolicyKeys;
-        typedef BlockRadixSortUpsweepPolicy <128, CUB_MAX(1, 15 / SCALE_FACTOR_4B), LOAD_DEFAULT, ALT_RADIX_BITS> AltUpsweepPolicyPairs;
+        typedef AgentRadixSortUpsweepPolicy <128, CUB_MAX(1, 15 / SCALE_FACTOR_4B), LOAD_DEFAULT, ALT_RADIX_BITS> AltUpsweepPolicyKeys;
+        typedef AgentRadixSortUpsweepPolicy <128, CUB_MAX(1, 15 / SCALE_FACTOR_4B), LOAD_DEFAULT, ALT_RADIX_BITS> AltUpsweepPolicyPairs;
         typedef typename If<KEYS_ONLY, AltUpsweepPolicyKeys, AltUpsweepPolicyPairs>::Type AltUpsweepPolicy;
 
         // ScanPolicy
-        typedef BlockScanSweepPolicy <256, 4, BLOCK_LOAD_VECTORIZE, false, LOAD_DEFAULT, BLOCK_STORE_VECTORIZE, false, BLOCK_SCAN_WARP_SCANS> ScanPolicy;
+        typedef AgentScanPolicy <256, 4, BLOCK_LOAD_VECTORIZE, LOAD_DEFAULT, BLOCK_STORE_VECTORIZE, BLOCK_SCAN_WARP_SCANS> ScanPolicy;
 
         // DownsweepPolicy
-        typedef BlockRadixSortDownsweepPolicy <64, CUB_MAX(1, 19 / SCALE_FACTOR_4B), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, PRIMARY_RADIX_BITS> DownsweepPolicyKeys;
-        typedef BlockRadixSortDownsweepPolicy <64, CUB_MAX(1, 19 / SCALE_FACTOR_4B), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, PRIMARY_RADIX_BITS> DownsweepPolicyPairs;
+        typedef AgentRadixSortDownsweepPolicy <64, CUB_MAX(1, 19 / SCALE_FACTOR_4B), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, PRIMARY_RADIX_BITS> DownsweepPolicyKeys;
+        typedef AgentRadixSortDownsweepPolicy <64, CUB_MAX(1, 19 / SCALE_FACTOR_4B), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, PRIMARY_RADIX_BITS> DownsweepPolicyPairs;
         typedef typename If<KEYS_ONLY, DownsweepPolicyKeys, DownsweepPolicyPairs>::Type DownsweepPolicy;
 
         // Alternate DownsweepPolicy for ALT_RADIX_BITS-bit passes
-        typedef BlockRadixSortDownsweepPolicy <128, CUB_MAX(1, 15 / SCALE_FACTOR_4B), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, ALT_RADIX_BITS> AltDownsweepPolicyKeys;
-        typedef BlockRadixSortDownsweepPolicy <128, CUB_MAX(1, 15 / SCALE_FACTOR_4B), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, ALT_RADIX_BITS> AltDownsweepPolicyPairs;
+        typedef AgentRadixSortDownsweepPolicy <128, CUB_MAX(1, 15 / SCALE_FACTOR_4B), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, ALT_RADIX_BITS> AltDownsweepPolicyKeys;
+        typedef AgentRadixSortDownsweepPolicy <128, CUB_MAX(1, 15 / SCALE_FACTOR_4B), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, ALT_RADIX_BITS> AltDownsweepPolicyPairs;
         typedef typename If<KEYS_ONLY, AltDownsweepPolicyKeys, AltDownsweepPolicyPairs>::Type AltDownsweepPolicy;
     };
 
@@ -343,19 +388,19 @@ struct DeviceRadixSortDispatch
         };
 
         // UpsweepPolicy
-        typedef BlockRadixSortUpsweepPolicy <64, CUB_MAX(1, 9 / SCALE_FACTOR_4B), LOAD_DEFAULT, PRIMARY_RADIX_BITS> PrimaryUpsweepPolicy;
+        typedef AgentRadixSortUpsweepPolicy <64, CUB_MAX(1, 9 / SCALE_FACTOR_4B), LOAD_DEFAULT, PRIMARY_RADIX_BITS> UpsweepPolicy;
 
         // Alternate UpsweepPolicy for ALT_RADIX_BITS-bit passes
-        typedef BlockRadixSortUpsweepPolicy <64, CUB_MAX(1, 9 / SCALE_FACTOR_4B), LOAD_DEFAULT, ALT_RADIX_BITS> AltUpsweepPolicy;
+        typedef AgentRadixSortUpsweepPolicy <64, CUB_MAX(1, 9 / SCALE_FACTOR_4B), LOAD_DEFAULT, ALT_RADIX_BITS> AltUpsweepPolicy;
 
         // ScanPolicy
-        typedef BlockScanSweepPolicy <256, 4, BLOCK_LOAD_VECTORIZE, false, LOAD_DEFAULT, BLOCK_STORE_VECTORIZE, false, BLOCK_SCAN_RAKING_MEMOIZE> ScanPolicy;
+        typedef AgentScanPolicy <256, 4, BLOCK_LOAD_VECTORIZE, LOAD_DEFAULT, BLOCK_STORE_VECTORIZE, BLOCK_SCAN_RAKING_MEMOIZE> ScanPolicy;
 
         // DownsweepPolicy
-        typedef BlockRadixSortDownsweepPolicy <64, CUB_MAX(1, 9 / SCALE_FACTOR_4B), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, PRIMARY_RADIX_BITS> DownsweepPolicy;
+        typedef AgentRadixSortDownsweepPolicy <64, CUB_MAX(1, 9 / SCALE_FACTOR_4B), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, PRIMARY_RADIX_BITS> DownsweepPolicy;
 
         // Alternate DownsweepPolicy for ALT_RADIX_BITS-bit passes
-        typedef BlockRadixSortDownsweepPolicy <64, CUB_MAX(1, 9 / SCALE_FACTOR_4B), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, ALT_RADIX_BITS> AltDownsweepPolicy;
+        typedef AgentRadixSortDownsweepPolicy <64, CUB_MAX(1, 9 / SCALE_FACTOR_4B), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, ALT_RADIX_BITS> AltDownsweepPolicy;
     };
 
 
@@ -363,7 +408,10 @@ struct DeviceRadixSortDispatch
      * Tuning policies of current PTX compiler pass
      ******************************************************************************/
 
-#if (CUB_PTX_ARCH >= 350)
+#if (CUB_PTX_ARCH >= 520)
+    typedef Policy520 PtxPolicy;
+
+#elif (CUB_PTX_ARCH >= 350)
     typedef Policy350 PtxPolicy;
 
 #elif (CUB_PTX_ARCH >= 300)
@@ -381,7 +429,7 @@ struct DeviceRadixSortDispatch
 #endif
 
     // "Opaque" policies (whose parameterizations aren't reflected in the type signature)
-    struct PtxUpsweepPolicy         : PtxPolicy::PrimaryUpsweepPolicy {};
+    struct PtxUpsweepPolicy         : PtxPolicy::UpsweepPolicy {};
     struct PtxAltUpsweepPolicy      : PtxPolicy::AltUpsweepPolicy {};
     struct PtxScanPolicy            : PtxPolicy::ScanPolicy {};
     struct PtxDownsweepPolicy       : PtxPolicy::DownsweepPolicy {};
@@ -418,7 +466,7 @@ struct DeviceRadixSortDispatch
     {
         cudaError_t error;
         do {
-            if (CubDebug(error = upsweep_config.template         InitUpsweepPolicy<typename Policy::PrimaryUpsweepPolicy>(         sm_version, sm_count, upsweep_kernel))) break;
+            if (CubDebug(error = upsweep_config.template         InitUpsweepPolicy<typename Policy::UpsweepPolicy>(         sm_version, sm_count, upsweep_kernel))) break;
             if (CubDebug(error = alt_upsweep_config.template     InitUpsweepPolicy<typename Policy::AltUpsweepPolicy>(      sm_version, sm_count, alt_upsweep_kernel))) break;
             if (CubDebug(error = scan_config.template            InitScanPolicy<typename Policy::ScanPolicy>(               sm_version, sm_count, scan_kernel))) break;
             if (CubDebug(error = downsweep_config.template       InitDownsweepPolicy<typename Policy::DownsweepPolicy>(     sm_version, sm_count, downsweep_kernel))) break;
@@ -450,7 +498,7 @@ struct DeviceRadixSortDispatch
         KernelConfig            &alt_downsweep_config,
         UpsweepKernelPtr        upsweep_kernel,
         UpsweepKernelPtr        alt_upsweep_kernel,
-        ScanKernelPtr          scan_kernel,
+        ScanKernelPtr           scan_kernel,
         DownsweepKernelPtr      downsweep_kernel,
         DownsweepKernelPtr      alt_downsweep_kernel)
     {
@@ -474,7 +522,11 @@ struct DeviceRadixSortDispatch
 
         // We're on the host, so lookup and initialize the kernel dispatch configurations with the policies that match the device's PTX version
         cudaError_t error;
-        if (ptx_version >= 350)
+        if (ptx_version >= 520)
+        {
+            error = InitConfigs<Policy520>(sm_version, sm_count, upsweep_config, alt_upsweep_config, scan_config, downsweep_config, alt_downsweep_config, upsweep_kernel, alt_upsweep_kernel, scan_kernel, downsweep_kernel, alt_downsweep_kernel);
+        }
+        else if (ptx_version >= 350)
         {
             error = InitConfigs<Policy350>(sm_version, sm_count, upsweep_config, alt_upsweep_config, scan_config, downsweep_config, alt_downsweep_config, upsweep_kernel, alt_upsweep_kernel, scan_kernel, downsweep_kernel, alt_downsweep_kernel);
         }
diff --git a/cub/device/dispatch/device_reduce_dispatch.cuh b/cub/device/dispatch/dispatch_reduce.cuh
similarity index 96%
rename from cub/device/dispatch/device_reduce_dispatch.cuh
rename to cub/device/dispatch/dispatch_reduce.cuh
index a5969fcb65..bd34a8052a 100644
--- a/cub/device/dispatch/device_reduce_dispatch.cuh
+++ b/cub/device/dispatch/dispatch_reduce.cuh
@@ -37,7 +37,7 @@
 #include <stdio.h>
 #include <iterator>
 
-#include "../../block_sweep/block_reduce_sweep.cuh"
+#include "../../agent/agent_reduce.cuh"
 #include "../../iterator/constant_input_iterator.cuh"
 #include "../../thread/thread_operators.cuh"
 #include "../../grid/grid_even_share.cuh"
@@ -61,12 +61,12 @@ namespace cub {
  * Reduce region kernel entry point (multi-block).  Computes privatized reductions, one per thread block.
  */
 template <
-    typename                BlockReduceSweepPolicy,     ///< Parameterized BlockReduceSweepPolicy tuning policy type
+    typename                AgentReducePolicy,     ///< Parameterized AgentReducePolicy tuning policy type
     typename                InputIteratorT,             ///< Random-access input iterator type for reading input items \iterator
     typename                OutputIteratorT,            ///< Output iterator type for recording the reduced aggregate \iterator
     typename                OffsetT,                    ///< Signed integer type for global offsets
     typename                ReductionOp>                ///< Binary reduction functor type having member <tt>T operator()(const T &a, const T &b)</tt>
-__launch_bounds__ (int(BlockReduceSweepPolicy::BLOCK_THREADS))
+__launch_bounds__ (int(AgentReducePolicy::BLOCK_THREADS))
 __global__ void DeviceReduceSweepKernel(
     InputIteratorT          d_in,                       ///< [in] Pointer to the input sequence of data items
     OutputIteratorT         d_out,                      ///< [out] Pointer to the output aggregate
@@ -79,21 +79,21 @@ __global__ void DeviceReduceSweepKernel(
     typedef typename std::iterator_traits<InputIteratorT>::value_type T;
 
     // Thread block type for reducing input tiles
-    typedef BlockReduceSweep<BlockReduceSweepPolicy, InputIteratorT, OffsetT, ReductionOp> BlockReduceSweepT;
+    typedef AgentReduce<AgentReducePolicy, InputIteratorT, OffsetT, ReductionOp> AgentReduceT;
 
     // Block-wide aggregate
     T block_aggregate;
 
     // Shared memory storage
-    __shared__ typename BlockReduceSweepT::TempStorage temp_storage;
+    __shared__ typename AgentReduceT::TempStorage temp_storage;
 
     // Consume input tiles
-    BlockReduceSweepT(temp_storage, d_in, reduction_op).ConsumeRange(
+    AgentReduceT(temp_storage, d_in, reduction_op).ConsumeRange(
         num_items,
         even_share,
         queue,
         block_aggregate,
-        Int2Type<BlockReduceSweepPolicy::GRID_MAPPING>());
+        Int2Type<AgentReducePolicy::GRID_MAPPING>());
 
     // Output result
     if (threadIdx.x == 0)
@@ -107,12 +107,12 @@ __global__ void DeviceReduceSweepKernel(
  * Reduce a single tile kernel entry point (single-block).  Can be used to aggregate privatized threadblock reductions from a previous multi-block reduction pass.
  */
 template <
-    typename                BlockReduceSweepPolicy,     ///< Parameterized BlockReduceSweepPolicy tuning policy type
+    typename                AgentReducePolicy,     ///< Parameterized AgentReducePolicy tuning policy type
     typename                InputIteratorT,             ///< Random-access input iterator type for reading input items \iterator
     typename                OutputIteratorT,            ///< Output iterator type for recording the reduced aggregate \iterator
     typename                OffsetT,                    ///< Signed integer type for global offsets
     typename                ReductionOp>                ///< Binary reduction functor type having member <tt>T operator()(const T &a, const T &b)</tt>
-__launch_bounds__ (int(BlockReduceSweepPolicy::BLOCK_THREADS), 1)
+__launch_bounds__ (int(AgentReducePolicy::BLOCK_THREADS), 1)
 __global__ void SingleReduceSweepKernel(
     InputIteratorT          d_in,                       ///< [in] Pointer to the input sequence of data items
     OutputIteratorT         d_out,                      ///< [out] Pointer to the output aggregate
@@ -123,16 +123,16 @@ __global__ void SingleReduceSweepKernel(
     typedef typename std::iterator_traits<InputIteratorT>::value_type T;
 
     // Thread block type for reducing input tiles
-    typedef BlockReduceSweep<BlockReduceSweepPolicy, InputIteratorT, OffsetT, ReductionOp> BlockReduceSweepT;
+    typedef AgentReduce<AgentReducePolicy, InputIteratorT, OffsetT, ReductionOp> AgentReduceT;
 
     // Block-wide aggregate
     T block_aggregate;
 
     // Shared memory storage
-    __shared__ typename BlockReduceSweepT::TempStorage temp_storage;
+    __shared__ typename AgentReduceT::TempStorage temp_storage;
 
     // Consume input tiles
-    BlockReduceSweepT(temp_storage, d_in, reduction_op).ConsumeRange(
+    AgentReduceT(temp_storage, d_in, reduction_op).ConsumeRange(
         OffsetT(0),
         OffsetT(num_items),
         block_aggregate);
@@ -159,7 +159,7 @@ template <
     typename OutputIteratorT,   ///< Output iterator type for recording the reduced aggregate \iterator
     typename OffsetT,           ///< Signed integer type for global offsets
     typename ReductionOp>       ///< Binary reduction functor type having member <tt>T operator()(const T &a, const T &b)</tt>
-struct DeviceReduceDispatch
+struct DispatchReduce
 {
     /******************************************************************************
      * Types and constants
@@ -188,7 +188,7 @@ struct DeviceReduceDispatch
             SMALL_NOMINAL_VECTOR_LOAD_LENGTH    = 4,
             SMALL_VECTOR_LOAD_LENGTH            = CUB_MIN(CUB_MIN(4, SMALL_ITEMS_PER_THREAD), CUB_MAX(1, (SMALL_NOMINAL_VECTOR_LOAD_LENGTH * 1 / sizeof(T)))),
         };
-        typedef BlockReduceSweepPolicy<
+        typedef AgentReducePolicy<
                 128,                                ///< Threads per thread block
                 SMALL_ITEMS_PER_THREAD,             ///< Items per thread per tile of input
                 SMALL_VECTOR_LOAD_LENGTH,           ///< Number of items per vectorized load
@@ -205,7 +205,7 @@ struct DeviceReduceDispatch
             NOMINAL_VECTOR_LOAD_LENGTH          = 4,
             VECTOR_LOAD_LENGTH                  = CUB_MIN(CUB_MIN(4, ITEMS_PER_THREAD), CUB_MAX(1, (NOMINAL_VECTOR_LOAD_LENGTH * 4 / sizeof(T)))),
         };
-        typedef BlockReduceSweepPolicy<
+        typedef AgentReducePolicy<
                 256,                                ///< Threads per thread block
                 ITEMS_PER_THREAD,                   ///< Items per thread per tile of input
                 VECTOR_LOAD_LENGTH,                 ///< Number of items per vectorized load
@@ -227,7 +227,7 @@ struct DeviceReduceDispatch
             SINGLE_NOMINAL_VECTOR_LOAD_LENGTH   = 1,
             SINGLE_VECTOR_LOAD_LENGTH           = CUB_MIN(CUB_MIN(4, ITEMS_PER_THREAD), CUB_MAX(1, (SINGLE_NOMINAL_VECTOR_LOAD_LENGTH * 4 / sizeof(T)))),
         };
-        typedef BlockReduceSweepPolicy<
+        typedef AgentReducePolicy<
                 256,                                ///< Threads per thread block
                 SINGLE_ITEMS_PER_THREAD,            ///< Items per thread per tile of input
                 SINGLE_VECTOR_LOAD_LENGTH,          ///< Number of items per vectorized load
@@ -248,7 +248,7 @@ struct DeviceReduceDispatch
             NOMINAL_VECTOR_LOAD_LENGTH          = 1,
             VECTOR_LOAD_LENGTH                  = CUB_MIN(CUB_MIN(4, ITEMS_PER_THREAD), CUB_MAX(1, (NOMINAL_VECTOR_LOAD_LENGTH * 4 / sizeof(T)))),
         };
-        typedef BlockReduceSweepPolicy<
+        typedef AgentReducePolicy<
                 256,                                ///< Threads per thread block
                 ITEMS_PER_THREAD,                   ///< Items per thread per tile of input
                 VECTOR_LOAD_LENGTH,                 ///< Number of items per vectorized load
@@ -265,7 +265,7 @@ struct DeviceReduceDispatch
             SINGLE_NOMINAL_VECTOR_LOAD_LENGTH   = 4,
             SINGLE_VECTOR_LOAD_LENGTH           = CUB_MIN(CUB_MIN(4, ITEMS_PER_THREAD), CUB_MAX(1, (SINGLE_NOMINAL_VECTOR_LOAD_LENGTH * 4 / sizeof(T)))),
         };
-        typedef BlockReduceSweepPolicy<
+        typedef AgentReducePolicy<
                 256,                                ///< Threads per thread block
                 SINGLE_ITEMS_PER_THREAD,            ///< Items per thread per tile of input
                 SINGLE_VECTOR_LOAD_LENGTH,          ///< Number of items per vectorized load
@@ -286,7 +286,7 @@ struct DeviceReduceDispatch
             SMALL_NOMINAL_VECTOR_LOAD_LENGTH    = 4,
             SMALL_VECTOR_LOAD_LENGTH            = CUB_MIN(CUB_MIN(4, SMALL_ITEMS_PER_THREAD), CUB_MAX(1, (SMALL_NOMINAL_VECTOR_LOAD_LENGTH * 1 / sizeof(T)))),
         };
-        typedef BlockReduceSweepPolicy<
+        typedef AgentReducePolicy<
                 192,                                ///< Threads per thread block
                 SMALL_ITEMS_PER_THREAD,             ///< Items per thread per tile of input
                 SMALL_VECTOR_LOAD_LENGTH,           ///< Number of items per vectorized load
@@ -305,7 +305,7 @@ struct DeviceReduceDispatch
             NOMINAL_VECTOR_LOAD_LENGTH          = 4,
             VECTOR_LOAD_LENGTH                  = CUB_MIN(CUB_MIN(4, ITEMS_PER_THREAD), CUB_MAX(1, (NOMINAL_VECTOR_LOAD_LENGTH * 4 / sizeof(T)))),
         };
-        typedef BlockReduceSweepPolicy<
+        typedef AgentReducePolicy<
                 128,                                ///< Threads per thread block
                 ITEMS_PER_THREAD,                   ///< Items per thread per tile of input
                 VECTOR_LOAD_LENGTH,                 ///< Number of items per vectorized load
@@ -327,7 +327,7 @@ struct DeviceReduceDispatch
             SINGLE_NOMINAL_VECTOR_LOAD_LENGTH   = 1,
             SINGLE_VECTOR_LOAD_LENGTH           = CUB_MIN(CUB_MIN(4, ITEMS_PER_THREAD), CUB_MAX(1, (SINGLE_NOMINAL_VECTOR_LOAD_LENGTH * 4 / sizeof(T)))),
         };
-        typedef BlockReduceSweepPolicy<
+        typedef AgentReducePolicy<
                 192,                                ///< Threads per thread block
                 SINGLE_ITEMS_PER_THREAD,            ///< Items per thread per tile of input
                 SINGLE_VECTOR_LOAD_LENGTH,          ///< Number of items per vectorized load
@@ -348,7 +348,7 @@ struct DeviceReduceDispatch
             NOMINAL_VECTOR_LOAD_LENGTH          = 2,
             VECTOR_LOAD_LENGTH                  = CUB_MIN(CUB_MIN(4, ITEMS_PER_THREAD), CUB_MAX(1, (NOMINAL_VECTOR_LOAD_LENGTH * 4 / sizeof(T)))),
         };
-        typedef BlockReduceSweepPolicy<
+        typedef AgentReducePolicy<
                 128,                                ///< Threads per thread block
                 ITEMS_PER_THREAD,                   ///< Items per thread per tile of input
                 VECTOR_LOAD_LENGTH,                 ///< Number of items per vectorized load
@@ -365,7 +365,7 @@ struct DeviceReduceDispatch
             SINGLE_NOMINAL_VECTOR_LOAD_LENGTH   = 2,
             SINGLE_VECTOR_LOAD_LENGTH           = CUB_MIN(CUB_MIN(4, ITEMS_PER_THREAD), CUB_MAX(1, (SINGLE_NOMINAL_VECTOR_LOAD_LENGTH * 4 / sizeof(T)))),
         };
-        typedef BlockReduceSweepPolicy<
+        typedef AgentReducePolicy<
                 32,                                 ///< Threads per thread block
                 SINGLE_ITEMS_PER_THREAD,            ///< Items per thread per tile of input
                 SINGLE_VECTOR_LOAD_LENGTH,          ///< Number of items per vectorized load
@@ -387,7 +387,7 @@ struct DeviceReduceDispatch
         };
 
         // RangeReducePolicy
-        typedef BlockReduceSweepPolicy<
+        typedef AgentReducePolicy<
                 128,                                ///< Threads per thread block
                 ITEMS_PER_THREAD,                   ///< Items per thread per tile of input
                 VECTOR_LOAD_LENGTH,                 ///< Number of items per vectorized load
@@ -404,7 +404,7 @@ struct DeviceReduceDispatch
             SINGLE_NOMINAL_VECTOR_LOAD_LENGTH   = 4,
             SINGLE_VECTOR_LOAD_LENGTH           = CUB_MIN(CUB_MIN(4, ITEMS_PER_THREAD), CUB_MAX(1, (SINGLE_NOMINAL_VECTOR_LOAD_LENGTH * 4 / sizeof(T)))),
         };
-        typedef BlockReduceSweepPolicy<
+        typedef AgentReducePolicy<
                 32,                                 ///< Threads per thread block
                 SINGLE_ITEMS_PER_THREAD,            ///< Items per thread per tile of input
                 SINGLE_VECTOR_LOAD_LENGTH,          ///< Number of items per vectorized load
diff --git a/cub/device/dispatch/device_reduce_by_key_dispatch.cuh b/cub/device/dispatch/dispatch_reduce_by_key.cuh
similarity index 92%
rename from cub/device/dispatch/device_reduce_by_key_dispatch.cuh
rename to cub/device/dispatch/dispatch_reduce_by_key.cuh
index 7f31eacfae..3f9458821c 100644
--- a/cub/device/dispatch/device_reduce_by_key_dispatch.cuh
+++ b/cub/device/dispatch/dispatch_reduce_by_key.cuh
@@ -38,7 +38,7 @@
 #include <iterator>
 
 #include "device_scan_dispatch.cuh"
-#include "../../block_sweep/block_reduce_by_key_sweep.cuh"
+#include "../../agent/agent_reduce_by_key.cuh"
 #include "../../thread/thread_operators.cuh"
 #include "../../grid/grid_queue.cuh"
 #include "../../util_device.cuh"
@@ -58,7 +58,7 @@ namespace cub {
  * Multi-block reduce-by-key sweep kernel entry point
  */
 template <
-    typename            BlockReduceSweepByKeyPolicy,            ///< Parameterized BlockReduceSweepByKeyPolicy tuning policy type
+    typename            AgentReduceByKeyPolicy,            ///< Parameterized AgentReduceByKeyPolicy tuning policy type
     typename            KeysInputIteratorT,                     ///< Random-access input iterator type for keys
     typename            UniqueOutputIteratorT,                  ///< Random-access output iterator type for keys
     typename            ValuesInputIteratorT,                   ///< Random-access input iterator type for values
@@ -68,7 +68,7 @@ template <
     typename            EqualityOp,                             ///< Key equality operator type
     typename            ReductionOp,                            ///< Value reduction operator type
     typename            OffsetT>                                ///< Signed integer type for global offsets
-__launch_bounds__ (int(BlockReduceSweepByKeyPolicy::BLOCK_THREADS))
+__launch_bounds__ (int(AgentReduceByKeyPolicy::BLOCK_THREADS))
 __global__ void DeviceReduceByKeySweepKernel(
     KeysInputIteratorT          d_keys_in,                      ///< [in] Pointer to the input sequence of keys
     UniqueOutputIteratorT       d_unique_out,                   ///< [out] Pointer to the output sequence of unique keys (one key per run)
@@ -83,21 +83,21 @@ __global__ void DeviceReduceByKeySweepKernel(
     GridQueue<int>              queue)                          ///< [in] Drain queue descriptor for dynamically mapping tile data onto thread blocks
 {
     // Thread block type for reducing tiles of value segments
-    typedef BlockReduceSweepByKey<
-        BlockReduceSweepByKeyPolicy,
+    typedef AgentReduceByKey<
+        AgentReduceByKeyPolicy,
         KeysInputIteratorT,
         UniqueOutputIteratorT,
         ValuesInputIteratorT,
         AggregatesOutputIteratorT,
         EqualityOp,
         ReductionOp,
-        OffsetT>BlockReduceSweepByKeyT;
+        OffsetT>AgentReduceByKeyT;
 
-    // Shared memory for BlockReduceSweepByKey
-    __shared__ typename BlockReduceSweepByKeyT::TempStorage temp_storage;
+    // Shared memory for AgentReduceByKey
+    __shared__ typename AgentReduceByKeyT::TempStorage temp_storage;
 
     // Process tiles
-    BlockReduceSweepByKeyT(temp_storage, d_keys_in, d_unique_out, d_values_in, d_aggregates_out, equality_op, reduction_op, num_items).ConsumeRange(
+    AgentReduceByKeyT(temp_storage, d_keys_in, d_unique_out, d_values_in, d_aggregates_out, equality_op, reduction_op, num_items).ConsumeRange(
         num_tiles,
         queue,
         tile_status,
@@ -123,7 +123,7 @@ template <
     typename    EqualityOp,                     ///< Key equality operator type
     typename    ReductionOp,                    ///< Value reduction operator type
     typename    OffsetT>                        ///< Signed integer type for global offsets
-struct DeviceReduceByKeyDispatch
+struct DispatchReduceByKey
 {
     /******************************************************************************
      * Types and constants
@@ -158,7 +158,7 @@ struct DeviceReduceByKeyDispatch
             ITEMS_PER_THREAD            = (MAX_INPUT_BYTES <= 8) ? 8 : CUB_MIN(NOMINAL_4B_ITEMS_PER_THREAD, CUB_MAX(1, ((NOMINAL_4B_ITEMS_PER_THREAD * 8) + COMBINED_INPUT_BYTES - 1) / COMBINED_INPUT_BYTES)),
         };
 
-        typedef BlockReduceSweepByKeyPolicy<
+        typedef AgentReduceByKeyPolicy<
                 128,
                 ITEMS_PER_THREAD,
                 BLOCK_LOAD_DIRECT,
@@ -176,7 +176,7 @@ struct DeviceReduceByKeyDispatch
             ITEMS_PER_THREAD            = CUB_MIN(NOMINAL_4B_ITEMS_PER_THREAD, CUB_MAX(1, ((NOMINAL_4B_ITEMS_PER_THREAD * 8) + COMBINED_INPUT_BYTES - 1) / COMBINED_INPUT_BYTES)),
         };
 
-        typedef BlockReduceSweepByKeyPolicy<
+        typedef AgentReduceByKeyPolicy<
                 128,
                 ITEMS_PER_THREAD,
                 BLOCK_LOAD_WARP_TRANSPOSE,
@@ -194,7 +194,7 @@ struct DeviceReduceByKeyDispatch
             ITEMS_PER_THREAD            = CUB_MIN(NOMINAL_4B_ITEMS_PER_THREAD, CUB_MAX(1, ((NOMINAL_4B_ITEMS_PER_THREAD * 8) + COMBINED_INPUT_BYTES - 1) / COMBINED_INPUT_BYTES)),
         };
 
-        typedef BlockReduceSweepByKeyPolicy<
+        typedef AgentReduceByKeyPolicy<
                 128,
                 ITEMS_PER_THREAD,
                 BLOCK_LOAD_WARP_TRANSPOSE,
@@ -212,7 +212,7 @@ struct DeviceReduceByKeyDispatch
             ITEMS_PER_THREAD            = CUB_MIN(NOMINAL_4B_ITEMS_PER_THREAD, CUB_MAX(1, ((NOMINAL_4B_ITEMS_PER_THREAD * 8) + COMBINED_INPUT_BYTES - 1) / COMBINED_INPUT_BYTES)),
         };
 
-        typedef BlockReduceSweepByKeyPolicy<
+        typedef AgentReduceByKeyPolicy<
                 128,
                 ITEMS_PER_THREAD,
                 BLOCK_LOAD_WARP_TRANSPOSE,
@@ -230,7 +230,7 @@ struct DeviceReduceByKeyDispatch
             ITEMS_PER_THREAD            = CUB_MIN(NOMINAL_4B_ITEMS_PER_THREAD, CUB_MAX(1, (NOMINAL_4B_ITEMS_PER_THREAD * 8) / COMBINED_INPUT_BYTES)),
         };
 
-        typedef BlockReduceSweepByKeyPolicy<
+        typedef AgentReduceByKeyPolicy<
                 64,
                 ITEMS_PER_THREAD,
                 BLOCK_LOAD_WARP_TRANSPOSE,
@@ -313,7 +313,7 @@ struct DeviceReduceByKeyDispatch
 
 
     /**
-     * Kernel kernel dispatch configuration.  Mirrors the constants within BlockReduceSweepByKeyPolicy.
+     * Kernel kernel dispatch configuration.  Mirrors the constants within AgentReduceByKeyPolicy.
      */
     struct KernelConfig
     {
@@ -323,15 +323,15 @@ struct DeviceReduceByKeyDispatch
         bool                    two_phase_scatter;
         BlockScanAlgorithm      scan_algorithm;
 
-        template <typename BlockReduceSweepByKeyPolicy>
+        template <typename AgentReduceByKeyPolicy>
         CUB_RUNTIME_FUNCTION __forceinline__
         void Init()
         {
-            block_threads               = BlockReduceSweepByKeyPolicy::BLOCK_THREADS;
-            items_per_thread            = BlockReduceSweepByKeyPolicy::ITEMS_PER_THREAD;
-            load_policy                 = BlockReduceSweepByKeyPolicy::LOAD_ALGORITHM;
-            two_phase_scatter           = BlockReduceSweepByKeyPolicy::TWO_PHASE_SCATTER;
-            scan_algorithm              = BlockReduceSweepByKeyPolicy::SCAN_ALGORITHM;
+            block_threads               = AgentReduceByKeyPolicy::BLOCK_THREADS;
+            items_per_thread            = AgentReduceByKeyPolicy::ITEMS_PER_THREAD;
+            load_policy                 = AgentReduceByKeyPolicy::LOAD_ALGORITHM;
+            two_phase_scatter           = AgentReduceByKeyPolicy::TWO_PHASE_SCATTER;
+            scan_algorithm              = AgentReduceByKeyPolicy::SCAN_ALGORITHM;
         }
 
         CUB_RUNTIME_FUNCTION __forceinline__
diff --git a/cub/device/dispatch/device_rle_dispatch.cuh b/cub/device/dispatch/dispatch_rle.cuh
similarity index 92%
rename from cub/device/dispatch/device_rle_dispatch.cuh
rename to cub/device/dispatch/dispatch_rle.cuh
index 2801ec04df..43107728df 100644
--- a/cub/device/dispatch/device_rle_dispatch.cuh
+++ b/cub/device/dispatch/dispatch_rle.cuh
@@ -38,7 +38,7 @@
 #include <iterator>
 
 #include "device_scan_dispatch.cuh"
-#include "../../block_sweep/block_rle_sweep.cuh"
+#include "../../agent/agent_rle.cuh"
 #include "../../thread/thread_operators.cuh"
 #include "../../grid/grid_queue.cuh"
 #include "../../util_device.cuh"
@@ -63,7 +63,7 @@ namespace cub {
  * Otherwise performs discontinuity selection (keep unique)
  */
 template <
-    typename            BlockRleSweepPolicy,        ///< Parameterized BlockRleSweepPolicy tuning policy type
+    typename            AgentRlePolicy,        ///< Parameterized AgentRlePolicy tuning policy type
     typename            InputIteratorT,             ///< Random-access input iterator type for reading input items \iterator
     typename            OffsetsOutputIteratorT,     ///< Random-access output iterator type for writing run-offset values \iterator
     typename            LengthsOutputIteratorT,     ///< Random-access output iterator type for writing run-length values \iterator
@@ -71,7 +71,7 @@ template <
     typename            ScanTileState,              ///< Tile status interface type
     typename            EqualityOp,                 ///< T equality operator type
     typename            OffsetT>                    ///< Signed integer type for global offsets
-__launch_bounds__ (int(BlockRleSweepPolicy::BLOCK_THREADS))
+__launch_bounds__ (int(AgentRlePolicy::BLOCK_THREADS))
 __global__ void DeviceRleSweepKernel(
     InputIteratorT              d_in,               ///< [in] Pointer to input sequence of data items
     OffsetsOutputIteratorT      d_offsets_out,      ///< [out] Pointer to output sequence of run-offsets
@@ -84,19 +84,19 @@ __global__ void DeviceRleSweepKernel(
     GridQueue<int>              queue)              ///< [in] Drain queue descriptor for dynamically mapping tile data onto thread blocks
 {
     // Thread block type for selecting data from input tiles
-    typedef BlockRleSweep<
-        BlockRleSweepPolicy,
+    typedef AgentRle<
+        AgentRlePolicy,
         InputIteratorT,
         OffsetsOutputIteratorT,
         LengthsOutputIteratorT,
         EqualityOp,
-        OffsetT> BlockRleSweepT;
+        OffsetT> AgentRleT;
 
-    // Shared memory for BlockRleSweep
-    __shared__ typename BlockRleSweepT::TempStorage temp_storage;
+    // Shared memory for AgentRle
+    __shared__ typename AgentRleT::TempStorage temp_storage;
 
     // Process tiles
-    BlockRleSweepT(temp_storage, d_in, d_offsets_out, d_lengths_out, equality_op, num_items).ConsumeRange(
+    AgentRleT(temp_storage, d_in, d_offsets_out, d_lengths_out, equality_op, num_items).ConsumeRange(
         num_tiles,
         queue,
         tile_status,
@@ -153,7 +153,7 @@ struct DeviceRleDispatch
             ITEMS_PER_THREAD            = CUB_MIN(NOMINAL_4B_ITEMS_PER_THREAD, CUB_MAX(1, (NOMINAL_4B_ITEMS_PER_THREAD * 4 / sizeof(T)))),
         };
 
-        typedef BlockRleSweepPolicy<
+        typedef AgentRlePolicy<
                 96,
                 ITEMS_PER_THREAD,
                 BLOCK_LOAD_DIRECT,
@@ -171,7 +171,7 @@ struct DeviceRleDispatch
             ITEMS_PER_THREAD            = CUB_MIN(NOMINAL_4B_ITEMS_PER_THREAD, CUB_MAX(1, (NOMINAL_4B_ITEMS_PER_THREAD * 4 / sizeof(T)))),
         };
 
-        typedef BlockRleSweepPolicy<
+        typedef AgentRlePolicy<
                 256,
                 ITEMS_PER_THREAD,
                 BLOCK_LOAD_WARP_TRANSPOSE,
@@ -189,7 +189,7 @@ struct DeviceRleDispatch
             ITEMS_PER_THREAD            = CUB_MIN(NOMINAL_4B_ITEMS_PER_THREAD, CUB_MAX(1, (NOMINAL_4B_ITEMS_PER_THREAD * 4 / sizeof(T)))),
         };
 
-        typedef BlockRleSweepPolicy<
+        typedef AgentRlePolicy<
                 128,
                 ITEMS_PER_THREAD,
                 BLOCK_LOAD_WARP_TRANSPOSE,
@@ -207,7 +207,7 @@ struct DeviceRleDispatch
             ITEMS_PER_THREAD            = CUB_MIN(NOMINAL_4B_ITEMS_PER_THREAD, CUB_MAX(1, (NOMINAL_4B_ITEMS_PER_THREAD * 4 / sizeof(T)))),
         };
 
-        typedef BlockRleSweepPolicy<
+        typedef AgentRlePolicy<
                 64,
                 ITEMS_PER_THREAD,
                 BLOCK_LOAD_WARP_TRANSPOSE,
@@ -225,7 +225,7 @@ struct DeviceRleDispatch
             ITEMS_PER_THREAD            = CUB_MIN(NOMINAL_4B_ITEMS_PER_THREAD, CUB_MAX(1, (NOMINAL_4B_ITEMS_PER_THREAD * 4 / sizeof(T)))),
         };
 
-        typedef BlockRleSweepPolicy<
+        typedef AgentRlePolicy<
                 256,
                 ITEMS_PER_THREAD,
                 BLOCK_LOAD_WARP_TRANSPOSE,
@@ -308,7 +308,7 @@ struct DeviceRleDispatch
 
 
     /**
-     * Kernel kernel dispatch configuration.  Mirrors the constants within BlockRleSweepPolicy.
+     * Kernel kernel dispatch configuration.  Mirrors the constants within AgentRlePolicy.
      */
     struct KernelConfig
     {
@@ -318,15 +318,15 @@ struct DeviceRleDispatch
         bool                    store_warp_time_slicing;
         BlockScanAlgorithm      scan_algorithm;
 
-        template <typename BlockRleSweepPolicy>
+        template <typename AgentRlePolicy>
         CUB_RUNTIME_FUNCTION __forceinline__
         void Init()
         {
-            block_threads               = BlockRleSweepPolicy::BLOCK_THREADS;
-            items_per_thread            = BlockRleSweepPolicy::ITEMS_PER_THREAD;
-            load_policy                 = BlockRleSweepPolicy::LOAD_ALGORITHM;
-            store_warp_time_slicing     = BlockRleSweepPolicy::STORE_WARP_TIME_SLICING;
-            scan_algorithm              = BlockRleSweepPolicy::SCAN_ALGORITHM;
+            block_threads               = AgentRlePolicy::BLOCK_THREADS;
+            items_per_thread            = AgentRlePolicy::ITEMS_PER_THREAD;
+            load_policy                 = AgentRlePolicy::LOAD_ALGORITHM;
+            store_warp_time_slicing     = AgentRlePolicy::STORE_WARP_TIME_SLICING;
+            scan_algorithm              = AgentRlePolicy::SCAN_ALGORITHM;
         }
 
         CUB_RUNTIME_FUNCTION __forceinline__
diff --git a/cub/device/dispatch/device_scan_dispatch.cuh b/cub/device/dispatch/dispatch_scan.cuh
similarity index 91%
rename from cub/device/dispatch/device_scan_dispatch.cuh
rename to cub/device/dispatch/dispatch_scan.cuh
index 372d01b753..a5cc8dc14e 100644
--- a/cub/device/dispatch/device_scan_dispatch.cuh
+++ b/cub/device/dispatch/dispatch_scan.cuh
@@ -37,7 +37,7 @@
 #include <stdio.h>
 #include <iterator>
 
-#include "../../block_sweep/block_scan_sweep.cuh"
+#include "../../agent/agent_scan.cuh"
 #include "../../thread/thread_operators.cuh"
 #include "../../grid/grid_queue.cuh"
 #include "../../util_debug.cuh"
@@ -79,14 +79,14 @@ __global__ void DeviceScanInitKernel(
  * Scan kernel entry point (multi-block)
  */
 template <
-    typename            BlockScanSweepPolicy,       ///< Parameterized BlockScanSweepPolicy tuning policy type
+    typename            AgentScanPolicy,       ///< Parameterized AgentScanPolicy tuning policy type
     typename            InputIteratorT,             ///< Random-access input iterator type for reading scan inputs \iterator
     typename            OutputIteratorT,            ///< Random-access output iterator type for writing scan outputs \iterator
     typename            ScanTileState,              ///< Tile status interface type
     typename            ScanOp,                     ///< Binary scan functor type having member <tt>T operator()(const T &a, const T &b)</tt>
     typename            Identity,                   ///< Identity value type (cub::NullType for inclusive scans)
     typename            OffsetT>                    ///< Signed integer type for global offsets
-__launch_bounds__ (int(BlockScanSweepPolicy::BLOCK_THREADS))
+__launch_bounds__ (int(AgentScanPolicy::BLOCK_THREADS))
 __global__ void DeviceScanSweepKernel(
     InputIteratorT      d_in,                       ///< Input data
     OutputIteratorT     d_out,                      ///< Output data
@@ -97,19 +97,19 @@ __global__ void DeviceScanSweepKernel(
     GridQueue<OffsetT>  queue)                      ///< Drain queue descriptor for dynamically mapping tile data onto thread blocks
 {
     // Thread block type for scanning input tiles
-    typedef BlockScanSweep<
-        BlockScanSweepPolicy,
+    typedef AgentScan<
+        AgentScanPolicy,
         InputIteratorT,
         OutputIteratorT,
         ScanOp,
         Identity,
-        OffsetT> BlockScanSweepT;
+        OffsetT> AgentScanT;
 
-    // Shared memory for BlockScanSweep
-    __shared__ typename BlockScanSweepT::TempStorage temp_storage;
+    // Shared memory for AgentScan
+    __shared__ typename AgentScanT::TempStorage temp_storage;
 
     // Process tiles
-    BlockScanSweepT(temp_storage, d_in, d_out, scan_op, identity).ConsumeRange(
+    AgentScanT(temp_storage, d_in, d_out, scan_op, identity).ConsumeRange(
         num_items,
         queue,
         tile_status);
@@ -131,7 +131,7 @@ template <
     typename ScanOp,             ///< Binary scan functor type having member <tt>T operator()(const T &a, const T &b)</tt>
     typename Identity,           ///< Identity value type (cub::NullType for inclusive scans)
     typename OffsetT>            ///< Signed integer type for global offsets
-struct DeviceScanDispatch
+struct DispatchScan
 {
     enum
     {
@@ -158,14 +158,12 @@ struct DeviceScanDispatch
         };
 
         // GTX Titan: 29.5B items/s (232.4 GB/s) @ 48M 32-bit T
-        typedef BlockScanSweepPolicy<
+        typedef AgentScanPolicy<
                 128,
                 ITEMS_PER_THREAD,
                 BLOCK_LOAD_DIRECT,
-                false,
                 LOAD_LDG,
-                BLOCK_STORE_WARP_TRANSPOSE,
-                true,
+                BLOCK_STORE_WARP_TRANSPOSE_TIMESLICED,
                 BLOCK_SCAN_RAKING_MEMOIZE>
             RangeScanPolicy;
     };
@@ -178,14 +176,12 @@ struct DeviceScanDispatch
             ITEMS_PER_THREAD            = CUB_MIN(NOMINAL_4B_ITEMS_PER_THREAD, CUB_MAX(1, (NOMINAL_4B_ITEMS_PER_THREAD * 4 / sizeof(T)))),
         };
 
-        typedef BlockScanSweepPolicy<
+        typedef AgentScanPolicy<
                 256,
                 ITEMS_PER_THREAD,
                 BLOCK_LOAD_WARP_TRANSPOSE,
-                false,
                 LOAD_DEFAULT,
                 BLOCK_STORE_WARP_TRANSPOSE,
-                false,
                 BLOCK_SCAN_RAKING_MEMOIZE>
             RangeScanPolicy;
     };
@@ -199,14 +195,12 @@ struct DeviceScanDispatch
         };
 
         // GTX 580: 20.3B items/s (162.3 GB/s) @ 48M 32-bit T
-        typedef BlockScanSweepPolicy<
+        typedef AgentScanPolicy<
                 128,
                 ITEMS_PER_THREAD,
                 BLOCK_LOAD_WARP_TRANSPOSE,
-                false,
                 LOAD_DEFAULT,
                 BLOCK_STORE_WARP_TRANSPOSE,
-                false,
                 BLOCK_SCAN_RAKING_MEMOIZE>
             RangeScanPolicy;
     };
@@ -219,14 +213,12 @@ struct DeviceScanDispatch
             ITEMS_PER_THREAD            = CUB_MIN(NOMINAL_4B_ITEMS_PER_THREAD, CUB_MAX(1, (NOMINAL_4B_ITEMS_PER_THREAD * 4 / sizeof(T)))),
         };
 
-        typedef BlockScanSweepPolicy<
+        typedef AgentScanPolicy<
                 96,
                 ITEMS_PER_THREAD,
                 BLOCK_LOAD_WARP_TRANSPOSE,
-                false,
                 LOAD_DEFAULT,
                 BLOCK_STORE_WARP_TRANSPOSE,
-                false,
                 BLOCK_SCAN_RAKING_MEMOIZE>
             RangeScanPolicy;
     };
@@ -239,14 +231,12 @@ struct DeviceScanDispatch
             ITEMS_PER_THREAD            = CUB_MIN(NOMINAL_4B_ITEMS_PER_THREAD, CUB_MAX(1, (NOMINAL_4B_ITEMS_PER_THREAD * 4 / sizeof(T)))),
         };
 
-        typedef BlockScanSweepPolicy<
+        typedef AgentScanPolicy<
                 64,
                 ITEMS_PER_THREAD,
                 BLOCK_LOAD_WARP_TRANSPOSE,
-                true,
                 LOAD_DEFAULT,
                 BLOCK_STORE_WARP_TRANSPOSE,
-                true,
                 BLOCK_SCAN_WARP_SCANS>
             RangeScanPolicy;
     };
@@ -324,7 +314,7 @@ struct DeviceScanDispatch
 
 
     /**
-     * Kernel kernel dispatch configuration.  Mirrors the constants within BlockScanSweepPolicy.
+     * Kernel kernel dispatch configuration.  Mirrors the constants within AgentScanPolicy.
      */
     struct KernelConfig
     {
@@ -334,15 +324,15 @@ struct DeviceScanDispatch
         BlockStoreAlgorithm     store_policy;
         BlockScanAlgorithm      scan_algorithm;
 
-        template <typename BlockScanSweepPolicy>
+        template <typename AgentScanPolicy>
         CUB_RUNTIME_FUNCTION __forceinline__
         void Init()
         {
-            block_threads               = BlockScanSweepPolicy::BLOCK_THREADS;
-            items_per_thread            = BlockScanSweepPolicy::ITEMS_PER_THREAD;
-            load_policy                 = BlockScanSweepPolicy::LOAD_ALGORITHM;
-            store_policy                = BlockScanSweepPolicy::STORE_ALGORITHM;
-            scan_algorithm              = BlockScanSweepPolicy::SCAN_ALGORITHM;
+            block_threads               = AgentScanPolicy::BLOCK_THREADS;
+            items_per_thread            = AgentScanPolicy::ITEMS_PER_THREAD;
+            load_policy                 = AgentScanPolicy::LOAD_ALGORITHM;
+            store_policy                = AgentScanPolicy::STORE_ALGORITHM;
+            scan_algorithm              = AgentScanPolicy::SCAN_ALGORITHM;
         }
 
         CUB_RUNTIME_FUNCTION __forceinline__
diff --git a/cub/device/dispatch/device_select_dispatch.cuh b/cub/device/dispatch/dispatch_select_if.cuh
similarity index 92%
rename from cub/device/dispatch/device_select_dispatch.cuh
rename to cub/device/dispatch/dispatch_select_if.cuh
index e3f0eb1911..74b6d524bb 100644
--- a/cub/device/dispatch/device_select_dispatch.cuh
+++ b/cub/device/dispatch/dispatch_select_if.cuh
@@ -38,7 +38,7 @@
 #include <iterator>
 
 #include "device_scan_dispatch.cuh"
-#include "../../block_sweep/block_select_sweep.cuh"
+#include "../../agent/agent_select_if.cuh"
 #include "../../thread/thread_operators.cuh"
 #include "../../grid/grid_queue.cuh"
 #include "../../util_device.cuh"
@@ -62,7 +62,7 @@ namespace cub {
  * Otherwise performs discontinuity selection (keep unique)
  */
 template <
-    typename            BlockSelectSweepPolicy,     ///< Parameterized BlockSelectSweepPolicy tuning policy type
+    typename            AgentSelectIfPolicy,     ///< Parameterized AgentSelectIfPolicy tuning policy type
     typename            InputIteratorT,             ///< Random-access input iterator type for reading input items
     typename            FlagsInputIteratorT,              ///< Random-access input iterator type for reading selection flags (NullType* if a selection functor or discontinuity flagging is to be used for selection)
     typename            SelectedOutputIteratorT,            ///< Random-access output iterator type for writing selected items
@@ -72,7 +72,7 @@ template <
     typename            EqualityOp,                 ///< Equality operator type (NullType if selection functor or selection flags is to be used for selection)
     typename            OffsetT,                    ///< Signed integer type for global offsets
     bool                KEEP_REJECTS>               ///< Whether or not we push rejected items to the back of the output
-__launch_bounds__ (int(BlockSelectSweepPolicy::BLOCK_THREADS))
+__launch_bounds__ (int(AgentSelectIfPolicy::BLOCK_THREADS))
 __global__ void DeviceSelectSweepKernel(
     InputIteratorT      d_in,                       ///< [in] Pointer to the input sequence of data items
     FlagsInputIteratorT       d_flags,                    ///< [in] Pointer to the input sequence of selection flags
@@ -86,21 +86,21 @@ __global__ void DeviceSelectSweepKernel(
     GridQueue<int>      queue)                      ///< [in] Drain queue descriptor for dynamically mapping tile data onto thread blocks
 {
     // Thread block type for selecting data from input tiles
-    typedef BlockSelectSweep<
-        BlockSelectSweepPolicy,
+    typedef AgentSelectIf<
+        AgentSelectIfPolicy,
         InputIteratorT,
         FlagsInputIteratorT,
         SelectedOutputIteratorT,
         SelectOp,
         EqualityOp,
         OffsetT,
-        KEEP_REJECTS> BlockSelectSweepT;
+        KEEP_REJECTS> AgentSelectIfT;
 
-    // Shared memory for BlockSelectSweep
-    __shared__ typename BlockSelectSweepT::TempStorage temp_storage;
+    // Shared memory for AgentSelectIf
+    __shared__ typename AgentSelectIfT::TempStorage temp_storage;
 
     // Process tiles
-    BlockSelectSweepT(temp_storage, d_in, d_flags, d_selected_out, select_op, equality_op, num_items).ConsumeRange(
+    AgentSelectIfT(temp_storage, d_in, d_flags, d_selected_out, select_op, equality_op, num_items).ConsumeRange(
         num_tiles,
         queue,
         tile_status,
@@ -126,7 +126,7 @@ template <
     typename    EqualityOp,                     ///< Equality operator type (NullType if selection functor or selection flags is to be used for selection)
     typename    OffsetT,                        ///< Signed integer type for global offsets
     bool        KEEP_REJECTS>                   ///< Whether or not we push rejected items to the back of the output
-struct DeviceSelectDispatch
+struct DispatchSelectIf
 {
     /******************************************************************************
      * Types and constants
@@ -159,7 +159,7 @@ struct DeviceSelectDispatch
             ITEMS_PER_THREAD            = CUB_MIN(NOMINAL_4B_ITEMS_PER_THREAD, CUB_MAX(1, (NOMINAL_4B_ITEMS_PER_THREAD * 4 / sizeof(T)))),
         };
 
-        typedef BlockSelectSweepPolicy<
+        typedef AgentSelectIfPolicy<
                 96,
                 ITEMS_PER_THREAD,
                 BLOCK_LOAD_DIRECT,
@@ -177,7 +177,7 @@ struct DeviceSelectDispatch
             ITEMS_PER_THREAD            = CUB_MIN(NOMINAL_4B_ITEMS_PER_THREAD, CUB_MAX(1, (NOMINAL_4B_ITEMS_PER_THREAD * 4 / sizeof(T)))),
         };
 
-        typedef BlockSelectSweepPolicy<
+        typedef AgentSelectIfPolicy<
                 256,
                 ITEMS_PER_THREAD,
                 BLOCK_LOAD_WARP_TRANSPOSE,
@@ -195,7 +195,7 @@ struct DeviceSelectDispatch
             ITEMS_PER_THREAD            = CUB_MIN(NOMINAL_4B_ITEMS_PER_THREAD, CUB_MAX(1, (NOMINAL_4B_ITEMS_PER_THREAD * 4 / sizeof(T)))),
         };
 
-        typedef BlockSelectSweepPolicy<
+        typedef AgentSelectIfPolicy<
                 128,
                 ITEMS_PER_THREAD,
                 BLOCK_LOAD_WARP_TRANSPOSE,
@@ -213,7 +213,7 @@ struct DeviceSelectDispatch
             ITEMS_PER_THREAD            = CUB_MIN(NOMINAL_4B_ITEMS_PER_THREAD, CUB_MAX(1, (NOMINAL_4B_ITEMS_PER_THREAD * 4 / sizeof(T)))),
         };
 
-        typedef BlockSelectSweepPolicy<
+        typedef AgentSelectIfPolicy<
                 64,
                 ITEMS_PER_THREAD,
                 BLOCK_LOAD_WARP_TRANSPOSE,
@@ -231,7 +231,7 @@ struct DeviceSelectDispatch
             ITEMS_PER_THREAD            = CUB_MIN(NOMINAL_4B_ITEMS_PER_THREAD, CUB_MAX(1, (NOMINAL_4B_ITEMS_PER_THREAD * 4 / sizeof(T)))),
         };
 
-        typedef BlockSelectSweepPolicy<
+        typedef AgentSelectIfPolicy<
                 256,
                 ITEMS_PER_THREAD,
                 BLOCK_LOAD_WARP_TRANSPOSE,
@@ -314,7 +314,7 @@ struct DeviceSelectDispatch
 
 
     /**
-     * Kernel kernel dispatch configuration.  Mirrors the constants within BlockSelectSweepPolicy.
+     * Kernel kernel dispatch configuration.  Mirrors the constants within AgentSelectIfPolicy.
      */
     struct KernelConfig
     {
@@ -324,15 +324,15 @@ struct DeviceSelectDispatch
         bool                    store_warp_time_slicing;
         BlockScanAlgorithm      scan_algorithm;
 
-        template <typename BlockSelectSweepPolicy>
+        template <typename AgentSelectIfPolicy>
         CUB_RUNTIME_FUNCTION __forceinline__
         void Init()
         {
-            block_threads               = BlockSelectSweepPolicy::BLOCK_THREADS;
-            items_per_thread            = BlockSelectSweepPolicy::ITEMS_PER_THREAD;
-            load_policy                 = BlockSelectSweepPolicy::LOAD_ALGORITHM;
-            store_warp_time_slicing     = BlockSelectSweepPolicy::STORE_WARP_TIME_SLICING;
-            scan_algorithm              = BlockSelectSweepPolicy::SCAN_ALGORITHM;
+            block_threads               = AgentSelectIfPolicy::BLOCK_THREADS;
+            items_per_thread            = AgentSelectIfPolicy::ITEMS_PER_THREAD;
+            load_policy                 = AgentSelectIfPolicy::LOAD_ALGORITHM;
+            store_warp_time_slicing     = AgentSelectIfPolicy::STORE_WARP_TIME_SLICING;
+            scan_algorithm              = AgentSelectIfPolicy::SCAN_ALGORITHM;
         }
 
         CUB_RUNTIME_FUNCTION __forceinline__
diff --git a/cub/device/dispatch/device_spmv_dispatch.cuh b/cub/device/dispatch/dispatch_spmv.cuh
similarity index 98%
rename from cub/device/dispatch/device_spmv_dispatch.cuh
rename to cub/device/dispatch/dispatch_spmv.cuh
index 1121ee0446..b979b04f70 100644
--- a/cub/device/dispatch/device_spmv_dispatch.cuh
+++ b/cub/device/dispatch/dispatch_spmv.cuh
@@ -38,7 +38,7 @@
 #include <iterator>
 #include <limits>
 
-#include "../../block_sweep/block_spmv_sweep.cuh"
+#include "../../agent/agent_spmv.cuh"
 #include "../../util_debug.cuh"
 #include "../../util_device.cuh"
 #include "../../thread/thread_search.cuh"
@@ -61,11 +61,11 @@ namespace cub {
  * Spmv privatized sweep kernel entry point (multi-block).  Computes privatized histograms, one per thread block.
  */
 template <
-    typename    BlockSpmvSweepPolicyT,      ///< Parameterized BlockSpmvSweepPolicy tuning policy type
+    typename    AgentSpmvPolicyT,      ///< Parameterized AgentSpmvPolicy tuning policy type
     typename    VertexT,                    ///< Integer type for vertex identifiers
     typename    ValueT,                     ///< Matrix and vector value type
     typename    OffsetT>                    ///< Signed integer type for sequence offsets
-__launch_bounds__ (int(BlockSpmvSweepPolicyT::BLOCK_THREADS))
+__launch_bounds__ (int(AgentSpmvPolicyT::BLOCK_THREADS))
 __global__ void DeviceSpmvSweepKernel(
     ValueT*     d_matrix_values,            ///< [in] Pointer to the array of \p num_nonzeros values of the corresponding nonzero elements of matrix <b>A</b>.
     OffsetT*    d_matrix_row_offsets,       ///< [in] Pointer to the array of \p m + 1 offsets demarcating the start of every row in \p d_matrix_column_indices and \p d_matrix_values (with the final entry being equal to \p num_nonzeros)
@@ -79,17 +79,17 @@ __global__ void DeviceSpmvSweepKernel(
     int         num_nonzeros)               ///< [in] number of nonzero elements of matrix <b>A</b>.
 {
     // Thread block type for compositing input tiles
-    typedef BlockSpmvSweep<
-            BlockSpmvSweepPolicyT,
+    typedef AgentSpmv<
+            AgentSpmvPolicyT,
             VertexT,
             ValueT,
             OffsetT>
-        BlockSpmvSweepT;
+        AgentSpmvT;
 
-    // Shared memory for BlockSpmvSweep
-    __shared__ typename BlockSpmvSweepT::TempStorage temp_storage;
+    // Shared memory for AgentSpmv
+    __shared__ typename AgentSpmvT::TempStorage temp_storage;
 
-    BlockSpmvSweepT block_sweep(
+    AgentSpmvT agent(
         temp_storage,
         d_samples,
         num_output_bins_wrapper.array,
@@ -100,7 +100,7 @@ __global__ void DeviceSpmvSweepKernel(
         privatized_decode_op_wrapper.array);
 
     // Consume input
-    block_sweep.ConsumeTiles(
+    agent.ConsumeTiles(
         num_row_pixels,
         num_rows,
         row_stride_samples,
@@ -295,7 +295,7 @@ struct DeviceSpmvDispatch
     struct Policy110
     {
         // SpmvSweepPolicy
-        typedef BlockSpmvSweepPolicy<
+        typedef AgentSpmvPolicy<
                 512,
                 2,
                 (NUM_CHANNELS == 1) ? BLOCK_LOAD_VECTORIZE : BLOCK_LOAD_DIRECT,
@@ -310,7 +310,7 @@ struct DeviceSpmvDispatch
     struct Policy200
     {
         // SpmvSweepPolicy
-        typedef BlockSpmvSweepPolicy<
+        typedef AgentSpmvPolicy<
                 (NUM_CHANNELS == 1) ? 256 : 128,
                 (NUM_CHANNELS == 1) ? 8 : 3,
                 (NUM_CHANNELS == 1) ? BLOCK_LOAD_VECTORIZE : BLOCK_LOAD_WARP_TRANSPOSE,
@@ -328,7 +328,7 @@ struct DeviceSpmvDispatch
     struct Policy350
     {
         // SpmvSweepPolicy
-        typedef BlockSpmvSweepPolicy<
+        typedef AgentSpmvPolicy<
                 128,
                 (NUM_CHANNELS == 1) ? 8 : 7,
                 (NUM_CHANNELS == 1) ? BLOCK_LOAD_VECTORIZE : BLOCK_LOAD_DIRECT,
@@ -343,7 +343,7 @@ struct DeviceSpmvDispatch
     struct Policy500
     {
         // SpmvSweepPolicy
-        typedef BlockSpmvSweepPolicy<
+        typedef AgentSpmvPolicy<
                 256,
                 8,
                 (NUM_CHANNELS == 1) ? BLOCK_LOAD_VECTORIZE : BLOCK_LOAD_DIRECT,
diff --git a/test/test_device_radix_sort.cu b/test/test_device_radix_sort.cu
index 14dd22f429..0c62c600f2 100644
--- a/test/test_device_radix_sort.cu
+++ b/test/test_device_radix_sort.cu
@@ -883,6 +883,9 @@ int main(int argc, char** argv)
     Test<CUB, unsigned int, NullType, false> (num_items, RANDOM, entropy_reduction, 0, bits, CUB_TYPE_STRING(unsigned int));
     Test<CUB, unsigned long long, NullType, false> (num_items, RANDOM, entropy_reduction, 0, bits, CUB_TYPE_STRING(unsigned int));
 
+    Test<CUB, unsigned int, unsigned int, false> (num_items, RANDOM, entropy_reduction, 0, bits, CUB_TYPE_STRING(unsigned int));
+    Test<CUB, unsigned long long, unsigned int, false> (num_items, RANDOM, entropy_reduction, 0, bits, CUB_TYPE_STRING(unsigned int));
+
 #elif defined(QUICK_TEST)
 
     // Compile/run quick tests