diff --git a/CHANGE_LOG.TXT b/CHANGE_LOG.TXT index 8c0b9a89ad..d0e0398204 100644 --- a/CHANGE_LOG.TXT +++ b/CHANGE_LOG.TXT @@ -1,25 +1,39 @@ //----------------------------------------------------------------------------- -1.3.0 03/03/2014 +1.3.1 05/22/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: + - 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. - Bug fixes: - Fixed bug in cub::WarpScan (which affected cub::BlockScan and cub::DeviceScan) where incorrect results (e.g., NAN) would often be @@ -34,7 +48,7 @@ //----------------------------------------------------------------------------- -1.2.3 03/03/2014 +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 ArgIndexInputIterator documentation @@ -43,18 +57,25 @@ 1.2.2 03/03/2014 - New features: - - Added device-wide reduce-by-key (DeviceReduce::ReduceByKey, DeviceReduce::RunLengthEncode) - Added MS VC++ project solutions for device-wide and block-wide examples - Performance: - - Improved DeviceScan, DeviceSelect, DevicePartition 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-portabiltiy plots for many device-wide primitives to doc + - 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: - - Fixed bug where inclusion of Thrust headers in a certain order prevented CUB device-wide primitives from working properly + - 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 diff --git a/cub/block/block_radix_rank.cuh b/cub/block/block_radix_rank.cuh index dc808a06cf..4b5a6a7615 100644 --- a/cub/block/block_radix_rank.cuh +++ b/cub/block/block_radix_rank.cuh @@ -198,13 +198,17 @@ private: UnsignedBits (&keys)[KEYS_PER_THREAD], // Key to decode DigitCounter (&thread_prefixes)[KEYS_PER_THREAD], // Prefix counter value (out parameter) DigitCounter* (&digit_counters)[KEYS_PER_THREAD], // Counter smem offset (out parameter) - int current_bit) // The least-significant bit position of the current digit to extract + int current_bit, // The least-significant bit position of the current digit to extract + int num_bits) // The number of bits in the current digit { + // Get digit + UnsignedBits digit = BFE(keys[COUNT], current_bit, num_bits); + // Get sub-counter - UnsignedBits sub_counter = BFE(keys[COUNT], current_bit + LOG_COUNTER_LANES, LOG_PACKING_RATIO); + UnsignedBits sub_counter = digit >> LOG_COUNTER_LANES; // Get counter lane - UnsignedBits counter_lane = BFE(keys[COUNT], current_bit, LOG_COUNTER_LANES); + UnsignedBits counter_lane = digit & (COUNTER_LANES - 1); if (DESCENDING) { @@ -222,7 +226,7 @@ private: *digit_counters[COUNT] = thread_prefixes[COUNT] + 1; // Iterate next key - Iterate::DecodeKeys(cta, keys, thread_prefixes, digit_counters, current_bit); + Iterate::DecodeKeys(cta, keys, thread_prefixes, digit_counters, current_bit, num_bits); } @@ -253,7 +257,9 @@ private: UnsignedBits (&keys)[KEYS_PER_THREAD], DigitCounter (&thread_prefixes)[KEYS_PER_THREAD], DigitCounter* (&digit_counters)[KEYS_PER_THREAD], - int current_bit) {} + int current_bit, // The least-significant bit position of the current digit to extract + int num_bits) // The number of bits in the current digit + {} // UpdateRanks @@ -261,7 +267,8 @@ private: static __device__ __forceinline__ void UpdateRanks( int (&ranks)[KEYS_PER_THREAD], DigitCounter (&thread_prefixes)[KEYS_PER_THREAD], - DigitCounter *(&digit_counters)[KEYS_PER_THREAD]) {} + DigitCounter *(&digit_counters)[KEYS_PER_THREAD]) + {} }; @@ -416,7 +423,8 @@ public: __device__ __forceinline__ void RankKeys( UnsignedBits (&keys)[KEYS_PER_THREAD], ///< [in] Keys for this tile int (&ranks)[KEYS_PER_THREAD], ///< [out] For each key, the local rank within the tile - int current_bit) ///< [in] The least-significant bit position of the current digit to extract + int current_bit, ///< [in] The least-significant bit position of the current digit to extract + int num_bits) ///< [in] The number of bits in the current digit { DigitCounter thread_prefixes[KEYS_PER_THREAD]; // For each key, the count of previous keys in this tile having the same digit DigitCounter* digit_counters[KEYS_PER_THREAD]; // For each key, the byte-offset of its corresponding digit counter in smem @@ -425,7 +433,7 @@ public: ResetCounters(); // Decode keys and update digit counters - Iterate<0, KEYS_PER_THREAD>::DecodeKeys(*this, keys, thread_prefixes, digit_counters, current_bit); + Iterate<0, KEYS_PER_THREAD>::DecodeKeys(*this, keys, thread_prefixes, digit_counters, current_bit, num_bits); __syncthreads(); @@ -449,10 +457,11 @@ public: UnsignedBits (&keys)[KEYS_PER_THREAD], ///< [in] Keys for this tile int (&ranks)[KEYS_PER_THREAD], ///< [out] For each key, the local rank within the tile (out parameter) int current_bit, ///< [in] The least-significant bit position of the current digit to extract + int num_bits, ///< [in] The number of bits in the current digit int &inclusive_digit_prefix) ///< [out] The incluisve prefix sum for the digit threadIdx.x { // Rank keys - RankKeys(keys, ranks, current_bit); + RankKeys(keys, ranks, current_bit, num_bits); // Get the inclusive and exclusive digit totals corresponding to the calling thread. if ((BLOCK_THREADS == RADIX_DIGITS) || (linear_tid < RADIX_DIGITS)) diff --git a/cub/block/block_radix_sort.cuh b/cub/block/block_radix_sort.cuh index a415f3d973..36006b337c 100644 --- a/cub/block/block_radix_sort.cuh +++ b/cub/block/block_radix_sort.cuh @@ -220,12 +220,14 @@ private: UnsignedBits (&unsigned_keys)[ITEMS_PER_THREAD], int (&ranks)[ITEMS_PER_THREAD], int begin_bit, + int pass_bits, Int2Type is_descending) { AscendingBlockRadixRank(temp_storage.asending_ranking_storage).RankKeys( unsigned_keys, ranks, - begin_bit); + begin_bit, + pass_bits); } /// Rank keys (specialized for descending sort) @@ -299,9 +301,11 @@ private: // Radix sorting passes while (true) { + int pass_bits = CUB_MIN(RADIX_BITS, end_bit - begin_bit); + // Rank the blocked keys int ranks[ITEMS_PER_THREAD]; - RankKeys(unsigned_keys, ranks, begin_bit, is_descending); + RankKeys(unsigned_keys, ranks, begin_bit, pass_bits, is_descending); begin_bit += RADIX_BITS; __syncthreads(); @@ -349,9 +353,11 @@ private: // Radix sorting passes while (true) { + int pass_bits = CUB_MIN(RADIX_BITS, end_bit - begin_bit); + // Rank the blocked keys int ranks[ITEMS_PER_THREAD]; - RankKeys(unsigned_keys, ranks, begin_bit, is_descending); + RankKeys(unsigned_keys, ranks, begin_bit, pass_bits, is_descending); begin_bit += RADIX_BITS; __syncthreads(); diff --git a/cub/block_range/block_range_radix_sort_downsweep.cuh b/cub/block_range/block_range_radix_sort_downsweep.cuh index 50546a5b79..4141315ed9 100644 --- a/cub/block_range/block_range_radix_sort_downsweep.cuh +++ b/cub/block_range/block_range_radix_sort_downsweep.cuh @@ -242,6 +242,9 @@ struct BlockRangeRadixSortDownsweep // The least-significant bit position of the current digit to extract int current_bit; + // Number of bits in current digit + int num_bits; + // Whether to short-ciruit bool short_circuit; @@ -261,7 +264,7 @@ struct BlockRangeRadixSortDownsweep #pragma unroll for (int KEY = 0; KEY < ITEMS_PER_THREAD; KEY++) { - UnsignedBits digit = BFE(twiddled_keys[KEY], current_bit, RADIX_BITS); + UnsignedBits digit = BFE(twiddled_keys[KEY], current_bit, num_bits); // Lookup base digit offset from shared memory relative_bin_offsets[KEY] = temp_storage.relative_bin_offsets[digit]; @@ -522,6 +525,7 @@ struct BlockRangeRadixSortDownsweep twiddled_keys, ranks, current_bit, + num_bits, inclusive_digit_prefix); // Update global scatter base offsets for each digit @@ -639,7 +643,8 @@ struct BlockRangeRadixSortDownsweep Key *d_keys_out, Value *d_values_in, Value *d_values_out, - int current_bit) + int current_bit, + int num_bits) : temp_storage(temp_storage.Alias()), bin_offset(bin_offset), @@ -648,6 +653,7 @@ struct BlockRangeRadixSortDownsweep d_values_in(d_values_in), d_values_out(d_values_out), current_bit(current_bit), + num_bits(num_bits), short_circuit(false) {} @@ -663,14 +669,16 @@ struct BlockRangeRadixSortDownsweep Key *d_keys_out, Value *d_values_in, Value *d_values_out, - int current_bit) + int current_bit, + int num_bits) : temp_storage(temp_storage.Alias()), d_keys_in(reinterpret_cast(d_keys_in)), d_keys_out(reinterpret_cast(d_keys_out)), d_values_in(d_values_in), d_values_out(d_values_out), - current_bit(current_bit) + current_bit(current_bit), + num_bits(num_bits) { // Load digit bin offsets (each of the first RADIX_DIGITS threads will load an offset for that digit) if (threadIdx.x < RADIX_DIGITS) diff --git a/cub/block_range/block_range_radix_sort_upsweep.cuh b/cub/block_range/block_range_radix_sort_upsweep.cuh index efb2f7bd30..faadbd3f47 100644 --- a/cub/block_range/block_range_radix_sort_upsweep.cuh +++ b/cub/block_range/block_range_radix_sort_upsweep.cuh @@ -171,6 +171,9 @@ struct BlockRangeRadixSortUpsweep // The least-significant bit position of the current digit to extract int current_bit; + // Number of bits in current digit + int num_bits; + //--------------------------------------------------------------------- @@ -214,15 +217,17 @@ struct BlockRangeRadixSortUpsweep // Perform transform op UnsignedBits converted_key = Traits::TwiddleIn(key); - // Add in sub-counter offset - UnsignedBits sub_counter = BFE(converted_key, current_bit, LOG_PACKING_RATIO); + // Extract current digit bits + UnsignedBits digit = BFE(converted_key, current_bit, num_bits); + + // Get sub-counter offset + UnsignedBits sub_counter = digit & (PACKING_RATIO - 1); - // Add in row offset - UnsignedBits row_offset = BFE(converted_key, current_bit + LOG_PACKING_RATIO, LOG_COUNTER_LANES); + // Get row offset + UnsignedBits row_offset = digit >> LOG_PACKING_RATIO; // Increment counter temp_storage.digit_counters[row_offset][threadIdx.x][sub_counter]++; - } @@ -372,11 +377,13 @@ struct BlockRangeRadixSortUpsweep __device__ __forceinline__ BlockRangeRadixSortUpsweep( TempStorage &temp_storage, Key *d_keys_in, - int current_bit) + int current_bit, + int num_bits) : temp_storage(temp_storage.Alias()), d_keys_in(reinterpret_cast(d_keys_in)), - current_bit(current_bit) + current_bit(current_bit), + num_bits(num_bits) {} diff --git a/cub/device/dispatch/device_radix_sort_dispatch.cuh b/cub/device/dispatch/device_radix_sort_dispatch.cuh index bde3b9dbdb..028a5684ec 100644 --- a/cub/device/dispatch/device_radix_sort_dispatch.cuh +++ b/cub/device/dispatch/device_radix_sort_dispatch.cuh @@ -63,12 +63,13 @@ template < bool DESCENDING, ///< Whether or not the sorted-order is high-to-low typename Key, ///< Key type typename Offset> ///< Signed integer type for global offsets -__launch_bounds__ (int(BlockRangeRadixSortUpsweepPolicy::BLOCK_THREADS)) +__launch_bounds__ (int(BlockRangeRadixSortUpsweepPolicy::BLOCK_THREADS), 1) __global__ void RadixSortUpsweepKernel( Key *d_keys, ///< [in] Input keys buffer Offset *d_spine, ///< [out] Privatized (per block) digit histograms (striped, i.e., 0s counts from each block, then 1s counts from each block, etc.) Offset num_items, ///< [in] Total number of input data items int current_bit, ///< [in] Bit position of current radix digit + int num_bits, ///< [in] Number of bits of current radix digit bool first_pass, ///< [in] Whether this is the first digit pass GridEvenShare even_share) ///< [in] Even-share descriptor for mapping an equal number of tiles onto each thread block { @@ -82,7 +83,7 @@ __global__ void RadixSortUpsweepKernel( even_share.BlockInit(); Offset bin_count; - BlockRangeRadixSortUpsweepT(temp_storage, d_keys, current_bit).ProcessRegion( + BlockRangeRadixSortUpsweepT(temp_storage, d_keys, current_bit, num_bits).ProcessRegion( even_share.block_offset, even_share.block_end, bin_count); @@ -136,23 +137,24 @@ __global__ void RadixSortScanKernel( * Downsweep pass kernel entry point (multi-block). Scatters keys (and values) into corresponding bins for the current digit place. */ template < - typename BlockRangeRadixSortDownsweepPolicy, ///< Parameterizable tuning policy type for cub::BlockRangeRadixSortUpsweep abstraction - bool DESCENDING, ///< Whether or not the sorted-order is high-to-low - typename Key, ///< Key type - typename Value, ///< Value type - typename Offset> ///< Signed integer type for global offsets -__launch_bounds__ (int(BlockRangeRadixSortDownsweepPolicy::BLOCK_THREADS)) + typename BlockRangeRadixSortDownsweepPolicy, ///< Parameterizable tuning policy type for cub::BlockRangeRadixSortUpsweep abstraction + bool DESCENDING, ///< Whether or not the sorted-order is high-to-low + typename Key, ///< Key type + typename Value, ///< Value type + typename Offset> ///< Signed integer type for global offsets +__launch_bounds__ (int(BlockRangeRadixSortDownsweepPolicy::BLOCK_THREADS), 1) __global__ void RadixSortDownsweepKernel( - Key *d_keys_in, ///< [in] Input keys ping buffer - Key *d_keys_out, ///< [in] Output keys pong buffer - Value *d_values_in, ///< [in] Input values ping buffer - Value *d_values_out, ///< [in] Output values pong buffer - Offset *d_spine, ///< [in] Scan of privatized (per block) digit histograms (striped, i.e., 0s counts from each block, then 1s counts from each block, etc.) - Offset num_items, ///< [in] Total number of input data items - int current_bit, ///< [in] Bit position of current radix digit - bool first_pass, ///< [in] Whether this is the first digit pass - bool last_pass, ///< [in] Whether this is the last digit pass - GridEvenShare even_share) ///< [in] Even-share descriptor for mapping an equal number of tiles onto each thread block + Key *d_keys_in, ///< [in] Input keys ping buffer + Key *d_keys_out, ///< [in] Output keys pong buffer + Value *d_values_in, ///< [in] Input values ping buffer + Value *d_values_out, ///< [in] Output values pong buffer + Offset *d_spine, ///< [in] Scan of privatized (per block) digit histograms (striped, i.e., 0s counts from each block, then 1s counts from each block, etc.) + Offset num_items, ///< [in] Total number of input data items + int current_bit, ///< [in] Bit position of current radix digit + int num_bits, ///< [in] Number of bits of current radix digit + bool first_pass, ///< [in] Whether this is the first digit pass + bool last_pass, ///< [in] Whether this is the last digit pass + GridEvenShare even_share) ///< [in] Even-share descriptor for mapping an equal number of tiles onto each thread block { // Parameterize BlockRangeRadixSortDownsweep type for the current configuration typedef BlockRangeRadixSortDownsweep BlockRangeRadixSortDownsweepT; @@ -164,7 +166,7 @@ __global__ void RadixSortDownsweepKernel( even_share.BlockInit(); // Process input tiles - BlockRangeRadixSortDownsweepT(temp_storage, num_items, d_spine, d_keys_in, d_keys_out, d_values_in, d_values_out, current_bit).ProcessRegion( + BlockRangeRadixSortDownsweepT(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); } @@ -209,7 +211,7 @@ struct DeviceRadixSortDispatch typedef typename If::Type AltUpsweepPolicy; // ScanPolicy - typedef BlockRangeScanPolicy <1024, 4, BLOCK_LOAD_VECTORIZE, false, LOAD_DEFAULT, BLOCK_STORE_VECTORIZE, false, BLOCK_SCAN_RAKING_MEMOIZE> ScanPolicy; + typedef BlockRangeScanPolicy <1024, 4, BLOCK_LOAD_VECTORIZE, false, LOAD_DEFAULT, BLOCK_STORE_VECTORIZE, false, BLOCK_SCAN_WARP_SCANS> ScanPolicy; // Primary DownsweepPolicy typedef BlockRangeRadixSortDownsweepPolicy <64, CUB_MAX(1, 18 / SCALE_FACTOR), BLOCK_LOAD_DIRECT, LOAD_LDG, false, true, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeEightByte, RADIX_BITS> DownsweepPolicyKeys; @@ -280,13 +282,13 @@ struct DeviceRadixSortDispatch typedef BlockRangeScanPolicy <512, 4, BLOCK_LOAD_VECTORIZE, false, LOAD_DEFAULT, BLOCK_STORE_VECTORIZE, false, BLOCK_SCAN_RAKING_MEMOIZE> ScanPolicy; // DownsweepPolicy - typedef BlockRangeRadixSortDownsweepPolicy <64, CUB_MAX(1, 18 / SCALE_FACTOR), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, false, BLOCK_SCAN_RAKING_MEMOIZE, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, RADIX_BITS> DownsweepPolicyKeys; - typedef BlockRangeRadixSortDownsweepPolicy <128, CUB_MAX(1, 13 / SCALE_FACTOR), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, false, BLOCK_SCAN_RAKING_MEMOIZE, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, RADIX_BITS> DownsweepPolicyPairs; + typedef BlockRangeRadixSortDownsweepPolicy <64, CUB_MAX(1, 18 / SCALE_FACTOR), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, RADIX_BITS> DownsweepPolicyKeys; + typedef BlockRangeRadixSortDownsweepPolicy <128, CUB_MAX(1, 13 / SCALE_FACTOR), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, RADIX_BITS> DownsweepPolicyPairs; typedef typename If::Type DownsweepPolicy; // Alternate DownsweepPolicy for (RADIX_BITS-1)-bit passes - typedef BlockRangeRadixSortDownsweepPolicy <64, CUB_MAX(1, 18 / SCALE_FACTOR), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, false, BLOCK_SCAN_RAKING_MEMOIZE, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, RADIX_BITS - 1> AltDownsweepPolicyKeys; - typedef BlockRangeRadixSortDownsweepPolicy <128, CUB_MAX(1, 13 / SCALE_FACTOR), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, false, BLOCK_SCAN_RAKING_MEMOIZE, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, RADIX_BITS - 1> AltDownsweepPolicyPairs; + typedef BlockRangeRadixSortDownsweepPolicy <64, CUB_MAX(1, 18 / SCALE_FACTOR), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, RADIX_BITS - 1> AltDownsweepPolicyKeys; + typedef BlockRangeRadixSortDownsweepPolicy <128, CUB_MAX(1, 13 / SCALE_FACTOR), BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, false, false, BLOCK_SCAN_WARP_SCANS, RADIX_SORT_SCATTER_TWO_PHASE, cudaSharedMemBankSizeFourByte, RADIX_BITS - 1> AltDownsweepPolicyPairs; typedef typename If::Type AltDownsweepPolicy; }; @@ -650,6 +652,8 @@ struct DeviceRadixSortDispatch int current_bit = begin_bit; while (current_bit < end_bit) { + int num_bits = CUB_MIN(end_bit - current_bit, downsweep_config.radix_bits); + #if (CUB_PTX_ARCH == 0) // Update smem config if necessary if (current_smem_config != upsweep_config.smem_config) @@ -670,6 +674,7 @@ struct DeviceRadixSortDispatch d_spine, num_items, current_bit, + num_bits, (current_bit == begin_bit), even_share); @@ -716,6 +721,7 @@ struct DeviceRadixSortDispatch d_spine, num_items, current_bit, + num_bits, (current_bit == begin_bit), (current_bit + downsweep_config.radix_bits >= end_bit), even_share); diff --git a/cub/device/dispatch/device_reduce_by_key_dispatch.cuh b/cub/device/dispatch/device_reduce_by_key_dispatch.cuh index 0a56560ebf..81c028e19e 100644 --- a/cub/device/dispatch/device_reduce_by_key_dispatch.cuh +++ b/cub/device/dispatch/device_reduce_by_key_dispatch.cuh @@ -64,7 +64,7 @@ template < typename ValueInputIterator, ///< Random-access input iterator type for values typename ValueOutputIterator, ///< Random-access output iterator type for values typename NumSegmentsIterator, ///< Output iterator type for recording number of segments encountered - typename ScanTileState, ///< Tile status interface type + typename ScanTileState, ///< Tile status interface type typename EqualityOp, ///< Key equality operator type typename ReductionOp, ///< Value reduction operator type typename Offset> ///< Signed integer type for global offsets @@ -75,7 +75,7 @@ __global__ void ReduceByKeyRegionKernel( ValueInputIterator d_values_in, ///< [in] Pointer to consecutive runs of input values ValueOutputIterator d_values_out, ///< [in] Pointer to output value aggregates (one aggregate per run) NumSegmentsIterator d_num_segments, ///< [in] Pointer to total number of runs - ScanTileState tile_status, ///< [in] Tile status interface + ScanTileState tile_status, ///< [in] Tile status interface EqualityOp equality_op, ///< [in] Key equality operator ReductionOp reduction_op, ///< [in] Value reduction operator Offset num_items, ///< [in] Total number of items to select from diff --git a/cub/warp/warp_reduce.cuh b/cub/warp/warp_reduce.cuh index 3f25a80b41..1cd3fe0cff 100644 --- a/cub/warp/warp_reduce.cuh +++ b/cub/warp/warp_reduce.cuh @@ -215,7 +215,7 @@ public: /** - * \brief Computes a warp-wide sum in each active warp. The output is valid in warp lane0. + * \brief Computes a warp-wide sum in the calling warp. The output is valid in warp lane0. * * \smemreuse * @@ -255,9 +255,9 @@ public: } /** - * \brief Computes a partially-full warp-wide sum in each active warp. The output is valid in warp lane0. + * \brief Computes a partially-full warp-wide sum in the calling warp. The output is valid in warp lane0. * - * All threads in each logical warp must agree on the same value for \p valid_items. Otherwise the result is undefined. + * All threads across the calling warp must agree on the same value for \p valid_items. Otherwise the result is undefined. * * \smemreuse * @@ -309,7 +309,7 @@ public: /** - * \brief Computes a segmented sum in each active warp where segments are defined by head-flags. The sum of each segment is returned to the first lane in that segment (which always includes lane0). + * \brief Computes a segmented sum in the calling warp where segments are defined by head-flags. The sum of each segment is returned to the first lane in that segment (which always includes lane0). * * \smemreuse * @@ -357,7 +357,7 @@ public: /** - * \brief Computes a segmented sum in each active warp where segments are defined by tail-flags. The sum of each segment is returned to the first lane in that segment (which always includes lane0). + * \brief Computes a segmented sum in the calling warp where segments are defined by tail-flags. The sum of each segment is returned to the first lane in that segment (which always includes lane0). * * \smemreuse * @@ -411,7 +411,7 @@ public: //@{ /** - * \brief Computes a warp-wide reduction in each active warp using the specified binary reduction functor. The output is valid in warp lane0. + * \brief Computes a warp-wide reduction in the calling warp using the specified binary reduction functor. The output is valid in warp lane0. * * Supports non-commutative reduction operators * @@ -457,9 +457,9 @@ public: } /** - * \brief Computes a partially-full warp-wide reduction in each active warp using the specified binary reduction functor. The output is valid in warp lane0. + * \brief Computes a partially-full warp-wide reduction in the calling warp using the specified binary reduction functor. The output is valid in warp lane0. * - * All threads in each logical warp must agree on the same value for \p valid_items. Otherwise the result is undefined. + * All threads across the calling warp must agree on the same value for \p valid_items. Otherwise the result is undefined. * * Supports non-commutative reduction operators * @@ -516,7 +516,7 @@ public: /** - * \brief Computes a segmented reduction in each active warp where segments are defined by head-flags. The reduction of each segment is returned to the first lane in that segment (which always includes lane0). + * \brief Computes a segmented reduction in the calling warp where segments are defined by head-flags. The reduction of each segment is returned to the first lane in that segment (which always includes lane0). * * Supports non-commutative reduction operators * @@ -567,7 +567,7 @@ public: /** - * \brief Computes a segmented reduction in each active warp where segments are defined by tail-flags. The reduction of each segment is returned to the first lane in that segment (which always includes lane0). + * \brief Computes a segmented reduction in the calling warp where segments are defined by tail-flags. The reduction of each segment is returned to the first lane in that segment (which always includes lane0). * * Supports non-commutative reduction operators * diff --git a/cub/warp/warp_scan.cuh b/cub/warp/warp_scan.cuh index b7e2c93c0a..ba604ec311 100644 --- a/cub/warp/warp_scan.cuh +++ b/cub/warp/warp_scan.cuh @@ -216,7 +216,7 @@ public: /** - * \brief Computes an inclusive prefix sum in each logical warp. + * \brief Computes an inclusive prefix sum across the calling warp. * * \smemreuse * @@ -257,7 +257,7 @@ public: /** - * \brief Computes an inclusive prefix sum in each logical warp. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. + * \brief Computes an inclusive prefix sum across the calling warp. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. * * The \p warp_aggregate is undefined in threads other than warp-lane0. * @@ -302,7 +302,7 @@ public: /** - * \brief Computes an inclusive prefix sum in each logical warp. Instead of using 0 as the warp-wide prefix, the call-back functor \p warp_prefix_op is invoked to provide the "seed" value that logically prefixes the warp's scan inputs. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. + * \brief Computes an inclusive prefix sum across the calling warp. Instead of using 0 as the warp-wide prefix, the call-back functor \p warp_prefix_op is invoked to provide the "seed" value that logically prefixes the warp's scan inputs. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. * * The \p warp_aggregate is undefined in threads other than warp-lane0. * @@ -415,7 +415,7 @@ private: InternalWarpScan(temp_storage).Scan(input, inclusive_output, exclusive_output, identity, cub::Sum()); } - /// Computes an exclusive prefix sum in each logical warp. + /// Computes an exclusive prefix sum across the calling warp. __device__ __forceinline__ void ExclusiveSum(T input, T &output, Int2Type is_integer) { // Compute exclusive warp scan from inclusive warp scan @@ -424,7 +424,7 @@ private: output = inclusive - input; } - /// Computes an exclusive prefix sum in each logical warp. Specialized for non-integer types. + /// Computes an exclusive prefix sum across the calling warp. Specialized for non-integer types. __device__ __forceinline__ void ExclusiveSum(T input, T &output, Int2Type is_integer) { // Delegate to regular scan for non-integer types (because we won't be able to use subtraction) @@ -432,7 +432,7 @@ private: ExclusiveScan(input, output, identity, cub::Sum()); } - /// Computes an exclusive prefix sum in each logical warp. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. + /// Computes an exclusive prefix sum across the calling warp. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. __device__ __forceinline__ void ExclusiveSum(T input, T &output, T &warp_aggregate, Int2Type is_integer) { // Compute exclusive warp scan from inclusive warp scan @@ -441,7 +441,7 @@ private: output = inclusive - input; } - /// Computes an exclusive prefix sum in each logical warp. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. Specialized for non-integer types. + /// Computes an exclusive prefix sum across the calling warp. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. Specialized for non-integer types. __device__ __forceinline__ void ExclusiveSum(T input, T &output, T &warp_aggregate, Int2Type is_integer) { // Delegate to regular scan for non-integer types (because we won't be able to use subtraction) @@ -449,7 +449,7 @@ private: ExclusiveScan(input, output, identity, cub::Sum(), warp_aggregate); } - /// Computes an exclusive prefix sum in each logical warp. Instead of using 0 as the warp-wide prefix, the call-back functor \p warp_prefix_op is invoked to provide the "seed" value that logically prefixes the warp's scan inputs. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. + /// Computes an exclusive prefix sum across the calling warp. Instead of using 0 as the warp-wide prefix, the call-back functor \p warp_prefix_op is invoked to provide the "seed" value that logically prefixes the warp's scan inputs. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. template __device__ __forceinline__ void ExclusiveSum(T input, T &output, T &warp_aggregate, WarpPrefixCallbackOp &warp_prefix_op, Int2Type is_integer) { @@ -459,7 +459,7 @@ private: output = inclusive - input; } - /// Computes an exclusive prefix sum in each logical warp. Instead of using 0 as the warp-wide prefix, the call-back functor \p warp_prefix_op is invoked to provide the "seed" value that logically prefixes the warp's scan inputs. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. Specialized for non-integer types. + /// Computes an exclusive prefix sum across the calling warp. Instead of using 0 as the warp-wide prefix, the call-back functor \p warp_prefix_op is invoked to provide the "seed" value that logically prefixes the warp's scan inputs. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. Specialized for non-integer types. template __device__ __forceinline__ void ExclusiveSum(T input, T &output, T &warp_aggregate, WarpPrefixCallbackOp &warp_prefix_op, Int2Type is_integer) { @@ -478,7 +478,7 @@ public: /** - * \brief Computes an exclusive prefix sum in each logical warp. + * \brief Computes an exclusive prefix sum across the calling warp. * * This operation assumes the value of obtained by the T's default * constructor (or by zero-initialization if no user-defined default @@ -525,7 +525,7 @@ public: /** - * \brief Computes an exclusive prefix sum in each logical warp. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. + * \brief Computes an exclusive prefix sum across the calling warp. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. * * This operation assumes the value of obtained by the T's default * constructor (or by zero-initialization if no user-defined default @@ -573,7 +573,7 @@ public: /** - * \brief Computes an exclusive prefix sum in each logical warp. Instead of using 0 as the warp-wide prefix, the call-back functor \p warp_prefix_op is invoked to provide the "seed" value that logically prefixes the warp's scan inputs. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. + * \brief Computes an exclusive prefix sum across the calling warp. Instead of using 0 as the warp-wide prefix, the call-back functor \p warp_prefix_op is invoked to provide the "seed" value that logically prefixes the warp's scan inputs. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. * * This operation assumes the value of obtained by the T's default * constructor (or by zero-initialization if no user-defined default @@ -668,7 +668,7 @@ public: //@{ /** - * \brief Computes an inclusive prefix sum using the specified binary scan functor in each logical warp. + * \brief Computes an inclusive prefix scan using the specified binary scan functor across the calling warp. * * Supports non-commutative scan operators. * @@ -715,7 +715,7 @@ public: /** - * \brief Computes an inclusive prefix sum using the specified binary scan functor in each logical warp. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. + * \brief Computes an inclusive prefix scan using the specified binary scan functor across the calling warp. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. * * Supports non-commutative scan operators. * @@ -767,7 +767,7 @@ public: /** - * \brief Computes an inclusive prefix sum using the specified binary scan functor in each logical warp. The call-back functor \p warp_prefix_op is invoked to provide the "seed" value that logically prefixes the warp's scan inputs. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. + * \brief Computes an inclusive prefix scan using the specified binary scan functor across the calling warp. The call-back functor \p warp_prefix_op is invoked to provide the "seed" value that logically prefixes the warp's scan inputs. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. * * The \p warp_prefix_op functor must implement a member function T operator()(T warp_aggregate). * The functor's input parameter \p warp_aggregate is the same value also returned by the scan operation. @@ -873,7 +873,7 @@ public: //@{ /** - * \brief Computes an exclusive prefix scan using the specified binary scan functor in each logical warp. + * \brief Computes an exclusive prefix scan using the specified binary scan functor across the calling warp. * * Supports non-commutative scan operators. * @@ -922,7 +922,7 @@ public: /** - * \brief Computes an exclusive prefix scan using the specified binary scan functor in each logical warp. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. + * \brief Computes an exclusive prefix scan using the specified binary scan functor across the calling warp. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. * * Supports non-commutative scan operators. * @@ -974,7 +974,7 @@ public: /** - * \brief Computes an exclusive prefix scan using the specified binary scan functor in each logical warp. The call-back functor \p warp_prefix_op is invoked to provide the "seed" value that logically prefixes the warp's scan inputs. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. + * \brief Computes an exclusive prefix scan using the specified binary scan functor across the calling warp. The call-back functor \p warp_prefix_op is invoked to provide the "seed" value that logically prefixes the warp's scan inputs. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. * * The \p warp_prefix_op functor must implement a member function T operator()(T warp_aggregate). * The functor's input parameter \p warp_aggregate is the same value also returned by the scan operation. @@ -1083,7 +1083,7 @@ public: /** - * \brief Computes an exclusive prefix scan using the specified binary scan functor in each logical warp. Because no identity value is supplied, the \p output computed for warp-lane0 is undefined. + * \brief Computes an exclusive prefix scan using the specified binary scan functor across the calling warp. Because no identity value is supplied, the \p output computed for warp-lane0 is undefined. * * Supports non-commutative scan operators. * @@ -1116,7 +1116,7 @@ public: * Suppose the set of input \p thread_data across the block of threads is {0, -1, 2, -3, ..., 126, -127}. * The corresponding output \p thread_data in the first warp would be * ?, 0, 0, 2, ..., 28, 30, the output for the second warp would be ?, 32, 32, 34, ..., 60, 62, etc. - * (The output \p thread_data in each warp lane0 is undefined.) + * (The output \p thread_data in warp lane0 is undefined.) * * \tparam ScanOp [inferred] Binary scan operator type having member T operator()(const T &a, const T &b) */ @@ -1132,7 +1132,7 @@ public: /** - * \brief Computes an exclusive prefix scan using the specified binary scan functor in each logical warp. Because no identity value is supplied, the \p output computed for warp-lane0 is undefined. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. + * \brief Computes an exclusive prefix scan using the specified binary scan functor across the calling warp. Because no identity value is supplied, the \p output computed for warp-lane0 is undefined. Also provides every thread with the warp-wide \p warp_aggregate of all inputs. * * Supports non-commutative scan operators. * @@ -1166,7 +1166,7 @@ public: * Suppose the set of input \p thread_data across the block of threads is {0, -1, 2, -3, ..., 126, -127}. * The corresponding output \p thread_data in the first warp would be * ?, 0, 0, 2, ..., 28, 30, the output for the second warp would be ?, 32, 32, 34, ..., 60, 62, etc. - * (The output \p thread_data in each warp lane0 is undefined.) Furthermore, \p warp_aggregate would be assigned \p 30 for threads in the first warp, \p 62 for threads + * (The output \p thread_data in warp lane0 is undefined.) Furthermore, \p warp_aggregate would be assigned \p 30 for threads in the first warp, \p 62 for threads * in the second warp, etc. * * \tparam ScanOp [inferred] Binary scan operator type having member T operator()(const T &a, const T &b) @@ -1183,7 +1183,7 @@ public: /** - * \brief Computes an exclusive prefix scan using the specified binary scan functor in each logical warp. The \p warp_prefix_op value from thread-thread-lane0 is applied to all scan outputs. Also computes the warp-wide \p warp_aggregate of all inputs for thread-thread-lane0. + * \brief Computes an exclusive prefix scan using the specified binary scan functor across the calling warp. The \p warp_prefix_op value from thread-thread-lane0 is applied to all scan outputs. Also computes the warp-wide \p warp_aggregate of all inputs for thread-thread-lane0. * * The \p warp_prefix_op functor must implement a member function T operator()(T warp_aggregate)}. * The functor's input parameter \p warp_aggregate is the same value also returned by the scan operation. @@ -1288,7 +1288,48 @@ public: *********************************************************************/ //@{ - /// Combination scan with identity + /** + * \brief Computes both inclusive and exclusive prefix sums across the calling warp. + * + * This operation assumes the value of obtained by the T's default + * constructor (or by zero-initialization if no user-defined default + * constructor exists) is suitable as the identity value "zero" for + * addition. + * + * \smemreuse + * + * \par Snippet + * The code snippet below illustrates four concurrent warp-wide prefix sums within a block of + * 128 threads (one per each of the 32-thread warps). + * \par + * \code + * #include + * + * __global__ void ExampleKernel(...) + * { + * // Specialize WarpScan for type int + * typedef cub::WarpScan WarpScan; + * + * // Allocate WarpScan shared memory for 4 warps + * __shared__ typename WarpScan::TempStorage temp_storage[4]; + * + * // Obtain one input item per thread + * int thread_data = ... + * + * // Compute in|exclusive warp-wide prefix sums + * int inclusive_partial, exclusive_partial; + * int warp_id = threadIdx.x / 32; + * WarpScan(temp_storage[warp_id]).Sum(thread_data, inclusive_partial, exclusive_partial); + * + * \endcode + * \par + * Suppose the set of input \p thread_data across the block of threads is {1, 1, 1, 1, ...}. + * The corresponding output \p inclusive_partial in each of the four warps of threads will be + * 1, 2, 3, ..., 32}. + * The corresponding output \p exclusive_partial in each of the four warps of threads will be + * 0, 1, 2, ..., 31}. + * + */ __device__ __forceinline__ void Sum( T input, ///< [in] Calling thread's input item. T &inclusive_output, ///< [out] Calling thread's inclusive-scan output item. @@ -1297,7 +1338,47 @@ public: Sum(input, inclusive_output, exclusive_output, Int2Type()); } - /// Combination scan with identity + + /** + * \brief Computes both inclusive and exclusive prefix scans using the specified binary scan functor across the calling warp. + * + * Supports non-commutative scan operators. + * + * \smemreuse + * + * \par Snippet + * The code snippet below illustrates four concurrent warp-wide prefix max scans within a block of + * 128 threads (one per each of the 32-thread warps). + * \par + * \code + * #include + * + * __global__ void ExampleKernel(...) + * { + * // Specialize WarpScan for type int + * typedef cub::WarpScan WarpScan; + * + * // Allocate WarpScan shared memory for 4 warps + * __shared__ typename WarpScan::TempStorage temp_storage[4]; + * + * // Obtain one input item per thread + * int thread_data = ... + * + * // Compute inclusive warp-wide prefix max scans + * int warp_id = threadIdx.x / 32; + * int inclusive_partial, exclusive_partial; + * WarpScan(temp_storage[warp_id]).Scan(thread_data, inclusive_partial, exclusive_partial, INT_MIN, cub::Max()); + * + * \endcode + * \par + * Suppose the set of input \p thread_data across the block of threads is {0, -1, 2, -3, ..., 126, -127}. + * The corresponding output \p inclusive_partial in the first warp would be + * 0, 0, 2, 2, ..., 30, 30, the output for the second warp would be 32, 32, 34, 34, ..., 62, 62, etc. + * The corresponding output \p exclusive_partial in the first warp would be + * INT_MIN, 0, 0, 2, ..., 28, 30, the output for the second warp would be 30, 32, 32, 34, ..., 60, 62, etc. + * + * \tparam ScanOp [inferred] Binary scan operator type having member T operator()(const T &a, const T &b) + */ template __device__ __forceinline__ void Scan( T input, ///< [in] Calling thread's input item. @@ -1309,7 +1390,47 @@ public: InternalWarpScan(temp_storage).Scan(input, inclusive_output, exclusive_output, identity, scan_op); } - /// Combination scan with without identity + + /** + * \brief Computes both inclusive and exclusive prefix scans using the specified binary scan functor across the calling warp. Because no identity value is supplied, the \p exclusive_output computed for warp-lane0 is undefined. + * + * Supports non-commutative scan operators. + * + * \smemreuse + * + * \par Snippet + * The code snippet below illustrates four concurrent warp-wide exclusive prefix max scans within a block of + * 128 threads (one per each of the 32-thread warps). + * \par + * \code + * #include + * + * __global__ void ExampleKernel(...) + * { + * // Specialize WarpScan for type int + * typedef cub::WarpScan WarpScan; + * + * // Allocate WarpScan shared memory for 4 warps + * __shared__ typename WarpScan::TempStorage temp_storage[4]; + * + * // Obtain one input item per thread + * int thread_data = ... + * + * // Compute exclusive warp-wide prefix max scans + * int inclusive_partial, exclusive_partial; + * WarpScan(temp_storage[warp_id]).Scan(thread_data, inclusive_partial, exclusive_partial, cub::Max()); + * + * \endcode + * \par + * Suppose the set of input \p thread_data across the block of threads is {0, -1, 2, -3, ..., 126, -127}. + * The corresponding output \p inclusive_partial in the first warp would be + * 0, 0, 2, 2, ..., 30, 30, the output for the second warp would be 32, 32, 34, 34, ..., 62, 62, etc. + * The corresponding output \p exclusive_partial in the first warp would be + * ?, 0, 0, 2, ..., 28, 30, the output for the second warp would be ?, 32, 32, 34, ..., 60, 62, etc. + * (The output \p thread_data in warp lane0 is undefined.) + * + * \tparam ScanOp [inferred] Binary scan operator type having member T operator()(const T &a, const T &b) + */ template __device__ __forceinline__ void Scan( T input, ///< [in] Calling thread's input item. diff --git a/docs/download_cub.html b/docs/download_cub.html index bce34c6d19..68ae20f4aa 100644 --- a/docs/download_cub.html +++ b/docs/download_cub.html @@ -37,14 +37,14 @@
If your download doesn't start in 3s:

