Experiment

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