- Fixes and tests for BlockHistogram

- Doc updates to v1.0.2

Former-commit-id: 526f484

Author: dumerrill

CHANGE_LOG.TXT

@@ -1,5 +1,12 @@
 //-----------------------------------------------------------------------------
 
+1.0.2    08/23/2013
+    - Corrections to code snippet examples for BlockLoad, BlockStore, and BlockDiscontinuity
+    - Cleaned up unnecessary/missing header includes.  You can now safely #inlude a specific .cuh (instead of cub.cuh)
+    - Bug/compilation fixes for BlockHistogram 
+
+//-----------------------------------------------------------------------------
+
 1.0.1    08/08/2013
     - New collective interface idiom (specialize::construct::invoke).
     - Added best-in-class DeviceRadixSort.  Implements short-circuiting for homogenous digit passes.

README.md

@@ -1,7 +1,7 @@
 <hr>
 <h3>About CUB</h3>
 
-Current release: v1.0.1 (August 8, 2013)
+Current release: v1.0.2 (August 23, 2013)
 
 We recommend the [CUB Project Website](http://nvlabs.github.com/cub) and the [cub-users discussion forum](http://groups.google.com/group/cub-users) for further information and examples.
 

cub/block/block_histogram.cuh

@@ -306,7 +306,7 @@ public:
             histogram[histo_offset + linear_tid] = 0;
         }
         // Finish up with guarded initialization if necessary
-        if ((histo_offset < BLOCK_THREADS) && (histo_offset + linear_tid < BINS))
+        if ((BINS % BLOCK_THREADS != 0) && (histo_offset + linear_tid < BINS))
         {
             histogram[histo_offset + linear_tid] = 0;
         }

cub/block/specializations/block_histogram_sort.cuh

@@ -152,7 +152,7 @@ struct BlockHistogramSort
             temp_storage.run_end[histo_offset + linear_tid] = TILE_SIZE;
         }
         // Finish up with guarded initialization if necessary
-        if ((histo_offset < BLOCK_THREADS) && (histo_offset + linear_tid < BINS))
+        if ((BINS % BLOCK_THREADS != 0) && (histo_offset + linear_tid < BINS))
         {
             temp_storage.run_begin[histo_offset + linear_tid] = TILE_SIZE;
             temp_storage.run_end[histo_offset + linear_tid] = TILE_SIZE;
@@ -182,7 +182,7 @@ struct BlockHistogramSort
             histogram[thread_offset] += count;
         }
         // Finish up with guarded composition if necessary
-        if ((histo_offset < BLOCK_THREADS) && (histo_offset + linear_tid < BINS))
+        if ((BINS % BLOCK_THREADS != 0) && (histo_offset + linear_tid < BINS))
         {
             int thread_offset = histo_offset + linear_tid;
             HistoCounter count = temp_storage.run_end[thread_offset] - temp_storage.run_begin[thread_offset];

cub/device/block/specializations/block_histo_tiles_satomic.cuh

@@ -127,7 +127,7 @@ struct BlockHistogramTilesSharedAtomic
                 this->temp_storage.histograms[CHANNEL][histo_offset + threadIdx.x] = 0;
             }
             // Finish up with guarded initialization if necessary
-            if ((histo_offset < BLOCK_THREADS) && (histo_offset + threadIdx.x < BINS))
+            if ((BINS % BLOCK_THREADS != 0) && (histo_offset + threadIdx.x < BINS))
             {
                 this->temp_storage.histograms[CHANNEL][histo_offset + threadIdx.x] = 0;
             }
@@ -222,7 +222,7 @@ struct BlockHistogramTilesSharedAtomic
                 d_out_histograms[CHANNEL][channel_offset + histo_offset + threadIdx.x] = temp_storage.histograms[CHANNEL][histo_offset + threadIdx.x];
             }
             // Finish up with guarded initialization if necessary
