Fix reduce by key tile state for Pascal

gevtushenko · Jun 13, 2023 · 7040b5f · 7040b5f
1 parent 01f8003
commit 7040b5f
Showing 1 changed file with 7 additions and 2 deletions.
diff --git a/cub/agent/single_pass_scan_operators.cuh b/cub/agent/single_pass_scan_operators.cuh
@@ -194,7 +194,12 @@ struct no_delay_constructor_t
 {
   struct delay_t
   {
-    __device__ __forceinline__ void operator()() {}
+    __device__ __forceinline__ void operator()() 
+    {
+      NV_IF_TARGET(NV_PROVIDES_SM_70,
+                  (),
+                  (__threadfence_block();));
+    }
   };
 
   __device__ __forceinline__ no_delay_constructor_t(unsigned int /* seed */)
@@ -997,7 +1002,7 @@ struct ReduceByKeyScanTileState<ValueT, KeyT, true>
     /**
      * Wait for the corresponding tile to become non-invalid
      */
-    template <class DelayT = detail::fixed_delay_constructor_t<350, 450>> 
+    template <class DelayT = detail::fixed_delay_constructor_t<350, 450>::delay_t> 
     __device__ __forceinline__ void WaitForValid(
         int                     tile_idx,
         StatusWord              &status,