- Remove smem-mode optimizations because they can cause uintentional/

undesirable device-wide synchronization Former-commit-id: 9a14645
GaryShen2008 · Jul 30, 2014 · 109e0f7 · 109e0f7
1 parent 470beb5
commit 109e0f7
Show file tree

Hide file tree

Showing 9 changed files with 274 additions and 293 deletions.
diff --git a/cub/device/device_radix_sort.cuh b/cub/device/device_radix_sort.cuh
@@ -155,7 +155,7 @@ struct DeviceRadixSort
         // Signed integer type for global offsets
         typedef int Offset;
 
-        return DeviceRadixSortDispatch<false, Key, Value, Offset>::Dispatch(
+        return DeviceRadixSortDispatch<false, false, Key, Value, Offset>::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_keys,
@@ -238,7 +238,7 @@ struct DeviceRadixSort
         // Signed integer type for global offsets
         typedef int Offset;
 
-        return DeviceRadixSortDispatch<true, Key, Value, Offset>::Dispatch(
+        return DeviceRadixSortDispatch<true, false, Key, Value, Offset>::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_keys,
@@ -319,7 +319,7 @@ struct DeviceRadixSort
         // Null value type
         DoubleBuffer<NullType> d_values;
 
-        return DeviceRadixSortDispatch<false, Key, NullType, Offset>::Dispatch(
+        return DeviceRadixSortDispatch<false, false, Key, NullType, Offset>::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_keys,
@@ -396,7 +396,7 @@ struct DeviceRadixSort
         // Null value type
         DoubleBuffer<NullType> d_values;
 
-        return DeviceRadixSortDispatch<true, Key, NullType, Offset>::Dispatch(
+        return DeviceRadixSortDispatch<true, false, Key, NullType, Offset>::Dispatch(
             d_temp_storage,
             temp_storage_bytes,
             d_keys,

diff --git a/cub/device/dispatch/device_histogram_dispatch.cuh b/cub/device/dispatch/device_histogram_dispatch.cuh
@@ -412,7 +412,7 @@ struct DeviceHistogramDispatch
                 GridQueue<int>::AllocationSize()                                                // bytes needed for grid queue descriptor
             };
 
-            // Alias the temporary allocations from the single storage blob (or set the necessary size of the blob)
+            // Alias the temporary allocations from the single storage blob (or compute the necessary size of the blob)
             if (CubDebug(error = AliasTemporaries(d_temp_storage, temp_storage_bytes, allocations, allocation_sizes))) break;
             if (d_temp_storage == NULL)
             {

diff --git a/cub/device/dispatch/device_radix_sort_dispatch.cuh b/cub/device/dispatch/device_radix_sort_dispatch.cuh
diff --git a/cub/device/dispatch/device_reduce_by_key_dispatch.cuh b/cub/device/dispatch/device_reduce_by_key_dispatch.cuh
@@ -412,7 +412,7 @@ struct DeviceReduceByKeyDispatch
             if (CubDebug(error = ScanTileState::AllocationSize(num_tiles, allocation_sizes[0]))) break;    // bytes needed for tile status descriptors
             allocation_sizes[1] = GridQueue<int>::AllocationSize();                                             // bytes needed for grid queue descriptor
 
-            // Compute allocation pointers into the single storage blob (or set the necessary size of the blob)
+            // Compute allocation pointers into the single storage blob (or compute the necessary size of the blob)
             void* allocations[2];
             if (CubDebug(error = AliasTemporaries(d_temp_storage, temp_storage_bytes, allocations, allocation_sizes))) break;
             if (d_temp_storage == NULL)

diff --git a/cub/device/dispatch/device_reduce_dispatch.cuh b/cub/device/dispatch/device_reduce_dispatch.cuh
@@ -607,7 +607,7 @@ struct DeviceReduceDispatch
                     GridQueue<int>::AllocationSize()        // bytes needed for grid queue descriptor
                 };
 
-                // Alias the temporary allocations from the single storage blob (or set the necessary size of the blob)
+                // Alias the temporary allocations from the single storage blob (or compute the necessary size of the blob)
                 if (CubDebug(error = AliasTemporaries(d_temp_storage, temp_storage_bytes, allocations, allocation_sizes))) break;
                 if (d_temp_storage == NULL)
                 {

diff --git a/cub/device/dispatch/device_rle_dispatch.cuh b/cub/device/dispatch/device_rle_dispatch.cuh
@@ -405,7 +405,7 @@ struct DeviceRleDispatch
             if (CubDebug(error = ScanTileState::AllocationSize(num_tiles, allocation_sizes[0]))) break;    // bytes needed for tile status descriptors
             allocation_sizes[1] = GridQueue<int>::AllocationSize();                                             // bytes needed for grid queue descriptor
 
-            // Compute allocation pointers into the single storage blob (or set the necessary size of the blob)
+            // Compute allocation pointers into the single storage blob (or compute the necessary size of the blob)
             void* allocations[2];
             if (CubDebug(error = AliasTemporaries(d_temp_storage, temp_storage_bytes, allocations, allocation_sizes))) break;
             if (d_temp_storage == NULL)

diff --git a/cub/device/dispatch/device_scan_dispatch.cuh b/cub/device/dispatch/device_scan_dispatch.cuh
@@ -416,7 +416,7 @@ struct DeviceScanDispatch
             if (CubDebug(error = ScanTileState::AllocationSize(num_tiles, allocation_sizes[0]))) break;    // bytes needed for tile status descriptors
             allocation_sizes[1] = GridQueue<int>::AllocationSize();                                             // bytes needed for grid queue descriptor
 
-            // Compute allocation pointers into the single storage blob (or set the necessary size of the blob)
+            // Compute allocation pointers into the single storage blob (or compute the necessary size of the blob)
             void* allocations[2];
             if (CubDebug(error = AliasTemporaries(d_temp_storage, temp_storage_bytes, allocations, allocation_sizes))) break;
             if (d_temp_storage == NULL)

diff --git a/cub/device/dispatch/device_select_dispatch.cuh b/cub/device/dispatch/device_select_dispatch.cuh
@@ -409,7 +409,7 @@ struct DeviceSelectDispatch
             if (CubDebug(error = ScanTileState::AllocationSize(num_tiles, allocation_sizes[0]))) break;    // bytes needed for tile status descriptors
             allocation_sizes[1] = GridQueue<int>::AllocationSize();                                             // bytes needed for grid queue descriptor
 
-            // Compute allocation pointers into the single storage blob (or set the necessary size of the blob)
+            // Compute allocation pointers into the single storage blob (or compute the necessary size of the blob)
             void* allocations[2];
             if (CubDebug(error = AliasTemporaries(d_temp_storage, temp_storage_bytes, allocations, allocation_sizes))) break;
             if (d_temp_storage == NULL)

diff --git a/cub/util_type.cuh b/cub/util_type.cuh
@@ -624,6 +624,10 @@ struct DoubleBuffer
 
     /// \brief Return pointer to the currently valid buffer
     __host__ __device__ __forceinline__ T* Current() { return d_buffers[selector]; }
+
+    /// \brief Return pointer to the currently invalid buffer
+    __host__ __device__ __forceinline__ T* Invalid() { return d_buffers[selector ^ 1]; }
+
 };