- -Download CUB! + +Download CUB!
diff --git a/docs/html/CHANGE_LOG.TXT b/docs/html/CHANGE_LOG.TXT index 8c0b9a89ad..d0e0398204 100644 --- a/docs/html/CHANGE_LOG.TXT +++ b/docs/html/CHANGE_LOG.TXT @@ -1,25 +1,39 @@ //----------------------------------------------------------------------------- -1.3.0 03/03/2014 +1.3.1 05/22/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: + - 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. - Bug fixes: - Fixed bug in cub::WarpScan (which affected cub::BlockScan and cub::DeviceScan) where incorrect results (e.g., NAN) would often be @@ -34,7 +48,7 @@ //----------------------------------------------------------------------------- -1.2.3 03/03/2014 +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 ArgIndexInputIterator documentation @@ -43,18 +57,25 @@ 1.2.2 03/03/2014 - New features: - - Added device-wide reduce-by-key (DeviceReduce::ReduceByKey, DeviceReduce::RunLengthEncode) - Added MS VC++ project solutions for device-wide and block-wide examples - Performance: - - Improved DeviceScan, DeviceSelect, DevicePartition 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-portabiltiy plots for many device-wide primitives to doc + - 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: - - Fixed bug where inclusion of Thrust headers in a certain order prevented CUB device-wide primitives from working properly + - 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 diff --git a/docs/html/annotated.html b/docs/html/annotated.html index bcbb2df3e2..93df51ac7c 100644 --- a/docs/html/annotated.html +++ b/docs/html/annotated.html @@ -228,7 +228,7 @@
-Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
diff --git a/docs/html/arg__index__input__iterator_8cuh.html b/docs/html/arg__index__input__iterator_8cuh.html index f42768c60f..a933b092f2 100644 --- a/docs/html/arg__index__input__iterator_8cuh.html +++ b/docs/html/arg__index__input__iterator_8cuh.html @@ -134,7 +134,7 @@
-Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
diff --git a/docs/html/arg__index__input__iterator_8cuh_source.html b/docs/html/arg__index__input__iterator_8cuh_source.html index 674927b7e1..f00a2c65f2 100644 --- a/docs/html/arg__index__input__iterator_8cuh_source.html +++ b/docs/html/arg__index__input__iterator_8cuh_source.html @@ -288,7 +288,7 @@
-Generated on Mon May 12 2014 17:09:51 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/block__discontinuity_8cuh.html b/docs/html/block__discontinuity_8cuh.html index b4a535455f..c9fb340986 100644 --- a/docs/html/block__discontinuity_8cuh.html +++ b/docs/html/block__discontinuity_8cuh.html @@ -138,7 +138,7 @@
-Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
diff --git a/docs/html/block__discontinuity_8cuh_source.html b/docs/html/block__discontinuity_8cuh_source.html index 8a3c12a61d..328f36edec 100644 --- a/docs/html/block__discontinuity_8cuh_source.html +++ b/docs/html/block__discontinuity_8cuh_source.html @@ -396,7 +396,7 @@
-Generated on Mon May 12 2014 17:09:51 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/block__exchange_8cuh.html b/docs/html/block__exchange_8cuh.html index 1de817f019..f37700506e 100644 --- a/docs/html/block__exchange_8cuh.html +++ b/docs/html/block__exchange_8cuh.html @@ -140,7 +140,7 @@
-Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
diff --git a/docs/html/block__exchange_8cuh_source.html.REMOVED.git-id b/docs/html/block__exchange_8cuh_source.html.REMOVED.git-id index 84a1445a09..8d436f3048 100644 --- a/docs/html/block__exchange_8cuh_source.html.REMOVED.git-id +++ b/docs/html/block__exchange_8cuh_source.html.REMOVED.git-id @@ -1 +1 @@ -58a08649e45c847ba4b2d22ebe22e8d46e63a836 \ No newline at end of file +4a654c1a08f97d2377e10c5fa8ac6f5b7a7c555f \ No newline at end of file diff --git a/docs/html/block__histogram_8cuh.html b/docs/html/block__histogram_8cuh.html index 00fa478728..3cf782f5fc 100644 --- a/docs/html/block__histogram_8cuh.html +++ b/docs/html/block__histogram_8cuh.html @@ -149,7 +149,7 @@
-Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
diff --git a/docs/html/block__histogram_8cuh_source.html b/docs/html/block__histogram_8cuh_source.html index 483d7424fa..4b43184a10 100644 --- a/docs/html/block__histogram_8cuh_source.html +++ b/docs/html/block__histogram_8cuh_source.html @@ -295,7 +295,7 @@
-Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/block__load_8cuh.html b/docs/html/block__load_8cuh.html index a9e11abc76..0ad934ec2a 100644 --- a/docs/html/block__load_8cuh.html +++ b/docs/html/block__load_8cuh.html @@ -204,7 +204,7 @@
-Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
diff --git a/docs/html/block__load_8cuh_source.html.REMOVED.git-id b/docs/html/block__load_8cuh_source.html.REMOVED.git-id index 4aa4cc533e..f577873aa8 100644 --- a/docs/html/block__load_8cuh_source.html.REMOVED.git-id +++ b/docs/html/block__load_8cuh_source.html.REMOVED.git-id @@ -1 +1 @@ -d179a051d3db874e3c1f771bea14100b788e43d5 \ No newline at end of file +fccac91152c0a9bd47d74f65f996881b13901088 \ No newline at end of file diff --git a/docs/html/block__radix__sort_8cuh.html b/docs/html/block__radix__sort_8cuh.html index d8046dec84..9aab574d63 100644 --- a/docs/html/block__radix__sort_8cuh.html +++ b/docs/html/block__radix__sort_8cuh.html @@ -141,7 +141,7 @@
-Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
diff --git a/docs/html/block__radix__sort_8cuh_source.html b/docs/html/block__radix__sort_8cuh_source.html index d60cf6a179..bfe502519b 100644 --- a/docs/html/block__radix__sort_8cuh_source.html +++ b/docs/html/block__radix__sort_8cuh_source.html @@ -237,289 +237,295 @@
220  UnsignedBits (&unsigned_keys)[ITEMS_PER_THREAD],
221  int (&ranks)[ITEMS_PER_THREAD],
222  int begin_bit,
-
223  Int2Type<false> is_descending)
-
224  {
-
225  AscendingBlockRadixRank(temp_storage.asending_ranking_storage).RankKeys(
-
226  unsigned_keys,
-
227  ranks,
-
228  begin_bit);
-
229  }
-
230 
-
232  __device__ __forceinline__ void RankKeys(
-
233  UnsignedBits (&unsigned_keys)[ITEMS_PER_THREAD],
-
234  int (&ranks)[ITEMS_PER_THREAD],
-
235  int begin_bit,
-
236  Int2Type<true> is_descending)
-
237  {
-
238  DescendingBlockRadixRank(temp_storage.descending_ranking_storage).RankKeys(
-
239  unsigned_keys,
-
240  ranks,
-
241  begin_bit);
-
242  }
-
243 
-
245  __device__ __forceinline__ void ExchangeValues(
-
246  Value (&values)[ITEMS_PER_THREAD],
-
247  int (&ranks)[ITEMS_PER_THREAD],
-
248  Int2Type<false> is_keys_only,
-
249  Int2Type<true> is_blocked)
-
250  {
-
251  __syncthreads();
-
252 
-
253  // Exchange values through shared memory in blocked arrangement
-
254  BlockExchangeValues(temp_storage.exchange_values).ScatterToBlocked(values, ranks);
-
255  }
-
256 
-
258  __device__ __forceinline__ void ExchangeValues(
-
259  Value (&values)[ITEMS_PER_THREAD],
-
260  int (&ranks)[ITEMS_PER_THREAD],
-
261  Int2Type<false> is_keys_only,
-
262  Int2Type<false> is_blocked)
-
263  {
-
264  __syncthreads();
-
265 
-
266  // Exchange values through shared memory in blocked arrangement
-
267  BlockExchangeValues(temp_storage.exchange_values).ScatterToStriped(values, ranks);
-
268  }
-
269 
-
271  template <int IS_BLOCKED>
-
272  __device__ __forceinline__ void ExchangeValues(
-
273  Value (&values)[ITEMS_PER_THREAD],
-
274  int (&ranks)[ITEMS_PER_THREAD],
-
275  Int2Type<true> is_keys_only,
-
276  Int2Type<IS_BLOCKED> is_blocked)
-
277  {}
-
278 
-
280  template <int DESCENDING, int KEYS_ONLY>
-
281  __device__ __forceinline__ void SortBlocked(
-
282  Key (&keys)[ITEMS_PER_THREAD],
-
283  Value (&values)[ITEMS_PER_THREAD],
-
284  int begin_bit,
-
285  int end_bit,
-
286  Int2Type<DESCENDING> is_descending,
-
287  Int2Type<KEYS_ONLY> is_keys_only)
-
288  {
-
289  UnsignedBits (&unsigned_keys)[ITEMS_PER_THREAD] =
-
290  reinterpret_cast<UnsignedBits (&)[ITEMS_PER_THREAD]>(keys);
-
291 
-
292  // Twiddle bits if necessary
-
293  #pragma unroll
-
294  for (int KEY = 0; KEY < ITEMS_PER_THREAD; KEY++)
-
295  {
-
296  unsigned_keys[KEY] = KeyTraits::TwiddleIn(unsigned_keys[KEY]);
-
297  }
-
298 
-
299  // Radix sorting passes
-
300  while (true)
-
301  {
-
302  // Rank the blocked keys
-
303  int ranks[ITEMS_PER_THREAD];
-
304  RankKeys(unsigned_keys, ranks, begin_bit, is_descending);
-
305  begin_bit += RADIX_BITS;
-
306 
-
307  __syncthreads();
-
308 
-
309  // Exchange keys through shared memory in blocked arrangement
-
310  BlockExchangeKeys(temp_storage.exchange_keys).ScatterToBlocked(keys, ranks);
-
311 
-
312  // Exchange values through shared memory in blocked arrangement
-
313  ExchangeValues(values, ranks, is_keys_only, Int2Type<true>());
-
314 
-
315  // Quit if done
-
316  if (begin_bit >= end_bit) break;
-
317 
-
318  __syncthreads();
-
319  }
-
320 
-
321  // Untwiddle bits if necessary
-
322  #pragma unroll
-
323  for (int KEY = 0; KEY < ITEMS_PER_THREAD; KEY++)
-
324  {
-
325  unsigned_keys[KEY] = KeyTraits::TwiddleOut(unsigned_keys[KEY]);
-
326  }
-
327  }
-
328 
-
330  template <int DESCENDING, int KEYS_ONLY>
-
331  __device__ __forceinline__ void SortBlockedToStriped(
-
332  Key (&keys)[ITEMS_PER_THREAD],
-
333  Value (&values)[ITEMS_PER_THREAD],
-
334  int begin_bit,
-
335  int end_bit,
-
336  Int2Type<DESCENDING> is_descending,
-
337  Int2Type<KEYS_ONLY> is_keys_only)
-
338  {
-
339  UnsignedBits (&unsigned_keys)[ITEMS_PER_THREAD] =
-
340  reinterpret_cast<UnsignedBits (&)[ITEMS_PER_THREAD]>(keys);
-
341 
-
342  // Twiddle bits if necessary
-
343  #pragma unroll
-
344  for (int KEY = 0; KEY < ITEMS_PER_THREAD; KEY++)
-
345  {
-
346  unsigned_keys[KEY] = KeyTraits::TwiddleIn(unsigned_keys[KEY]);
-
347  }
-
348 
-
349  // Radix sorting passes
-
350  while (true)
-
351  {
-
352  // Rank the blocked keys
-
353  int ranks[ITEMS_PER_THREAD];
-
354  RankKeys(unsigned_keys, ranks, begin_bit, is_descending);
-
355  begin_bit += RADIX_BITS;
-
356 
-
357  __syncthreads();
-
358 
-
359  // Check if this is the last pass
-
360  if (begin_bit >= end_bit)
-
361  {
-
362  // Last pass exchanges keys through shared memory in striped arrangement
-
363  BlockExchangeKeys(temp_storage.exchange_keys).ScatterToStriped(keys, ranks);
+
223  int pass_bits,
+
224  Int2Type<false> is_descending)
+
225  {
+
226  AscendingBlockRadixRank(temp_storage.asending_ranking_storage).RankKeys(
+
227  unsigned_keys,
+
228  ranks,
+
229  begin_bit,
+
230  pass_bits);
+
231  }
+
232 
+
234  __device__ __forceinline__ void RankKeys(
+
235  UnsignedBits (&unsigned_keys)[ITEMS_PER_THREAD],
+
236  int (&ranks)[ITEMS_PER_THREAD],
+
237  int begin_bit,
+
238  Int2Type<true> is_descending)
+
239  {
+
240  DescendingBlockRadixRank(temp_storage.descending_ranking_storage).RankKeys(
+
241  unsigned_keys,
+
242  ranks,
+
243  begin_bit);
+
244  }
+
245 
+
247  __device__ __forceinline__ void ExchangeValues(
+
248  Value (&values)[ITEMS_PER_THREAD],
+
249  int (&ranks)[ITEMS_PER_THREAD],
+
250  Int2Type<false> is_keys_only,
+
251  Int2Type<true> is_blocked)
+
252  {
+
253  __syncthreads();
+
254 
+
255  // Exchange values through shared memory in blocked arrangement
+
256  BlockExchangeValues(temp_storage.exchange_values).ScatterToBlocked(values, ranks);
+
257  }
+
258 
+
260  __device__ __forceinline__ void ExchangeValues(
+
261  Value (&values)[ITEMS_PER_THREAD],
+
262  int (&ranks)[ITEMS_PER_THREAD],
+
263  Int2Type<false> is_keys_only,
+
264  Int2Type<false> is_blocked)
+
265  {
+
266  __syncthreads();
+
267 
+
268  // Exchange values through shared memory in blocked arrangement
+
269  BlockExchangeValues(temp_storage.exchange_values).ScatterToStriped(values, ranks);
+
270  }
+
271 
+
273  template <int IS_BLOCKED>
+
274  __device__ __forceinline__ void ExchangeValues(
+
275  Value (&values)[ITEMS_PER_THREAD],
+
276  int (&ranks)[ITEMS_PER_THREAD],
+
277  Int2Type<true> is_keys_only,
+
278  Int2Type<IS_BLOCKED> is_blocked)
+
279  {}
+
280 
+
282  template <int DESCENDING, int KEYS_ONLY>
+
283  __device__ __forceinline__ void SortBlocked(
+
284  Key (&keys)[ITEMS_PER_THREAD],
+
285  Value (&values)[ITEMS_PER_THREAD],
+
286  int begin_bit,
+
287  int end_bit,
+
288  Int2Type<DESCENDING> is_descending,
+
289  Int2Type<KEYS_ONLY> is_keys_only)
+
290  {
+
291  UnsignedBits (&unsigned_keys)[ITEMS_PER_THREAD] =
+
292  reinterpret_cast<UnsignedBits (&)[ITEMS_PER_THREAD]>(keys);
+
293 
+
294  // Twiddle bits if necessary
+
295  #pragma unroll
+
296  for (int KEY = 0; KEY < ITEMS_PER_THREAD; KEY++)
+
297  {
+
298  unsigned_keys[KEY] = KeyTraits::TwiddleIn(unsigned_keys[KEY]);
+
299  }
+
300 
+
301  // Radix sorting passes
+
302  while (true)
+
303  {
+
304  int pass_bits = CUB_MIN(RADIX_BITS, end_bit - begin_bit);
+
305 
+
306  // Rank the blocked keys
+
307  int ranks[ITEMS_PER_THREAD];
+
308  RankKeys(unsigned_keys, ranks, begin_bit, pass_bits, is_descending);
+
309  begin_bit += RADIX_BITS;
+
310 
+
311  __syncthreads();
+
312 
+
313  // Exchange keys through shared memory in blocked arrangement
+
314  BlockExchangeKeys(temp_storage.exchange_keys).ScatterToBlocked(keys, ranks);
+
315 
+
316  // Exchange values through shared memory in blocked arrangement
+
317  ExchangeValues(values, ranks, is_keys_only, Int2Type<true>());
+
318 
+
319  // Quit if done
+
320  if (begin_bit >= end_bit) break;
+
321 
+
322  __syncthreads();
+
323  }
+
324 
+
325  // Untwiddle bits if necessary
+
326  #pragma unroll
+
327  for (int KEY = 0; KEY < ITEMS_PER_THREAD; KEY++)
+
328  {
+
329  unsigned_keys[KEY] = KeyTraits::TwiddleOut(unsigned_keys[KEY]);
+
330  }
+
331  }
+
332 
+
334  template <int DESCENDING, int KEYS_ONLY>
+
335  __device__ __forceinline__ void SortBlockedToStriped(
+
336  Key (&keys)[ITEMS_PER_THREAD],
+
337  Value (&values)[ITEMS_PER_THREAD],
+
338  int begin_bit,
+
339  int end_bit,
+
340  Int2Type<DESCENDING> is_descending,
+
341  Int2Type<KEYS_ONLY> is_keys_only)
+
342  {
+
343  UnsignedBits (&unsigned_keys)[ITEMS_PER_THREAD] =
+
344  reinterpret_cast<UnsignedBits (&)[ITEMS_PER_THREAD]>(keys);
+
345 
+
346  // Twiddle bits if necessary
+
347  #pragma unroll
+
348  for (int KEY = 0; KEY < ITEMS_PER_THREAD; KEY++)
+
349  {
+
350  unsigned_keys[KEY] = KeyTraits::TwiddleIn(unsigned_keys[KEY]);
+
351  }
+
352 
+
353  // Radix sorting passes
+
354  while (true)
+
355  {
+
356  int pass_bits = CUB_MIN(RADIX_BITS, end_bit - begin_bit);
+
357 
+
358  // Rank the blocked keys
+
359  int ranks[ITEMS_PER_THREAD];
+
360  RankKeys(unsigned_keys, ranks, begin_bit, pass_bits, is_descending);
+
361  begin_bit += RADIX_BITS;
+
362 
+
363  __syncthreads();
364 
-
365  // Last pass exchanges through shared memory in striped arrangement
-
366  ExchangeValues(values, ranks, is_keys_only, Int2Type<false>());
-
367 
-
368  // Quit
-
369  break;
-
370  }
-
371 
-
372  // Exchange keys through shared memory in blocked arrangement
-
373  BlockExchangeKeys(temp_storage.exchange_keys).ScatterToBlocked(keys, ranks);
-
374 
-
375  // Exchange values through shared memory in blocked arrangement
-
376  ExchangeValues(values, ranks, is_keys_only, Int2Type<true>());
+
365  // Check if this is the last pass
+
366  if (begin_bit >= end_bit)
+
367  {
+
368  // Last pass exchanges keys through shared memory in striped arrangement
+
369  BlockExchangeKeys(temp_storage.exchange_keys).ScatterToStriped(keys, ranks);
+
370 
+
371  // Last pass exchanges through shared memory in striped arrangement
+
372  ExchangeValues(values, ranks, is_keys_only, Int2Type<false>());
+
373 
+
374  // Quit
+
375  break;
+
376  }
377 
-
378  __syncthreads();
-
379  }
+
378  // Exchange keys through shared memory in blocked arrangement
+
379  BlockExchangeKeys(temp_storage.exchange_keys).ScatterToBlocked(keys, ranks);
380 
-
381  // Untwiddle bits if necessary
-
382  #pragma unroll
-
383  for (int KEY = 0; KEY < ITEMS_PER_THREAD; KEY++)
-
384  {
-
385  unsigned_keys[KEY] = KeyTraits::TwiddleOut(unsigned_keys[KEY]);
-
386  }
-
387  }
-
388 
-
389 
-
390 
-
391 public:
-
392 
-
394  struct TempStorage : Uninitialized<_TempStorage> {};
+
381  // Exchange values through shared memory in blocked arrangement
+
382  ExchangeValues(values, ranks, is_keys_only, Int2Type<true>());
+
383 
+
384  __syncthreads();
+
385  }
+
386 
+
387  // Untwiddle bits if necessary
+
388  #pragma unroll
+
389  for (int KEY = 0; KEY < ITEMS_PER_THREAD; KEY++)
+
390  {
+
391  unsigned_keys[KEY] = KeyTraits::TwiddleOut(unsigned_keys[KEY]);
+
392  }
+
393  }
+
394 
395 
396 
-
397  /******************************************************************/
+
397 public:
+
398 
+
400  struct TempStorage : Uninitialized<_TempStorage> {};
401 
-
405  __device__ __forceinline__ BlockRadixSort()
-
406  :
-
407  temp_storage(PrivateStorage()),
-
408  linear_tid(RowMajorTid(BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z))
-
409  {}
-
410 
-
411 
-
415  __device__ __forceinline__ BlockRadixSort(
-
416  TempStorage &temp_storage)
-
417  :
-
418  temp_storage(temp_storage.Alias()),
-
419  linear_tid(RowMajorTid(BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z))
-
420  {}
-
421 
-
422 
-
424  /******************************************************************/
+
402 
+
403  /******************************************************************/
+
407 
+
411  __device__ __forceinline__ BlockRadixSort()
+
412  :
+
413  temp_storage(PrivateStorage()),
+
414  linear_tid(RowMajorTid(BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z))
+
415  {}
+
416 
+
417 
+
421  __device__ __forceinline__ BlockRadixSort(
+
422  TempStorage &temp_storage)
+
423  :
+
424  temp_storage(temp_storage.Alias()),
+
425  linear_tid(RowMajorTid(BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z))
+
426  {}
+
427 
428 
-
466  __device__ __forceinline__ void Sort(
-
467  Key (&keys)[ITEMS_PER_THREAD],
-
468  int begin_bit = 0,
-
469  int end_bit = sizeof(Key) * 8)
-
470  {
-
471  NullType values[ITEMS_PER_THREAD];
-
472 
-
473  SortBlocked(keys, values, begin_bit, end_bit, Int2Type<false>(), Int2Type<KEYS_ONLY>());
-
474  }
-
475 
-
476 
-
521  __device__ __forceinline__ void Sort(
-
522  Key (&keys)[ITEMS_PER_THREAD],
-
523  Value (&values)[ITEMS_PER_THREAD],
-
524  int begin_bit = 0,
-
525  int end_bit = sizeof(Key) * 8)
-
526  {
-
527  SortBlocked(keys, values, begin_bit, end_bit, Int2Type<false>(), Int2Type<KEYS_ONLY>());
-
528  }
-
529 
-
567  __device__ __forceinline__ void SortDescending(
-
568  Key (&keys)[ITEMS_PER_THREAD],
-
569  int begin_bit = 0,
-
570  int end_bit = sizeof(Key) * 8)
-
571  {
-
572  NullType values[ITEMS_PER_THREAD];
-
573 
-
574  SortBlocked(keys, values, begin_bit, end_bit, Int2Type<true>(), Int2Type<KEYS_ONLY>());
-
575  }
-
576 
-
577 
-
622  __device__ __forceinline__ void SortDescending(
-
623  Key (&keys)[ITEMS_PER_THREAD],
-
624  Value (&values)[ITEMS_PER_THREAD],
-
625  int begin_bit = 0,
-
626  int end_bit = sizeof(Key) * 8)
-
627  {
-
628  SortBlocked(keys, values, begin_bit, end_bit, Int2Type<true>(), Int2Type<KEYS_ONLY>());
-
629  }
-
630 
-
631 
-
633  /******************************************************************/
+
430  /******************************************************************/
+
434 
+
472  __device__ __forceinline__ void Sort(
+
473  Key (&keys)[ITEMS_PER_THREAD],
+
474  int begin_bit = 0,
+
475  int end_bit = sizeof(Key) * 8)
+
476  {
+
477  NullType values[ITEMS_PER_THREAD];
+
478 
+
479  SortBlocked(keys, values, begin_bit, end_bit, Int2Type<false>(), Int2Type<KEYS_ONLY>());
+
480  }
+
481 
+
482 
+
527  __device__ __forceinline__ void Sort(
+
528  Key (&keys)[ITEMS_PER_THREAD],
+
529  Value (&values)[ITEMS_PER_THREAD],
+
530  int begin_bit = 0,
+
531  int end_bit = sizeof(Key) * 8)
+
532  {
+
533  SortBlocked(keys, values, begin_bit, end_bit, Int2Type<false>(), Int2Type<KEYS_ONLY>());
+
534  }
+
535 
+
573  __device__ __forceinline__ void SortDescending(
+
574  Key (&keys)[ITEMS_PER_THREAD],
+
575  int begin_bit = 0,
+
576  int end_bit = sizeof(Key) * 8)
+
577  {
+
578  NullType values[ITEMS_PER_THREAD];
+
579 
+
580  SortBlocked(keys, values, begin_bit, end_bit, Int2Type<true>(), Int2Type<KEYS_ONLY>());
+
581  }
+
582 
+
583 
+
628  __device__ __forceinline__ void SortDescending(
+
629  Key (&keys)[ITEMS_PER_THREAD],
+
630  Value (&values)[ITEMS_PER_THREAD],
+
631  int begin_bit = 0,
+
632  int end_bit = sizeof(Key) * 8)
+
633  {
+
634  SortBlocked(keys, values, begin_bit, end_bit, Int2Type<true>(), Int2Type<KEYS_ONLY>());
+
635  }
+
636 
637 
-
638 
-
677  __device__ __forceinline__ void SortBlockedToStriped(
-
678  Key (&keys)[ITEMS_PER_THREAD],
-
679  int begin_bit = 0,
-
680  int end_bit = sizeof(Key) * 8)
-
681  {
-
682  NullType values[ITEMS_PER_THREAD];
-
683 
-
684  SortBlockedToStriped(keys, values, begin_bit, end_bit, Int2Type<false>(), Int2Type<KEYS_ONLY>());
-
685  }
-
686 
-
687 
-
732  __device__ __forceinline__ void SortBlockedToStriped(
-
733  Key (&keys)[ITEMS_PER_THREAD],
-
734  Value (&values)[ITEMS_PER_THREAD],
-
735  int begin_bit = 0,
-
736  int end_bit = sizeof(Key) * 8)
-
737  {
-
738  SortBlockedToStriped(keys, values, begin_bit, end_bit, Int2Type<false>(), Int2Type<KEYS_ONLY>());
-
739  }
-
740 
-
741 
-
780  __device__ __forceinline__ void SortDescendingBlockedToStriped(
-
781  Key (&keys)[ITEMS_PER_THREAD],
-
782  int begin_bit = 0,
-
783  int end_bit = sizeof(Key) * 8)
-
784  {
-
785  NullType values[ITEMS_PER_THREAD];
-
786 
-
787  SortBlockedToStriped(keys, values, begin_bit, end_bit, Int2Type<true>(), Int2Type<KEYS_ONLY>());
-
788  }
-
789 
-
790 
-
835  __device__ __forceinline__ void SortDescendingBlockedToStriped(
-
836  Key (&keys)[ITEMS_PER_THREAD],
-
837  Value (&values)[ITEMS_PER_THREAD],
-
838  int begin_bit = 0,
-
839  int end_bit = sizeof(Key) * 8)
-
840  {
-
841  SortBlockedToStriped(keys, values, begin_bit, end_bit, Int2Type<true>(), Int2Type<KEYS_ONLY>());
-
842  }
-
843 
-
844 
-
846 
-
847 };
-
848 
-
853 } // CUB namespace
-
854 CUB_NS_POSTFIX // Optional outer namespace(s)
-
855 
+
639  /******************************************************************/
+
643 
+
644 
+
683  __device__ __forceinline__ void SortBlockedToStriped(
+
684  Key (&keys)[ITEMS_PER_THREAD],
+
685  int begin_bit = 0,
+
686  int end_bit = sizeof(Key) * 8)
+
687  {
+
688  NullType values[ITEMS_PER_THREAD];
+
689 
+
690  SortBlockedToStriped(keys, values, begin_bit, end_bit, Int2Type<false>(), Int2Type<KEYS_ONLY>());
+
691  }
+
692 
+
693 
+
738  __device__ __forceinline__ void SortBlockedToStriped(
+
739  Key (&keys)[ITEMS_PER_THREAD],
+
740  Value (&values)[ITEMS_PER_THREAD],
+
741  int begin_bit = 0,
+
742  int end_bit = sizeof(Key) * 8)
+
743  {
+
744  SortBlockedToStriped(keys, values, begin_bit, end_bit, Int2Type<false>(), Int2Type<KEYS_ONLY>());
+
745  }
+
746 
+
747 
+
786  __device__ __forceinline__ void SortDescendingBlockedToStriped(
+
787  Key (&keys)[ITEMS_PER_THREAD],
+
788  int begin_bit = 0,
+
789  int end_bit = sizeof(Key) * 8)
+
790  {
+
791  NullType values[ITEMS_PER_THREAD];
+
792 
+
793  SortBlockedToStriped(keys, values, begin_bit, end_bit, Int2Type<true>(), Int2Type<KEYS_ONLY>());
+
794  }
+
795 
+
796 
+
841  __device__ __forceinline__ void SortDescendingBlockedToStriped(
+
842  Key (&keys)[ITEMS_PER_THREAD],
+
843  Value (&values)[ITEMS_PER_THREAD],
+
844  int begin_bit = 0,
+
845  int end_bit = sizeof(Key) * 8)
+
846  {
+
847  SortBlockedToStriped(keys, values, begin_bit, end_bit, Int2Type<true>(), Int2Type<KEYS_ONLY>());
+
848  }
+
849 
+
850 
+
852 
+
853 };
+
854 
+
859 } // CUB namespace
+
860 CUB_NS_POSTFIX // Optional outer namespace(s)
+
861 
-Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/block__reduce_8cuh.html b/docs/html/block__reduce_8cuh.html index 58ca2b1e50..54cf646833 100644 --- a/docs/html/block__reduce_8cuh.html +++ b/docs/html/block__reduce_8cuh.html @@ -151,7 +151,7 @@
-Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
diff --git a/docs/html/block__reduce_8cuh_source.html b/docs/html/block__reduce_8cuh_source.html index e0d9f678a6..c6297bddff 100644 --- a/docs/html/block__reduce_8cuh_source.html +++ b/docs/html/block__reduce_8cuh_source.html @@ -328,7 +328,7 @@
-Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/block__scan_8cuh.html b/docs/html/block__scan_8cuh.html index 9a1abcd5ea..31c00fe4e1 100644 --- a/docs/html/block__scan_8cuh.html +++ b/docs/html/block__scan_8cuh.html @@ -151,7 +151,7 @@
-Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
diff --git a/docs/html/block__scan_8cuh_source.html.REMOVED.git-id b/docs/html/block__scan_8cuh_source.html.REMOVED.git-id index 905f3cc657..1e6ba7872a 100644 --- a/docs/html/block__scan_8cuh_source.html.REMOVED.git-id +++ b/docs/html/block__scan_8cuh_source.html.REMOVED.git-id @@ -1 +1 @@ -841d80cb75904c0b2bf77ef7e64f72f0f68b7141 \ No newline at end of file +402ec7b45869fdc77cee768440435e3bcc85847b \ No newline at end of file diff --git a/docs/html/block__store_8cuh.html b/docs/html/block__store_8cuh.html index 520ccb0aba..31dd545df2 100644 --- a/docs/html/block__store_8cuh.html +++ b/docs/html/block__store_8cuh.html @@ -192,7 +192,7 @@
-Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
diff --git a/docs/html/block__store_8cuh_source.html b/docs/html/block__store_8cuh_source.html index 451d2e2807..3482b2cb8f 100644 --- a/docs/html/block__store_8cuh_source.html +++ b/docs/html/block__store_8cuh_source.html @@ -617,7 +617,7 @@
-Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/cache__modified__input__iterator_8cuh.html b/docs/html/cache__modified__input__iterator_8cuh.html index eebaea2ccb..2e1525d258 100644 --- a/docs/html/cache__modified__input__iterator_8cuh.html +++ b/docs/html/cache__modified__input__iterator_8cuh.html @@ -133,7 +133,7 @@
-Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
diff --git a/docs/html/cache__modified__input__iterator_8cuh_source.html b/docs/html/cache__modified__input__iterator_8cuh_source.html index 484a209d7a..00f552a9ba 100644 --- a/docs/html/cache__modified__input__iterator_8cuh_source.html +++ b/docs/html/cache__modified__input__iterator_8cuh_source.html @@ -279,7 +279,7 @@
-Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/cache__modified__output__iterator_8cuh.html b/docs/html/cache__modified__output__iterator_8cuh.html index 1684846daa..4701ec9255 100644 --- a/docs/html/cache__modified__output__iterator_8cuh.html +++ b/docs/html/cache__modified__output__iterator_8cuh.html @@ -133,7 +133,7 @@
-Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
diff --git a/docs/html/cache__modified__output__iterator_8cuh_source.html b/docs/html/cache__modified__output__iterator_8cuh_source.html index ec14de3a13..311be0e013 100644 --- a/docs/html/cache__modified__output__iterator_8cuh_source.html +++ b/docs/html/cache__modified__output__iterator_8cuh_source.html @@ -288,7 +288,7 @@
-Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_arg_index_input_iterator-members.html b/docs/html/classcub_1_1_arg_index_input_iterator-members.html index ab0e4bee39..60469a8b52 100644 --- a/docs/html/classcub_1_1_arg_index_input_iterator-members.html +++ b/docs/html/classcub_1_1_arg_index_input_iterator-members.html @@ -128,7 +128,7 @@
-Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_arg_index_input_iterator.html b/docs/html/classcub_1_1_arg_index_input_iterator.html index 318464e7e5..e0774ba975 100644 --- a/docs/html/classcub_1_1_arg_index_input_iterator.html +++ b/docs/html/classcub_1_1_arg_index_input_iterator.html @@ -315,7 +315,7 @@

