# Using \_\_syncthreads() while using two shared arrays

**URL:** <https://forums.developer.nvidia.com/t/using-syncthreads-while-using-two-shared-arrays/210082>\
**Category:** CUDA Programming and Performance\
**Tags:** cuda\
**Created:** [April 1, 2022, 6:16pm UTC](https://forums.developer.nvidia.com/t/using-syncthreads-while-using-two-shared-arrays/210082 "2022-04-01T18:16:04Z")\
**Posts on this page:** 2\
**Page:** 1

<div class="post-metadata">

**Author:** ![sajal](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@sajal](https://forums.developer.nvidia.com/u/sajal)\
**Post date:** [April 1, 2022, 6:16pm UTC](https://forums.developer.nvidia.com/t/using-syncthreads-while-using-two-shared-arrays/210082/1 "2022-04-01T18:16:04Z")

</div>

My CUDA kernel looks somewhat like this. I am using two shared arrays. I am doing some sort of loop tiling inside the kernel. So, the first shared array needs to be updated with outer loop’s (block1) iteration, while the second shared array needs to be update with the iteration of the inner loop (block2).

However, it looks like I am not getting first shared array’s data right, while the second array’s data looks correct. I am using only \_\_syncthreads after updating the second array in the hope that every threads will need to wait till this point. Do I need to use the \_\_syncthreads() after the first array update (the line where I commented out the \_\_syncthreads()). Using this here slows down the whole process significantly.

What should I do?

```auto
extern __shared__ unsigned long long int shared_matrix[];
extern __shared__ unsigned long long int shared_matrix2[];

for(int block1 = block_index; block1 < num_blocks; block1++){
   int i3_start = block1 * 32;
   for(int i = 0; i < group_size; i++) {
       shared_matrix[localIdx * group_size + i] = global_matrix[(block1 * 32 + localIdx) * group_size + i];
   }
   //__syncthreads(); Should I need this line since I have a following syncthreads later?
   for(int block2 = block1; block2 < num_blocks; block2++) {
       int i4_start = block2 * 32;
       
       for(int i = 0; i < group_size; i++) {
           shared_matrix2[localIdx * group_size + i] = global_matrix[(block1 * 32 + localIdx) * group_size + i];
       }
       __syncthreads();
       
       for(int i3 = i3_start; i3 < i3_max; i3++){
           for(int i4 = i4_start; i4 < i4_max; i4++) {
               f(shared_matrix[g(i3)], shared_matrix2[g(i4)]);

           }
       }
    }
}

```

---

<div class="post-metadata">

**Author:** ![Robert\_Crovella](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/robert_crovella/32/14043_2.png) [@Robert\_Crovella](https://forums.developer.nvidia.com/u/Robert_Crovella)\
**Post date:** [April 1, 2022, 6:49pm UTC](https://forums.developer.nvidia.com/t/using-syncthreads-while-using-two-shared-arrays/210082/2 "2022-04-01T18:49:45Z")

</div>

- its frequently the case that with a loop around shared load/usage cycles, at least two `__syncthreads()` statements are needed. The reads and writes often need to be protected from each other, _as the loop proceeds_. Since you have multiple shared arrays and multiple sets of loops, its quite possible you might need more than two.
- the `compute-sanitizer --racecheck` [tool](https://docs.nvidia.com/compute-sanitizer/ComputeSanitizer/index.html) is very useful for this type of work. You might want to learn how to use it.
- I wouldn’t be able to make definitive statements based on incomplete snippets.