title: Different results each run in CUDA
body: <p>I'm programming in cuda the next:</p>
  
  <ol>
  <li>I get a matrix with image values (d_Data)</li>
  <li>I copy the matrix in shared memory (tile)</li>
  <li>I get a basic pixel difference (pixel(i)-pixel(i+1)) (d_diff)</li>
  <li>If difference is a specific value for example 0, A number 1 is set i a matrix (d_A) in each position where the number 0 was found in d_diff. This is to be able to get the frequency of 0's in the original differences matrix.</li>
  <li>Parallel cumulative sum.</li>
  <li>Frequency result go to the frequencies vector.</li>
  </ol>
  
  <p>Checking step by step, every thing goes as expected until the cumulative sum. When I launch the code the software-calculated value is 104347 but sometimes from CUDA I get a nan result, other times I get any number for example 2425. The very strange is if I persists running the kernel 20 or 30 times, the value becomes the expected 104347 :S.</p>
  
  <p>I'm using for each matrix:</p>
  
  <pre><code>h_Data  = (float *)malloc(data_size);
  h_diff  = (float *)malloc(data_size);
  h_A         = (float *)malloc(data_size);
  </code></pre>
  
  <p>and </p>
  
  <pre><code> cudaFree(d_A);
  cudaFree(d_diff);
  cudaFree(d_Av);
  </code></pre>
  
  <p>so I don't understand why the code is closer and closer to the correct result when i run enough times. By the way, when the correct value is reached, it doesn't move any more no matter how many times i run the code.</p>
  
  <p>The code: </p>
  
  <pre><code> __global__ void spam(float *d_Data, float *d_diff, float *d_A, int dw, int dh, float *d_Av){
  
  long bx = blockIdx.x;  long by = blockIdx.y;
  long tx = threadIdx.x; long ty = threadIdx.y;
  
  
  // Identify the row and column of the Pd element to work on
  long Row = by * TILE_WIDTH + ty;
  long Col = bx * TILE_WIDTH + tx;
  long tid = Row*dw+Col;
  long i=512*512;
  long r = MASK_DIM/2;
  long s = 0;
  
   __shared__ int tile[BLOCK_WIDTH][BLOCK_WIDTH];
  
  for (int k=0; k&lt;=8; k++)
       d_Av[k]=0; 
  
  
      if(tid &lt; dw*dh)
      {
  
     // to shared memory.
                                            tile[ty + r][tx + r]=d_Data[Row*dw+Col];
          if (Col-r &gt;=0)                    tile[ty + r]  [tx] = d_Data[Row*dw+Col-r];
          if (Col+r &lt;dw)                    tile[ty + r]  [tx + 2*r] = d_Data[Row*dw+Col+r];
          if (Row-r &gt;=0)                    tile[ty]      [tx + r] = d_Data[(Row - r)*dw + Col];
          if (Row+r &lt;dw)                    tile[ty + 2*r][tx + r] = d_Data[(Row + r)*dw + Col];
          if (Row - r &gt;= 0 &amp;&amp; Col - r &gt;= 0) tile[ty]      [tx] = d_Data[(Row-r)*dw+Col-r];
          if(Row - r &gt;= 0 &amp;&amp; Col + r &lt; dw)  tile[ty]      [tx + 2*r] = d_Data[(Row-r)*dw+Col+r];
          if (Row + r &lt; dw &amp;&amp; Col - r &gt;= 0) tile[ty + 2*r][tx] = d_Data[(Row+r)*dw+Col-r];
          if(Row + r &lt;dw &amp;&amp; Col + r &lt; dw)   tile[ty + 2*r][tx + 2*r] = d_Data[(Row-r)*dw+Col+r]; 
  
          //Calculates the difference matrix
         d_diff[tid] = (tile[ty + r][tx +r] - tile[ty + r][tx + r + 1]);
  
  
          d_A[tid]=0;
  
         //Set a 1 in each position in d_A where 0 was found in d_diff.
          if (d_diff[tid] == 0)
          { d_A[tid]=1;}
          __syncthreads();
  
          //cumulative sum to get the frecuency of value 0 in d_diff.  // The error is HERE
        for (s = (i/2); s&gt;=1; s=s/2) {
              if (tid &lt; s)
              {   d_A[tid] += d_A[tid+s];
              }
          } 
  
         // set the frequency value in frequencies vector.
          d_Av[0] = d_A[0];
  
  }} // END IF tid &lt; dw*dh
  </code></pre>
  
  <p>Any idea is welcome :D</p>