Constructor & Destructor Documentation

-Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_block_discontinuity-members.html b/docs/html/classcub_1_1_block_discontinuity-members.html index da2167cc8b..95b581fd03 100644 --- a/docs/html/classcub_1_1_block_discontinuity-members.html +++ b/docs/html/classcub_1_1_block_discontinuity-members.html @@ -114,7 +114,7 @@
-Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_block_discontinuity.html b/docs/html/classcub_1_1_block_discontinuity.html index 417c61c31d..f8ff42d146 100644 --- a/docs/html/classcub_1_1_block_discontinuity.html +++ b/docs/html/classcub_1_1_block_discontinuity.html @@ -637,7 +637,7 @@

Member Function Documentation

-Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_block_exchange-members.html b/docs/html/classcub_1_1_block_exchange-members.html index 71b76889c3..bae3a384a2 100644 --- a/docs/html/classcub_1_1_block_exchange-members.html +++ b/docs/html/classcub_1_1_block_exchange-members.html @@ -118,7 +118,7 @@
-Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_block_exchange.html b/docs/html/classcub_1_1_block_exchange.html index 27e2a4ec04..fe6d1b912d 100644 --- a/docs/html/classcub_1_1_block_exchange.html +++ b/docs/html/classcub_1_1_block_exchange.html @@ -749,7 +749,7 @@

Member Function Documentation

-Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_block_histogram-members.html b/docs/html/classcub_1_1_block_histogram-members.html index e4cf4d6775..aa498ac87b 100644 --- a/docs/html/classcub_1_1_block_histogram-members.html +++ b/docs/html/classcub_1_1_block_histogram-members.html @@ -113,7 +113,7 @@
-Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_block_histogram.html b/docs/html/classcub_1_1_block_histogram.html index eb99ca4aa6..ec96bd3a78 100644 --- a/docs/html/classcub_1_1_block_histogram.html +++ b/docs/html/classcub_1_1_block_histogram.html @@ -479,7 +479,7 @@

Member Function Documentation

-Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_block_load-members.html b/docs/html/classcub_1_1_block_load-members.html index 9e89752286..72fd06c841 100644 --- a/docs/html/classcub_1_1_block_load-members.html +++ b/docs/html/classcub_1_1_block_load-members.html @@ -113,7 +113,7 @@
-Generated on Mon May 12 2014 17:09:55 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_block_load.html b/docs/html/classcub_1_1_block_load.html index 9703e164b1..d391968ad9 100644 --- a/docs/html/classcub_1_1_block_load.html +++ b/docs/html/classcub_1_1_block_load.html @@ -464,7 +464,7 @@

Member Function Documentation

-Generated on Mon May 12 2014 17:09:55 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_block_radix_sort-members.html b/docs/html/classcub_1_1_block_radix_sort-members.html index 6b081f7229..9c1cfd83c7 100644 --- a/docs/html/classcub_1_1_block_radix_sort-members.html +++ b/docs/html/classcub_1_1_block_radix_sort-members.html @@ -118,7 +118,7 @@
-Generated on Mon May 12 2014 17:09:55 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_block_radix_sort.html b/docs/html/classcub_1_1_block_radix_sort.html index db7633a1ae..239d58e3dc 100644 --- a/docs/html/classcub_1_1_block_radix_sort.html +++ b/docs/html/classcub_1_1_block_radix_sort.html @@ -252,7 +252,7 @@

Constructor & Destructor Documentation

-

Definition at line 415 of file block_radix_sort.cuh.

+

Definition at line 421 of file block_radix_sort.cuh.

@@ -331,7 +331,7 @@

Member Function Documentation

-

Definition at line 466 of file block_radix_sort.cuh.

+

Definition at line 472 of file block_radix_sort.cuh.

@@ -418,7 +418,7 @@

Member Function Documentation

-

Definition at line 521 of file block_radix_sort.cuh.

+

Definition at line 527 of file block_radix_sort.cuh.

@@ -496,7 +496,7 @@

Member Function Documentation

-

Definition at line 567 of file block_radix_sort.cuh.

+

Definition at line 573 of file block_radix_sort.cuh.

@@ -583,7 +583,7 @@

Member Function Documentation

-

Definition at line 622 of file block_radix_sort.cuh.

+

Definition at line 628 of file block_radix_sort.cuh.

@@ -661,7 +661,7 @@

Member Function Documentation

-

Definition at line 677 of file block_radix_sort.cuh.

+

Definition at line 683 of file block_radix_sort.cuh.

@@ -748,7 +748,7 @@

Member Function Documentation

-

Definition at line 732 of file block_radix_sort.cuh.

+

Definition at line 738 of file block_radix_sort.cuh.

@@ -826,7 +826,7 @@

Member Function Documentation

-

Definition at line 780 of file block_radix_sort.cuh.

+

Definition at line 786 of file block_radix_sort.cuh.

@@ -913,7 +913,7 @@

Member Function Documentation

-

Definition at line 835 of file block_radix_sort.cuh.

+

Definition at line 841 of file block_radix_sort.cuh.

@@ -924,7 +924,7 @@

Member Function Documentation

-Generated on Mon May 12 2014 17:09:55 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_block_reduce-members.html b/docs/html/classcub_1_1_block_reduce-members.html index 42744c516c..135cadb5cd 100644 --- a/docs/html/classcub_1_1_block_reduce-members.html +++ b/docs/html/classcub_1_1_block_reduce-members.html @@ -116,7 +116,7 @@
-Generated on Mon May 12 2014 17:09:55 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_block_reduce.html b/docs/html/classcub_1_1_block_reduce.html index 18d2eba84d..b2a1ab67a1 100644 --- a/docs/html/classcub_1_1_block_reduce.html +++ b/docs/html/classcub_1_1_block_reduce.html @@ -708,7 +708,7 @@

Member Function Documentation

-Generated on Mon May 12 2014 17:09:55 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_block_scan-members.html b/docs/html/classcub_1_1_block_scan-members.html index 64db347bfb..60b76f86ad 100644 --- a/docs/html/classcub_1_1_block_scan-members.html +++ b/docs/html/classcub_1_1_block_scan-members.html @@ -134,7 +134,7 @@
-Generated on Mon May 12 2014 17:09:55 for CUB by   +Generated on Thu May 22 2014 16:10:32 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_block_scan.html.REMOVED.git-id b/docs/html/classcub_1_1_block_scan.html.REMOVED.git-id index 3bca5c5ffd..88101d7ad8 100644 --- a/docs/html/classcub_1_1_block_scan.html.REMOVED.git-id +++ b/docs/html/classcub_1_1_block_scan.html.REMOVED.git-id @@ -1 +1 @@ -10fdfe8a82d8f43e9798bef7ae7b7bd50a268d34 \ No newline at end of file +8af2a8768b7eef0b0eab9d379aecff4374aa070f \ No newline at end of file diff --git a/docs/html/classcub_1_1_block_store-members.html b/docs/html/classcub_1_1_block_store-members.html index b3a877b086..d1030f4725 100644 --- a/docs/html/classcub_1_1_block_store-members.html +++ b/docs/html/classcub_1_1_block_store-members.html @@ -112,7 +112,7 @@
-Generated on Mon May 12 2014 17:09:56 for CUB by   +Generated on Thu May 22 2014 16:10:32 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_block_store.html b/docs/html/classcub_1_1_block_store.html index b4083bb930..bdf0c85684 100644 --- a/docs/html/classcub_1_1_block_store.html +++ b/docs/html/classcub_1_1_block_store.html @@ -391,7 +391,7 @@

Member Function Documentation

-Generated on Mon May 12 2014 17:09:56 for CUB by   +Generated on Thu May 22 2014 16:10:32 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_cache_modified_input_iterator-members.html b/docs/html/classcub_1_1_cache_modified_input_iterator-members.html index d88cbd5d35..c2e7d44319 100644 --- a/docs/html/classcub_1_1_cache_modified_input_iterator-members.html +++ b/docs/html/classcub_1_1_cache_modified_input_iterator-members.html @@ -128,7 +128,7 @@
-Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_cache_modified_input_iterator.html b/docs/html/classcub_1_1_cache_modified_input_iterator.html index cff87a2c09..e15f412519 100644 --- a/docs/html/classcub_1_1_cache_modified_input_iterator.html +++ b/docs/html/classcub_1_1_cache_modified_input_iterator.html @@ -297,7 +297,7 @@

Constructor & Destructor Documentation

-Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_cache_modified_output_iterator-members.html b/docs/html/classcub_1_1_cache_modified_output_iterator-members.html index 43088bd27e..a3eb09841f 100644 --- a/docs/html/classcub_1_1_cache_modified_output_iterator-members.html +++ b/docs/html/classcub_1_1_cache_modified_output_iterator-members.html @@ -127,7 +127,7 @@
-Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_cache_modified_output_iterator.html b/docs/html/classcub_1_1_cache_modified_output_iterator.html index 2be7bc0d7c..c3840740a4 100644 --- a/docs/html/classcub_1_1_cache_modified_output_iterator.html +++ b/docs/html/classcub_1_1_cache_modified_output_iterator.html @@ -297,7 +297,7 @@

Constructor & Destructor Documentation

-Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_constant_input_iterator-members.html b/docs/html/classcub_1_1_constant_input_iterator-members.html index f9be292f2c..b348c924a9 100644 --- a/docs/html/classcub_1_1_constant_input_iterator-members.html +++ b/docs/html/classcub_1_1_constant_input_iterator-members.html @@ -128,7 +128,7 @@
-Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_constant_input_iterator.html b/docs/html/classcub_1_1_constant_input_iterator.html index 2e13ebb4bc..fb50595c87 100644 --- a/docs/html/classcub_1_1_constant_input_iterator.html +++ b/docs/html/classcub_1_1_constant_input_iterator.html @@ -302,7 +302,7 @@

Constructor & Destructor Documentation

-Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_counting_input_iterator-members.html b/docs/html/classcub_1_1_counting_input_iterator-members.html index 1a24f5c3cc..f9b6da0bc8 100644 --- a/docs/html/classcub_1_1_counting_input_iterator-members.html +++ b/docs/html/classcub_1_1_counting_input_iterator-members.html @@ -128,7 +128,7 @@
-Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_counting_input_iterator.html b/docs/html/classcub_1_1_counting_input_iterator.html index e3bef85273..dbcd2265c6 100644 --- a/docs/html/classcub_1_1_counting_input_iterator.html +++ b/docs/html/classcub_1_1_counting_input_iterator.html @@ -290,7 +290,7 @@

Constructor & Destructor Documentation

-Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_tex_obj_input_iterator-members.html b/docs/html/classcub_1_1_tex_obj_input_iterator-members.html index e651d1cdef..555b356551 100644 --- a/docs/html/classcub_1_1_tex_obj_input_iterator-members.html +++ b/docs/html/classcub_1_1_tex_obj_input_iterator-members.html @@ -130,7 +130,7 @@
-Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_tex_obj_input_iterator.html b/docs/html/classcub_1_1_tex_obj_input_iterator.html index 9ec75cacc9..f677c5d716 100644 --- a/docs/html/classcub_1_1_tex_obj_input_iterator.html +++ b/docs/html/classcub_1_1_tex_obj_input_iterator.html @@ -328,7 +328,7 @@

Member Function Documentation

-Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_tex_ref_input_iterator-members.html b/docs/html/classcub_1_1_tex_ref_input_iterator-members.html index d7e1c7121a..1bd05eb3e3 100644 --- a/docs/html/classcub_1_1_tex_ref_input_iterator-members.html +++ b/docs/html/classcub_1_1_tex_ref_input_iterator-members.html @@ -130,7 +130,7 @@
-Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_tex_ref_input_iterator.html b/docs/html/classcub_1_1_tex_ref_input_iterator.html index 6e858724be..47aa452440 100644 --- a/docs/html/classcub_1_1_tex_ref_input_iterator.html +++ b/docs/html/classcub_1_1_tex_ref_input_iterator.html @@ -333,7 +333,7 @@

Member Function Documentation

-Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_transform_input_iterator-members.html b/docs/html/classcub_1_1_transform_input_iterator-members.html index eae3f64ed2..7f332285c6 100644 --- a/docs/html/classcub_1_1_transform_input_iterator-members.html +++ b/docs/html/classcub_1_1_transform_input_iterator-members.html @@ -128,7 +128,7 @@
-Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_transform_input_iterator.html b/docs/html/classcub_1_1_transform_input_iterator.html index c5d3465887..9004c3f1b0 100644 --- a/docs/html/classcub_1_1_transform_input_iterator.html +++ b/docs/html/classcub_1_1_transform_input_iterator.html @@ -320,7 +320,7 @@

Constructor & Destructor Documentation

-Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_warp_reduce-members.html b/docs/html/classcub_1_1_warp_reduce-members.html index e99a3fbe6c..7a346bb99b 100644 --- a/docs/html/classcub_1_1_warp_reduce-members.html +++ b/docs/html/classcub_1_1_warp_reduce-members.html @@ -117,7 +117,7 @@
-Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_warp_reduce.html b/docs/html/classcub_1_1_warp_reduce.html index 3970966c4f..fc7a4473ea 100644 --- a/docs/html/classcub_1_1_warp_reduce.html +++ b/docs/html/classcub_1_1_warp_reduce.html @@ -200,35 +200,35 @@  
Summation reductions
__device__ __forceinline__ T Sum (T input) - Computes a warp-wide sum in each active warp. The output is valid in warp lane0. More...
+ Computes a warp-wide sum in the calling warp. The output is valid in warp lane0. More...
  __device__ __forceinline__ T Sum (T input, int valid_items) - Computes a partially-full warp-wide sum in each active warp. The output is valid in warp lane0. More...
+ Computes a partially-full warp-wide sum in the calling warp. The output is valid in warp lane0. More...
  template<typename Flag > __device__ __forceinline__ T HeadSegmentedSum (T input, Flag head_flag) - Computes a segmented sum in each active warp where segments are defined by head-flags. The sum of each segment is returned to the first lane in that segment (which always includes lane0). More...
+ Computes a segmented sum in the calling warp where segments are defined by head-flags. The sum of each segment is returned to the first lane in that segment (which always includes lane0). More...
  template<typename Flag > __device__ __forceinline__ T TailSegmentedSum (T input, Flag tail_flag) - Computes a segmented sum in each active warp where segments are defined by tail-flags. The sum of each segment is returned to the first lane in that segment (which always includes lane0). More...
+ Computes a segmented sum in the calling warp where segments are defined by tail-flags. The sum of each segment is returned to the first lane in that segment (which always includes lane0). More...
 
Generic reductions
template<typename ReductionOp > __device__ __forceinline__ T Reduce (T input, ReductionOp reduction_op) - Computes a warp-wide reduction in each active warp using the specified binary reduction functor. The output is valid in warp lane0. More...
+ Computes a warp-wide reduction in the calling warp using the specified binary reduction functor. The output is valid in warp lane0. More...
  template<typename ReductionOp > __device__ __forceinline__ T Reduce (T input, ReductionOp reduction_op, int valid_items) - Computes a partially-full warp-wide reduction in each active warp using the specified binary reduction functor. The output is valid in warp lane0. More...
+ Computes a partially-full warp-wide reduction in the calling warp using the specified binary reduction functor. The output is valid in warp lane0. More...
  template<typename ReductionOp , typename Flag > __device__ __forceinline__ T HeadSegmentedReduce (T input, Flag head_flag, ReductionOp reduction_op) - Computes a segmented reduction in each active warp where segments are defined by head-flags. The reduction of each segment is returned to the first lane in that segment (which always includes lane0). More...
+ Computes a segmented reduction in the calling warp where segments are defined by head-flags. The reduction of each segment is returned to the first lane in that segment (which always includes lane0). More...
  template<typename ReductionOp , typename Flag > __device__ __forceinline__ T TailSegmentedReduce (T input, Flag tail_flag, ReductionOp reduction_op) - Computes a segmented reduction in each active warp where segments are defined by tail-flags. The reduction of each segment is returned to the first lane in that segment (which always includes lane0). More...
