Fix

gevtushenko · May 9, 2023 · dee9014 · dee9014
1 parent 9376f3f
commit dee9014
Showing 1 changed file with 12 additions and 6 deletions.
diff --git a/cub/agent/single_pass_scan_operators.cuh b/cub/agent/single_pass_scan_operators.cuh
@@ -555,10 +555,13 @@ struct ClusterTilePrefixCallbackOp
               {
                 TileDescriptor tile_descriptor;
 
-                do
+                if (src_cta < cta_rank)
                 {
-                  LoadTileDescriptor(src_cta, tile_descriptor);
-                } while (tile_descriptor.status == SCAN_TILE_INVALID);
+                  do
+                  {
+                    LoadTileDescriptor(src_cta, tile_descriptor);
+                  } while (tile_descriptor.status == SCAN_TILE_INVALID);
+                }
                 exclusive_prefix = Reduce(cta_rank, src_cta, tile_descriptor.value);
 
                 if (__shfl_sync(CUB_DETAIL_CLUSTER_WARP_MASK,
@@ -586,10 +589,13 @@ struct ClusterTilePrefixCallbackOp
               {
                 TileDescriptor tile_descriptor;
                 StatusWord target_status = src_cta == 0 ? SCAN_TILE_INCLUSIVE : SCAN_TILE_PARTIAL;
-                do
+                if (src_cta < cta_rank)
                 {
-                  LoadTileDescriptor(src_cta, tile_descriptor);
-                } while (tile_descriptor.status != target_status);
+                  do
+                  {
+                    LoadTileDescriptor(src_cta, tile_descriptor);
+                  } while (tile_descriptor.status != target_status);
+                }
                 exclusive_prefix = Reduce(cta_rank, src_cta, tile_descriptor.value);
               }
             }