From 9a07e28ecb4351ff29a02e2aff09be84620b559b Mon Sep 17 00:00:00 2001
From: Samyak Gangwal <sam.gangwal97@gmail.com>
Date: Fri, 14 Jun 2024 14:11:01 -0400
Subject: [PATCH] Fixed data races in aidw-cuda, mixbench-cuda, word2vec-cuda

---
 src/aidw-cuda/main.cu     | 3 +++
 src/mixbench-cuda/main.cu | 4 ++++
 src/word2vec-cuda/cbow.cu | 1 +
 3 files changed, 8 insertions(+)

diff --git a/src/aidw-cuda/main.cu b/src/aidw-cuda/main.cu
index 23f81a678..ec52cc324 100644
--- a/src/aidw-cuda/main.cu
+++ b/src/aidw-cuda/main.cu
@@ -143,6 +143,9 @@ void AIDW_Kernel_Tiled(
       dist = (six_s * six_s + siy_s * siy_s);
       t = 1.f / powf(dist, alpha);  sum_dn += t;  sum_up += t * sdz[e];
     }
+  
+  __syncthreads();
+
   }
   iz[tid] = sum_up / sum_dn;
 }
diff --git a/src/mixbench-cuda/main.cu b/src/mixbench-cuda/main.cu
index 7bbc50aee..d3c89d2cd 100644
--- a/src/mixbench-cuda/main.cu
+++ b/src/mixbench-cuda/main.cu
@@ -35,6 +35,8 @@ __global__ void benchmark_func(float *g_data, const int blockdim,
         tmps[j] = tmps[j]*tmps[j]+seed;
     }
 
+    __syncthreads();
+
     // Multiply add reduction
     float sum = 0.f;
     #pragma unroll
@@ -44,6 +46,8 @@ __global__ void benchmark_func(float *g_data, const int blockdim,
     #pragma unroll
     for(int j=0; j<granularity; j++)
       g_data[idx+k*big_stride] = sum;
+
+    __syncthreads();
   }
 }
 
diff --git a/src/word2vec-cuda/cbow.cu b/src/word2vec-cuda/cbow.cu
index ce8a1295f..eea5bd49b 100644
--- a/src/word2vec-cuda/cbow.cu
+++ b/src/word2vec-cuda/cbow.cu
@@ -167,6 +167,7 @@ void device_cbow(
           }
       }
     }// End for sentence_idx
+    __syncthreads();
 
     // Update d_random
     if (idInWarp == 0 ) d_random[sentence_position] = next_random;