From 35a752c58076bf6842a504a49aad228df62ff043 Mon Sep 17 00:00:00 2001
From: Georgy Evtushenko <evtushenko.georgy@gmail.com>
Date: Fri, 26 May 2023 10:39:48 +0400
Subject: [PATCH] Strange

---
 cub/agent/single_pass_scan_operators.cuh | 13 ++++++++++++-
 1 file changed, 12 insertions(+), 1 deletion(-)

diff --git a/cub/agent/single_pass_scan_operators.cuh b/cub/agent/single_pass_scan_operators.cuh
index f461dd3f2..6b2135f1b 100644
--- a/cub/agent/single_pass_scan_operators.cuh
+++ b/cub/agent/single_pass_scan_operators.cuh
@@ -143,7 +143,18 @@ __device__ __forceinline__ void delay()
 template <unsigned int GridThreshold = 500>
 __device__ __forceinline__ void delay(int ns)
 {
-  NV_IF_TARGET(NV_PROVIDES_SM_70, (__nanosleep(ns);), ((void)ns;));
+  NV_IF_TARGET(NV_PROVIDES_SM_70,
+               (if (ns > 0) 
+                {
+                  if (gridDim.x < GridThreshold) 
+                  {
+                    __threadfence_block();
+                  }
+                  else 
+                  {
+                    __nanosleep(ns); 
+                  }
+                }));
 }
 
 template <unsigned int GridThreshold = 500>