Fix for bug in cub::DeviceReduce where reductions of small problems

(small enough to only dispatch a single threadblock) would run in the default stream (stream zero) regardless of whether an alternate stream was specified Former-commit-id: 99f91b9
GaryShen2008 · Jul 29, 2014 · 26b1d58 · 26b1d58
1 parent 9785d22
commit 26b1d58
Show file tree

Hide file tree

Showing 5 changed files with 29 additions and 8 deletions.
diff --git a/CHANGE_LOG.TXT b/CHANGE_LOG.TXT
@@ -1,5 +1,14 @@
 //-----------------------------------------------------------------------------
 
+1.3.2    07/28/2014
+    - Bug fixes: 
+        - Fix for cub::DeviceReduce where reductions of small problems 
+          (small enough to only dispatch a single threadblock) would run in 
+          the default stream (stream zero) regardless of whether an alternate
+          stream was specified.  
+
+//-----------------------------------------------------------------------------
+
 1.3.1    05/23/2014
     - Bug fixes: 
         - Workaround for a benign WAW race warning reported by cuda-memcheck

diff --git a/README.md b/README.md
@@ -1,7 +1,7 @@
 <hr>
 <h3>About CUB</h3>
 
-Current release: v1.3.1 (May 23, 2014)
+Current release: v1.3.2 (July 28, 2014)
 
 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.
 
@@ -92,7 +92,7 @@ to contact us if you would like such things.)  See the
  
 | Date | Version |
 | ---- | ------- |
-| 05/23/2014 | [CUB v1.3.1](https://github.com/NVlabs/cub/archive/1.3.1.zip) |
+| 05/23/2014 | [CUB v1.3.2](https://github.com/NVlabs/cub/archive/1.3.2.zip) |
 | 04/01/2014 | [CUB v1.2.3](https://github.com/NVlabs/cub/archive/1.2.3.zip) |
 | 12/10/2013 | [CUB v1.1.1](https://github.com/NVlabs/cub/archive/1.1.1.zip) |
 | 08/08/2013 | [CUB v1.0.1](https://github.com/NVlabs/cub/archive/1.0.1.zip) |

diff --git a/cub/device/dispatch/device_reduce_dispatch.cuh b/cub/device/dispatch/device_reduce_dispatch.cuh
@@ -542,7 +542,7 @@ struct DeviceReduceDispatch
                     single_reduce_sweep_config.block_threads, (long long) stream, single_reduce_sweep_config.items_per_thread);
 
                 // Invoke single_reduce_sweep_kernel
-                single_reduce_sweep_kernel<<<1, single_reduce_sweep_config.block_threads>>>(
+                single_reduce_sweep_kernel<<<1, single_reduce_sweep_config.block_threads, 0, stream>>>(
                     d_in,
                     d_out,
                     num_items,

diff --git a/docs/download_cub.html b/docs/download_cub.html
@@ -37,14 +37,14 @@
 </head>
 
 <body 
-	onload="downloadURL('https://github.com/NVlabs/cub/archive/1.3.1.zip');" 
+	onload="downloadURL('https://github.com/NVlabs/cub/archive/1.3.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.3.1.zip"><img src="download-icon.png" style="position:relative; bottom:-10px; border:0px;"/></a>
-<a href="https://github.com/NVlabs/cub/archive/1.3.1.zip"><em>Download CUB!</em></a>
+<a href="https://github.com/NVlabs/cub/archive/1.3.2.zip"><img src="download-icon.png" style="position:relative; bottom:-10px; border:0px;"/></a>
+<a href="https://github.com/NVlabs/cub/archive/1.3.2.zip"><em>Download CUB!</em></a>
 </center>
 
 </body>

diff --git a/docs/mainpage.dox b/docs/mainpage.dox
@@ -58,7 +58,7 @@
 
 <a href="download_cub.html"><img src="download-icon.png" style="position:relative; bottom:-10px; border:0px;"/></a>
 &nbsp;&nbsp;
-<a href="download_cub.html"><em><b>Download CUB v1.3.1</b></em></a>
+<a href="download_cub.html"><em><b>Download CUB v1.3.2</b></em></a>
 
 </td><td>
 
@@ -522,9 +522,21 @@ The following table enumerates prior feature releases as well as update versions
 for the current feature release.
 <table>
 
+<tr><td style="white-space: nowrap; vertical-align:text-top;">
+07/28/2014<br>
+[<b>CUB v1.3.2</b>](download_cub.html)
+</td><td style="vertical-align:text-top;">
+  - Bug fixes: 
+    - Fix for bug in cub::DeviceReduce where reductions of small problems 
+      (small enough to only dispatch a single threadblock) would run in 
+      the default stream (stream zero) regardless of whether an alternate
+      stream was specified.  
+  - See the [change-log](CHANGE_LOG.TXT) for further details, including bug-fixes
+</td></tr>
+
 <tr><td style="white-space: nowrap; vertical-align:text-top;">
 05/23/2014<br>
-[<b>CUB v1.3.1</b>](download_cub.html)
+<b>CUB v1.3.1</b>
 </td><td style="vertical-align:text-top;">
   - Bug fixes: 
     - Workaround for a benign WAW race warning reported by cuda-memcheck