Add jitter

gevtushenko · May 25, 2023 · 2ed91b2 · 2ed91b2
1 parent 837855f
commit 2ed91b2
Showing 1 changed file with 16 additions and 1 deletion.
diff --git a/cub/agent/single_pass_scan_operators.cuh b/cub/agent/single_pass_scan_operators.cuh
@@ -154,6 +154,20 @@ __device__ __forceinline__ void delay_or_prevent_hoisting(int ns)
                ((void)ns; __threadfence_block();));
 }
 
+struct rng_t
+{
+    static constexpr unsigned int a = 16807; 
+    static constexpr unsigned int c = 0; 
+    static constexpr unsigned int m = 1u << 31;
+
+    unsigned int seed;
+
+    __device__ unsigned int operator()(unsigned int min, unsigned int max)
+    {
+        return (seed = (a * seed + c) % m) % (max + 1 - min) + min;
+    }
+};
+
 }
 
 /**
@@ -314,11 +328,12 @@ struct ScanTileState<T, true>
           tile_descriptor = reinterpret_cast<TileDescriptor&>(alias);
         }
 
+        detail::rng_t rng{static_cast<unsigned int>(tile_idx)};
         int delay = CUB_DETAIL_INITIAL_L2_BACKOFF_NS;
 
         while (WARP_ANY((tile_descriptor.status == SCAN_TILE_INVALID), 0xffffffff))
         {   
-          detail::delay_or_prevent_hoisting(delay <<= 1);
+          detail::delay_or_prevent_hoisting(rng(0, delay <<= 1));
           TxnWord alias = detail::load_relaxed(d_tile_descriptors + TILE_STATUS_PADDING + tile_idx);
           tile_descriptor = reinterpret_cast<TileDescriptor&>(alias);
         }