1.8.0    02/15/2018
    - API change: change to the interfaces of cub::ShuffleIndex, cub::ShuffleUp, and 
      cub::ShuffleDown to better compute the PTX shfl control constant for 
      logical warps smaller than 32 threads 
    - Bug fixes: 
        - Issue #112: Bug in WarpScan's broadcast of warp-wide aggregate for 
          logical warps < 32 threads 
        
          		  
//-----------------------------------------------------------------------------

1.7.5    02/08/2018
    - Added radix-sorting support for __half keys
    - Updated sorting policies for improved 8b-key performance 
    - Bug fixes: 
        - Syntax tweaks to mollify Clang
        - Issue #127: DeviceRunLengthEncode::Encode returns wrong results 
        - Issue #128: 7-bit sorting passes fail for sm61 w/ large-values
        
          		  
//-----------------------------------------------------------------------------

1.7.4    09/20/2017
    - Bug fixes: 
        - Issue #114: Can't pair non-trivially-constructible values in radix sort
        - Issue #115: WarpReduce segmented reduction broken in CUDA 9 for logical warp sizes < 32 
          		  
//-----------------------------------------------------------------------------

1.7.3    08/28/2017
    - Bug fixes: 
        - Issue #110: DeviceHistogram null-pointer exception bug for iterator inputs
          		  
//-----------------------------------------------------------------------------

1.7.2    08/26/2017
    - Bug fixes: 
        - Issue #104: Device-wide reduction is now "run-to-run" deterministic for 
          pseudo-associative reduction operators (like floating point addition)
          		  
//-----------------------------------------------------------------------------

1.7.1    08/18/2017
    - Updated Volta radix sorting tuning policies 
    - Bug fixes: 
        - Issue #104 (uint64_t warp-reduce broken for cub 1.7.0 on cuda 8 and older)
        - Issue #103 (Can't mix Thrust 9.0 and CUB)
        - Issue #102 (CUB pulls in windows.h which defines min/max macros that conflict with std::min/std::max)
        - Issue #99 (Radix sorting crashes NVCC on Windows 10 for SM52)
        - Issue #98 (cuda-memcheck: --tool initcheck failed with lineOfSight)
        - Issue #94 (Git clone size)
        - Issue #93 (accept iterators for segment offsets)
        - Issue #87 (CUB uses anonymous unions which is not valid C++)
        - Issue #44 (Check for C++ 11 should be changed that Visual Studio 2013 is also recognized as C++ 11 capable)
          		  
//-----------------------------------------------------------------------------

1.7.0    06/07/2017
    - Compatible with CUDA9 and SM7.x (Volta) independent thread scheduling 
    - API change: remove cub::WarpAll() and cub::WarpAny().  These functions served to 
      emulate __all and __any functionality for SM1.x devices, which did not have those 
      operations.  However, the SM1.x devices are now deprecated in CUDA, and the 
      interfaces of the these two functions are now lacking the lane-mask needed 
      for collectives to run on Volta SMs having independent thread scheduling. 
    - Bug fixes: 
        - Issue #86 Incorrect results with ReduceByKey
          		  
//-----------------------------------------------------------------------------

1.6.4    12/06/2016
    - Updated sm_5x, sm_6x tuning policies for radix sorting (3.5B and 3.4B 
      32b keys/s on TitanX and GTX 1080, respectively)
    - Bug fixes: 
        - Restore fence work-around for scan (reduce-by-key, etc.) hangs 
          in CUDA 8.5
        - Issue 65: DeviceSegmentedRadixSort should allow inputs to have 
          pointer-to-const type 
        - Mollify Clang device-side warnings
        - Remove out-dated VC project files
          		  
//-----------------------------------------------------------------------------