-            if ((histo_offset < BLOCK_THREADS) && (histo_offset + threadIdx.x < BINS))
+            if ((BINS % BLOCK_THREADS != 0) && (histo_offset + threadIdx.x < BINS))
             {
                 d_out_histograms[CHANNEL][channel_offset + histo_offset + threadIdx.x] = temp_storage.histograms[CHANNEL][histo_offset + threadIdx.x];
             }

cub/util_allocator.cuh

@@ -33,9 +33,12 @@
 
 #pragma once
 
+#ifndef __CUDA_ARCH__
+    #include <set>              // NVCC (EDG, really) takes FOREVER to compile std::map
+    #include <map>
+#endif
+
 #include <math.h>
-#include <set>
-#include <map>
 
 #include "util_namespace.cuh"
 #include "util_debug.cuh"
@@ -199,6 +202,8 @@ struct CachingDeviceAllocator
     /// BlockDescriptor comparator function interface
     typedef bool (*Compare)(const BlockDescriptor &, const BlockDescriptor &);
 
+#ifndef __CUDA_ARCH__   // Only define STL container members in host code
+
     /// Set type for cached blocks (ordered by size)
     typedef std::multiset<BlockDescriptor, Compare> CachedBlocks;
 
@@ -208,6 +213,7 @@ struct CachingDeviceAllocator
     /// Map type of device ordinals to the number of cached bytes cached by each device
     typedef std::map<int, size_t> GpuCachedBytes;
 
+#endif // __CUDA_ARCH__
 
     //---------------------------------------------------------------------
     // Fields

docs/Doxyfile

@@ -467,7 +467,7 @@ INLINE_INFO            = YES
 # alphabetically by member name. If set to NO the members will appear in
 # declaration order.
 
-SORT_MEMBER_DOCS       = NO
+SORT_MEMBER_DOCS       = YES
 
 # If the SORT_BRIEF_DOCS tag is set to YES then doxygen will sort the
 # brief documentation of file, namespace and class members alphabetically
@@ -816,7 +816,7 @@ FILTER_SOURCE_PATTERNS =
 # Note: To get rid of all source code in the generated output, make sure also
 # VERBATIM_HEADERS is set to NO.
 
-SOURCE_BROWSER         = NO
+SOURCE_BROWSER         = YES
 
 # Setting the INLINE_SOURCES tag to YES will include the body
 # of functions and classes directly in the documentation.

docs/download_cub.html

@@ -37,14 +37,14 @@
 </head>
 
 <body 
-	onload="downloadURL('https://github.com/NVlabs/cub/archive/1.0.1.zip');" 
+	onload="downloadURL('https://github.com/NVlabs/cub/archive/1.0.2.zip');" 
 	style="color: rgb(102, 102, 102); font-family: Helvetica, arial, freesans, clean, sans-serif; font-size: 13px; font-style: normal; font-variant: normal; font-weight: 300; height: 18px;">
 
 <center>
 If your download doesn't start in 3s:
 <br><br>
-<a href="https://github.com/NVlabs/cub/archive/1.0.1.zip"><img src="download-icon.png" style="position:relative; bottom:-10px; border:0px;"/></a>
-<a href="https://github.com/NVlabs/cub/archive/1.0.1.zip"><em>Download CUB!</em></a>
+<a href="https://github.com/NVlabs/cub/archive/1.0.2.zip"><img src="download-icon.png" style="position:relative; bottom:-10px; border:0px;"/></a>
+<a href="https://github.com/NVlabs/cub/archive/1.0.2.zip"><em>Download CUB!</em></a>
 </center>
 
 </body>

docs/extra_stylesheet.css

@@ -267,7 +267,6 @@ table.memname tbody tr:first-child td:nth-child(2) {
 }
 
 table.memname tbody tr td {
-		line-height: normal;
 }
 
 table.memname tbody tr td.paramtype {
@@ -422,6 +421,8 @@ td.memname, td.paramtype, div.memtemplate {
 }
 
 .memtemplate {
+		padding-bottom: 10px;
+		line-height: 180%;
 		font-size: 85%;
 		color: cadetblue;
 }