Not that good

gevtushenko · Jun 7, 2023 · 3534daa · 3534daa
1 parent 0cf1181
commit 3534daa
Showing 1 changed file with 7 additions and 13 deletions.
diff --git a/cub/device/dispatch/dispatch_scan.cuh b/cub/device/dispatch/dispatch_scan.cuh
@@ -214,7 +214,7 @@ struct sm90_tuning
   static constexpr int threads = 128;
   static constexpr int items   = 15;
 
-  using delay_constructor = detail::exponential_backoff_constructor_t<32, 1140>;
+  using delay_constructor = detail::default_delay_constructor_t<AccumT>;
 };
 
 template <class AccumT>
@@ -235,28 +235,22 @@ struct sm90_tuning<AccumT, 2, true>
   using delay_constructor = detail::fixed_delay_constructor_t<376, 1125>;
 };
 
-template <class AccumT, bool Primitive>
-struct sm90_tuning<AccumT, 4, Primitive>
+template <class AccumT>
+struct sm90_tuning<AccumT, 4, true>
 {
   static constexpr int threads = 128;
   static constexpr int items   = 24;
 
-  using delay_constructor =
-    cub::detail::conditional_t<Primitive,
-                               detail::fixed_delay_constructor_t<648, 1245>,
-                               detail::exponential_backoff_constructor_t<64, 1245>>;
+  using delay_constructor = detail::fixed_delay_constructor_t<648, 1245>;
 };
 
-template <class AccumT, bool Primitive>
-struct sm90_tuning<AccumT, 8, Primitive>
+template <class AccumT>
+struct sm90_tuning<AccumT, 8, true>
 {
   static constexpr int threads = 224;
   static constexpr int items   = 24;
 
-  using delay_constructor =
-    cub::detail::conditional_t<Primitive,
-                               detail::fixed_delay_constructor_t<632, 1290>,
-                               detail::exponential_backoff_constructor_t<64, 1290>>;
+  using delay_constructor = detail::fixed_delay_constructor_t<632, 1290>;
 };
 
 #if CUB_IS_INT128_ENABLED