1.6.3    11/20/2016
    - API change: BlockLoad and BlockStore are now templated by the local
      data type, instead of the Iterator type.  This allows for output iterators
      having \p void as their \p value_type (e.g., discard iterators).
    - Updated GP100 tuning policies for radix sorting (6.2B 32b keys/s)
    - Bug fixes: 
        - Issue #74: Warpreduce executes reduction operator for out-of-bounds items
        - Issue #72 (cub:InequalityWrapper::operator() should be non-const)
        - Issue #71 (KeyVairPair won't work if Key has non-trivial ctor)
		- Issue #70 1.5.3 breaks BlockScan API.  Retroactively reversioned
		  from v1.5.3 -> v1.6 to appropriately indicate API change.
		- Issue #69 cub::BlockStore::Store doesn't compile if OutputIteratorT::value_type != T  
        - Issue #68 (cub::TilePrefixCallbackOp::WarpReduce doesn't permit ptx 
          arch specialization)
		- Improved support for Win32 platforms (warnings, alignment, etc)
		  
//-----------------------------------------------------------------------------

1.6.2 (was 1.5.5)    10/25/2016
    - Updated Pascal tuning policies for radix sorting
    - Bug fixes: 
        - Fix for arm64 compilation of caching allocator

//-----------------------------------------------------------------------------

1.6.1 (was 1.5.4)    10/14/2016
    - Bug fixes: 
        - Fix for radix sorting bug introduced by scan refactorization

//-----------------------------------------------------------------------------

1.6.0 (was 1.5.3)    10/11/2016
    - API change: Device/block/warp-wide exclusive scans have been revised to now 
      accept an "initial value" (instead of an "identity value") for seeding the 
      computation with an arbitrary prefix.  
    - API change: Device-wide reductions and scans can now have input sequence types that are 
      different from output sequence types (as long as they are coercible)
      value") for seeding the computation with an arbitrary prefix
    - Reduce repository size (move doxygen binary to doc repository)
    - Minor reductions in block-scan instruction count
    - Bug fixes: 
        - Issue #55: warning in cub/device/dispatch/dispatch_reduce_by_key.cuh 
        - Issue #59: cub::DeviceScan::ExclusiveSum can't prefix sum of float into double
        - Issue #58: Infinite loop in cub::CachingDeviceAllocator::NearestPowerOf
        - Issue #47: Caching allocator needs to clean up cuda error upon successful retry 
        - Issue #46: Very high amount of needed memory from the cub::DeviceHistogram::HistogramEven routine
        - Issue #45: Caching Device Allocator fails with debug output enabled
        - Fix for generic-type reduce-by-key warpscan (sm3.x and newer)

//-----------------------------------------------------------------------------

1.5.2    03/21/2016
	- Improved medium-size scan performance for sm5x (Maxwell)
    - Refactored caching allocator for device memory
   		- Spends less time locked
		- Failure to allocate a block from the runtime will retry once after
		  freeing cached allocations
		- Now respects max-bin (issue where blocks in excess of max-bin were
		  still being retained in free cache)
		- Uses C++11 mutex when available
    - Bug fixes: 
        - Fix for generic-type reduce-by-key warpscan (sm3.x and newer)
          
//-----------------------------------------------------------------------------

1.5.1    12/28/2015
    - Bug fixes: 
        - Fix for incorrect DeviceRadixSort output for some small problems on 
          Maxwell SM52 architectures
        - Fix for macro redefinition warnings when compiling with Thrust sort
          
//-----------------------------------------------------------------------------

1.5.0    12/14/2015
    - New Features:
        - Added new segmented device-wide operations for device-wide sort and 
          reduction primitives.
    - Bug fixes: 
        - Fix for Git Issue 36 (Compilation error with GCC 4.8.4 nvcc 7.0.27) and
          Forums thread (ThreadLoad generates compiler errors when loading from 
          pointer-to-const)
        - Fix for Git Issue 29 (DeviceRadixSort::SortKeys<bool> yields compiler 
          errors)
        - Fix for Git Issue 26 (CUDA error: misaligned address after 
          cub::DeviceRadixSort::SortKeys())
        - Fix for incorrect/crash on 0-length problems, e.g., Git Issue 25 (Floating 
          point exception (core dumped) during cub::DeviceRadixSort::SortKeys)
        - Fix for CUDA 7.5 issues on SM 5.2 with SHFL-based warp-scan and warp-reduction 
          on non-primitive data types (e.g., user-defined structs)
        - Fix for small radix sorting problems where 0 temporary bytes were 
          required and users code was invoking malloc(0) on some systems where
          that returns NULL.  (Impl assumed was asking for size again and was not 
          running the sort.)
          
