Move __syncthreads after thread copy

spcl · Dec 12, 2023 · f5cd14b · f5cd14b
1 parent 8b99971
commit f5cd14b
Showing 1 changed file with 9 additions and 9 deletions.
diff --git a/dace/runtime/include/dace/cuda/copy.cuh b/dace/runtime/include/dace/cuda/copy.cuh
@@ -743,9 +743,6 @@ namespace dace
 
         static DACE_DFI void Copy(const T *smem, int src_xstride, T *ptr, int dst_xstride)
         {
-           if (!ASYNC)
-                __syncthreads();
-
             // Linear thread ID
             int ltid = GetLinearTID<BLOCK_WIDTH, BLOCK_HEIGHT, BLOCK_DEPTH>();
 
@@ -759,14 +756,14 @@ namespace dace
                 *(ptr + (ltid + WRITES*BLOCK_SIZE)* dst_xstride) =
                     *(smem + (ltid + WRITES * BLOCK_SIZE) * src_xstride);
             }
+
+            if (!ASYNC)
+                __syncthreads();
         }
 
         template <typename WCR>
         static DACE_DFI void Accum(const T *smem, int src_xstride, T *ptr, int dst_xstride, WCR wcr)
         {
-            if (!ASYNC)
-                __syncthreads();
-
             // Linear thread ID
             int ltid = GetLinearTID<BLOCK_WIDTH, BLOCK_HEIGHT, BLOCK_DEPTH>();
 
@@ -783,14 +780,14 @@ namespace dace
                         ptr + (ltid + WRITES * BLOCK_SIZE)* dst_xstride,
                         *(smem + (ltid + WRITES * BLOCK_SIZE) * src_xstride));
             }
+
+            if (!ASYNC)
+                __syncthreads();
         }
 
         template <ReductionType REDTYPE>
         static DACE_DFI void Accum(const T *smem, int src_xstride, T *ptr, int dst_xstride)
         {
-            if (!ASYNC)
-                __syncthreads();
-
             // Linear thread ID
             int ltid = GetLinearTID<BLOCK_WIDTH, BLOCK_HEIGHT, BLOCK_DEPTH>();
 
@@ -807,6 +804,9 @@ namespace dace
                         ptr + (ltid + WRITES*BLOCK_SIZE)* dst_xstride,
                         *(smem + (ltid + WRITES * BLOCK_SIZE) * src_xstride));
             }
+
+            if (!ASYNC)
+                __syncthreads();
         }
     };