+ Computes a segmented reduction in the calling warp where segments are defined by tail-flags. The reduction of each segment is returned to the first lane in that segment (which always includes lane0). More...
 

Constructor & Destructor Documentation

@@ -293,7 +293,7 @@

Member Function Documentation

-

Computes a warp-wide sum in each active warp. The output is valid in warp lane0.

+

Computes a warp-wide sum in the calling warp. The output is valid in warp lane0.

A subsequent __syncthreads() threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., temp_storage) is to be reused or repurposed.

Snippet
The code snippet below illustrates four concurrent warp sum reductions within a block of 128 threads (one per each of the 32-thread warps).
#include <cub/cub.cuh>
@@ -359,8 +359,8 @@

Member Function Documentation

-

Computes a partially-full warp-wide sum in each active warp. The output is valid in warp lane0.

-

All threads in each logical warp must agree on the same value for valid_items. Otherwise the result is undefined.

+

Computes a partially-full warp-wide sum in the calling warp. The output is valid in warp lane0.

+

All threads across the calling warp must agree on the same value for valid_items. Otherwise the result is undefined.

A subsequent __syncthreads() threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., temp_storage) is to be reused or repurposed.

Snippet
The code snippet below illustrates a sum reduction within a single, partially-full block of 32 threads (one warp).
#include <cub/cub.cuh>
@@ -431,7 +431,7 @@

Member Function Documentation

-

Computes a segmented sum in each active warp where segments are defined by head-flags. The sum of each segment is returned to the first lane in that segment (which always includes lane0).

+

Computes a segmented sum in the calling warp where segments are defined by head-flags. The sum of each segment is returned to the first lane in that segment (which always includes lane0).

A subsequent __syncthreads() threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., temp_storage) is to be reused or repurposed.

Snippet
The code snippet below illustrates a head-segmented warp sum reduction within a block of 32 threads (one warp).
#include <cub/cub.cuh>
@@ -507,7 +507,7 @@

Member Function Documentation

-

Computes a segmented sum in each active warp where segments are defined by tail-flags. The sum of each segment is returned to the first lane in that segment (which always includes lane0).

+

Computes a segmented sum in the calling warp where segments are defined by tail-flags. The sum of each segment is returned to the first lane in that segment (which always includes lane0).

A subsequent __syncthreads() threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., temp_storage) is to be reused or repurposed.

Snippet
The code snippet below illustrates a tail-segmented warp sum reduction within a block of 32 threads (one warp).
#include <cub/cub.cuh>
@@ -583,7 +583,7 @@

Member Function Documentation

-

Computes a warp-wide reduction in each active warp using the specified binary reduction functor. The output is valid in warp lane0.

+

Computes a warp-wide reduction in the calling warp using the specified binary reduction functor. The output is valid in warp lane0.

Supports non-commutative reduction operators

A subsequent __syncthreads() threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., temp_storage) is to be reused or repurposed.

Snippet
The code snippet below illustrates four concurrent warp max reductions within a block of 128 threads (one per each of the 32-thread warps).
@@ -666,8 +666,8 @@

Member Function Documentation

-

Computes a partially-full warp-wide reduction in each active warp using the specified binary reduction functor. The output is valid in warp lane0.

-

All threads in each logical warp must agree on the same value for valid_items. Otherwise the result is undefined.

+

Computes a partially-full warp-wide reduction in the calling warp using the specified binary reduction functor. The output is valid in warp lane0.

+

All threads across the calling warp must agree on the same value for valid_items. Otherwise the result is undefined.

Supports non-commutative reduction operators

A subsequent __syncthreads() threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., temp_storage) is to be reused or repurposed.

Snippet
The code snippet below illustrates a max reduction within a single, partially-full block of 32 threads (one warp).
@@ -752,7 +752,7 @@

Member Function Documentation

-

Computes a segmented reduction in each active warp where segments are defined by head-flags. The reduction of each segment is returned to the first lane in that segment (which always includes lane0).

+

Computes a segmented reduction in the calling warp where segments are defined by head-flags. The reduction of each segment is returned to the first lane in that segment (which always includes lane0).

Supports non-commutative reduction operators

A subsequent __syncthreads() threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., temp_storage) is to be reused or repurposed.

Snippet
The code snippet below illustrates a head-segmented warp max reduction within a block of 32 threads (one warp).
@@ -836,7 +836,7 @@

Member Function Documentation

-

Computes a segmented reduction in each active warp where segments are defined by tail-flags. The reduction of each segment is returned to the first lane in that segment (which always includes lane0).

+

Computes a segmented reduction in the calling warp where segments are defined by tail-flags. The reduction of each segment is returned to the first lane in that segment (which always includes lane0).

Supports non-commutative reduction operators

A subsequent __syncthreads() threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., temp_storage) is to be reused or repurposed.

Snippet
The code snippet below illustrates a tail-segmented warp max reduction within a block of 32 threads (one warp).
@@ -885,7 +885,7 @@

Member Function Documentation

-Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:32 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_warp_scan-members.html b/docs/html/classcub_1_1_warp_scan-members.html index 3bd39d86bb..3835296648 100644 --- a/docs/html/classcub_1_1_warp_scan-members.html +++ b/docs/html/classcub_1_1_warp_scan-members.html @@ -119,12 +119,15 @@ InclusiveSum(T input, T &output)cub::WarpScan< T, LOGICAL_WARP_THREADS, PTX_ARCH >inline InclusiveSum(T input, T &output, T &warp_aggregate)cub::WarpScan< T, LOGICAL_WARP_THREADS, PTX_ARCH >inline InclusiveSum(T input, T &output, T &warp_aggregate, WarpPrefixCallbackOp &warp_prefix_op)cub::WarpScan< T, LOGICAL_WARP_THREADS, PTX_ARCH >inline - WarpScan(TempStorage &temp_storage)cub::WarpScan< T, LOGICAL_WARP_THREADS, PTX_ARCH >inline + Scan(T input, T &inclusive_output, T &exclusive_output, T identity, ScanOp scan_op)cub::WarpScan< T, LOGICAL_WARP_THREADS, PTX_ARCH >inline + Scan(T input, T &inclusive_output, T &exclusive_output, ScanOp scan_op)cub::WarpScan< T, LOGICAL_WARP_THREADS, PTX_ARCH >inline + Sum(T input, T &inclusive_output, T &exclusive_output)cub::WarpScan< T, LOGICAL_WARP_THREADS, PTX_ARCH >inline + WarpScan(TempStorage &temp_storage)cub::WarpScan< T, LOGICAL_WARP_THREADS, PTX_ARCH >inline
-Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:32 for CUB by   doxygen 1.8.4
diff --git a/docs/html/classcub_1_1_warp_scan.html.REMOVED.git-id b/docs/html/classcub_1_1_warp_scan.html.REMOVED.git-id index 4c4d2987d8..6419915653 100644 --- a/docs/html/classcub_1_1_warp_scan.html.REMOVED.git-id +++ b/docs/html/classcub_1_1_warp_scan.html.REMOVED.git-id @@ -1 +1 @@ -9e41b78e4ff4d0a1687556fc9149d95462015984 \ No newline at end of file +ddf0ca2f1ff17f995e27a3f33a2347c8c74037fc \ No newline at end of file diff --git a/docs/html/classes.html b/docs/html/classes.html index 95ac1a90cf..6d1e3c4028 100644 --- a/docs/html/classes.html +++ b/docs/html/classes.html @@ -141,7 +141,7 @@
-Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
diff --git a/docs/html/constant__input__iterator_8cuh.html b/docs/html/constant__input__iterator_8cuh.html index 353fe357d2..46f505c2d9 100644 --- a/docs/html/constant__input__iterator_8cuh.html +++ b/docs/html/constant__input__iterator_8cuh.html @@ -132,7 +132,7 @@
-Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
diff --git a/docs/html/constant__input__iterator_8cuh_source.html b/docs/html/constant__input__iterator_8cuh_source.html index b2133c6364..9b1cf509d3 100644 --- a/docs/html/constant__input__iterator_8cuh_source.html +++ b/docs/html/constant__input__iterator_8cuh_source.html @@ -282,7 +282,7 @@
-Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/counting__input__iterator_8cuh.html b/docs/html/counting__input__iterator_8cuh.html index 2fef4e7dc7..6d950308a2 100644 --- a/docs/html/counting__input__iterator_8cuh.html +++ b/docs/html/counting__input__iterator_8cuh.html @@ -133,7 +133,7 @@
-Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
diff --git a/docs/html/counting__input__iterator_8cuh_source.html b/docs/html/counting__input__iterator_8cuh_source.html index 7cc841e276..7e59aa7c84 100644 --- a/docs/html/counting__input__iterator_8cuh_source.html +++ b/docs/html/counting__input__iterator_8cuh_source.html @@ -277,7 +277,7 @@
-Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/cub_8cuh.html b/docs/html/cub_8cuh.html index dcefe31e10..3ea9cd9902 100644 --- a/docs/html/cub_8cuh.html +++ b/docs/html/cub_8cuh.html @@ -137,7 +137,7 @@
-Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
diff --git a/docs/html/cub_8cuh_source.html b/docs/html/cub_8cuh_source.html index 61dfae5f7e..1dacd93a99 100644 --- a/docs/html/cub_8cuh_source.html +++ b/docs/html/cub_8cuh_source.html @@ -195,7 +195,7 @@
-Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/device__histogram_8cuh.html b/docs/html/device__histogram_8cuh.html index 97fe8ade68..980cd13fcd 100644 --- a/docs/html/device__histogram_8cuh.html +++ b/docs/html/device__histogram_8cuh.html @@ -136,7 +136,7 @@
-Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
diff --git a/docs/html/device__histogram_8cuh_source.html b/docs/html/device__histogram_8cuh_source.html index 0ce3592056..3551f9e62e 100644 --- a/docs/html/device__histogram_8cuh_source.html +++ b/docs/html/device__histogram_8cuh_source.html @@ -399,7 +399,7 @@
-Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/device__partition_8cuh.html b/docs/html/device__partition_8cuh.html index e3216f4587..2d9f3c68dd 100644 --- a/docs/html/device__partition_8cuh.html +++ b/docs/html/device__partition_8cuh.html @@ -136,7 +136,7 @@
-Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
diff --git a/docs/html/device__partition_8cuh_source.html b/docs/html/device__partition_8cuh_source.html index 5809b3f624..e026f34d66 100644 --- a/docs/html/device__partition_8cuh_source.html +++ b/docs/html/device__partition_8cuh_source.html @@ -226,7 +226,7 @@
-Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/device__radix__sort_8cuh.html b/docs/html/device__radix__sort_8cuh.html index 9c41da78b7..68b503c5ff 100644 --- a/docs/html/device__radix__sort_8cuh.html +++ b/docs/html/device__radix__sort_8cuh.html @@ -136,7 +136,7 @@
-Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
diff --git a/docs/html/device__radix__sort_8cuh_source.html b/docs/html/device__radix__sort_8cuh_source.html index 1a74478763..91f113d268 100644 --- a/docs/html/device__radix__sort_8cuh_source.html +++ b/docs/html/device__radix__sort_8cuh_source.html @@ -278,7 +278,7 @@
-Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/device__reduce_8cuh.html b/docs/html/device__reduce_8cuh.html index 9de0471011..a7dd4284bc 100644 --- a/docs/html/device__reduce_8cuh.html +++ b/docs/html/device__reduce_8cuh.html @@ -136,7 +136,7 @@
-Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
diff --git a/docs/html/device__reduce_8cuh_source.html b/docs/html/device__reduce_8cuh_source.html index c2f22608af..ee054cb8ed 100644 --- a/docs/html/device__reduce_8cuh_source.html +++ b/docs/html/device__reduce_8cuh_source.html @@ -440,7 +440,7 @@
-Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/device__scan_8cuh.html b/docs/html/device__scan_8cuh.html index 3d7edca17e..68565ee06e 100644 --- a/docs/html/device__scan_8cuh.html +++ b/docs/html/device__scan_8cuh.html @@ -136,7 +136,7 @@
-Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
diff --git a/docs/html/device__scan_8cuh_source.html b/docs/html/device__scan_8cuh_source.html index 466df73006..5e4a7f843a 100644 --- a/docs/html/device__scan_8cuh_source.html +++ b/docs/html/device__scan_8cuh_source.html @@ -285,7 +285,7 @@
-Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/device__select_8cuh.html b/docs/html/device__select_8cuh.html index e67542f10c..aebcfb2100 100644 --- a/docs/html/device__select_8cuh.html +++ b/docs/html/device__select_8cuh.html @@ -136,7 +136,7 @@
-Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
diff --git a/docs/html/device__select_8cuh_source.html b/docs/html/device__select_8cuh_source.html index b3841f68e0..1ef8a22db0 100644 --- a/docs/html/device__select_8cuh_source.html +++ b/docs/html/device__select_8cuh_source.html @@ -261,7 +261,7 @@
-Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/dir_011e1c944d88f71be72e1e24a5fda7cf.html b/docs/html/dir_011e1c944d88f71be72e1e24a5fda7cf.html index a08065f1af..b763afd320 100644 --- a/docs/html/dir_011e1c944d88f71be72e1e24a5fda7cf.html +++ b/docs/html/dir_011e1c944d88f71be72e1e24a5fda7cf.html @@ -119,7 +119,7 @@
-Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
diff --git a/docs/html/dir_18fc672d63781b5a743137aee24ff656.html b/docs/html/dir_18fc672d63781b5a743137aee24ff656.html index 02360faa7f..762c3ffee9 100644 --- a/docs/html/dir_18fc672d63781b5a743137aee24ff656.html +++ b/docs/html/dir_18fc672d63781b5a743137aee24ff656.html @@ -115,7 +115,7 @@
-Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
diff --git a/docs/html/dir_80932b4cec52750ff92b1a1912314cf5.html b/docs/html/dir_80932b4cec52750ff92b1a1912314cf5.html index dcdba06be0..a6858691ba 100644 --- a/docs/html/dir_80932b4cec52750ff92b1a1912314cf5.html +++ b/docs/html/dir_80932b4cec52750ff92b1a1912314cf5.html @@ -119,7 +119,7 @@
-Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
diff --git a/docs/html/dir_bb50a5ef59f19d030d06415663184d05.html b/docs/html/dir_bb50a5ef59f19d030d06415663184d05.html index 729eb6b7f2..c332be2ae0 100644 --- a/docs/html/dir_bb50a5ef59f19d030d06415663184d05.html +++ b/docs/html/dir_bb50a5ef59f19d030d06415663184d05.html @@ -109,7 +109,7 @@
-Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
diff --git a/docs/html/dir_cb3a671affffe7eeb3fdf5ae58e42cc8.html b/docs/html/dir_cb3a671affffe7eeb3fdf5ae58e42cc8.html index 44252b7618..ed1025221d 100644 --- a/docs/html/dir_cb3a671affffe7eeb3fdf5ae58e42cc8.html +++ b/docs/html/dir_cb3a671affffe7eeb3fdf5ae58e42cc8.html @@ -107,7 +107,7 @@
-Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
diff --git a/docs/html/dir_d583f216f1aafe19404e836b0c097ad2.html b/docs/html/dir_d583f216f1aafe19404e836b0c097ad2.html index 4f024e59d1..bfeec08949 100644 --- a/docs/html/dir_d583f216f1aafe19404e836b0c097ad2.html +++ b/docs/html/dir_d583f216f1aafe19404e836b0c097ad2.html @@ -130,7 +130,7 @@
-Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
diff --git a/docs/html/download_cub.html b/docs/html/download_cub.html index bce34c6d19..68ae20f4aa 100644 --- a/docs/html/download_cub.html +++ b/docs/html/download_cub.html @@ -37,14 +37,14 @@
If your download doesn't start in 3s:

- -Download CUB! + +Download CUB!
diff --git a/docs/html/example_block_radix_sort_8cu-example.html b/docs/html/example_block_radix_sort_8cu-example.html index 75c686482c..ed5b7f1d56 100644 --- a/docs/html/example_block_radix_sort_8cu-example.html +++ b/docs/html/example_block_radix_sort_8cu-example.html @@ -398,7 +398,7 @@
-Generated on Mon May 12 2014 17:09:51 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/example_block_reduce_8cu-example.html b/docs/html/example_block_reduce_8cu-example.html index b5bdf1d763..d343d28949 100644 --- a/docs/html/example_block_reduce_8cu-example.html +++ b/docs/html/example_block_reduce_8cu-example.html @@ -365,7 +365,7 @@
-Generated on Mon May 12 2014 17:09:51 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/example_block_scan_8cu-example.html b/docs/html/example_block_scan_8cu-example.html index 53d93bf07f..cc7130969f 100644 --- a/docs/html/example_block_scan_8cu-example.html +++ b/docs/html/example_block_scan_8cu-example.html @@ -414,7 +414,7 @@
-Generated on Mon May 12 2014 17:09:51 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/example_device_histogram_8cu-example.html b/docs/html/example_device_histogram_8cu-example.html index d9c169d139..9570858070 100644 --- a/docs/html/example_device_histogram_8cu-example.html +++ b/docs/html/example_device_histogram_8cu-example.html @@ -312,7 +312,7 @@
-Generated on Mon May 12 2014 17:09:51 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/example_device_partition_flagged_8cu-example.html b/docs/html/example_device_partition_flagged_8cu-example.html index 671171ebc6..a50c44a10f 100644 --- a/docs/html/example_device_partition_flagged_8cu-example.html +++ b/docs/html/example_device_partition_flagged_8cu-example.html @@ -318,7 +318,7 @@
-Generated on Mon May 12 2014 17:09:51 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/example_device_partition_if_8cu-example.html b/docs/html/example_device_partition_if_8cu-example.html index 0a83f05141..b857e78b13 100644 --- a/docs/html/example_device_partition_if_8cu-example.html +++ b/docs/html/example_device_partition_if_8cu-example.html @@ -329,7 +329,7 @@
-Generated on Mon May 12 2014 17:09:51 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/example_device_radix_sort_8cu-example.html b/docs/html/example_device_radix_sort_8cu-example.html index 25c7b71863..8fe62c81c1 100644 --- a/docs/html/example_device_radix_sort_8cu-example.html +++ b/docs/html/example_device_radix_sort_8cu-example.html @@ -311,7 +311,7 @@
-Generated on Mon May 12 2014 17:09:51 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/example_device_reduce_8cu-example.html b/docs/html/example_device_reduce_8cu-example.html index bec6d5cb2d..3c3b072287 100644 --- a/docs/html/example_device_reduce_8cu-example.html +++ b/docs/html/example_device_reduce_8cu-example.html @@ -266,7 +266,7 @@
-Generated on Mon May 12 2014 17:09:51 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/example_device_scan_8cu-example.html b/docs/html/example_device_scan_8cu-example.html index c431023eff..25d0afe628 100644 --- a/docs/html/example_device_scan_8cu-example.html +++ b/docs/html/example_device_scan_8cu-example.html @@ -272,7 +272,7 @@
-Generated on Mon May 12 2014 17:09:51 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/example_device_select_flagged_8cu-example.html b/docs/html/example_device_select_flagged_8cu-example.html index 350b1fdeb4..6bed07f0e5 100644 --- a/docs/html/example_device_select_flagged_8cu-example.html +++ b/docs/html/example_device_select_flagged_8cu-example.html @@ -318,7 +318,7 @@
-Generated on Mon May 12 2014 17:09:51 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/example_device_select_if_8cu-example.html b/docs/html/example_device_select_if_8cu-example.html index 2090a00b2f..0140fba2df 100644 --- a/docs/html/example_device_select_if_8cu-example.html +++ b/docs/html/example_device_select_if_8cu-example.html @@ -327,7 +327,7 @@
-Generated on Mon May 12 2014 17:09:51 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/example_device_select_unique_8cu-example.html b/docs/html/example_device_select_unique_8cu-example.html index 8da117c1b5..24a8166e1e 100644 --- a/docs/html/example_device_select_unique_8cu-example.html +++ b/docs/html/example_device_select_unique_8cu-example.html @@ -307,7 +307,7 @@
-Generated on Mon May 12 2014 17:09:51 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
diff --git a/docs/html/examples.html b/docs/html/examples.html index c94ca6e741..03ad1dd0c9 100644 --- a/docs/html/examples.html +++ b/docs/html/examples.html @@ -121,7 +121,7 @@
-Generated on Mon May 12 2014 17:09:58 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
diff --git a/docs/html/files.html b/docs/html/files.html index 0d4f2a2d0d..2ea363efc4 100644 --- a/docs/html/files.html +++ b/docs/html/files.html @@ -139,7 +139,7 @@
-Generated on Mon May 12 2014 17:09:58 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
diff --git a/docs/html/functions.html b/docs/html/functions.html index 4dfc22aefa..6ead986c12 100644 --- a/docs/html/functions.html +++ b/docs/html/functions.html @@ -149,7 +149,7 @@

- a -

-Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
diff --git a/docs/html/functions_0x74.html b/docs/html/functions_0x74.html index 372e93975f..190c11c1b0 100644 --- a/docs/html/functions_0x74.html +++ b/docs/html/functions_0x74.html @@ -155,7 +155,7 @@

- t -

-Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
diff --git a/docs/html/functions_func_0x74.html b/docs/html/functions_func_0x74.html index e61cdae9f0..6622971228 100644 --- a/docs/html/functions_func_0x74.html +++ b/docs/html/functions_func_0x74.html @@ -146,7 +146,7 @@