//-----------------------------------------------------------------------------

1.4.1    04/13/2015
    - Bug fixes: 
        - Fixes for CUDA 7.0 issues with SHFL-based warp-scan and warp-reduction 
          on non-primitive data types (e.g., user-defined structs)
        - Fixes for minor CUDA 7.0 performance regressions in cub::DeviceScan,
          DeviceReduceByKey
        - Fixes to allow cub::DeviceRadixSort and cub::BlockRadixSort on bool types
        - Remove requirement for callers to define the CUB_CDP macro 
          when invoking CUB device-wide rountines using CUDA dynamic parallelism
        - Fix for headers not being included in the proper order (or missing includes)
          for some block-wide functions
          
//-----------------------------------------------------------------------------

1.4.0    03/18/2015
    - New Features:
		- Support and performance tuning for new Maxwell GPU architectures
        - Updated cub::DeviceHistogram implementation that provides the same 
          "histogram-even" and "histogram-range" functionality as IPP/NPP.
          Provides extremely fast and, perhaps more importantly, very 
          uniform performance response across diverse real-world datasets, 
          including pathological (homogeneous) sample distributions (resilience)
        - New cub::DeviceSpmv methods for multiplying sparse matrices by 
          dense vectors, load-balanced using a merge-based parallel decomposition.
        - New cub::DeviceRadixSort sorting entry-points that always return
          the sorted output into the specified buffer (as opposed to the 
          cub::DoubleBuffer in which it could end up in either buffer)
        - New cub::DeviceRunLengthEncode::NonTrivialRuns for finding the starting 
          offsets and lengths of all non-trivial runs (i.e., length > 1) of keys in 
          a given sequence.  (Useful for top-down partitioning algorithms like 
          MSD sorting of very-large keys.)
          
//-----------------------------------------------------------------------------

1.3.2    07/28/2014
    - Bug fixes: 
        - Fix for cub::DeviceReduce where reductions of small problems 
          (small enough to only dispatch a single thread block) would run in 
          the default stream (stream zero) regardless of whether an alternate
          stream was specified.  
          
//-----------------------------------------------------------------------------

1.3.1    05/23/2014
    - Bug fixes: 
        - Workaround for a benign WAW race warning reported by cuda-memcheck
          in BlockScan specialized for BLOCK_SCAN_WARP_SCANS algorithm.
        - Fix for bug in DeviceRadixSort where the algorithm may sort more 
          key bits than the caller specified (up to the nearest radix digit).
        - Fix for ~3% DeviceRadixSort performance regression on Kepler and 
          Fermi that was introduced in v1.3.0.  

//-----------------------------------------------------------------------------

1.3.0    05/12/2014
    - New features:
        - CUB's collective (block-wide, warp-wide) primitives underwent a minor 
          interface refactoring:
            - To provide the appropriate support for multidimensional thread blocks,
              The interfaces for collective classes are now template-parameterized 
              by X, Y, and Z block dimensions (with BLOCK_DIM_Y and BLOCK_DIM_Z being 
              optional, and BLOCK_DIM_X replacing BLOCK_THREADS).  Furthermore, the 
              constructors that accept remapped linear thread-identifiers have been 
              removed: all primitives now assume a row-major thread-ranking for 
              multidimensional thread blocks.  
            - To allow the host program (compiled by the host-pass) to 
              accurately determine the device-specific storage requirements for 
              a given collective (compiled for each device-pass), the interfaces 
              for collective classes are now (optionally) template-parameterized 
              by the desired PTX compute capability. This is useful when 
              aliasing collective storage to shared memory that has been 
              allocated dynamically by the host at the kernel call site.   
            - Most CUB programs having typical 1D usage should not require any 
              changes to accomodate these updates.
        - Added new "combination" WarpScan methods for efficiently computing 
          both inclusive and exclusive prefix scans (and sums).
    - Bug fixes: 
        - Fixed bug in cub::WarpScan (which affected cub::BlockScan and 
          cub::DeviceScan) where incorrect results (e.g., NAN) would often be 
          returned when parameterized for floating-point types (fp32, fp64).
        - Workaround-fix for ptxas error when compiling with with -G flag on Linux 
          (for debug instrumentation) 
        - Misc. workaround-fixes for certain scan scenarios (using custom 
          scan operators) where code compiled for SM1x is run on newer 
          GPUs of higher compute-capability: the compiler could not tell
          which memory space was being used collective operations and was 
          mistakenly using global ops instead of shared ops. 

