Half-done

gevtushenko · May 6, 2023 · f88c6c8 · f88c6c8
1 parent 9c2acdf
commit f88c6c8
Show file tree

Hide file tree

Showing 3 changed files with 73 additions and 30 deletions.
diff --git a/benchmarks/README.md b/benchmarks/README.md
@@ -0,0 +1,42 @@
+## Benchmarks 
+
+### Ready
+
+- radix sort
+  - keys
+  - pairs 
+- merge sort
+  - keys
+  - pairs 
+- reduce
+  - sum
+  - max
+- scan
+  - sum
+  - max
+
+### In Progress
+
+- select
+  - flagged
+  - if
+
+### TODO
+
+- adjacent difference 
+  - left
+  - right
+- histogram
+  - even
+  - range
+- partition
+  - flagged
+  - if
+- rle
+  - encode
+  - non trivial runs
+- reduce
+  - by key
+- scan
+  - by key
+- segmented
diff --git a/benchmarks/bench/select/flagged.cu b/benchmarks/bench/select/flagged.cu
@@ -0,0 +1 @@
+#include <cub/device/device_reduce.cuh>
diff --git a/cub/device/dispatch/dispatch_select_if.cuh b/cub/device/dispatch/dispatch_select_if.cuh
@@ -107,6 +107,31 @@ __global__ void DeviceSelectSweepKernel(
 }
 
 
+namespace detail 
+{
+
+template <class InputT, bool MayAlias>
+struct device_select_policy_hub
+{
+    struct Policy350 : ChainedPolicy<350, Policy350, Policy350>
+    {
+      static constexpr int NOMINAL_4B_ITEMS_PER_THREAD = 10;
+
+      static constexpr int ITEMS_PER_THREAD =
+        CUB_MIN(NOMINAL_4B_ITEMS_PER_THREAD,
+                CUB_MAX(1, (NOMINAL_4B_ITEMS_PER_THREAD * 4 / sizeof(InputT))));
+
+      using SelectIfPolicyT = AgentSelectIfPolicy<128,
+                                                  ITEMS_PER_THREAD,
+                                                  BLOCK_LOAD_DIRECT,
+                                                  MayAlias ? LOAD_CA : LOAD_LDG,
+                                                  BLOCK_SCAN_WARP_SCANS>;
+    };
+
+    using MaxPolicy = Policy350;
+};
+
+} // detail
 
 
 /******************************************************************************
@@ -125,8 +150,10 @@ template <
     typename    EqualityOpT,                    ///< Equality operator type (NullType if selection functor or selection flags is to be used for selection)
     typename    OffsetT,                        ///< Signed integer type for global offsets
     bool        KEEP_REJECTS,                   ///< Whether or not we push rejected items to the back of the output
-    bool        MayAlias = false>                   
-struct DispatchSelectIf
+    bool        MayAlias = false,
+    typename    SelectedPolicy = 
+        detail::device_select_policy_hub<cub::detail::value_t<InputIteratorT>, MayAlias>>
+struct DispatchSelectIf : SelectedPolicy
 {
     /******************************************************************************
      * Types and constants
@@ -147,35 +174,8 @@ struct DispatchSelectIf
     typedef ScanTileState<OffsetT> ScanTileStateT;
 
 
-    /******************************************************************************
-     * Tuning policies
-     ******************************************************************************/
-
-    /// SM35
-    struct Policy350
-    {
-        enum {
-            NOMINAL_4B_ITEMS_PER_THREAD = 10,
-            ITEMS_PER_THREAD            = CUB_MIN(NOMINAL_4B_ITEMS_PER_THREAD, CUB_MAX(1, (NOMINAL_4B_ITEMS_PER_THREAD * 4 / sizeof(InputT)))),
-        };
-
-        typedef AgentSelectIfPolicy<
-                128,
-                ITEMS_PER_THREAD,
-                BLOCK_LOAD_DIRECT,
-                MayAlias ? LOAD_CA : LOAD_LDG,
-                BLOCK_SCAN_WARP_SCANS>
-            SelectIfPolicyT;
-    };
-
-    /******************************************************************************
-     * Tuning policies of current PTX compiler pass
-     ******************************************************************************/
-
-    typedef Policy350 PtxPolicy;
-
     // "Opaque" policies (whose parameterizations aren't reflected in the type signature)
-    struct PtxSelectIfPolicyT : PtxPolicy::SelectIfPolicyT {};
+    struct PtxSelectIfPolicyT : SelectedPolicy::SelectIfPolicyT {};
 
 
     /******************************************************************************