- t -

    -Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
    diff --git a/docs/html/functions_func_0x75.html b/docs/html/functions_func_0x75.html index 3c008c3f71..61d112af06 100644 --- a/docs/html/functions_func_0x75.html +++ b/docs/html/functions_func_0x75.html @@ -138,7 +138,7 @@

    - u -

      -Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
      diff --git a/docs/html/functions_func_0x77.html b/docs/html/functions_func_0x77.html index 9a3f796472..e4048adf74 100644 --- a/docs/html/functions_func_0x77.html +++ b/docs/html/functions_func_0x77.html @@ -140,7 +140,7 @@

      - w -

        -Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
        diff --git a/docs/html/functions_func_0x7e.html b/docs/html/functions_func_0x7e.html index 402337dc78..e516879382 100644 --- a/docs/html/functions_func_0x7e.html +++ b/docs/html/functions_func_0x7e.html @@ -134,7 +134,7 @@

        - ~ -

          -Generated on Mon May 12 2014 17:09:58 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
          diff --git a/docs/html/functions_rela.html b/docs/html/functions_rela.html index 5acfae273c..512d381bc4 100644 --- a/docs/html/functions_rela.html +++ b/docs/html/functions_rela.html @@ -118,7 +118,7 @@
          -Generated on Mon May 12 2014 17:09:58 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
          diff --git a/docs/html/functions_type.html b/docs/html/functions_type.html index f96277bb3e..d58e959300 100644 --- a/docs/html/functions_type.html +++ b/docs/html/functions_type.html @@ -233,7 +233,7 @@

          - v -

            -Generated on Mon May 12 2014 17:09:58 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
            diff --git a/docs/html/functions_vars.html b/docs/html/functions_vars.html index ea0f332dd6..0900baf663 100644 --- a/docs/html/functions_vars.html +++ b/docs/html/functions_vars.html @@ -136,7 +136,7 @@
            -Generated on Mon May 12 2014 17:09:58 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
            diff --git a/docs/html/globals.html b/docs/html/globals.html index 222a27f43d..bafd9f0e9b 100644 --- a/docs/html/globals.html +++ b/docs/html/globals.html @@ -168,7 +168,7 @@
            -Generated on Mon May 12 2014 17:09:58 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
            diff --git a/docs/html/globals_defs.html b/docs/html/globals_defs.html index 48dd363737..37b21d8e37 100644 --- a/docs/html/globals_defs.html +++ b/docs/html/globals_defs.html @@ -168,7 +168,7 @@
            -Generated on Mon May 12 2014 17:09:58 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
            diff --git a/docs/html/group___block_module.html b/docs/html/group___block_module.html index 8a06823258..8ecdeb1df3 100644 --- a/docs/html/group___block_module.html +++ b/docs/html/group___block_module.html @@ -165,7 +165,7 @@
            -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
            diff --git a/docs/html/group___device_module.html b/docs/html/group___device_module.html index 81133acc55..5213b03415 100644 --- a/docs/html/group___device_module.html +++ b/docs/html/group___device_module.html @@ -149,7 +149,7 @@
            -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
            diff --git a/docs/html/group___primitive_module.html b/docs/html/group___primitive_module.html index b0798134a6..3d91ea6c38 100644 --- a/docs/html/group___primitive_module.html +++ b/docs/html/group___primitive_module.html @@ -107,7 +107,7 @@
            -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
            diff --git a/docs/html/group___util_io.html b/docs/html/group___util_io.html index 84fae1a54e..594619d9d3 100644 --- a/docs/html/group___util_io.html +++ b/docs/html/group___util_io.html @@ -1459,7 +1459,7 @@

            Function Documentation

            -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
            diff --git a/docs/html/group___util_iterator.html b/docs/html/group___util_iterator.html index 74cb41df67..3f3a6135a1 100644 --- a/docs/html/group___util_iterator.html +++ b/docs/html/group___util_iterator.html @@ -125,7 +125,7 @@
            -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
            diff --git a/docs/html/group___util_mgmt.html b/docs/html/group___util_mgmt.html index 0ad68b1fc1..faf00e7152 100644 --- a/docs/html/group___util_mgmt.html +++ b/docs/html/group___util_mgmt.html @@ -628,7 +628,7 @@

            Function Documentation

            -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
            diff --git a/docs/html/group___util_module.html b/docs/html/group___util_module.html index cc6d868aaa..7aca61afb8 100644 --- a/docs/html/group___util_module.html +++ b/docs/html/group___util_module.html @@ -193,7 +193,7 @@
            -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
            diff --git a/docs/html/group___util_ptx.html b/docs/html/group___util_ptx.html index 0977b8c94d..96fd5b499e 100644 --- a/docs/html/group___util_ptx.html +++ b/docs/html/group___util_ptx.html @@ -216,7 +216,7 @@

            Function Documentation

            -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
            diff --git a/docs/html/group___warp_module.html b/docs/html/group___warp_module.html index 847abb8be0..c1e2956c9f 100644 --- a/docs/html/group___warp_module.html +++ b/docs/html/group___warp_module.html @@ -337,7 +337,7 @@

            Function Documentation

            -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
            diff --git a/docs/html/hierarchy.html b/docs/html/hierarchy.html index 7d05145578..93c5b2015a 100644 --- a/docs/html/hierarchy.html +++ b/docs/html/hierarchy.html @@ -235,7 +235,7 @@
            -Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
            diff --git a/docs/html/index.html b/docs/html/index.html index cf1be933e9..22065f4c5a 100644 --- a/docs/html/index.html +++ b/docs/html/index.html @@ -115,7 +115,7 @@    -Download CUB v1.3.0 +Download CUB v1.3.1 @@ -393,23 +393,34 @@

            (8) Recent News

            + +
            05/22/2014
            + CUB v1.3.1
              +
            • New features:
                +
              • Added new "combination scan" methods to cub::WarpScan for efficiently computing both inclusive and exclusive prefix scans (and sums).
                • +
              +
              • +
            • Bug fixes: +
              • +
            • See the change-log for further details, including bug-fixes
              • +
            +

            +
            05/12/2014
            CUB v1.3.0
              -
            • 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.
                • -
              -
              • -
            • 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.
                • -
              -
            • See the change-log for further details

            @@ -497,7 +508,7 @@

            -Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
            diff --git a/docs/html/modules.html b/docs/html/modules.html index a2eb96b847..b96c42e77a 100644 --- a/docs/html/modules.html +++ b/docs/html/modules.html @@ -108,7 +108,7 @@
            -Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
            diff --git a/docs/html/namespacecub.html.REMOVED.git-id b/docs/html/namespacecub.html.REMOVED.git-id index 0f32aeab4a..88b9eaaf78 100644 --- a/docs/html/namespacecub.html.REMOVED.git-id +++ b/docs/html/namespacecub.html.REMOVED.git-id @@ -1 +1 @@ -aebfb4926d86d84f2378d7aec2a69840c65fdb7b \ No newline at end of file +21c2a6f52695280b17d031e348725c8f29808759 \ No newline at end of file diff --git a/docs/html/namespacemembers.html b/docs/html/namespacemembers.html index f683aa939e..b5be82240f 100644 --- a/docs/html/namespacemembers.html +++ b/docs/html/namespacemembers.html @@ -374,7 +374,7 @@

            - w -

              -Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
              diff --git a/docs/html/namespacemembers_enum.html b/docs/html/namespacemembers_enum.html index 129b705b95..665cbaf823 100644 --- a/docs/html/namespacemembers_enum.html +++ b/docs/html/namespacemembers_enum.html @@ -131,7 +131,7 @@
              -Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
              diff --git a/docs/html/namespacemembers_eval.html b/docs/html/namespacemembers_eval.html index 49ba91d5d4..45f97dcd2a 100644 --- a/docs/html/namespacemembers_eval.html +++ b/docs/html/namespacemembers_eval.html @@ -194,7 +194,7 @@
              -Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
              diff --git a/docs/html/namespacemembers_func.html b/docs/html/namespacemembers_func.html index 83da4ba82c..ed03d3f32f 100644 --- a/docs/html/namespacemembers_func.html +++ b/docs/html/namespacemembers_func.html @@ -258,7 +258,7 @@

              - w -

                -Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/namespaces.html b/docs/html/namespaces.html index 52b71d2a43..12cc6cd0bb 100644 --- a/docs/html/namespaces.html +++ b/docs/html/namespaces.html @@ -106,7 +106,7 @@
                -Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/search/all_73.js b/docs/html/search/all_73.js index 7050060ebb..5b9e4d403f 100644 --- a/docs/html/search/all_73.js +++ b/docs/html/search/all_73.js @@ -1,5 +1,6 @@ var searchData= [ + ['scan',['Scan',['../classcub_1_1_warp_scan.html#a32a247bea83254e7c3f3ae4ed8a1d5c7',1,'cub::WarpScan::Scan(T input, T &inclusive_output, T &exclusive_output, T identity, ScanOp scan_op)'],['../classcub_1_1_warp_scan.html#a3266a375b79e44f77087cec512e0c1a3',1,'cub::WarpScan::Scan(T input, T &inclusive_output, T &exclusive_output, ScanOp scan_op)']]], ['scattertoblocked',['ScatterToBlocked',['../classcub_1_1_block_exchange.html#a756e7903c5369261fcf6139e5db52824',1,'cub::BlockExchange']]], ['scattertostriped',['ScatterToStriped',['../classcub_1_1_block_exchange.html#a2ba02edda6bfed9327cd4ee1748aa678',1,'cub::BlockExchange::ScatterToStriped(T items[ITEMS_PER_THREAD], Offset ranks[ITEMS_PER_THREAD])'],['../classcub_1_1_block_exchange.html#a171c90da1f2d5572b22acf059e0b06ea',1,'cub::BlockExchange::ScatterToStriped(T items[ITEMS_PER_THREAD], Offset ranks[ITEMS_PER_THREAD], ValidFlag is_valid[ITEMS_PER_THREAD])']]], ['scattertostripedguarded',['ScatterToStripedGuarded',['../classcub_1_1_block_exchange.html#a0232bb6ad4bdfe1f1578e1ba8d175beb',1,'cub::BlockExchange']]], @@ -36,6 +37,6 @@ var searchData= ['storedirectstriped',['StoreDirectStriped',['../group___util_io.html#gafa774cd981172f96137620c868e628ba',1,'cub::StoreDirectStriped(int linear_tid, OutputIterator block_itr, T(&items)[ITEMS_PER_THREAD])'],['../group___util_io.html#gac9ab95ee7528c9acb3f92c2dd0fc7992',1,'cub::StoreDirectStriped(int linear_tid, OutputIterator block_itr, T(&items)[ITEMS_PER_THREAD], int valid_items)']]], ['storedirectwarpstriped',['StoreDirectWarpStriped',['../group___util_io.html#ga8ffe6de6fbf7c5a617b863c460b0de48',1,'cub::StoreDirectWarpStriped(int linear_tid, OutputIterator block_itr, T(&items)[ITEMS_PER_THREAD])'],['../group___util_io.html#gae201cc0d0cc452e5500a5f28c7e4a234',1,'cub::StoreDirectWarpStriped(int linear_tid, OutputIterator block_itr, T(&items)[ITEMS_PER_THREAD], int valid_items)']]], ['stripedtoblocked',['StripedToBlocked',['../classcub_1_1_block_exchange.html#a2855471bbbcc4d66ac6a29d35a040e0c',1,'cub::BlockExchange']]], - ['sum',['Sum',['../structcub_1_1_sum.html',1,'cub']]], - ['sum',['Sum',['../classcub_1_1_block_reduce.html#a7632bd9c8950dd6a3528ca99fa3f0890',1,'cub::BlockReduce::Sum(T input)'],['../classcub_1_1_block_reduce.html#ac5d4591d9513f08b180d4112cb0c4c51',1,'cub::BlockReduce::Sum(T(&inputs)[ITEMS_PER_THREAD])'],['../classcub_1_1_block_reduce.html#a33ddffdde07275ab0c4e1bf61b0d9409',1,'cub::BlockReduce::Sum(T input, int num_valid)'],['../structcub_1_1_device_reduce.html#aaf68d747ee676df2b5b5ff016ee6c371',1,'cub::DeviceReduce::Sum()'],['../classcub_1_1_warp_reduce.html#abe4aeeabf8859a7582a0b5858b84ee7a',1,'cub::WarpReduce::Sum(T input)'],['../classcub_1_1_warp_reduce.html#ad9c4a8d85a7795cf220713f362c36f30',1,'cub::WarpReduce::Sum(T input, int valid_items)']]] + ['sum',['Sum',['../classcub_1_1_block_reduce.html#a7632bd9c8950dd6a3528ca99fa3f0890',1,'cub::BlockReduce::Sum(T input)'],['../classcub_1_1_block_reduce.html#ac5d4591d9513f08b180d4112cb0c4c51',1,'cub::BlockReduce::Sum(T(&inputs)[ITEMS_PER_THREAD])'],['../classcub_1_1_block_reduce.html#a33ddffdde07275ab0c4e1bf61b0d9409',1,'cub::BlockReduce::Sum(T input, int num_valid)'],['../structcub_1_1_device_reduce.html#aaf68d747ee676df2b5b5ff016ee6c371',1,'cub::DeviceReduce::Sum()'],['../classcub_1_1_warp_scan.html#a25bd83f795e88b9260ec2bcbf846fb20',1,'cub::WarpScan::Sum()'],['../classcub_1_1_warp_reduce.html#abe4aeeabf8859a7582a0b5858b84ee7a',1,'cub::WarpReduce::Sum(T input)'],['../classcub_1_1_warp_reduce.html#ad9c4a8d85a7795cf220713f362c36f30',1,'cub::WarpReduce::Sum(T input, int valid_items)']]], + ['sum',['Sum',['../structcub_1_1_sum.html',1,'cub']]] ]; diff --git a/docs/html/search/all_76.js b/docs/html/search/all_76.js index 3f5e1f4a55..021b8d7f85 100644 --- a/docs/html/search/all_76.js +++ b/docs/html/search/all_76.js @@ -1,5 +1,5 @@ var searchData= [ - ['value',['Value',['../structcub_1_1_key_value_pair.html#a9fd385872c09fd3757e9ba59b2754955',1,'cub::KeyValuePair::Value()'],['../structcub_1_1_item_offset_pair.html#a7faea6eea84a2e5bb4250bd78e765685',1,'cub::ItemOffsetPair::value()'],['../structcub_1_1_key_value_pair.html#a468bef1440a66f45bd9b5193594bf1a4',1,'cub::KeyValuePair::value()']]], + ['value',['value',['../structcub_1_1_item_offset_pair.html#a7faea6eea84a2e5bb4250bd78e765685',1,'cub::ItemOffsetPair::value()'],['../structcub_1_1_key_value_pair.html#a468bef1440a66f45bd9b5193594bf1a4',1,'cub::KeyValuePair::value()'],['../structcub_1_1_key_value_pair.html#a9fd385872c09fd3757e9ba59b2754955',1,'cub::KeyValuePair::Value()']]], ['value_5ftype',['value_type',['../classcub_1_1_arg_index_input_iterator.html#ada4b2b51d9e2957c4839590c7ac135c3',1,'cub::ArgIndexInputIterator::value_type()'],['../classcub_1_1_cache_modified_input_iterator.html#aaaaa98e306495a1e4107144f97e35737',1,'cub::CacheModifiedInputIterator::value_type()'],['../classcub_1_1_cache_modified_output_iterator.html#a5137ea14538ce430886974945005cfef',1,'cub::CacheModifiedOutputIterator::value_type()'],['../classcub_1_1_constant_input_iterator.html#acde3cbd43e78ba2f8a8435cacee84990',1,'cub::ConstantInputIterator::value_type()'],['../classcub_1_1_counting_input_iterator.html#acbf0b49c06d27338f1b2e2ed5b1faece',1,'cub::CountingInputIterator::value_type()'],['../classcub_1_1_tex_obj_input_iterator.html#ad7daad59aaab9b38c1da340831bb9521',1,'cub::TexObjInputIterator::value_type()'],['../classcub_1_1_tex_ref_input_iterator.html#ae4525990cb440826692f7e99abb41ee0',1,'cub::TexRefInputIterator::value_type()'],['../classcub_1_1_transform_input_iterator.html#a71f31c0245f2a927e12d469e91287960',1,'cub::TransformInputIterator::value_type()']]] ]; diff --git a/docs/html/search/functions_73.js b/docs/html/search/functions_73.js index 2ba6e181a6..ea0f06e70b 100644 --- a/docs/html/search/functions_73.js +++ b/docs/html/search/functions_73.js @@ -1,5 +1,6 @@ var searchData= [ + ['scan',['Scan',['../classcub_1_1_warp_scan.html#a32a247bea83254e7c3f3ae4ed8a1d5c7',1,'cub::WarpScan::Scan(T input, T &inclusive_output, T &exclusive_output, T identity, ScanOp scan_op)'],['../classcub_1_1_warp_scan.html#a3266a375b79e44f77087cec512e0c1a3',1,'cub::WarpScan::Scan(T input, T &inclusive_output, T &exclusive_output, ScanOp scan_op)']]], ['scattertoblocked',['ScatterToBlocked',['../classcub_1_1_block_exchange.html#a756e7903c5369261fcf6139e5db52824',1,'cub::BlockExchange']]], ['scattertostriped',['ScatterToStriped',['../classcub_1_1_block_exchange.html#a2ba02edda6bfed9327cd4ee1748aa678',1,'cub::BlockExchange::ScatterToStriped(T items[ITEMS_PER_THREAD], Offset ranks[ITEMS_PER_THREAD])'],['../classcub_1_1_block_exchange.html#a171c90da1f2d5572b22acf059e0b06ea',1,'cub::BlockExchange::ScatterToStriped(T items[ITEMS_PER_THREAD], Offset ranks[ITEMS_PER_THREAD], ValidFlag is_valid[ITEMS_PER_THREAD])']]], ['scattertostripedguarded',['ScatterToStripedGuarded',['../classcub_1_1_block_exchange.html#a0232bb6ad4bdfe1f1578e1ba8d175beb',1,'cub::BlockExchange']]], @@ -27,5 +28,5 @@ var searchData= ['storedirectstriped',['StoreDirectStriped',['../group___util_io.html#gafa774cd981172f96137620c868e628ba',1,'cub::StoreDirectStriped(int linear_tid, OutputIterator block_itr, T(&items)[ITEMS_PER_THREAD])'],['../group___util_io.html#gac9ab95ee7528c9acb3f92c2dd0fc7992',1,'cub::StoreDirectStriped(int linear_tid, OutputIterator block_itr, T(&items)[ITEMS_PER_THREAD], int valid_items)']]], ['storedirectwarpstriped',['StoreDirectWarpStriped',['../group___util_io.html#ga8ffe6de6fbf7c5a617b863c460b0de48',1,'cub::StoreDirectWarpStriped(int linear_tid, OutputIterator block_itr, T(&items)[ITEMS_PER_THREAD])'],['../group___util_io.html#gae201cc0d0cc452e5500a5f28c7e4a234',1,'cub::StoreDirectWarpStriped(int linear_tid, OutputIterator block_itr, T(&items)[ITEMS_PER_THREAD], int valid_items)']]], ['stripedtoblocked',['StripedToBlocked',['../classcub_1_1_block_exchange.html#a2855471bbbcc4d66ac6a29d35a040e0c',1,'cub::BlockExchange']]], - ['sum',['Sum',['../classcub_1_1_block_reduce.html#a7632bd9c8950dd6a3528ca99fa3f0890',1,'cub::BlockReduce::Sum(T input)'],['../classcub_1_1_block_reduce.html#ac5d4591d9513f08b180d4112cb0c4c51',1,'cub::BlockReduce::Sum(T(&inputs)[ITEMS_PER_THREAD])'],['../classcub_1_1_block_reduce.html#a33ddffdde07275ab0c4e1bf61b0d9409',1,'cub::BlockReduce::Sum(T input, int num_valid)'],['../structcub_1_1_device_reduce.html#aaf68d747ee676df2b5b5ff016ee6c371',1,'cub::DeviceReduce::Sum()'],['../classcub_1_1_warp_reduce.html#abe4aeeabf8859a7582a0b5858b84ee7a',1,'cub::WarpReduce::Sum(T input)'],['../classcub_1_1_warp_reduce.html#ad9c4a8d85a7795cf220713f362c36f30',1,'cub::WarpReduce::Sum(T input, int valid_items)']]] + ['sum',['Sum',['../classcub_1_1_block_reduce.html#a7632bd9c8950dd6a3528ca99fa3f0890',1,'cub::BlockReduce::Sum(T input)'],['../classcub_1_1_block_reduce.html#ac5d4591d9513f08b180d4112cb0c4c51',1,'cub::BlockReduce::Sum(T(&inputs)[ITEMS_PER_THREAD])'],['../classcub_1_1_block_reduce.html#a33ddffdde07275ab0c4e1bf61b0d9409',1,'cub::BlockReduce::Sum(T input, int num_valid)'],['../structcub_1_1_device_reduce.html#aaf68d747ee676df2b5b5ff016ee6c371',1,'cub::DeviceReduce::Sum()'],['../classcub_1_1_warp_scan.html#a25bd83f795e88b9260ec2bcbf846fb20',1,'cub::WarpScan::Sum()'],['../classcub_1_1_warp_reduce.html#abe4aeeabf8859a7582a0b5858b84ee7a',1,'cub::WarpReduce::Sum(T input)'],['../classcub_1_1_warp_reduce.html#ad9c4a8d85a7795cf220713f362c36f30',1,'cub::WarpReduce::Sum(T input, int valid_items)']]] ]; diff --git a/docs/html/structcub_1_1_arg_max-members.html b/docs/html/structcub_1_1_arg_max-members.html index 2c153ad368..a839e2fc81 100644 --- a/docs/html/structcub_1_1_arg_max-members.html +++ b/docs/html/structcub_1_1_arg_max-members.html @@ -109,7 +109,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_arg_max.html b/docs/html/structcub_1_1_arg_max.html index 4dc6180580..55b0dae93e 100644 --- a/docs/html/structcub_1_1_arg_max.html +++ b/docs/html/structcub_1_1_arg_max.html @@ -126,7 +126,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_arg_min-members.html b/docs/html/structcub_1_1_arg_min-members.html index c9754ebacc..97544ed0ab 100644 --- a/docs/html/structcub_1_1_arg_min-members.html +++ b/docs/html/structcub_1_1_arg_min-members.html @@ -109,7 +109,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_arg_min.html b/docs/html/structcub_1_1_arg_min.html index 1e1af1b22b..bfc6f79e54 100644 --- a/docs/html/structcub_1_1_arg_min.html +++ b/docs/html/structcub_1_1_arg_min.html @@ -126,7 +126,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_base_traits-members.html b/docs/html/structcub_1_1_base_traits-members.html index fa4fc6416e..8d4bc59d40 100644 --- a/docs/html/structcub_1_1_base_traits-members.html +++ b/docs/html/structcub_1_1_base_traits-members.html @@ -111,7 +111,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_base_traits.html b/docs/html/structcub_1_1_base_traits.html index ed026d28cd..4339d455d3 100644 --- a/docs/html/structcub_1_1_base_traits.html +++ b/docs/html/structcub_1_1_base_traits.html @@ -138,7 +138,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_block_discontinuity_1_1_temp_storage-members.html b/docs/html/structcub_1_1_block_discontinuity_1_1_temp_storage-members.html index 092f52c954..3057311087 100644 --- a/docs/html/structcub_1_1_block_discontinuity_1_1_temp_storage-members.html +++ b/docs/html/structcub_1_1_block_discontinuity_1_1_temp_storage-members.html @@ -112,7 +112,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_block_discontinuity_1_1_temp_storage.html b/docs/html/structcub_1_1_block_discontinuity_1_1_temp_storage.html index 955a4a0e72..13776e8fd3 100644 --- a/docs/html/structcub_1_1_block_discontinuity_1_1_temp_storage.html +++ b/docs/html/structcub_1_1_block_discontinuity_1_1_temp_storage.html @@ -154,7 +154,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_block_exchange_1_1_temp_storage-members.html b/docs/html/structcub_1_1_block_exchange_1_1_temp_storage-members.html index f0f39366f2..d03631e7ec 100644 --- a/docs/html/structcub_1_1_block_exchange_1_1_temp_storage-members.html +++ b/docs/html/structcub_1_1_block_exchange_1_1_temp_storage-members.html @@ -112,7 +112,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_block_exchange_1_1_temp_storage.html b/docs/html/structcub_1_1_block_exchange_1_1_temp_storage.html index 1584898cc4..e2ca2c21da 100644 --- a/docs/html/structcub_1_1_block_exchange_1_1_temp_storage.html +++ b/docs/html/structcub_1_1_block_exchange_1_1_temp_storage.html @@ -156,7 +156,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_block_histogram_1_1_temp_storage-members.html b/docs/html/structcub_1_1_block_histogram_1_1_temp_storage-members.html index 81349a3424..4946591fc4 100644 --- a/docs/html/structcub_1_1_block_histogram_1_1_temp_storage-members.html +++ b/docs/html/structcub_1_1_block_histogram_1_1_temp_storage-members.html @@ -112,7 +112,7 @@
                -Generated on Mon May 12 2014 17:09:55 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_block_histogram_1_1_temp_storage.html b/docs/html/structcub_1_1_block_histogram_1_1_temp_storage.html index 0d9f57c6b0..7a94adb7e4 100644 --- a/docs/html/structcub_1_1_block_histogram_1_1_temp_storage.html +++ b/docs/html/structcub_1_1_block_histogram_1_1_temp_storage.html @@ -157,7 +157,7 @@
                -Generated on Mon May 12 2014 17:09:55 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_block_load_1_1_load_internal_3_01_b_l_o_c_k___l_o_a_d___t_r_a_n_s_p_o_s_e_00_01_d_5ea8dad2df262b118ec77ecff8dc9dd3.html b/docs/html/structcub_1_1_block_load_1_1_load_internal_3_01_b_l_o_c_k___l_o_a_d___t_r_a_n_s_p_o_s_e_00_01_d_5ea8dad2df262b118ec77ecff8dc9dd3.html index 8b49ae5326..9a57a20f29 100644 --- a/docs/html/structcub_1_1_block_load_1_1_load_internal_3_01_b_l_o_c_k___l_o_a_d___t_r_a_n_s_p_o_s_e_00_01_d_5ea8dad2df262b118ec77ecff8dc9dd3.html +++ b/docs/html/structcub_1_1_block_load_1_1_load_internal_3_01_b_l_o_c_k___l_o_a_d___t_r_a_n_s_p_o_s_e_00_01_d_5ea8dad2df262b118ec77ecff8dc9dd3.html @@ -112,7 +112,7 @@
                -Generated on Mon May 12 2014 17:09:55 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_block_load_1_1_load_internal_3_01_b_l_o_c_k___l_o_a_d___t_r_a_n_s_p_o_s_e_00_01_d_u_m_m_y_01_4_1_1_temp_storage.html b/docs/html/structcub_1_1_block_load_1_1_load_internal_3_01_b_l_o_c_k___l_o_a_d___t_r_a_n_s_p_o_s_e_00_01_d_u_m_m_y_01_4_1_1_temp_storage.html index 30da1b12f5..fd16a16f30 100644 --- a/docs/html/structcub_1_1_block_load_1_1_load_internal_3_01_b_l_o_c_k___l_o_a_d___t_r_a_n_s_p_o_s_e_00_01_d_u_m_m_y_01_4_1_1_temp_storage.html +++ b/docs/html/structcub_1_1_block_load_1_1_load_internal_3_01_b_l_o_c_k___l_o_a_d___t_r_a_n_s_p_o_s_e_00_01_d_u_m_m_y_01_4_1_1_temp_storage.html @@ -159,7 +159,7 @@
                -Generated on Mon May 12 2014 17:09:55 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_block_load_1_1_load_internal_3_01_b_l_o_c_k___l_o_a_d___w_a_r_p___t_r_a_n_s_p_o_s_402c3164d23f1ec647db5dad06a54584.html b/docs/html/structcub_1_1_block_load_1_1_load_internal_3_01_b_l_o_c_k___l_o_a_d___w_a_r_p___t_r_a_n_s_p_o_s_402c3164d23f1ec647db5dad06a54584.html index 37c201d30f..0c52c814e9 100644 --- a/docs/html/structcub_1_1_block_load_1_1_load_internal_3_01_b_l_o_c_k___l_o_a_d___w_a_r_p___t_r_a_n_s_p_o_s_402c3164d23f1ec647db5dad06a54584.html +++ b/docs/html/structcub_1_1_block_load_1_1_load_internal_3_01_b_l_o_c_k___l_o_a_d___w_a_r_p___t_r_a_n_s_p_o_s_402c3164d23f1ec647db5dad06a54584.html @@ -159,7 +159,7 @@
                -Generated on Mon May 12 2014 17:09:55 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_block_load_1_1_load_internal_3_01_b_l_o_c_k___l_o_a_d___w_a_r_p___t_r_a_n_s_p_o_s_b58863673477c12a4e46def6747d1835.html b/docs/html/structcub_1_1_block_load_1_1_load_internal_3_01_b_l_o_c_k___l_o_a_d___w_a_r_p___t_r_a_n_s_p_o_s_b58863673477c12a4e46def6747d1835.html index 863100ece0..d1dfec4c11 100644 --- a/docs/html/structcub_1_1_block_load_1_1_load_internal_3_01_b_l_o_c_k___l_o_a_d___w_a_r_p___t_r_a_n_s_p_o_s_b58863673477c12a4e46def6747d1835.html +++ b/docs/html/structcub_1_1_block_load_1_1_load_internal_3_01_b_l_o_c_k___l_o_a_d___w_a_r_p___t_r_a_n_s_p_o_s_b58863673477c12a4e46def6747d1835.html @@ -112,7 +112,7 @@
                -Generated on Mon May 12 2014 17:09:55 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_block_load_1_1_temp_storage-members.html b/docs/html/structcub_1_1_block_load_1_1_temp_storage-members.html index 419bad30cb..bdec8c3993 100644 --- a/docs/html/structcub_1_1_block_load_1_1_temp_storage-members.html +++ b/docs/html/structcub_1_1_block_load_1_1_temp_storage-members.html @@ -112,7 +112,7 @@
                -Generated on Mon May 12 2014 17:09:55 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_block_load_1_1_temp_storage.html b/docs/html/structcub_1_1_block_load_1_1_temp_storage.html index 7a6844cec4..a59a4f59e0 100644 --- a/docs/html/structcub_1_1_block_load_1_1_temp_storage.html +++ b/docs/html/structcub_1_1_block_load_1_1_temp_storage.html @@ -157,7 +157,7 @@
                -Generated on Mon May 12 2014 17:09:55 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_block_radix_sort_1_1_temp_storage-members.html b/docs/html/structcub_1_1_block_radix_sort_1_1_temp_storage-members.html index adc85c3089..7df83dae72 100644 --- a/docs/html/structcub_1_1_block_radix_sort_1_1_temp_storage-members.html +++ b/docs/html/structcub_1_1_block_radix_sort_1_1_temp_storage-members.html @@ -112,7 +112,7 @@
                -Generated on Mon May 12 2014 17:09:55 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_block_radix_sort_1_1_temp_storage.html b/docs/html/structcub_1_1_block_radix_sort_1_1_temp_storage.html index 862df76df9..e159f09b03 100644 --- a/docs/html/structcub_1_1_block_radix_sort_1_1_temp_storage.html +++ b/docs/html/structcub_1_1_block_radix_sort_1_1_temp_storage.html @@ -120,7 +120,7 @@

                The operations exposed by BlockScan require a temporary memory allocation of this nested type for thread communication. This opaque storage can be allocated directly using the __shared__ keyword. Alternatively, it can be aliased to externally allocated memory (shared or global) or union'd with other storage allocation types to facilitate memory reuse.

                -

                Definition at line 394 of file block_radix_sort.cuh.

                +

                Definition at line 400 of file block_radix_sort.cuh.

                Inheritance diagram for cub::BlockRadixSort< Key, BLOCK_DIM_X, ITEMS_PER_THREAD, Value, RADIX_BITS, MEMOIZE_OUTER_SCAN, INNER_SCAN_ALGORITHM, SMEM_CONFIG, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::TempStorage:
                @@ -160,7 +160,7 @@
                -Generated on Mon May 12 2014 17:09:55 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_block_reduce_1_1_temp_storage-members.html b/docs/html/structcub_1_1_block_reduce_1_1_temp_storage-members.html index 8b31087517..48fd4cc3c0 100644 --- a/docs/html/structcub_1_1_block_reduce_1_1_temp_storage-members.html +++ b/docs/html/structcub_1_1_block_reduce_1_1_temp_storage-members.html @@ -112,7 +112,7 @@
                -Generated on Mon May 12 2014 17:09:55 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_block_reduce_1_1_temp_storage.html b/docs/html/structcub_1_1_block_reduce_1_1_temp_storage.html index 5347a92e40..6cbbd079d1 100644 --- a/docs/html/structcub_1_1_block_reduce_1_1_temp_storage.html +++ b/docs/html/structcub_1_1_block_reduce_1_1_temp_storage.html @@ -155,7 +155,7 @@
                -Generated on Mon May 12 2014 17:09:55 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_block_scan_1_1_temp_storage-members.html b/docs/html/structcub_1_1_block_scan_1_1_temp_storage-members.html index 5d92c7002f..e353966fa5 100644 --- a/docs/html/structcub_1_1_block_scan_1_1_temp_storage-members.html +++ b/docs/html/structcub_1_1_block_scan_1_1_temp_storage-members.html @@ -112,7 +112,7 @@
                -Generated on Mon May 12 2014 17:09:56 for CUB by   +Generated on Thu May 22 2014 16:10:32 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_block_scan_1_1_temp_storage.html b/docs/html/structcub_1_1_block_scan_1_1_temp_storage.html index 7bd8cc3d99..6dcdf13944 100644 --- a/docs/html/structcub_1_1_block_scan_1_1_temp_storage.html +++ b/docs/html/structcub_1_1_block_scan_1_1_temp_storage.html @@ -155,7 +155,7 @@
                -Generated on Mon May 12 2014 17:09:55 for CUB by   +Generated on Thu May 22 2014 16:10:32 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_block_store_1_1_store_internal_3_01_b_l_o_c_k___s_t_o_r_e___t_r_a_n_s_p_o_s_e_00_07dc8a0abd37bbc1ba3abf8440890a6c0.html b/docs/html/structcub_1_1_block_store_1_1_store_internal_3_01_b_l_o_c_k___s_t_o_r_e___t_r_a_n_s_p_o_s_e_00_07dc8a0abd37bbc1ba3abf8440890a6c0.html index 96c1bbcd89..f54a9583e5 100644 --- a/docs/html/structcub_1_1_block_store_1_1_store_internal_3_01_b_l_o_c_k___s_t_o_r_e___t_r_a_n_s_p_o_s_e_00_07dc8a0abd37bbc1ba3abf8440890a6c0.html +++ b/docs/html/structcub_1_1_block_store_1_1_store_internal_3_01_b_l_o_c_k___s_t_o_r_e___t_r_a_n_s_p_o_s_e_00_07dc8a0abd37bbc1ba3abf8440890a6c0.html @@ -112,7 +112,7 @@
                -Generated on Mon May 12 2014 17:09:56 for CUB by   +Generated on Thu May 22 2014 16:10:32 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_block_store_1_1_store_internal_3_01_b_l_o_c_k___s_t_o_r_e___t_r_a_n_s_p_o_s_e_00_09dfae03f13932c7dbdb41be30a5767ba.html b/docs/html/structcub_1_1_block_store_1_1_store_internal_3_01_b_l_o_c_k___s_t_o_r_e___t_r_a_n_s_p_o_s_e_00_09dfae03f13932c7dbdb41be30a5767ba.html index b311357a8d..dba56a4881 100644 --- a/docs/html/structcub_1_1_block_store_1_1_store_internal_3_01_b_l_o_c_k___s_t_o_r_e___t_r_a_n_s_p_o_s_e_00_09dfae03f13932c7dbdb41be30a5767ba.html +++ b/docs/html/structcub_1_1_block_store_1_1_store_internal_3_01_b_l_o_c_k___s_t_o_r_e___t_r_a_n_s_p_o_s_e_00_09dfae03f13932c7dbdb41be30a5767ba.html @@ -159,7 +159,7 @@
                -Generated on Mon May 12 2014 17:09:56 for CUB by   +Generated on Thu May 22 2014 16:10:32 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_block_store_1_1_store_internal_3_01_b_l_o_c_k___s_t_o_r_e___w_a_r_p___t_r_a_n_s_p_2a6ef8c29850a6a6dfbf5e4acbc50c1f.html b/docs/html/structcub_1_1_block_store_1_1_store_internal_3_01_b_l_o_c_k___s_t_o_r_e___w_a_r_p___t_r_a_n_s_p_2a6ef8c29850a6a6dfbf5e4acbc50c1f.html index 08eac3944a..1df4ea282b 100644 --- a/docs/html/structcub_1_1_block_store_1_1_store_internal_3_01_b_l_o_c_k___s_t_o_r_e___w_a_r_p___t_r_a_n_s_p_2a6ef8c29850a6a6dfbf5e4acbc50c1f.html +++ b/docs/html/structcub_1_1_block_store_1_1_store_internal_3_01_b_l_o_c_k___s_t_o_r_e___w_a_r_p___t_r_a_n_s_p_2a6ef8c29850a6a6dfbf5e4acbc50c1f.html @@ -112,7 +112,7 @@
                -Generated on Mon May 12 2014 17:09:56 for CUB by   +Generated on Thu May 22 2014 16:10:32 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_block_store_1_1_store_internal_3_01_b_l_o_c_k___s_t_o_r_e___w_a_r_p___t_r_a_n_s_p_8d170856b7ed1df0ed565731a681b449.html b/docs/html/structcub_1_1_block_store_1_1_store_internal_3_01_b_l_o_c_k___s_t_o_r_e___w_a_r_p___t_r_a_n_s_p_8d170856b7ed1df0ed565731a681b449.html index 1a149a7ee4..c082fc271a 100644 --- a/docs/html/structcub_1_1_block_store_1_1_store_internal_3_01_b_l_o_c_k___s_t_o_r_e___w_a_r_p___t_r_a_n_s_p_8d170856b7ed1df0ed565731a681b449.html +++ b/docs/html/structcub_1_1_block_store_1_1_store_internal_3_01_b_l_o_c_k___s_t_o_r_e___w_a_r_p___t_r_a_n_s_p_8d170856b7ed1df0ed565731a681b449.html @@ -159,7 +159,7 @@
                -Generated on Mon May 12 2014 17:09:56 for CUB by   +Generated on Thu May 22 2014 16:10:32 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_block_store_1_1_temp_storage-members.html b/docs/html/structcub_1_1_block_store_1_1_temp_storage-members.html index 9a4a0f5b4b..62e57e359f 100644 --- a/docs/html/structcub_1_1_block_store_1_1_temp_storage-members.html +++ b/docs/html/structcub_1_1_block_store_1_1_temp_storage-members.html @@ -112,7 +112,7 @@
                -Generated on Mon May 12 2014 17:09:56 for CUB by   +Generated on Thu May 22 2014 16:10:32 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_block_store_1_1_temp_storage.html b/docs/html/structcub_1_1_block_store_1_1_temp_storage.html index 12c0b134a1..b531b073c4 100644 --- a/docs/html/structcub_1_1_block_store_1_1_temp_storage.html +++ b/docs/html/structcub_1_1_block_store_1_1_temp_storage.html @@ -157,7 +157,7 @@
                -Generated on Mon May 12 2014 17:09:56 for CUB by   +Generated on Thu May 22 2014 16:10:32 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_caching_device_allocator-members.html b/docs/html/structcub_1_1_caching_device_allocator-members.html index 87d464f1fa..68749b3e6a 100644 --- a/docs/html/structcub_1_1_caching_device_allocator-members.html +++ b/docs/html/structcub_1_1_caching_device_allocator-members.html @@ -117,7 +117,7 @@
                -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_caching_device_allocator.html b/docs/html/structcub_1_1_caching_device_allocator.html index 2eea3456c8..13246d6563 100644 --- a/docs/html/structcub_1_1_caching_device_allocator.html +++ b/docs/html/structcub_1_1_caching_device_allocator.html @@ -281,7 +281,7 @@

                Constructor & Destructor Documentation

                -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_cast-members.html b/docs/html/structcub_1_1_cast-members.html index a89ea87777..f3045df1fe 100644 --- a/docs/html/structcub_1_1_cast-members.html +++ b/docs/html/structcub_1_1_cast-members.html @@ -109,7 +109,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_cast.html b/docs/html/structcub_1_1_cast.html index 6ace142407..293378be9a 100644 --- a/docs/html/structcub_1_1_cast.html +++ b/docs/html/structcub_1_1_cast.html @@ -129,7 +129,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_cub_vector.html b/docs/html/structcub_1_1_cub_vector.html index fdf03d40e7..9f90577453 100644 --- a/docs/html/structcub_1_1_cub_vector.html +++ b/docs/html/structcub_1_1_cub_vector.html @@ -117,7 +117,7 @@
                -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_device_histogram-members.html b/docs/html/structcub_1_1_device_histogram-members.html index da41b212dd..afe13f2cea 100644 --- a/docs/html/structcub_1_1_device_histogram-members.html +++ b/docs/html/structcub_1_1_device_histogram-members.html @@ -114,7 +114,7 @@
                -Generated on Mon May 12 2014 17:09:56 for CUB by   +Generated on Thu May 22 2014 16:10:32 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_device_histogram.html b/docs/html/structcub_1_1_device_histogram.html index ae3b6badce..a74e0eda41 100644 --- a/docs/html/structcub_1_1_device_histogram.html +++ b/docs/html/structcub_1_1_device_histogram.html @@ -951,7 +951,7 @@

                Member Function Documentation

                -Generated on Mon May 12 2014 17:09:56 for CUB by   +Generated on Thu May 22 2014 16:10:32 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_device_partition-members.html b/docs/html/structcub_1_1_device_partition-members.html index 6ebbe32d35..245c268e2c 100644 --- a/docs/html/structcub_1_1_device_partition-members.html +++ b/docs/html/structcub_1_1_device_partition-members.html @@ -110,7 +110,7 @@
                -Generated on Mon May 12 2014 17:09:56 for CUB by   +Generated on Thu May 22 2014 16:10:32 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_device_partition.html b/docs/html/structcub_1_1_device_partition.html index 857bbf7a16..ab5699a458 100644 --- a/docs/html/structcub_1_1_device_partition.html +++ b/docs/html/structcub_1_1_device_partition.html @@ -467,7 +467,7 @@

                Member Function Documentation

                -Generated on Mon May 12 2014 17:09:56 for CUB by   +Generated on Thu May 22 2014 16:10:32 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_device_radix_sort-members.html b/docs/html/structcub_1_1_device_radix_sort-members.html index 3be9aeefbd..f13be7d8de 100644 --- a/docs/html/structcub_1_1_device_radix_sort-members.html +++ b/docs/html/structcub_1_1_device_radix_sort-members.html @@ -112,7 +112,7 @@
                -Generated on Mon May 12 2014 17:09:56 for CUB by   +Generated on Thu May 22 2014 16:10:32 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_device_radix_sort.html b/docs/html/structcub_1_1_device_radix_sort.html index ec2df01b0e..20552a2576 100644 --- a/docs/html/structcub_1_1_device_radix_sort.html +++ b/docs/html/structcub_1_1_device_radix_sort.html @@ -699,7 +699,7 @@

                Member Function Documentation

                -Generated on Mon May 12 2014 17:09:56 for CUB by   +Generated on Thu May 22 2014 16:10:32 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_device_reduce-members.html b/docs/html/structcub_1_1_device_reduce-members.html index 0f756907e7..1d7f5e9517 100644 --- a/docs/html/structcub_1_1_device_reduce-members.html +++ b/docs/html/structcub_1_1_device_reduce-members.html @@ -116,7 +116,7 @@
                -Generated on Mon May 12 2014 17:09:56 for CUB by   +Generated on Thu May 22 2014 16:10:32 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_device_reduce.html b/docs/html/structcub_1_1_device_reduce.html index 5d63263d89..af78a95a14 100644 --- a/docs/html/structcub_1_1_device_reduce.html +++ b/docs/html/structcub_1_1_device_reduce.html @@ -1251,7 +1251,7 @@

                Member Function Documentation

                -Generated on Mon May 12 2014 17:09:56 for CUB by   +Generated on Thu May 22 2014 16:10:32 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_device_scan-members.html b/docs/html/structcub_1_1_device_scan-members.html index 183c119ac6..44cf72f85c 100644 --- a/docs/html/structcub_1_1_device_scan-members.html +++ b/docs/html/structcub_1_1_device_scan-members.html @@ -112,7 +112,7 @@
                -Generated on Mon May 12 2014 17:09:56 for CUB by   +Generated on Thu May 22 2014 16:10:32 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_device_scan.html b/docs/html/structcub_1_1_device_scan.html index 5b1853117c..9b4f1e46be 100644 --- a/docs/html/structcub_1_1_device_scan.html +++ b/docs/html/structcub_1_1_device_scan.html @@ -680,7 +680,7 @@

                Member Function Documentation

                -Generated on Mon May 12 2014 17:09:56 for CUB by   +Generated on Thu May 22 2014 16:10:32 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_device_select-members.html b/docs/html/structcub_1_1_device_select-members.html index 670a29fb1d..d56d7b055b 100644 --- a/docs/html/structcub_1_1_device_select-members.html +++ b/docs/html/structcub_1_1_device_select-members.html @@ -111,7 +111,7 @@
                -Generated on Mon May 12 2014 17:09:56 for CUB by   +Generated on Thu May 22 2014 16:10:32 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_device_select.html b/docs/html/structcub_1_1_device_select.html index 7399afb716..7673e958ca 100644 --- a/docs/html/structcub_1_1_device_select.html +++ b/docs/html/structcub_1_1_device_select.html @@ -630,7 +630,7 @@

                Member Function Documentation

                -Generated on Mon May 12 2014 17:09:56 for CUB by   +Generated on Thu May 22 2014 16:10:32 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_double_buffer-members.html b/docs/html/structcub_1_1_double_buffer-members.html index 6768386776..c499429fe5 100644 --- a/docs/html/structcub_1_1_double_buffer-members.html +++ b/docs/html/structcub_1_1_double_buffer-members.html @@ -113,7 +113,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_double_buffer.html b/docs/html/structcub_1_1_double_buffer.html index cd333fbb8c..07f9d7ee7b 100644 --- a/docs/html/structcub_1_1_double_buffer.html +++ b/docs/html/structcub_1_1_double_buffer.html @@ -197,7 +197,7 @@

                Constructor & Destructor Documentation

                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_equality-members.html b/docs/html/structcub_1_1_equality-members.html index 27e86c9ef0..c8d2ff00c3 100644 --- a/docs/html/structcub_1_1_equality-members.html +++ b/docs/html/structcub_1_1_equality-members.html @@ -109,7 +109,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_equality.html b/docs/html/structcub_1_1_equality.html index c17339ff4f..d0caf9d14d 100644 --- a/docs/html/structcub_1_1_equality.html +++ b/docs/html/structcub_1_1_equality.html @@ -125,7 +125,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_equals-members.html b/docs/html/structcub_1_1_equals-members.html index 0b1e1d30c0..c7fdeae9b8 100644 --- a/docs/html/structcub_1_1_equals-members.html +++ b/docs/html/structcub_1_1_equals-members.html @@ -110,7 +110,7 @@
                -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_equals.html b/docs/html/structcub_1_1_equals.html index a16674c95f..7162af5d75 100644 --- a/docs/html/structcub_1_1_equals.html +++ b/docs/html/structcub_1_1_equals.html @@ -128,7 +128,7 @@
                -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_if-members.html b/docs/html/structcub_1_1_if-members.html index cfa4daa4e6..10156a2ebb 100644 --- a/docs/html/structcub_1_1_if-members.html +++ b/docs/html/structcub_1_1_if-members.html @@ -109,7 +109,7 @@
                -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_if.html b/docs/html/structcub_1_1_if.html index f34d51aa9d..28ea210498 100644 --- a/docs/html/structcub_1_1_if.html +++ b/docs/html/structcub_1_1_if.html @@ -129,7 +129,7 @@
                -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_inequality-members.html b/docs/html/structcub_1_1_inequality-members.html index 0bbc9c0ece..3ecbc03b89 100644 --- a/docs/html/structcub_1_1_inequality-members.html +++ b/docs/html/structcub_1_1_inequality-members.html @@ -109,7 +109,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_inequality.html b/docs/html/structcub_1_1_inequality.html index 88c6fa9074..79295940df 100644 --- a/docs/html/structcub_1_1_inequality.html +++ b/docs/html/structcub_1_1_inequality.html @@ -125,7 +125,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_inequality_wrapper-members.html b/docs/html/structcub_1_1_inequality_wrapper-members.html index f40a1ad9bf..414a432711 100644 --- a/docs/html/structcub_1_1_inequality_wrapper-members.html +++ b/docs/html/structcub_1_1_inequality_wrapper-members.html @@ -111,7 +111,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_inequality_wrapper.html b/docs/html/structcub_1_1_inequality_wrapper.html index d23fe8c277..3a97c277ea 100644 --- a/docs/html/structcub_1_1_inequality_wrapper.html +++ b/docs/html/structcub_1_1_inequality_wrapper.html @@ -141,7 +141,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_int2_type-members.html b/docs/html/structcub_1_1_int2_type-members.html index 96758c9178..10d0b41bc0 100644 --- a/docs/html/structcub_1_1_int2_type-members.html +++ b/docs/html/structcub_1_1_int2_type-members.html @@ -109,7 +109,7 @@
                -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_int2_type.html b/docs/html/structcub_1_1_int2_type.html index 4a1c732ff7..233c9fa8eb 100644 --- a/docs/html/structcub_1_1_int2_type.html +++ b/docs/html/structcub_1_1_int2_type.html @@ -126,7 +126,7 @@
                -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_item_offset_pair-members.html b/docs/html/structcub_1_1_item_offset_pair-members.html index 2edf0f2147..65cd880310 100644 --- a/docs/html/structcub_1_1_item_offset_pair-members.html +++ b/docs/html/structcub_1_1_item_offset_pair-members.html @@ -114,7 +114,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_item_offset_pair.html b/docs/html/structcub_1_1_item_offset_pair.html index 7713cdb2de..aa422daf36 100644 --- a/docs/html/structcub_1_1_item_offset_pair.html +++ b/docs/html/structcub_1_1_item_offset_pair.html @@ -159,7 +159,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_key_value_pair-members.html b/docs/html/structcub_1_1_key_value_pair-members.html index 28fdea2328..28041c195f 100644 --- a/docs/html/structcub_1_1_key_value_pair-members.html +++ b/docs/html/structcub_1_1_key_value_pair-members.html @@ -113,7 +113,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_key_value_pair.html b/docs/html/structcub_1_1_key_value_pair.html index b279ac3d0e..3058ceec8b 100644 --- a/docs/html/structcub_1_1_key_value_pair.html +++ b/docs/html/structcub_1_1_key_value_pair.html @@ -153,7 +153,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_log2-members.html b/docs/html/structcub_1_1_log2-members.html index 4228f55cbd..5c45140b42 100644 --- a/docs/html/structcub_1_1_log2-members.html +++ b/docs/html/structcub_1_1_log2-members.html @@ -109,7 +109,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_log2.html b/docs/html/structcub_1_1_log2.html index f9ebd5b805..26a035c5c5 100644 --- a/docs/html/structcub_1_1_log2.html +++ b/docs/html/structcub_1_1_log2.html @@ -130,7 +130,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_max-members.html b/docs/html/structcub_1_1_max-members.html index 810a54e75a..474d2b1e5f 100644 --- a/docs/html/structcub_1_1_max-members.html +++ b/docs/html/structcub_1_1_max-members.html @@ -109,7 +109,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_max.html b/docs/html/structcub_1_1_max.html index a3c1bc963e..1846cb319c 100644 --- a/docs/html/structcub_1_1_max.html +++ b/docs/html/structcub_1_1_max.html @@ -125,7 +125,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_min-members.html b/docs/html/structcub_1_1_min-members.html index 4a8f2f73cb..adcd71e36c 100644 --- a/docs/html/structcub_1_1_min-members.html +++ b/docs/html/structcub_1_1_min-members.html @@ -109,7 +109,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_min.html b/docs/html/structcub_1_1_min.html index 5fe76148c7..7c6690edc6 100644 --- a/docs/html/structcub_1_1_min.html +++ b/docs/html/structcub_1_1_min.html @@ -125,7 +125,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_null_type.html b/docs/html/structcub_1_1_null_type.html index 80b7de067f..5752f3238e 100644 --- a/docs/html/structcub_1_1_null_type.html +++ b/docs/html/structcub_1_1_null_type.html @@ -112,7 +112,7 @@
                -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_numeric_traits-members.html b/docs/html/structcub_1_1_numeric_traits-members.html index 491a2899b2..c0b417bd15 100644 --- a/docs/html/structcub_1_1_numeric_traits-members.html +++ b/docs/html/structcub_1_1_numeric_traits-members.html @@ -111,7 +111,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_numeric_traits.html b/docs/html/structcub_1_1_numeric_traits.html index e7c8eace38..32e147a868 100644 --- a/docs/html/structcub_1_1_numeric_traits.html +++ b/docs/html/structcub_1_1_numeric_traits.html @@ -139,7 +139,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_power_of_two-members.html b/docs/html/structcub_1_1_power_of_two-members.html index 2e404ae8d0..41e091d15d 100644 --- a/docs/html/structcub_1_1_power_of_two-members.html +++ b/docs/html/structcub_1_1_power_of_two-members.html @@ -109,7 +109,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_power_of_two.html b/docs/html/structcub_1_1_power_of_two.html index 47028f30ea..8516dd2117 100644 --- a/docs/html/structcub_1_1_power_of_two.html +++ b/docs/html/structcub_1_1_power_of_two.html @@ -126,7 +126,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_sum-members.html b/docs/html/structcub_1_1_sum-members.html index 47219abb61..4342785205 100644 --- a/docs/html/structcub_1_1_sum-members.html +++ b/docs/html/structcub_1_1_sum-members.html @@ -109,7 +109,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_sum.html b/docs/html/structcub_1_1_sum.html index 0a07794f05..6f6f8e64c8 100644 --- a/docs/html/structcub_1_1_sum.html +++ b/docs/html/structcub_1_1_sum.html @@ -125,7 +125,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:31 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_traits-members.html b/docs/html/structcub_1_1_traits-members.html index 9ed0248c88..92ad4b6fa4 100644 --- a/docs/html/structcub_1_1_traits-members.html +++ b/docs/html/structcub_1_1_traits-members.html @@ -111,7 +111,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_traits.html b/docs/html/structcub_1_1_traits.html index cbf940880a..a03704f665 100644 --- a/docs/html/structcub_1_1_traits.html +++ b/docs/html/structcub_1_1_traits.html @@ -140,7 +140,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_uninitialized-members.html b/docs/html/structcub_1_1_uninitialized-members.html index 886454dc7d..964ff02db3 100644 --- a/docs/html/structcub_1_1_uninitialized-members.html +++ b/docs/html/structcub_1_1_uninitialized-members.html @@ -112,7 +112,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_uninitialized.html b/docs/html/structcub_1_1_uninitialized.html index 8ef04ccedc..68fb59603d 100644 --- a/docs/html/structcub_1_1_uninitialized.html +++ b/docs/html/structcub_1_1_uninitialized.html @@ -147,7 +147,7 @@
                -Generated on Mon May 12 2014 17:09:54 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_warp_reduce_1_1_temp_storage-members.html b/docs/html/structcub_1_1_warp_reduce_1_1_temp_storage-members.html index ad5591f721..438d69f321 100644 --- a/docs/html/structcub_1_1_warp_reduce_1_1_temp_storage-members.html +++ b/docs/html/structcub_1_1_warp_reduce_1_1_temp_storage-members.html @@ -112,7 +112,7 @@
                -Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_warp_reduce_1_1_temp_storage.html b/docs/html/structcub_1_1_warp_reduce_1_1_temp_storage.html index 9bbb3a7868..38a4c2b76b 100644 --- a/docs/html/structcub_1_1_warp_reduce_1_1_temp_storage.html +++ b/docs/html/structcub_1_1_warp_reduce_1_1_temp_storage.html @@ -152,7 +152,7 @@
                -Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:33 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_warp_scan_1_1_temp_storage-members.html b/docs/html/structcub_1_1_warp_scan_1_1_temp_storage-members.html index d5950756bd..cd87af4932 100644 --- a/docs/html/structcub_1_1_warp_scan_1_1_temp_storage-members.html +++ b/docs/html/structcub_1_1_warp_scan_1_1_temp_storage-members.html @@ -112,7 +112,7 @@
                -Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:32 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/structcub_1_1_warp_scan_1_1_temp_storage.html b/docs/html/structcub_1_1_warp_scan_1_1_temp_storage.html index 038213f6a2..baf46e939c 100644 --- a/docs/html/structcub_1_1_warp_scan_1_1_temp_storage.html +++ b/docs/html/structcub_1_1_warp_scan_1_1_temp_storage.html @@ -152,7 +152,7 @@
                -Generated on Mon May 12 2014 17:09:57 for CUB by   +Generated on Thu May 22 2014 16:10:32 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/tex__obj__input__iterator_8cuh.html b/docs/html/tex__obj__input__iterator_8cuh.html index 0beee7075a..866b400e5b 100644 --- a/docs/html/tex__obj__input__iterator_8cuh.html +++ b/docs/html/tex__obj__input__iterator_8cuh.html @@ -134,7 +134,7 @@
                -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/tex__obj__input__iterator_8cuh_source.html b/docs/html/tex__obj__input__iterator_8cuh_source.html index e3fe648003..61a4a4c769 100644 --- a/docs/html/tex__obj__input__iterator_8cuh_source.html +++ b/docs/html/tex__obj__input__iterator_8cuh_source.html @@ -339,7 +339,7 @@
                -Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/tex__ref__input__iterator_8cuh.html b/docs/html/tex__ref__input__iterator_8cuh.html index 672ac26e37..185417a26f 100644 --- a/docs/html/tex__ref__input__iterator_8cuh.html +++ b/docs/html/tex__ref__input__iterator_8cuh.html @@ -134,7 +134,7 @@
                -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/tex__ref__input__iterator_8cuh_source.html b/docs/html/tex__ref__input__iterator_8cuh_source.html index 28a4805bfe..e5b24993b5 100644 --- a/docs/html/tex__ref__input__iterator_8cuh_source.html +++ b/docs/html/tex__ref__input__iterator_8cuh_source.html @@ -389,7 +389,7 @@
                -Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/thread__load_8cuh.html b/docs/html/thread__load_8cuh.html index c925b70cd9..a9447296fd 100644 --- a/docs/html/thread__load_8cuh.html +++ b/docs/html/thread__load_8cuh.html @@ -152,7 +152,7 @@
                -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/thread__load_8cuh_source.html b/docs/html/thread__load_8cuh_source.html index ecb8057f0d..2594f3b42e 100644 --- a/docs/html/thread__load_8cuh_source.html +++ b/docs/html/thread__load_8cuh_source.html @@ -378,76 +378,84 @@
                359  typedef typename UnitWord<T>::VolatileWord VolatileWord; // Word type for memcopying
                360 
                361  const int VOLATILE_MULTIPLE = sizeof(T) / sizeof(VolatileWord);
                -
                362 
                -
                363  VolatileWord words[VOLATILE_MULTIPLE];
                -
                364 
                -
                365  IterateThreadLoad<0, VOLATILE_MULTIPLE>::Dereference(
                -
                366  reinterpret_cast<volatile VolatileWord*>(ptr),
                -
                367  words);
                -
                368 
                -
                369  return *reinterpret_cast<T*>(words);
                -
                370 
                -
                371 #endif // CUB_PTX_ARCH <= 130
                -
                372 }
                -
                373 
                -
                374 
                -
                378 template <typename T>
                -
                379 __device__ __forceinline__ T ThreadLoad(
                -
                380  T *ptr,
                -
                381  Int2Type<LOAD_VOLATILE> modifier,
                -
                382  Int2Type<true> is_pointer)
                -
                383 {
                -
                384  // Apply tags for partial-specialization
                -
                385  return ThreadLoadVolatilePointer(ptr, Int2Type<Traits<T>::PRIMITIVE>());
                -
                386 }
                -
                387 
                -
                388 
                -
                392 template <typename T, int MODIFIER>
                -
                393 __device__ __forceinline__ T ThreadLoad(
                -
                394  T *ptr,
                -
                395  Int2Type<MODIFIER> modifier,
                -
                396  Int2Type<true> is_pointer)
                -
                397 {
                -
                398  typedef typename UnitWord<T>::DeviceWord DeviceWord;
                -
                399 
                -
                400  const int DEVICE_MULTIPLE = sizeof(T) / sizeof(DeviceWord);
                -
                401 
                -
                402  DeviceWord words[DEVICE_MULTIPLE];
                -
                403 
                -
                404  IterateThreadLoad<0, DEVICE_MULTIPLE>::template Load<CacheLoadModifier(MODIFIER)>(
                -
                405  reinterpret_cast<DeviceWord*>(ptr),
                -
                406  words);
                +
                362 /*
                +
                363  VolatileWord words[VOLATILE_MULTIPLE];
                +
                364 
                +
                365  IterateThreadLoad<0, VOLATILE_MULTIPLE>::Dereference(
                +
                366  reinterpret_cast<volatile VolatileWord*>(ptr),
                +
                367  words);
                +
                368 
                +
                369  return *reinterpret_cast<T*>(words);
                +
                370 */
                +
                371 
                +
                372  T retval;
                +
                373  VolatileWord *words = reinterpret_cast<VolatileWord*>(&retval);
                +
                374  IterateThreadLoad<0, VOLATILE_MULTIPLE>::Dereference(
                +
                375  reinterpret_cast<volatile VolatileWord*>(ptr),
                +
                376  words);
                +
                377  return retval;
                +
                378 
                +
                379 #endif // CUB_PTX_ARCH <= 130
                +
                380 }
                +
                381 
                +
                382 
                +
                386 template <typename T>
                +
                387 __device__ __forceinline__ T ThreadLoad(
                +
                388  T *ptr,
                +
                389  Int2Type<LOAD_VOLATILE> modifier,
                +
                390  Int2Type<true> is_pointer)
                +
                391 {
                +
                392  // Apply tags for partial-specialization
                +
                393  return ThreadLoadVolatilePointer(ptr, Int2Type<Traits<T>::PRIMITIVE>());
                +
                394 }
                +
                395 
                +
                396 
                +
                400 template <typename T, int MODIFIER>
                +
                401 __device__ __forceinline__ T ThreadLoad(
                +
                402  T *ptr,
                +
                403  Int2Type<MODIFIER> modifier,
                +
                404  Int2Type<true> is_pointer)
                +
                405 {
                +
                406  typedef typename UnitWord<T>::DeviceWord DeviceWord;
                407 
                -
                408  return *reinterpret_cast<T*>(words);
                -
                409 }
                -
                410 
                +
                408  const int DEVICE_MULTIPLE = sizeof(T) / sizeof(DeviceWord);
                +
                409 
                +
                410  DeviceWord words[DEVICE_MULTIPLE];
                411 
                -
                415 template <
                -
                416  CacheLoadModifier MODIFIER,
                -
                417  typename InputIterator>
                -
                418 __device__ __forceinline__ typename std::iterator_traits<InputIterator>::value_type ThreadLoad(InputIterator itr)
                -
                419 {
                -
                420  // Apply tags for partial-specialization
                -
                421  return ThreadLoad(
                -
                422  itr,
                -
                423  Int2Type<MODIFIER>(),
                -
                424  Int2Type<IsPointer<InputIterator>::VALUE>());
                -
                425 }
                -
                426 
                -
                427 
                -
                428 
                -
                429 #endif // DOXYGEN_SHOULD_SKIP_THIS
                -
                430 
                -
                431  // end group UtilIo
                -
                433 
                +
                412  IterateThreadLoad<0, DEVICE_MULTIPLE>::template Load<CacheLoadModifier(MODIFIER)>(
                +
                413  reinterpret_cast<DeviceWord*>(ptr),
                +
                414  words);
                +
                415 
                +
                416  return *reinterpret_cast<T*>(words);
                +
                417 }
                +
                418 
                +
                419 
                +
                423 template <
                +
                424  CacheLoadModifier MODIFIER,
                +
                425  typename InputIterator>
                +
                426 __device__ __forceinline__ typename std::iterator_traits<InputIterator>::value_type ThreadLoad(InputIterator itr)
                +
                427 {
                +
                428  // Apply tags for partial-specialization
                +
                429  return ThreadLoad(
                +
                430  itr,
                +
                431  Int2Type<MODIFIER>(),
                +
                432  Int2Type<IsPointer<InputIterator>::VALUE>());
                +
                433 }
                434 
                -
                435 } // CUB namespace
                -
                436 CUB_NS_POSTFIX // Optional outer namespace(s)
                +
                435 
                +
                436 
                +
                437 #endif // DOXYGEN_SHOULD_SKIP_THIS
                +
                438 
                +
                439  // end group UtilIo
                +
                441 
                +
                442 
                +
                443 } // CUB namespace
                +
                444 CUB_NS_POSTFIX // Optional outer namespace(s)
                -Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/thread__operators_8cuh.html b/docs/html/thread__operators_8cuh.html index 98af22f3a8..c06228492f 100644 --- a/docs/html/thread__operators_8cuh.html +++ b/docs/html/thread__operators_8cuh.html @@ -154,7 +154,7 @@
                -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/thread__operators_8cuh_source.html b/docs/html/thread__operators_8cuh_source.html index eae181722e..c84448de95 100644 --- a/docs/html/thread__operators_8cuh_source.html +++ b/docs/html/thread__operators_8cuh_source.html @@ -260,7 +260,7 @@
                -Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/thread__store_8cuh.html b/docs/html/thread__store_8cuh.html index b06b0699b3..4e75827dd0 100644 --- a/docs/html/thread__store_8cuh.html +++ b/docs/html/thread__store_8cuh.html @@ -148,7 +148,7 @@
                -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/thread__store_8cuh_source.html b/docs/html/thread__store_8cuh_source.html index ad3c694e7f..bc726b71ca 100644 --- a/docs/html/thread__store_8cuh_source.html +++ b/docs/html/thread__store_8cuh_source.html @@ -353,71 +353,73 @@
                340  VolatileWord words[VOLATILE_MULTIPLE];
                341  *reinterpret_cast<T*>(words) = val;
                342 
                -
                343  IterateThreadStore<0, VOLATILE_MULTIPLE>::template Dereference(
                -
                344  reinterpret_cast<volatile VolatileWord*>(ptr),
                -
                345  words);
                -
                346 
                -
                347 #endif // CUB_PTX_ARCH <= 130
                -
                348 
                -
                349 }
                -
                350 
                -
                351 
                -
                355 template <typename T>
                -
                356 __device__ __forceinline__ void ThreadStore(
                -
                357  T *ptr,
                -
                358  T val,
                -
                359  Int2Type<STORE_VOLATILE> modifier,
                -
                360  Int2Type<true> is_pointer)
                -
                361 {
                -
                362  ThreadStoreVolatilePtr(ptr, val, Int2Type<Traits<T>::PRIMITIVE>());
                -
                363 }
                -
                364 
                -
                365 
                -
                369 template <typename T, int MODIFIER>
                -
                370 __device__ __forceinline__ void ThreadStore(
                -
                371  T *ptr,
                -
                372  T val,
                -
                373  Int2Type<MODIFIER> modifier,
                -
                374  Int2Type<true> is_pointer)
                -
                375 {
                -
                376  typedef typename UnitWord<T>::DeviceWord DeviceWord; // Word type for memcopying
                -
                377 
                -
                378  const int DEVICE_MULTIPLE = sizeof(T) / sizeof(DeviceWord);
                +
                343 // VolatileWord *words = reinterpret_cast<VolatileWord*>(&val);
                +
                344 
                +
                345  IterateThreadStore<0, VOLATILE_MULTIPLE>::template Dereference(
                +
                346  reinterpret_cast<volatile VolatileWord*>(ptr),
                +
                347  words);
                +
                348 
                +
                349 #endif // CUB_PTX_ARCH <= 130
                +
                350 
                +
                351 }
                +
                352 
                +
                353 
                +
                357 template <typename T>
                +
                358 __device__ __forceinline__ void ThreadStore(
                +
                359  T *ptr,
                +
                360  T val,
                +
                361  Int2Type<STORE_VOLATILE> modifier,
                +
                362  Int2Type<true> is_pointer)
                +
                363 {
                +
                364  ThreadStoreVolatilePtr(ptr, val, Int2Type<Traits<T>::PRIMITIVE>());
                +
                365 }
                +
                366 
                +
                367 
                +
                371 template <typename T, int MODIFIER>
                +
                372 __device__ __forceinline__ void ThreadStore(
                +
                373  T *ptr,
                +
                374  T val,
                +
                375  Int2Type<MODIFIER> modifier,
                +
                376  Int2Type<true> is_pointer)
                +
                377 {
                +
                378  typedef typename UnitWord<T>::DeviceWord DeviceWord; // Word type for memcopying
                379 
                -
                380  DeviceWord words[DEVICE_MULTIPLE];
                +
                380  const int DEVICE_MULTIPLE = sizeof(T) / sizeof(DeviceWord);
                381 
                -
                382  *reinterpret_cast<T*>(words) = val;
                +
                382  DeviceWord words[DEVICE_MULTIPLE];
                383 
                -
                384  IterateThreadStore<0, DEVICE_MULTIPLE>::template Store<CacheStoreModifier(MODIFIER)>(
                -
                385  reinterpret_cast<DeviceWord*>(ptr),
                -
                386  words);
                -
                387 }
                -
                388 
                -
                389 
                -
                393 template <CacheStoreModifier MODIFIER, typename OutputIterator, typename T>
                -
                394 __device__ __forceinline__ void ThreadStore(OutputIterator itr, T val)
                -
                395 {
                -
                396  ThreadStore(
                -
                397  itr,
                -
                398  val,
                -
                399  Int2Type<MODIFIER>(),
                -
                400  Int2Type<IsPointer<OutputIterator>::VALUE>());
                -
                401 }
                -
                402 
                -
                403 
                +
                384  *reinterpret_cast<T*>(words) = val;
                +
                385 
                +
                386  IterateThreadStore<0, DEVICE_MULTIPLE>::template Store<CacheStoreModifier(MODIFIER)>(
                +
                387  reinterpret_cast<DeviceWord*>(ptr),
                +
                388  words);
                +
                389 }
                +
                390 
                +
                391 
                +
                395 template <CacheStoreModifier MODIFIER, typename OutputIterator, typename T>
                +
                396 __device__ __forceinline__ void ThreadStore(OutputIterator itr, T val)
                +
                397 {
                +
                398  ThreadStore(
                +
                399  itr,
                +
                400  val,
                +
                401  Int2Type<MODIFIER>(),
                +
                402  Int2Type<IsPointer<OutputIterator>::VALUE>());
                +
                403 }
                404 
                -
                405 #endif // DOXYGEN_SHOULD_SKIP_THIS
                -
                406 
                -
                407  // end group UtilIo
                -
                409 
                -
                410 
                -
                411 } // CUB namespace
                -
                412 CUB_NS_POSTFIX // Optional outer namespace(s)
                +
                405 
                +
                406 
                +
                407 #endif // DOXYGEN_SHOULD_SKIP_THIS
                +
                408 
                +
                409  // end group UtilIo
                +
                411 
                +
                412 
                +
                413 } // CUB namespace
                +
                414 CUB_NS_POSTFIX // Optional outer namespace(s)
                -Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/transform__input__iterator_8cuh.html b/docs/html/transform__input__iterator_8cuh.html index 1356f2c50c..e05357789c 100644 --- a/docs/html/transform__input__iterator_8cuh.html +++ b/docs/html/transform__input__iterator_8cuh.html @@ -133,7 +133,7 @@
                -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/transform__input__iterator_8cuh_source.html b/docs/html/transform__input__iterator_8cuh_source.html index 347b13e44f..28adf74272 100644 --- a/docs/html/transform__input__iterator_8cuh_source.html +++ b/docs/html/transform__input__iterator_8cuh_source.html @@ -280,7 +280,7 @@
                -Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/util__allocator_8cuh_source.html b/docs/html/util__allocator_8cuh_source.html index 5a126ad4a6..bff10c07dd 100644 --- a/docs/html/util__allocator_8cuh_source.html +++ b/docs/html/util__allocator_8cuh_source.html @@ -675,7 +675,7 @@
                -Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/util__arch_8cuh.html b/docs/html/util__arch_8cuh.html index 05e4fbd28d..f4f02a641f 100644 --- a/docs/html/util__arch_8cuh.html +++ b/docs/html/util__arch_8cuh.html @@ -188,7 +188,7 @@
                -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/util__arch_8cuh_source.html b/docs/html/util__arch_8cuh_source.html index 9899f2bca8..8b95a6eaba 100644 --- a/docs/html/util__arch_8cuh_source.html +++ b/docs/html/util__arch_8cuh_source.html @@ -270,7 +270,7 @@
                -Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/util__debug_8cuh.html b/docs/html/util__debug_8cuh.html index f070f7c415..515fe3f265 100644 --- a/docs/html/util__debug_8cuh.html +++ b/docs/html/util__debug_8cuh.html @@ -151,7 +151,7 @@
                -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/util__debug_8cuh_source.html b/docs/html/util__debug_8cuh_source.html index 6c05c0c1b9..52c1d2c361 100644 --- a/docs/html/util__debug_8cuh_source.html +++ b/docs/html/util__debug_8cuh_source.html @@ -188,7 +188,7 @@
                -Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/util__device_8cuh.html b/docs/html/util__device_8cuh.html index 4e7dcc67d5..8e87c60625 100644 --- a/docs/html/util__device_8cuh.html +++ b/docs/html/util__device_8cuh.html @@ -142,7 +142,7 @@
                -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/util__device_8cuh_source.html b/docs/html/util__device_8cuh_source.html index b6de359b00..652680d3d8 100644 --- a/docs/html/util__device_8cuh_source.html +++ b/docs/html/util__device_8cuh_source.html @@ -413,7 +413,7 @@
                -Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/util__ptx_8cuh.html b/docs/html/util__ptx_8cuh.html index 274c94fae5..ecdfb7f019 100644 --- a/docs/html/util__ptx_8cuh.html +++ b/docs/html/util__ptx_8cuh.html @@ -228,7 +228,7 @@
                -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/util__ptx_8cuh_source.html b/docs/html/util__ptx_8cuh_source.html index d38d24acd6..1ed517faa9 100644 --- a/docs/html/util__ptx_8cuh_source.html +++ b/docs/html/util__ptx_8cuh_source.html @@ -514,7 +514,7 @@
                -Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:29 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/util__type_8cuh.html b/docs/html/util__type_8cuh.html index 9c1e9ac696..4dd099152c 100644 --- a/docs/html/util__type_8cuh.html +++ b/docs/html/util__type_8cuh.html @@ -182,7 +182,7 @@
                -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/util__type_8cuh_source.html.REMOVED.git-id b/docs/html/util__type_8cuh_source.html.REMOVED.git-id index 92c9a27cbf..913971646e 100644 --- a/docs/html/util__type_8cuh_source.html.REMOVED.git-id +++ b/docs/html/util__type_8cuh_source.html.REMOVED.git-id @@ -1 +1 @@ -dc4aa91876ad363c1904f266b038493ac3fa741a \ No newline at end of file +9bf38f8cdb2756a62c2088b15f6248c2c3bbfb35 \ No newline at end of file diff --git a/docs/html/warp__reduce_8cuh.html b/docs/html/warp__reduce_8cuh.html index 45b2d74353..f12fb74185 100644 --- a/docs/html/warp__reduce_8cuh.html +++ b/docs/html/warp__reduce_8cuh.html @@ -141,7 +141,7 @@
                -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/warp__reduce_8cuh_source.html b/docs/html/warp__reduce_8cuh_source.html index 3c93bc5b95..77a71d97f6 100644 --- a/docs/html/warp__reduce_8cuh_source.html +++ b/docs/html/warp__reduce_8cuh_source.html @@ -207,13 +207,13 @@
                211  /******************************************************************/
                215 
                216 
                -
                251  __device__ __forceinline__ T Sum(
                +
                251  __device__ __forceinline__ T Sum(
                252  T input)
                253  {
                254  return InternalWarpReduce(temp_storage).Sum<true, 1>(input, LOGICAL_WARP_THREADS);
                255  }
                256 
                -
                295  __device__ __forceinline__ T Sum(
                +
                295  __device__ __forceinline__ T Sum(
                296  T input,
                297  int valid_items)
                298  {
                @@ -231,21 +231,21 @@
                310 
                349  template <
                350  typename Flag>
                -
                351  __device__ __forceinline__ T HeadSegmentedSum(
                +
                351  __device__ __forceinline__ T HeadSegmentedSum(
                352  T input,
                353  Flag head_flag)
                354  {
                -
                355  return HeadSegmentedReduce(input, head_flag, cub::Sum());
                +
                355  return HeadSegmentedReduce(input, head_flag, cub::Sum());
                356  }
                357 
                358 
                396  template <
                397  typename Flag>
                -
                398  __device__ __forceinline__ T TailSegmentedSum(
                +
                398  __device__ __forceinline__ T TailSegmentedSum(
                399  T input,
                400  Flag tail_flag)
                401  {
                -
                402  return TailSegmentedReduce(input, tail_flag, cub::Sum());
                +
                402  return TailSegmentedReduce(input, tail_flag, cub::Sum());
                403  }
                404 
                405 
                @@ -253,7 +253,7 @@
                408  /******************************************************************/
                412 
                451  template <typename ReductionOp>
                -
                452  __device__ __forceinline__ T Reduce(
                +
                452  __device__ __forceinline__ T Reduce(
                453  T input,
                454  ReductionOp reduction_op)
                455  {
                @@ -261,7 +261,7 @@
                457  }
                458 
                500  template <typename ReductionOp>
                -
                501  __device__ __forceinline__ T Reduce(
                +
                501  __device__ __forceinline__ T Reduce(
                502  T input,
                503  ReductionOp reduction_op,
                504  int valid_items)
                @@ -281,7 +281,7 @@
                557  template <
                558  typename ReductionOp,
                559  typename Flag>
                -
                560  __device__ __forceinline__ T HeadSegmentedReduce(
                +
                560  __device__ __forceinline__ T HeadSegmentedReduce(
                561  T input,
                562  Flag head_flag,
                563  ReductionOp reduction_op)
                @@ -293,7 +293,7 @@
                608  template <
                609  typename ReductionOp,
                610  typename Flag>
                -
                611  __device__ __forceinline__ T TailSegmentedReduce(
                +
                611  __device__ __forceinline__ T TailSegmentedReduce(
                612  T input,
                613  Flag tail_flag,
                614  ReductionOp reduction_op)
                @@ -312,7 +312,7 @@
                -Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/warp__scan_8cuh.html b/docs/html/warp__scan_8cuh.html index 61ff9420c6..81881d28ca 100644 --- a/docs/html/warp__scan_8cuh.html +++ b/docs/html/warp__scan_8cuh.html @@ -141,7 +141,7 @@
                -Generated on Mon May 12 2014 17:09:53 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/html/warp__scan_8cuh_source.html b/docs/html/warp__scan_8cuh_source.html index 868a0f49a9..662332a033 100644 --- a/docs/html/warp__scan_8cuh_source.html +++ b/docs/html/warp__scan_8cuh_source.html @@ -202,32 +202,32 @@
                212  /******************************************************************/
                216 
                217 
                -
                251  __device__ __forceinline__ void InclusiveSum(
                +
                251  __device__ __forceinline__ void InclusiveSum(
                252  T input,
                253  T &output)
                254  {
                -
                255  InternalWarpScan(temp_storage).InclusiveSum(input, output);
                +
                255  InternalWarpScan(temp_storage).InclusiveScan(input, output, cub::Sum());
                256  }
                257 
                258 
                -
                295  __device__ __forceinline__ void InclusiveSum(
                +
                295  __device__ __forceinline__ void InclusiveSum(
                296  T input,
                297  T &output,
                298  T &warp_aggregate)
                299  {
                -
                300  InternalWarpScan(temp_storage).InclusiveSum(input, output, warp_aggregate);
                +
                300  InternalWarpScan(temp_storage).InclusiveScan(input, output, cub::Sum(), warp_aggregate);
                301  }
                302 
                303 
                379  template <typename WarpPrefixCallbackOp>
                -
                380  __device__ __forceinline__ void InclusiveSum(
                +
                380  __device__ __forceinline__ void InclusiveSum(
                381  T input,
                382  T &output,
                383  T &warp_aggregate,
                384  WarpPrefixCallbackOp &warp_prefix_op)
                385  {
                386  // Compute inclusive warp scan
                -
                387  InclusiveSum(input, output, warp_aggregate);
                +
                387  InclusiveSum(input, output, warp_aggregate);
                388 
                389  // Compute warp-wide prefix from aggregate, then broadcast to other lanes
                390  T prefix;
                @@ -241,240 +241,290 @@
                399 
                400 private:
                401 
                -
                403  __device__ __forceinline__ void ExclusiveSum(T input, T &output, Int2Type<true> is_integer)
                +
                403  __device__ __forceinline__ void Sum(T input, T &inclusive_output, T &exclusive_output, Int2Type<true> is_integer)
                404  {
                405  // Compute exclusive warp scan from inclusive warp scan
                -
                406  T inclusive;
                -
                407  InclusiveSum(input, inclusive);
                -
                408  output = inclusive - input;
                -
                409  }
                -
                410 
                -
                412  __device__ __forceinline__ void ExclusiveSum(T input, T &output, Int2Type<false> is_integer)
                -
                413  {
                -
                414  // Delegate to regular scan for non-integer types (because we won't be able to use subtraction)
                -
                415  T identity = ZeroInitialize<T>();
                -
                416  ExclusiveScan(input, output, identity, Sum());
                -
                417  }
                -
                418 
                -
                420  __device__ __forceinline__ void ExclusiveSum(T input, T &output, T &warp_aggregate, Int2Type<true> is_integer)
                -
                421  {
                -
                422  // Compute exclusive warp scan from inclusive warp scan
                -
                423  T inclusive;
                -
                424  InclusiveSum(input, inclusive, warp_aggregate);
                -
                425  output = inclusive - input;
                -
                426  }
                -
                427 
                -
                429  __device__ __forceinline__ void ExclusiveSum(T input, T &output, T &warp_aggregate, Int2Type<false> is_integer)
                -
                430  {
                -
                431  // Delegate to regular scan for non-integer types (because we won't be able to use subtraction)
                -
                432  T identity = ZeroInitialize<T>();
                -
                433  ExclusiveScan(input, output, identity, Sum(), warp_aggregate);
                -
                434  }
                -
                435 
                -
                437  template <typename WarpPrefixCallbackOp>
                -
                438  __device__ __forceinline__ void ExclusiveSum(T input, T &output, T &warp_aggregate, WarpPrefixCallbackOp &warp_prefix_op, Int2Type<true> is_integer)
                -
                439  {
                -
                440  // Compute exclusive warp scan from inclusive warp scan
                -
                441  T inclusive;
                -
                442  InclusiveSum(input, inclusive, warp_aggregate, warp_prefix_op);
                -
                443  output = inclusive - input;
                -
                444  }
                -
                445 
                -
                447  template <typename WarpPrefixCallbackOp>
                -
                448  __device__ __forceinline__ void ExclusiveSum(T input, T &output, T &warp_aggregate, WarpPrefixCallbackOp &warp_prefix_op, Int2Type<false> is_integer)
                -
                449  {
                -
                450  // Delegate to regular scan for non-integer types (because we won't be able to use subtraction)
                -
                451  T identity = ZeroInitialize<T>();
                -
                452  ExclusiveScan(input, output, identity, Sum(), warp_aggregate, warp_prefix_op);
                -
                453  }
                -
                454 
                -
                455 public:
                -
                456 
                -
                457 
                -
                458  /******************************************************************/
                -
                462 
                -
                463 
                -
                503  __device__ __forceinline__ void ExclusiveSum(
                -
                504  T input,
                -
                505  T &output)
                -
                506  {
                -
                507  ExclusiveSum(input, output, Int2Type<IS_INTEGER>());
                -
                508  }
                -
                509 
                -
                510 
                -
                550  __device__ __forceinline__ void ExclusiveSum(
                -
                551  T input,
                -
                552  T &output,
                -
                553  T &warp_aggregate)
                -
                554  {
                -
                555  ExclusiveSum(input, output, warp_aggregate, Int2Type<IS_INTEGER>());
                -
                556  }
                -
                557 
                -
                558 
                -
                637  template <typename WarpPrefixCallbackOp>
                -
                638  __device__ __forceinline__ void ExclusiveSum(
                -
                639  T input,
                -
                640  T &output,
                -
                641  T &warp_aggregate,
                -
                642  WarpPrefixCallbackOp &warp_prefix_op)
                -
                643  {
                -
                644  ExclusiveSum(input, output, warp_aggregate, warp_prefix_op, Int2Type<IS_INTEGER>());
                -
                645  }
                -
                646 
                -
                647 
                -
                649  /******************************************************************/
                -
                653 
                -
                691  template <typename ScanOp>
                -
                692  __device__ __forceinline__ void InclusiveScan(
                -
                693  T input,
                -
                694  T &output,
                -
                695  ScanOp scan_op)
                -
                696  {
                -
                697  InternalWarpScan(temp_storage).InclusiveScan(input, output, scan_op);
                -
                698  }
                -
                699 
                -
                700 
                -
                742  template <typename ScanOp>
                -
                743  __device__ __forceinline__ void InclusiveScan(
                -
                744  T input,
                -
                745  T &output,
                -
                746  ScanOp scan_op,
                -
                747  T &warp_aggregate)
                -
                748  {
                -
                749  InternalWarpScan(temp_storage).InclusiveScan(input, output, scan_op, warp_aggregate);
                -
                750  }
                -
                751 
                -
                752 
                -
                830  template <
                -
                831  typename ScanOp,
                -
                832  typename WarpPrefixCallbackOp>
                -
                833  __device__ __forceinline__ void InclusiveScan(
                -
                834  T input,
                -
                835  T &output,
                -
                836  ScanOp scan_op,
                -
                837  T &warp_aggregate,
                -
                838  WarpPrefixCallbackOp &warp_prefix_op)
                -
                839  {
                -
                840  // Compute inclusive warp scan
                -
                841  InclusiveScan(input, output, scan_op, warp_aggregate);
                -
                842 
                -
                843  // Compute warp-wide prefix from aggregate, then broadcast to other lanes
                -
                844  T prefix;
                -
                845  prefix = warp_prefix_op(warp_aggregate);
                -
                846  prefix = InternalWarpScan(temp_storage).Broadcast(prefix, 0);
                -
                847 
                -
                848  // Update output
                -
                849  output = scan_op(prefix, output);
                -
                850  }
                -
                851 
                -
                852 
                -
                854  /******************************************************************/
                +
                406  InclusiveSum(input, inclusive_output);
                +
                407  exclusive_output = inclusive_output - input;
                +
                408  }
                +
                409 
                +
                411  __device__ __forceinline__ void Sum(T input, T &inclusive_output, T &exclusive_output, Int2Type<false> is_integer)
                +
                412  {
                +
                413  // Delegate to regular scan for non-integer types (because we won't be able to use subtraction)
                +
                414  T identity = ZeroInitialize<T>();
                +
                415  InternalWarpScan(temp_storage).Scan(input, inclusive_output, exclusive_output, identity, cub::Sum());
                +
                416  }
                +
                417 
                +
                419  __device__ __forceinline__ void ExclusiveSum(T input, T &output, Int2Type<true> is_integer)
                +
                420  {
                +
                421  // Compute exclusive warp scan from inclusive warp scan
                +
                422  T inclusive;
                +
                423  InclusiveSum(input, inclusive);
                +
                424  output = inclusive - input;
                +
                425  }
                +
                426 
                +
                428  __device__ __forceinline__ void ExclusiveSum(T input, T &output, Int2Type<false> is_integer)
                +
                429  {
                +
                430  // Delegate to regular scan for non-integer types (because we won't be able to use subtraction)
                +
                431  T identity = ZeroInitialize<T>();
                +
                432  ExclusiveScan(input, output, identity, cub::Sum());
                +
                433  }
                +
                434 
                +
                436  __device__ __forceinline__ void ExclusiveSum(T input, T &output, T &warp_aggregate, Int2Type<true> is_integer)
                +
                437  {
                +
                438  // Compute exclusive warp scan from inclusive warp scan
                +
                439  T inclusive;
                +
                440  InclusiveSum(input, inclusive, warp_aggregate);
                +
                441  output = inclusive - input;
                +
                442  }
                +
                443 
                +
                445  __device__ __forceinline__ void ExclusiveSum(T input, T &output, T &warp_aggregate, Int2Type<false> is_integer)
                +
                446  {
                +
                447  // Delegate to regular scan for non-integer types (because we won't be able to use subtraction)
                +
                448  T identity = ZeroInitialize<T>();
                +
                449  ExclusiveScan(input, output, identity, cub::Sum(), warp_aggregate);
                +
                450  }
                +
                451 
                +
                453  template <typename WarpPrefixCallbackOp>
                +
                454  __device__ __forceinline__ void ExclusiveSum(T input, T &output, T &warp_aggregate, WarpPrefixCallbackOp &warp_prefix_op, Int2Type<true> is_integer)
                +
                455  {
                +
                456  // Compute exclusive warp scan from inclusive warp scan
                +
                457  T inclusive;
                +
                458  InclusiveSum(input, inclusive, warp_aggregate, warp_prefix_op);
                +
                459  output = inclusive - input;
                +
                460  }
                +
                461 
                +
                463  template <typename WarpPrefixCallbackOp>
                +
                464  __device__ __forceinline__ void ExclusiveSum(T input, T &output, T &warp_aggregate, WarpPrefixCallbackOp &warp_prefix_op, Int2Type<false> is_integer)
                +
                465  {
                +
                466  // Delegate to regular scan for non-integer types (because we won't be able to use subtraction)
                +
                467  T identity = ZeroInitialize<T>();
                +
                468  ExclusiveScan(input, output, identity, cub::Sum(), warp_aggregate, warp_prefix_op);
                +
                469  }
                +
                470 
                +
                471 public:
                +
                472 
                +
                473 
                +
                474  /******************************************************************/
                +
                478 
                +
                479 
                +
                519  __device__ __forceinline__ void ExclusiveSum(
                +
                520  T input,
                +
                521  T &output)
                +
                522  {
                +
                523  ExclusiveSum(input, output, Int2Type<IS_INTEGER>());
                +
                524  }
                +
                525 
                +
                526 
                +
                566  __device__ __forceinline__ void ExclusiveSum(
                +
                567  T input,
                +
                568  T &output,
                +
                569  T &warp_aggregate)
                +
                570  {
                +
                571  ExclusiveSum(input, output, warp_aggregate, Int2Type<IS_INTEGER>());
                +
                572  }
                +
                573 
                +
                574 
                +
                653  template <typename WarpPrefixCallbackOp>
                +
                654  __device__ __forceinline__ void ExclusiveSum(
                +
                655  T input,
                +
                656  T &output,
                +
                657  T &warp_aggregate,
                +
                658  WarpPrefixCallbackOp &warp_prefix_op)
                +
                659  {
                +
                660  ExclusiveSum(input, output, warp_aggregate, warp_prefix_op, Int2Type<IS_INTEGER>());
                +
                661  }
                +
                662 
                +
                663 
                +
                665  /******************************************************************/
                +
                669 
                +
                707  template <typename ScanOp>
                +
                708  __device__ __forceinline__ void InclusiveScan(
                +
                709  T input,
                +
                710  T &output,
                +
                711  ScanOp scan_op)
                +
                712  {
                +
                713  InternalWarpScan(temp_storage).InclusiveScan(input, output, scan_op);
                +
                714  }
                +
                715 
                +
                716 
                +
                758  template <typename ScanOp>
                +
                759  __device__ __forceinline__ void InclusiveScan(
                +
                760  T input,
                +
                761  T &output,
                +
                762  ScanOp scan_op,
                +
                763  T &warp_aggregate)
                +
                764  {
                +
                765  InternalWarpScan(temp_storage).InclusiveScan(input, output, scan_op, warp_aggregate);
                +
                766  }
                +
                767 
                +
                768 
                +
                846  template <
                +
                847  typename ScanOp,
                +
                848  typename WarpPrefixCallbackOp>
                +
                849  __device__ __forceinline__ void InclusiveScan(
                +
                850  T input,
                +
                851  T &output,
                +
                852  ScanOp scan_op,
                +
                853  T &warp_aggregate,
                +
                854  WarpPrefixCallbackOp &warp_prefix_op)
                +
                855  {
                +
                856  // Compute inclusive warp scan
                +
                857  InclusiveScan(input, output, scan_op, warp_aggregate);
                858 
                -
                896  template <typename ScanOp>
                -
                897  __device__ __forceinline__ void ExclusiveScan(
                -
                898  T input,
                -
                899  T &output,
                -
                900  T identity,
                -
                901  ScanOp scan_op)
                -
                902  {
                -
                903  InternalWarpScan(temp_storage).ExclusiveScan(input, output, identity, scan_op);
                -
                904  }
                -
                905 
                -
                906 
                -
                947  template <typename ScanOp>
                -
                948  __device__ __forceinline__ void ExclusiveScan(
                -
                949  T input,
                -
                950  T &output,
                -
                951  T identity,
                -
                952  ScanOp scan_op,
                -
                953  T &warp_aggregate)
                -
                954  {
                -
                955  InternalWarpScan(temp_storage).ExclusiveScan(input, output, identity, scan_op, warp_aggregate);
                -
                956  }
                -
                957 
                -
                958 
                -
                1036  template <
                -
                1037  typename ScanOp,
                -
                1038  typename WarpPrefixCallbackOp>
                -
                1039  __device__ __forceinline__ void ExclusiveScan(
                -
                1040  T input,
                -
                1041  T &output,
                -
                1042  T identity,
                -
                1043  ScanOp scan_op,
                -
                1044  T &warp_aggregate,
                -
                1045  WarpPrefixCallbackOp &warp_prefix_op)
                -
                1046  {
                -
                1047  // Exclusive warp scan
                -
                1048  ExclusiveScan(input, output, identity, scan_op, warp_aggregate);
                -
                1049 
                -
                1050  // Compute warp-wide prefix from aggregate, then broadcast to other lanes
                -
                1051  T prefix = warp_prefix_op(warp_aggregate);
                -
                1052  prefix = InternalWarpScan(temp_storage).Broadcast(prefix, 0);
                -
                1053 
                -
                1054  // Update output
                -
                1055  output = (lane_id == 0) ?
                -
                1056  prefix :
                -
                1057  scan_op(prefix, output);
                -
                1058  }
                -
                1059 
                -
                1060 
                -
                1062  /******************************************************************/
                +
                859  // Compute warp-wide prefix from aggregate, then broadcast to other lanes
                +
                860  T prefix;
                +
                861  prefix = warp_prefix_op(warp_aggregate);
                +
                862  prefix = InternalWarpScan(temp_storage).Broadcast(prefix, 0);
                +
                863 
                +
                864  // Update output
                +
                865  output = scan_op(prefix, output);
                +
                866  }
                +
                867 
                +
                868 
                +
                870  /******************************************************************/
                +
                874 
                +
                912  template <typename ScanOp>
                +
                913  __device__ __forceinline__ void ExclusiveScan(
                +
                914  T input,
                +
                915  T &output,
                +
                916  T identity,
                +
                917  ScanOp scan_op)
                +
                918  {
                +
                919  T inclusive_output;
                +
                920  InternalWarpScan(temp_storage).Scan(input, inclusive_output, output, identity, scan_op);
                +
                921  }
                +
                922 
                +
                923 
                +
                964  template <typename ScanOp>
                +
                965  __device__ __forceinline__ void ExclusiveScan(
                +
                966  T input,
                +
                967  T &output,
                +
                968  T identity,
                +
                969  ScanOp scan_op,
                +
                970  T &warp_aggregate)
                +
                971  {
                +
                972  InternalWarpScan(temp_storage).ExclusiveScan(input, output, identity, scan_op, warp_aggregate);
                +
                973  }
                +
                974 
                +
                975 
                +
                1053  template <
                +
                1054  typename ScanOp,
                +
                1055  typename WarpPrefixCallbackOp>
                +
                1056  __device__ __forceinline__ void ExclusiveScan(
                +
                1057  T input,
                +
                1058  T &output,
                +
                1059  T identity,
                +
                1060  ScanOp scan_op,
                +
                1061  T &warp_aggregate,
                +
                1062  WarpPrefixCallbackOp &warp_prefix_op)
                +
                1063  {
                +
                1064  // Exclusive warp scan
                +
                1065  ExclusiveScan(input, output, identity, scan_op, warp_aggregate);
                1066 
                -
                1067 
                -
                1106  template <typename ScanOp>
                -
                1107  __device__ __forceinline__ void ExclusiveScan(
                -
                1108  T input,
                -
                1109  T &output,
                -
                1110  ScanOp scan_op)
                -
                1111  {
                -
                1112  InternalWarpScan(temp_storage).ExclusiveScan(input, output, scan_op);
                -
                1113  }
                -
                1114 
                -
                1115 
                -
                1156  template <typename ScanOp>
                -
                1157  __device__ __forceinline__ void ExclusiveScan(
                -
                1158  T input,
                -
                1159  T &output,
                -
                1160  ScanOp scan_op,
                -
                1161  T &warp_aggregate)
                -
                1162  {
                -
                1163  InternalWarpScan(temp_storage).ExclusiveScan(input, output, scan_op, warp_aggregate);
                -
                1164  }
                -
                1165 
                -
                1166 
                -
                1244  template <
                -
                1245  typename ScanOp,
                -
                1246  typename WarpPrefixCallbackOp>
                -
                1247  __device__ __forceinline__ void ExclusiveScan(
                -
                1248  T input,
                -
                1249  T &output,
                -
                1250  ScanOp scan_op,
                -
                1251  T &warp_aggregate,
                -
                1252  WarpPrefixCallbackOp &warp_prefix_op)
                -
                1253  {
                -
                1254  // Exclusive warp scan
                -
                1255  ExclusiveScan(input, output, scan_op, warp_aggregate);
                -
                1256 
                -
                1257  // Compute warp-wide prefix from aggregate, then broadcast to other lanes
                -
                1258  T prefix = warp_prefix_op(warp_aggregate);
                -
                1259  prefix = InternalWarpScan(temp_storage).Broadcast(prefix, 0);
                -
                1260 
                -
                1261  // Update output with prefix
                -
                1262  output = (lane_id == 0) ?
                -
                1263  prefix :
                -
                1264  scan_op(prefix, output);
                -
                1265  }
                -
                1266 
                -
                1268 };
                -
                1269  // end group WarpModule
                -
                1271 
                -
                1272 } // CUB namespace
                -
                1273 CUB_NS_POSTFIX // Optional outer namespace(s)
                +
                1067  // Compute warp-wide prefix from aggregate, then broadcast to other lanes
                +
                1068  T prefix = warp_prefix_op(warp_aggregate);
                +
                1069  prefix = InternalWarpScan(temp_storage).Broadcast(prefix, 0);
                +
                1070 
                +
                1071  // Update output
                +
                1072  output = (lane_id == 0) ?
                +
                1073  prefix :
                +
                1074  scan_op(prefix, output);
                +
                1075  }
                +
                1076 
                +
                1077 
                +
                1079  /******************************************************************/
                +
                1083 
                +
                1084 
                +
                1123  template <typename ScanOp>
                +
                1124  __device__ __forceinline__ void ExclusiveScan(
                +
                1125  T input,
                +
                1126  T &output,
                +
                1127  ScanOp scan_op)
                +
                1128  {
                +
                1129  T inclusive_output;
                +
                1130  InternalWarpScan(temp_storage).Scan(input, inclusive_output, output, scan_op);
                +
                1131  }
                +
                1132 
                +
                1133 
                +
                1174  template <typename ScanOp>
                +
                1175  __device__ __forceinline__ void ExclusiveScan(
                +
                1176  T input,
                +
                1177  T &output,
                +
                1178  ScanOp scan_op,
                +
                1179  T &warp_aggregate)
                +
                1180  {
                +
                1181  InternalWarpScan(temp_storage).ExclusiveScan(input, output, scan_op, warp_aggregate);
                +
                1182  }
                +
                1183 
                +
                1184 
                +
                1262  template <
                +
                1263  typename ScanOp,
                +
                1264  typename WarpPrefixCallbackOp>
                +
                1265  __device__ __forceinline__ void ExclusiveScan(
                +
                1266  T input,
                +
                1267  T &output,
                +
                1268  ScanOp scan_op,
                +
                1269  T &warp_aggregate,
                +
                1270  WarpPrefixCallbackOp &warp_prefix_op)
                +
                1271  {
                +
                1272  // Exclusive warp scan
                +
                1273  ExclusiveScan(input, output, scan_op, warp_aggregate);
                +
                1274 
                +
                1275  // Compute warp-wide prefix from aggregate, then broadcast to other lanes
                +
                1276  T prefix = warp_prefix_op(warp_aggregate);
                +
                1277  prefix = InternalWarpScan(temp_storage).Broadcast(prefix, 0);
                +
                1278 
                +
                1279  // Update output with prefix
                +
                1280  output = (lane_id == 0) ?
                +
                1281  prefix :
                +
                1282  scan_op(prefix, output);
                +
                1283  }
                +
                1284 
                +
                1286  /******************************************************************/
                +
                1290 
                +
                1333  __device__ __forceinline__ void Sum(
                +
                1334  T input,
                +
                1335  T &inclusive_output,
                +
                1336  T &exclusive_output)
                +
                1337  {
                +
                1338  Sum(input, inclusive_output, exclusive_output, Int2Type<IS_INTEGER>());
                +
                1339  }
                +
                1340 
                +
                1341 
                +
                1382  template <typename ScanOp>
                +
                1383  __device__ __forceinline__ void Scan(
                +
                1384  T input,
                +
                1385  T &inclusive_output,
                +
                1386  T &exclusive_output,
                +
                1387  T identity,
                +
                1388  ScanOp scan_op)
                +
                1389  {
                +
                1390  InternalWarpScan(temp_storage).Scan(input, inclusive_output, exclusive_output, identity, scan_op);
                +
                1391  }
                +
                1392 
                +
                1393 
                +
                1434  template <typename ScanOp>
                +
                1435  __device__ __forceinline__ void Scan(
                +
                1436  T input,
                +
                1437  T &inclusive_output,
                +
                1438  T &exclusive_output,
                +
                1439  ScanOp scan_op)
                +
                1440  {
                +
                1441  InternalWarpScan(temp_storage).Scan(input, inclusive_output, exclusive_output, scan_op);
                +
                1442  }
                +
                1443 
                +
                1444 
                +
                1446 };
                +
                1447  // end group WarpModule
                +
                1449 
                +
                1450 } // CUB namespace
                +
                1451 CUB_NS_POSTFIX // Optional outer namespace(s)
                -Generated on Mon May 12 2014 17:09:52 for CUB by   +Generated on Thu May 22 2014 16:10:30 for CUB by   doxygen 1.8.4
                diff --git a/docs/mainpage.dox b/docs/mainpage.dox index f5d7c99266..e0ecaa2c06 100644 --- a/docs/mainpage.dox +++ b/docs/mainpage.dox @@ -58,7 +58,7 @@    -Download CUB v1.3.0 +Download CUB v1.3.1

            		@@ -510,39 +510,54 @@ for many of the algorithmic techniques used by CUB. + + diff --git a/test/Makefile b/test/Makefile index fdf915bd89..f907b0bd9d 100644 --- a/test/Makefile +++ b/test/Makefile @@ -258,6 +258,7 @@ ALL = link \ test_device_reduce \ test_device_histogram \ test_device_scan \ + test_device_radix_sort \ test_device_reduce_by_key\ test_device_select_unique \ test_device_select_flagged \ @@ -288,7 +289,7 @@ clean : # make all #------------------------------------------------------------------------------- -all : $(ALL) test_device_radix_sort +all : $(ALL) #------------------------------------------------------------------------------- diff --git a/test/test_device_radix_sort.cu b/test/test_device_radix_sort.cu index 2bd0324f7d..aab8291715 100644 --- a/test/test_device_radix_sort.cu +++ b/test/test_device_radix_sort.cu @@ -55,7 +55,6 @@ using namespace cub; bool g_verbose = false; int g_timing_iterations = 0; int g_repeat = 0; -int g_bits = -1; CachingDeviceAllocator g_allocator(true); //--------------------------------------------------------------------- @@ -325,24 +324,15 @@ struct Pair /** - * Initialize key-value sorting problem. + * Initialize key data */ -template -void Initialize( +template +void InitializeKeyBits( GenMode gen_mode, Key *h_keys, - Value *h_values, - Key **h_reference_keys, - Value **h_reference_values, int num_items, - int entropy_reduction, - int begin_bit, - int end_bit) + int entropy_reduction) { - const bool KEYS_ONLY = Equals::VALUE; - - Pair *h_pairs = new Pair[num_items]; - for (int i = 0; i < num_items; ++i) { if (gen_mode == RANDOM) { @@ -352,51 +342,48 @@ void Initialize( } else { h_keys[i] = i; } + } +} - if (h_values != NULL) - RandomBits(h_values[i]); +/** + * Initialize solution + */ +template +void InitializeSolution( + Key *h_keys, + int num_items, + int begin_bit, + int end_bit, + int *&h_reference_ranks, + Key *&h_reference_keys) +{ + Pair *h_pairs = new Pair[num_items]; + int num_bits = end_bit - begin_bit; + for (int i = 0; i < num_items; ++i) + { // Mask off unwanted portions - int num_bits = end_bit - begin_bit; - if ((begin_bit > 0) || (end_bit < sizeof(Key) * 8)) - { - unsigned long long base = 0; - memcpy(&base, &h_keys[i], sizeof(Key)); - base &= ((1ull << num_bits) - 1) << begin_bit; - memcpy(&h_keys[i], &base, sizeof(Key)); - } + unsigned long long base = 0; + memcpy(&base, &h_keys[i], sizeof(Key)); + base &= ((1ull << num_bits) - 1) << begin_bit; + memcpy(&h_pairs[i].key, &base, sizeof(Key)); - h_pairs[i].key = h_keys[i]; - h_pairs[i].value = h_values[i]; + h_pairs[i].value = i; } - if (g_verbose) - { - printf("Input keys:\n"); - DisplayResults(h_keys, num_items); - printf("\n\n"); - if (!KEYS_ONLY) - { - printf("Input values:\n"); - DisplayResults(h_values, num_items); - printf("\n\n"); - } - } - - + printf("\nSorting reference solution on CPU..."); fflush(stdout); if (DESCENDING) std::reverse(h_pairs, h_pairs + num_items); std::stable_sort(h_pairs, h_pairs + num_items); if (DESCENDING) std::reverse(h_pairs, h_pairs + num_items); + printf(" Done.\n"); fflush(stdout); - *h_reference_keys = new Key[num_items]; - *h_reference_values = (KEYS_ONLY) ? NULL : new Value[num_items]; + h_reference_ranks = new int[num_items]; + h_reference_keys = new Key[num_items]; for (int i = 0; i < num_items; ++i) { - (*h_reference_keys)[i] = h_pairs[i].key; - - if ((*h_reference_values) != NULL) - (*h_reference_values)[i] = h_pairs[i].value; + h_reference_ranks[i] = h_pairs[i].value; + h_reference_keys[i] = h_keys[h_pairs[i].value]; } delete[] h_pairs; @@ -404,54 +391,38 @@ void Initialize( - /** * Test DeviceRadixSort */ template < - Backend BACKEND, - typename Key, - typename Value, - bool DESCENDING> + Backend BACKEND, + bool DESCENDING, + typename Key, + typename Value> void Test( - int num_items, - GenMode gen_mode, - int entropy_reduction, - int begin_bit, - int end_bit, - char* type_string) + Key *h_keys, + Value *h_values, + int num_items, + int begin_bit, + int end_bit, + Key *h_reference_keys, + Value *h_reference_values) { const bool KEYS_ONLY = Equals::VALUE; - if (end_bit < 0) end_bit = sizeof(Key) * 8; - - printf("%s %s cub::DeviceRadixSort %d items, %s %d-byte keys %d-byte values, gen-mode %s, descending %d, entropy_reduction %d, begin_bit %d, end_bit %d\n", + printf("%s %s cub::DeviceRadixSort %d items, %d-byte keys %d-byte values, descending %d, begin_bit %d, end_bit %d\n", (BACKEND == CDP) ? "CDP CUB" : (BACKEND == THRUST) ? "Thrust" : "CUB", (KEYS_ONLY) ? "keys-only" : "key-value", - num_items, type_string, - (int) sizeof(Key), (KEYS_ONLY) ? 0 : (int) sizeof(Value), - (gen_mode == RANDOM) ? "RANDOM" : (gen_mode == INTEGER_SEED) ? "SEQUENTIAL" : "HOMOGENOUS", - DESCENDING, entropy_reduction, begin_bit, end_bit); + num_items, (int) sizeof(Key), (KEYS_ONLY) ? 0 : (int) sizeof(Value), + DESCENDING, begin_bit, end_bit); fflush(stdout); - // Allocate host arrays - Key *h_keys = new Key[num_items]; - Value *h_values = (KEYS_ONLY) ? NULL : new Value[num_items]; - - Key *h_reference_keys; - Value *h_reference_values; - - // Initialize problem and solution on host - Initialize( - gen_mode, - h_keys, - h_values, - &h_reference_keys, - &h_reference_values, - num_items, - entropy_reduction, - begin_bit, - end_bit); + if (g_verbose) + { + printf("Input keys:\n"); + DisplayResults(h_keys, num_items); + printf("\n\n"); + } // Allocate device arrays DoubleBuffer d_keys; @@ -488,21 +459,25 @@ void Test( // Run warmup/correctness iteration CubDebugExit(Dispatch(Int2Type(), Int2Type(), d_selector, d_temp_storage_bytes, d_cdp_error, d_temp_storage, temp_storage_bytes, d_keys, d_values, num_items, begin_bit, end_bit, 0, true)); + // Flush any stdout/stderr + fflush(stdout); + fflush(stderr); + // Check for correctness (and display results, if specified) + printf("Warmup done. Checking results:\n"); fflush(stdout); int compare = CompareDeviceResults(h_reference_keys, d_keys.Current(), num_items, true, g_verbose); - printf("\t Compare keys (selector %d): %s ", d_keys.selector, compare ? "FAIL" : "PASS"); + printf("\t Compare keys (selector %d): %s ", d_keys.selector, compare ? "FAIL" : "PASS"); fflush(stdout); if (!KEYS_ONLY) { int values_compare = CompareDeviceResults(h_reference_values, d_values.Current(), num_items, true, g_verbose); compare |= values_compare; - printf("\t Compare values (selector %d): %s ", d_values.selector, values_compare ? "FAIL" : "PASS"); + printf("\t Compare values (selector %d): %s ", d_values.selector, values_compare ? "FAIL" : "PASS"); fflush(stdout); } - // Flush any stdout/stderr - fflush(stdout); - fflush(stderr); - // Performance + if (g_timing_iterations) + printf("\nPerforming timing iterations:\n"); fflush(stdout); + GpuTimer gpu_timer; float elapsed_millis = 0.0f; for (int i = 0; i < g_timing_iterations; ++i) @@ -536,11 +511,6 @@ void Test( printf("\n\n"); // Cleanup - if (h_keys) delete[] h_keys; - if (h_reference_keys) delete[] h_reference_keys; - if (h_values) delete[] h_values; - if (h_reference_values) delete[] h_reference_values; - if (d_keys.d_buffers[0]) CubDebugExit(g_allocator.DeviceFree(d_keys.d_buffers[0])); if (d_keys.d_buffers[1]) CubDebugExit(g_allocator.DeviceFree(d_keys.d_buffers[1])); if (d_values.d_buffers[0]) CubDebugExit(g_allocator.DeviceFree(d_values.d_buffers[0])); @@ -554,119 +524,229 @@ void Test( AssertEquals(0, compare); } + +/** + * Test backend + */ +template +void TestBackend( + Key *h_keys, + int num_items, + int begin_bit, + int end_bit, + Key *h_reference_keys, + int *h_reference_ranks) +{ + const bool KEYS_ONLY = Equals::VALUE; + + Value *h_values = NULL; + Value *h_reference_values = NULL; + + if (!KEYS_ONLY) + { + h_values = new Value[num_items]; + h_reference_values = new Value[num_items]; + + for (int i = 0; i < num_items; ++i) + { + InitValue(INTEGER_SEED, h_values[i], i); + InitValue(INTEGER_SEED, h_reference_values[i], h_reference_ranks[i]); + } + } + + Test(h_keys, h_values, num_items, begin_bit, end_bit, h_reference_keys, h_reference_values); +#ifdef CUB_CDP + Test(h_keys, h_values, num_items, begin_bit, end_bit, h_reference_keys, h_reference_values); +#endif + + if (h_values) delete[] h_values; + if (h_reference_values) delete[] h_reference_values; +} + + + + +/** + * Test value type + */ +template +void TestValueTypes( + Key *h_keys, + int num_items, + int begin_bit, + int end_bit) +{ + // Initialize the solution + + int *h_reference_ranks = NULL; + Key *h_reference_keys = NULL; + InitializeSolution(h_keys, num_items, begin_bit, end_bit, h_reference_ranks, h_reference_keys); + + // Test value types + + TestBackend (h_keys, num_items, begin_bit, end_bit, h_reference_keys, h_reference_ranks); + + TestBackend (h_keys, num_items, begin_bit, end_bit, h_reference_keys, h_reference_ranks); + + if (!Equals::VALUE) + TestBackend (h_keys, num_items, begin_bit, end_bit, h_reference_keys, h_reference_ranks); + + if (!Equals::VALUE) + TestBackend(h_keys, num_items, begin_bit, end_bit, h_reference_keys, h_reference_ranks); + + TestBackend (h_keys, num_items, begin_bit, end_bit, h_reference_keys, h_reference_ranks); + + // Cleanup + + if (h_reference_ranks) delete[] h_reference_ranks; + if (h_reference_keys) delete[] h_reference_keys; +} + + + /** * Test ascending/descending */ -template < - Backend BACKEND, - typename Key, - typename Value> -void Test( - int num_items, - GenMode gen_mode, - int entropy_reduction, - int begin_bit, - int end_bit, - char* type_string) +template +void TestDirection( + Key *h_keys, + int num_items, + int begin_bit, + int end_bit) { - Test(num_items, gen_mode, entropy_reduction, begin_bit, end_bit, type_string); - Test(num_items, gen_mode, entropy_reduction, begin_bit, end_bit, type_string); + TestValueTypes(h_keys, num_items, begin_bit, end_bit); + TestValueTypes(h_keys, num_items, begin_bit, end_bit); } + /** - * Test problem generation + * Test different bit ranges */ -template < - Backend BACKEND, - typename Key, - typename Value> -void Test( - int num_items, - int begin_bit, - int end_bit, - char* type_string) +template +void TestBits( + Key *h_keys, + int num_items) { - for (int entropy_reduction = 0; entropy_reduction <= 6; entropy_reduction += 3) + if (Traits::CATEGORY == UNSIGNED_INTEGER) { - Test(num_items, RANDOM, entropy_reduction, begin_bit, end_bit, type_string); + // Don't test partial-word sorting for fp or signed types (the bit-flipping techniques get in the way) + printf("Testing key bits [%d,%d)\n", 3, 4); fflush(stdout); + TestDirection(h_keys, num_items, 3, 4); } - Test(num_items, UNIFORM, 0, begin_bit, end_bit, type_string); - Test(num_items, INTEGER_SEED, 0, begin_bit, end_bit, type_string); + printf("Testing key bits [%d,%d)\n", 0, sizeof(Key) * 8); fflush(stdout); + TestDirection(h_keys, num_items, 0, sizeof(Key) * 8); } + /** - * Test CDP and num items + * Test different (sub)lengths */ -template < - typename Key, - typename Value> -void Test( - int num_items, - int begin_bit, - int end_bit, - char* type_string) +template +void TestSizes( + Key *h_keys, + int max_items) { - Test(num_items, begin_bit, end_bit, type_string); + while (true) + { + TestBits(h_keys, max_items); -#ifdef CUB_CDP - Test(num_items, begin_bit, end_bit, type_string); -#endif + if (max_items == 1) + break; + + max_items = (max_items + 31) / 32; + } } /** - * Test CDP and num items + * Test key sampling distributions */ -template < - typename Key, - typename Value> -void TestItems( - int num_items, - int begin_bit, - int end_bit, +template +void TestGen( + int max_items, char* type_string) { - // Get ptx version - int ptx_version; - CubDebugExit(PtxVersion(ptx_version)); - - if (num_items < 0) + if (max_items < 0) { - Test(1, begin_bit, end_bit, type_string); - Test(32, begin_bit, end_bit, type_string); - Test(3203, begin_bit, end_bit, type_string); - Test(320003, begin_bit, end_bit, type_string); - if (ptx_version > 100) - Test(9000003, begin_bit, end_bit, type_string); - else - Test(5000003, begin_bit, end_bit, type_string); + int ptx_version; + CubDebugExit(PtxVersion(ptx_version)); + max_items = (ptx_version > 100) ? 9000003 : max_items = 5000003; } - else + + Key *h_keys = new Key[max_items]; + + for (int entropy_reduction = 0; entropy_reduction <= 6; entropy_reduction += 3) { - Test(num_items, begin_bit, end_bit, type_string); + printf("\nTesting random %s keys with entropy reduction factor %d\n", type_string, entropy_reduction); fflush(stdout); + InitializeKeyBits(RANDOM, h_keys, max_items, entropy_reduction); + TestSizes(h_keys, max_items); } + + printf("\nTesting uniform %s keys\n", type_string); fflush(stdout); + InitializeKeyBits(UNIFORM, h_keys, max_items, 0); + TestSizes(h_keys, max_items); + + printf("\nTesting natural number %s keys\n", type_string); fflush(stdout); + InitializeKeyBits(INTEGER_SEED, h_keys, max_items, 0); + TestSizes(h_keys, max_items); + + if (h_keys) delete[] h_keys; } -/** - * Test value type - */ -template -void TestItems( - int num_items, - int begin_bit, - int end_bit, - char* type_string) + +template < + Backend BACKEND, + typename Key, + typename Value, + bool DESCENDING> +void Test( + int num_items, + GenMode gen_mode, + int entropy_reduction, + int begin_bit, + int end_bit, + char *type_string) { - TestItems(num_items, begin_bit, end_bit, type_string); - TestItems(num_items, begin_bit, end_bit, type_string); - TestItems(num_items, begin_bit, end_bit, type_string); - TestItems(num_items, begin_bit, end_bit, type_string); - TestItems(num_items, begin_bit, end_bit, type_string); + const bool KEYS_ONLY = Equals::VALUE; + + Key *h_keys = new Key[num_items]; + int *h_reference_ranks = NULL; + Key *h_reference_keys = NULL; + Value *h_values = NULL; + Value *h_reference_values = NULL; + + if (end_bit < 0) + end_bit = sizeof(Key) * 8; + + InitializeKeyBits(gen_mode, h_keys, num_items, entropy_reduction); + InitializeSolution(h_keys, num_items, begin_bit, end_bit, h_reference_ranks, h_reference_keys); + + if (!KEYS_ONLY) + { + h_values = new Value[num_items]; + h_reference_values = new Value[num_items]; + + for (int i = 0; i < num_items; ++i) + { + InitValue(INTEGER_SEED, h_values[i], i); + InitValue(INTEGER_SEED, h_reference_values[i], h_reference_ranks[i]); + } + } + if (h_reference_ranks) delete[] h_reference_ranks; + + printf("\nTesting bits [%d,%d) of %s keys with gen-mode %d\n", begin_bit, end_bit, type_string, gen_mode); fflush(stdout); + Test(h_keys, h_values, num_items, begin_bit, end_bit, h_reference_keys, h_reference_values); + + if (h_keys) delete[] h_keys; + if (h_reference_keys) delete[] h_reference_keys; + if (h_values) delete[] h_values; + if (h_reference_values) delete[] h_reference_values; } + //--------------------------------------------------------------------- // Main //--------------------------------------------------------------------- @@ -676,6 +756,7 @@ void TestItems( */ int main(int argc, char** argv) { + int bits = -1; int num_items = -1; int entropy_reduction = 0; @@ -685,7 +766,7 @@ int main(int argc, char** argv) args.GetCmdLineArgument("n", num_items); args.GetCmdLineArgument("i", g_timing_iterations); args.GetCmdLineArgument("repeat", g_repeat); - args.GetCmdLineArgument("bits", g_bits); + args.GetCmdLineArgument("bits", bits); args.GetCmdLineArgument("entropy", entropy_reduction); // Print usage @@ -716,47 +797,47 @@ int main(int argc, char** argv) if (num_items < 0) num_items = 20000000; // Compare CUB and thrust on 32b keys-only - Test (num_items, RANDOM, entropy_reduction, 0, g_bits, CUB_TYPE_STRING(unsigned int)); - Test (num_items, RANDOM, entropy_reduction, 0, g_bits, CUB_TYPE_STRING(unsigned int)); + Test (num_items, RANDOM, entropy_reduction, 0, bits, CUB_TYPE_STRING(unsigned int)); + Test (num_items, RANDOM, entropy_reduction, 0, bits, CUB_TYPE_STRING(unsigned int)); // Compare CUB and thrust on 64b keys-only - Test (num_items, RANDOM, entropy_reduction, 0, g_bits, CUB_TYPE_STRING(unsigned long long)); - Test (num_items, RANDOM, entropy_reduction, 0, g_bits, CUB_TYPE_STRING(unsigned long long)); + Test (num_items, RANDOM, entropy_reduction, 0, bits, CUB_TYPE_STRING(unsigned long long)); + Test (num_items, RANDOM, entropy_reduction, 0, bits, CUB_TYPE_STRING(unsigned long long)); // Compare CUB and thrust on 32b key-value pairs - Test (num_items, RANDOM, entropy_reduction, 0, g_bits, CUB_TYPE_STRING(unsigned int)); - Test (num_items, RANDOM, entropy_reduction, 0, g_bits, CUB_TYPE_STRING(unsigned int)); + Test (num_items, RANDOM, entropy_reduction, 0, bits, CUB_TYPE_STRING(unsigned int)); + Test (num_items, RANDOM, entropy_reduction, 0, bits, CUB_TYPE_STRING(unsigned int)); // Compare CUB and thrust on 64b key-value pairs - Test (num_items, RANDOM, entropy_reduction, 0, g_bits, CUB_TYPE_STRING(unsigned long long)); - Test (num_items, RANDOM, entropy_reduction, 0, g_bits, CUB_TYPE_STRING(unsigned long long)); + Test (num_items, RANDOM, entropy_reduction, 0, bits, CUB_TYPE_STRING(unsigned long long)); + Test (num_items, RANDOM, entropy_reduction, 0, bits, CUB_TYPE_STRING(unsigned long long)); #else // Compile/run thorough tests for (int i = 0; i <= g_repeat; ++i) { - TestItems (num_items, 0, g_bits, CUB_TYPE_STRING(char)); - TestItems (num_items, 0, g_bits, CUB_TYPE_STRING(signed char)); - TestItems (num_items, 0, g_bits, CUB_TYPE_STRING(unsigned char)); + TestGen (num_items, CUB_TYPE_STRING(char)); + TestGen (num_items, CUB_TYPE_STRING(signed char)); + TestGen (num_items, CUB_TYPE_STRING(unsigned char)); - TestItems (num_items, 0, g_bits, CUB_TYPE_STRING(short)); - TestItems (num_items, 0, g_bits, CUB_TYPE_STRING(unsigned short)); + TestGen (num_items, CUB_TYPE_STRING(short)); + TestGen (num_items, CUB_TYPE_STRING(unsigned short)); - TestItems (num_items, 0, g_bits, CUB_TYPE_STRING(int)); - TestItems (num_items, 0, g_bits, CUB_TYPE_STRING(unsigned int)); + TestGen (num_items, CUB_TYPE_STRING(int)); + TestGen (num_items, CUB_TYPE_STRING(unsigned int)); - TestItems (num_items, 0, g_bits, CUB_TYPE_STRING(long)); - TestItems (num_items, 0, g_bits, CUB_TYPE_STRING(unsigned long)); + TestGen (num_items, CUB_TYPE_STRING(long)); + TestGen (num_items, CUB_TYPE_STRING(unsigned long)); - TestItems (num_items, 0, g_bits, CUB_TYPE_STRING(long long)); - TestItems (num_items, 0, g_bits, CUB_TYPE_STRING(unsigned long long)); + TestGen (num_items, CUB_TYPE_STRING(long long)); + TestGen (num_items, CUB_TYPE_STRING(unsigned long long)); - TestItems (num_items, 0, g_bits, CUB_TYPE_STRING(float)); + TestGen (num_items, CUB_TYPE_STRING(float)); if (ptx_version > 100) // Don't check doubles on PTX100 because they're down-converted - TestItems (num_items, 0, g_bits, CUB_TYPE_STRING(double)); + TestGen (num_items, CUB_TYPE_STRING(double)); } #endif diff --git a/test/test_util.h b/test/test_util.h index 6527f8f263..41ce8b5709 100644 --- a/test/test_util.h +++ b/test/test_util.h @@ -413,6 +413,12 @@ __host__ __device__ __forceinline__ void InitValue(GenMode gen_mode, T &value, i } +/** + * TestFoo test initialization + */ +__host__ __device__ __forceinline__ void InitValue(GenMode gen_mode, cub::NullType &value, int index = 0) +{} + /****************************************************************************** * Comparison and ostream operators ******************************************************************************/ @@ -943,15 +949,18 @@ int CompareResults(float* computed, float* reference, Offset len, bool verbose = { for (Offset i = 0; i < len; i++) { - float difference = std::abs(computed[i]-reference[i]); - float fraction = difference / std::abs(reference[i]); - - if (fraction > 0.0001) + if (computed[i] != reference[i]) { - if (verbose) std::cout << "INCORRECT: [" << i << "]: " - << CoutCast(computed[i]) << " != " - << CoutCast(reference[i]) << " (difference:" << difference << ", fraction: " << fraction << ")"; - return 1; + float difference = std::abs(computed[i]-reference[i]); + float fraction = difference / std::abs(reference[i]); + + if (fraction > 0.0001) + { + if (verbose) std::cout << "INCORRECT: [" << i << "]: " + << CoutCast(computed[i]) << " != " + << CoutCast(reference[i]) << " (difference:" << difference << ", fraction: " << fraction << ")"; + return 1; + } } } return 0; @@ -975,15 +984,18 @@ int CompareResults(double* computed, double* reference, Offset len, bool verbose { for (Offset i = 0; i < len; i++) { - double difference = std::abs(computed[i]-reference[i]); - double fraction = difference / std::abs(reference[i]); - - if (fraction > 0.0001) + if (computed[i] != reference[i]) { - if (verbose) std::cout << "INCORRECT: [" << i << "]: " - << CoutCast(computed[i]) << " != " - << CoutCast(reference[i]) << " (difference:" << difference << ", fraction: " << fraction << ")"; - return 1; + double difference = std::abs(computed[i]-reference[i]); + double fraction = difference / std::abs(reference[i]); + + if (fraction > 0.0001) + { + if (verbose) std::cout << "INCORRECT: [" << i << "]: " + << CoutCast(computed[i]) << " != " + << CoutCast(reference[i]) << " (difference:" << difference << ", fraction: " << fraction << ")"; + return 1; + } } } return 0; @@ -1018,7 +1030,7 @@ int CompareDeviceResults( bool display_data = false) { // Allocate array on host - T *h_data = (T*) malloc(num_items * sizeof(T)); + T *h_data = new T[num_items]; // Copy data back cudaMemcpy(h_data, d_data, sizeof(T) * num_items, cudaMemcpyDeviceToHost); @@ -1043,7 +1055,7 @@ int CompareDeviceResults( int retval = CompareResults(h_data, h_reference, num_items, verbose); // Cleanup - if (h_data) free(h_data); + if (h_data) delete[] h_data; return retval; } @@ -1062,8 +1074,8 @@ int CompareDeviceDeviceResults( bool display_data = false) { // Allocate array on host - T *h_reference = (T*) malloc(num_items * sizeof(T)); - T *h_data = (T*) malloc(num_items * sizeof(T)); + T *h_reference = new T[num_items]; + T *h_data = new T[num_items]; // Copy data back cudaMemcpy(h_reference, d_reference, sizeof(T) * num_items, cudaMemcpyDeviceToHost); @@ -1088,8 +1100,8 @@ int CompareDeviceDeviceResults( int retval = CompareResults(h_data, h_reference, num_items, verbose); // Cleanup - if (h_reference) free(h_reference); - if (h_data) free(h_data); + if (h_reference) delete[] h_reference; + if (h_data) delete[] h_data; return retval; }
            -05/12/2014
            -[CUB v1.3.0](download_cub.html) +05/22/2014
            +CUB v1.3.1             		- 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 \p BLOCK_DIM_Y and \p BLOCK_DIM_Z being - optional, and \p BLOCK_DIM_X replacing \p 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 scan" methods to cub::WarpScan 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. + - Workaround for a benign WAW race warning reported by cuda-memcheck + in cub::BlockScan specialized for cub::BLOCK_SCAN_WARP_SCANS algorithm. + - Fix for bug in cub::DeviceRadixSort where the algorithm may sort more + key bits than the caller specified (up to the nearest radix digit). + - Fix for a cub::DeviceRadixSort performance regression (~3%) on Kepler and + Fermi that was introduced in v1.3.0. + - See the [change-log](CHANGE_LOG.TXT) for further details, including bug-fixes +
            +05/12/2014
            +[CUB v1.3.0](download_cub.html) +            		 + - 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 \p BLOCK_DIM_Y and \p BLOCK_DIM_Z being + optional, and \p BLOCK_DIM_X replacing \p 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. + - 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. - See the [change-log](CHANGE_LOG.TXT) for further details