- Added sm52 scan tuning

- Added fast in-core processing for very small sorting problems Former-commit-id: 7702e4dd652a9cd43e93f168d3522dc9a06d1ebe
NVIDIA · Feb 23, 2015 · ce261ed · ce261ed
1 parent 6330f93
commit ce261ed
Show file tree

Hide file tree

Showing 6 changed files with 346 additions and 197 deletions.
diff --git a/cub/agent/agent_radix_sort_downsweep.cuh b/cub/agent/agent_radix_sort_downsweep.cuh
@@ -571,6 +571,9 @@ struct AgentRadixSortDownsweep
         GatherScatterValues<FULL_TILE>(values, relative_bin_offsets, ranks, block_offset, valid_items);
     }
 
+    //---------------------------------------------------------------------
+    // Copy shortcut
+    //---------------------------------------------------------------------
 
     /**
      * Copy tiles within the range of input
@@ -728,6 +731,7 @@ struct AgentRadixSortDownsweep
             }
         }
     }
+
 };
 
 

diff --git a/cub/agent/agent_scan.cuh b/cub/agent/agent_scan.cuh
@@ -367,7 +367,6 @@ struct AgentScan
         GridQueue<int>      queue,              ///< Queue descriptor for assigning tiles of work to thread blocks
         ScanTileState       &tile_status)       ///< Global list of tile status
     {
-#if (CUB_PTX_ARCH <= 130)
         // Blocks are launched in increasing order, so just assign one tile per block
 
         int     tile_idx        = (blockIdx.y * gridDim.x) + blockIdx.x;    // Current tile index
@@ -378,44 +377,6 @@ struct AgentScan
             ConsumeTile<false>(num_items, num_remaining, tile_idx, block_offset, tile_status);
         else if (num_remaining > 0)
             ConsumeTile<true>(num_items, num_remaining, tile_idx, block_offset, tile_status);
-
-#else
-        // Blocks may not be launched in increasing order, so work-steal tiles
-
-        // Get first tile index
-        if (threadIdx.x == 0)
-            temp_storage.tile_idx = queue.Drain(1);
-
-        __syncthreads();
-
-        int     tile_idx        = temp_storage.tile_idx;
-        OffsetT block_offset    = TILE_ITEMS * tile_idx;
-        OffsetT num_remaining   = num_items - block_offset;
-
-        while (num_remaining >= TILE_ITEMS)
-        {
-            // Consume full tile
-            ConsumeTile<false>(num_items, num_remaining, tile_idx, block_offset, tile_status);
-
-            // Get next tile
-            if (threadIdx.x == 0)
-                temp_storage.tile_idx = queue.Drain(1);
-
-            __syncthreads();
-
-            tile_idx        = temp_storage.tile_idx;
-            block_offset    = TILE_ITEMS * tile_idx;
-            num_remaining   = num_items - block_offset;
-        }
-
-        // Consume the last (and potentially partially-full) tile
-        if (num_remaining > 0)
-        {
-            ConsumeTile<true>(num_items, num_remaining, tile_idx, block_offset, tile_status);
-        }
-
-#endif
-
     }
 
 

diff --git a/cub/block/block_radix_sort.cuh b/cub/block/block_radix_sort.cuh
@@ -332,6 +332,10 @@ private:
         }
     }
 
+public:
+
+#ifndef DOXYGEN_SHOULD_SKIP_THIS    // Do not document
+
     /// Sort blocked -> striped arrangement
     template <int DESCENDING, int KEYS_ONLY>
     __device__ __forceinline__ void SortBlockedToStriped(
@@ -394,9 +398,7 @@ private:
         }
     }
 
-
-
-public:
+#endif // DOXYGEN_SHOULD_SKIP_THIS
 
     /// \smemstorage{BlockScan}
     struct TempStorage : Uninitialized<_TempStorage> {};