Hope that's it

gevtushenko · May 10, 2023 · a4a4d99 · a4a4d99
1 parent 0a50907
commit a4a4d99
Show file tree

Hide file tree

Showing 2 changed files with 2 additions and 2 deletions.
diff --git a/cub/agent/agent_scan.cuh b/cub/agent/agent_scan.cuh
@@ -357,8 +357,8 @@ struct AgentScan
     }
 
     // Wait for all threads in the cluster to finish loading / dsmem initialization
+    CTA_SYNC();
     cooperative_groups::cluster_group::barrier_wait(std::move(token));
-    __threadfence();
 
     // Perform tile scan
     if (tile_idx == 0)

diff --git a/cub/agent/single_pass_scan_operators.cuh b/cub/agent/single_pass_scan_operators.cuh
@@ -714,7 +714,7 @@ struct ClusterTilePrefixCallbackOp
         descriptor->status = StatusWord(SCAN_TILE_INVALID);
 
         if (threadIdx.x < CUB_DETAIL_CLUSTER_SIZE) {
-            lsmem_st_relaxed(temp_storage.dsmem + threadIdx.x, val);
+            temp_storage.dsmem[threadIdx.x] = val;
         }
     }
 };