From a4a4d99d6625c8e18b33d62a30327ed96286358e Mon Sep 17 00:00:00 2001
From: Georgy Evtushenko <evtushenko.georgy@gmail.com>
Date: Wed, 10 May 2023 16:31:08 +0400
Subject: [PATCH] Hope that's it

---
 cub/agent/agent_scan.cuh                 | 2 +-
 cub/agent/single_pass_scan_operators.cuh | 2 +-
 2 files changed, 2 insertions(+), 2 deletions(-)

diff --git a/cub/agent/agent_scan.cuh b/cub/agent/agent_scan.cuh
index 82b29912c..0ba6ca845 100644
--- 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
index 07ab66e90..52aaa9843 100644
--- 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;
         }
     }
 };