//-----------------------------------------------------------------------------

1.2.3    04/01/2014
    - Bug fixes: 
        - Fixed access violation bug in DeviceReduce::ReduceByKey for non-primitive value types
        - Fixed code-snippet bug in ArgIndexInputIteratorT documentation 

//-----------------------------------------------------------------------------

1.2.2    03/03/2014
    - New features:
        - Added MS VC++ project solutions for device-wide and block-wide examples 
    - Performance:
        - Added a third algorithmic variant of cub::BlockReduce for improved performance
          when using commutative operators (e.g., numeric addition)
    - Bug fixes: 
        - Fixed bug where inclusion of Thrust headers in a certain order prevented CUB device-wide primitives from working properly

//-----------------------------------------------------------------------------

1.2.0    02/25/2014
    - New features:
        - Added device-wide reduce-by-key (DeviceReduce::ReduceByKey, DeviceReduce::RunLengthEncode) 
    - Performance
        - Improved DeviceScan, DeviceSelect, DevicePartition performance
    - Documentation and testing:
        - Compatible with CUDA 6.0
        - Added performance-portability plots for many device-wide primitives to doc 
        - Update doc and tests to reflect iterator (in)compatibilities with CUDA 5.0 (and older) and Thrust 1.6 (and older).
    - Bug fixes 
        - Revised the operation of temporary tile status bookkeeping for DeviceScan (and similar) to be safe for current code run on future platforms (now uses proper fences)  
        - Fixed DeviceScan bug where Win32 alignment disagreements between host and device regarding user-defined data types would corrupt tile status
        - Fixed BlockScan bug where certain exclusive scans on custom data types for the BLOCK_SCAN_WARP_SCANS variant would return incorrect results for the first thread in the block
        - Added workaround for TexRefInputIteratorTto work with CUDA 6.0
    
//-----------------------------------------------------------------------------

1.1.1    12/11/2013
    - New features:
        - Added TexObjInputIteratorT, TexRefInputIteratorT, CacheModifiedInputIteratorT, and CacheModifiedOutputIterator types for loading & storing arbitrary types through the cache hierarchy.  Compatible with Thrust API. 
        - Added descending sorting to DeviceRadixSort and BlockRadixSort
        - Added min, max, arg-min, and arg-max to DeviceReduce
        - Added DeviceSelect (select-unique, select-if, and select-flagged)
        - Added DevicePartition (partition-if, partition-flagged)
        - Added generic cub::ShuffleUp(), cub::ShuffleDown(), and cub::ShuffleIndex() for warp-wide communication of arbitrary data types (SM3x+)
        - Added cub::MaxSmOccupancy() for accurately determining SM occupancy for any given kernel function pointer
    - Performance
        - Improved DeviceScan and DeviceRadixSort performance for older architectures (SM10-SM30)
    - Interface changes:
        - Refactored block-wide I/O (BlockLoad and BlockStore), removing cache-modifiers from their interfaces.  The CacheModifiedInputIteratorTand CacheModifiedOutputIterator should now be used with BlockLoad and BlockStore to effect that behavior.
        - Rename device-wide "stream_synchronous" param to "debug_synchronous" to avoid confusion about usage
    - Documentation and testing:
        - Added simple examples of device-wide methods
        - Improved doxygen documentation and example snippets
        - Improved test coverege to include up to 21,000 kernel variants and 851,000 unit tests (per architecture, per platform)
    - Bug fixes 
        - Fixed misc DeviceScan, BlockScan, DeviceReduce, and BlockReduce bugs when operating on non-primitive types for older architectures SM10-SM13
        - Fixed DeviceScan / WarpReduction bug: SHFL-based segmented reduction producting incorrect results for multi-word types (size > 4B) on Linux 
        - Fixed BlockScan bug: For warpscan-based scans, not all threads in the first warp were entering the prefix callback functor
        - Fixed DeviceRadixSort bug: race condition with key-value pairs for pre-SM35 architectures
        - Fixed DeviceRadixSort bug: incorrect bitfield-extract behavior with long keys on 64bit Linux
        - Fixed BlockDiscontinuity bug: complation error in for types other than int32/uint32
        - CDP (device-callable) versions of device-wide methods now report the same temporary storage allocation size requirement as their host-callable counterparts
     

