Strange

gevtushenko · May 26, 2023 · 35a752c · 35a752c
1 parent 2ed91b2
commit 35a752c
Showing 1 changed file with 12 additions and 1 deletion.
diff --git 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>