# \_\_syncthreads screwes calculation

**URL:** <https://forums.developer.nvidia.com/t/syncthreads-screwes-calculation/1894>\
**Category:** CUDA Programming and Performance\
**Created:** [November 22, 2007, 12:55pm UTC](https://forums.developer.nvidia.com/t/syncthreads-screwes-calculation/1894 "2007-11-22T12:55:05Z")\
**Posts on this page:** 3\
**Page:** 1

<div class="post-metadata">

**Author:** ![axolotl](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@axolotl](https://forums.developer.nvidia.com/u/axolotl)\
**Post date:** [November 22, 2007, 12:55pm UTC](https://forums.developer.nvidia.com/t/syncthreads-screwes-calculation/1894/1 "2007-11-22T12:55:05Z")

</div>

Hi,

I am using CUDA to develop a monte-carlo code. GPU kernel basically does pretty uniform caclulations and works fine, though slow. While improving performance I was trying different things and found out that adding \_\_syncthreads in the middle of the code leads to incorrect caclulation, with non-predictable outcome (though within certain range). Yes, it works fine without syncthreads and silently produces wrong result with it. There is no aliasing, I just want to load numbers into _shared_ before applying calculation to them

code snippet: s is grid.x, c is grid.y, H/U are _shared_ cuda\_flink is cuda\_array

{ n = neighbour\_table[8 \* i + dir]; { U[s\*3+c+16\*0] = cuda\_flink[8 \* i + dir].e  
[c][0].real; U[s\*3+c+16\*1] = cuda\_flink[8 \* i + dir].e[c][0].imag; U[s\*3+c+16\*2] =  
cuda\_flink[8 \* i + dir].e[c][1].real; U[s\*3+c+16\*3] = cuda\_flink[8 \* i + dir].e[c][  
1].imag; U[s\*3+c+16\*4] = cuda\_flink[8 \* i + dir].e[c][2].real; U[s\*3+c+16\*5] = cuda  
\_flink[8 \* i + dir].e[c][2].imag; }; };  
\_\_syncthreads();

```
H[s].h[0].c[c].real =src[n].c[c].d[0].real-src[n].c[c].d[3].imag;
H[s].h[0].c[c].imag =src[n].c[c].d[0].imag+src[n].c[c].d[3].real;
H[s].h[1].c[c].real =src[n].c[c].d[1].real-src[n].c[c].d[2].imag;
H[s].h[1].c[c].imag =src[n].c[c].d[1].imag+src[n].c[c].d[2].real;

```

usually you screw things up when forgetting syncthreads but not by adding.  
cuda 1.0 and 1.1 beta give same results.

any ideas?

thanks!

---

<div class="post-metadata">

**Author:** ![DenisR](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@DenisR](https://forums.developer.nvidia.com/u/DenisR)\
**Post date:** [November 22, 2007, 1:23pm UTC](https://forums.developer.nvidia.com/t/syncthreads-screwes-calculation/1894/2 "2007-11-22T13:23:51Z")

</div>

This is easier to read :

code snippet: s is grid.x, c is grid.y, H/U are _shared_ cuda\_flink is cuda\_array

```auto
{

 n = neighbour_table[8 * i + dir];

 {

  U[s*3+c+16*0] = cuda_flink[8 * i + dir].e[c][0].real; 

  U[s*3+c+16*1] = cuda_flink[8 * i + dir].e[c][0].imag;

  U[s*3+c+16*2] =cuda_flink[8 * i + dir].e[c][1].real;

  U[s*3+c+16*3] = cuda_flink[8 * i + dir].e[c][1].imag;

  U[s*3+c+16*4] = cuda_flink[8 * i + dir].e[c][2].real;

  U[s*3+c+16*5] = cuda_flink[8 * i + dir].e[c][2].imag;

 };

};

__syncthreads();

    H[s].h[0].c[c].real =src[n].c[c].d[0].real-src[n].c[c].d[3].imag;

    H[s].h[0].c[c].imag =src[n].c[c].d[0].imag+src[n].c[c].d[3].real;

    H[s].h[1].c[c].real =src[n].c[c].d[1].real-src[n].c[c].d[2].imag;

    H[s].h[1].c[c].imag =src[n].c[c].d[1].imag+src[n].c[c].d[2].real;

```

---

<div class="post-metadata">

**Author:** ![jordyvaneijk](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@jordyvaneijk](https://forums.developer.nvidia.com/u/jordyvaneijk)\
**Post date:** [November 22, 2007, 2:03pm UTC](https://forums.developer.nvidia.com/t/syncthreads-screwes-calculation/1894/3 "2007-11-22T14:03:38Z")

</div>

> [@](#):
>
> …
> 
> [snapback]283013[/snapback]

I Realy can’t tell you whats wrong but maybe it’s better next time to post only “one” thread instead of “three” ;)