//-----------------------------------------------------------------------------

1.0.2    08/23/2013
    - Corrections to code snippet examples for BlockLoad, BlockStore, and BlockDiscontinuity
    - Cleaned up unnecessary/missing header includes.  You can now safely #inlude a specific .cuh (instead of cub.cuh)
    - Bug/compilation fixes for BlockHistogram 

//-----------------------------------------------------------------------------

1.0.1    08/08/2013
    - New collective interface idiom (specialize::construct::invoke).
    - Added best-in-class DeviceRadixSort.  Implements short-circuiting for homogenous digit passes.
    - Added best-in-class DeviceScan.  Implements single-pass "adaptive-lookback" strategy.
    - Significantly improved documentation (with example code snippets) 
    - More extensive regression test suit for aggressively testing collective variants
    - Allow non-trially-constructed types (previously unions had prevented aliasing temporary storage of those types)
    - Improved support for Kepler SHFL (collective ops now use SHFL for types larger than 32b)
    - Better code generation for 64-bit addressing within BlockLoad/BlockStore
    - DeviceHistogram now supports histograms of arbitrary bins
    - Misc. fixes
      - Workarounds for SM10 codegen issues in uncommonly-used WarpScan/Reduce specializations
      - Updates to accommodate CUDA 5.5 dynamic parallelism   


//-----------------------------------------------------------------------------

0.9.4    05/07/2013

    - Fixed compilation errors for SM10-SM13
    - Fixed compilation errors for some WarpScan entrypoints on SM30+
    - Added block-wide histogram (BlockHistogram256)
    - Added device-wide histogram (DeviceHistogram256)
    - Added new BlockScan algorithm variant BLOCK_SCAN_RAKING_MEMOIZE, which 
      trades more register consumption for less shared memory I/O)
    - Updates to BlockRadixRank to use BlockScan (which improves performance
      on Kepler due to SHFL instruction)
    - Allow types other than C++ primitives to be used in WarpScan::*Sum methods 
      if they only have operator + overloaded.  (Previously they also required 
      to support assignment from int(0).) 
    - Update BlockReduce's BLOCK_REDUCE_WARP_REDUCTIONS algorithm to work even 
      when block size is not an even multiple of warp size
    - Added work management utility descriptors (GridQueue, GridEvenShare)
    - Refactoring of DeviceAllocator interface and CachingDeviceAllocator 
      implementation 
    - Misc. documentation updates and corrections. 
     
//-----------------------------------------------------------------------------

0.9.2    04/04/2013

    - Added WarpReduce.  WarpReduce uses the SHFL instruction when applicable. 
      BlockReduce now uses this WarpReduce instead of implementing its own.
    - Misc. fixes for 64-bit Linux compilation warnings and errors.
    - Misc. documentation updates and corrections. 

//-----------------------------------------------------------------------------

0.9.1    03/09/2013

    - Fix for ambiguity in BlockScan::Reduce() between generic reduction and 
      summation.  Summation entrypoints are now called ::Sum(), similar to the 
      convention in BlockScan.
    - Small edits to mainpage documentation and download tracking
    
//-----------------------------------------------------------------------------

0.9.0    03/07/2013    

    - Intial "preview" release.    CUB is the first durable, high-performance library 
      of cooperative block-level, warp-level, and thread-level primitives for CUDA 
      kernel programming.  More primitives and examples coming soon!