title: Cuda: data corruption in memory after kernel call
body: <p>Please look this code:</p>
  
  <pre><code>#include &lt;stdlib.h&gt;
  #include &lt;stdio.h&gt;
  
  int N, L, I;
  float * inputs;
  float * temp;
  
  // first kernel
  __global__ void mulKernel ( float * output, float * inputs)///, float * weights)
  {
     int idx = blockIdx.x * blockDim.x + threadIdx.x;
  
     output [idx] = inputs [idx] * 3;//weights [idx];
     //weights [idx] = 4;
  
     //__syncthreads();
  }
  
  //second kernel
  __global__ void sumKernel ( float * output, float * input)
  {
     int idx = blockIdx.x * blockDim.x + threadIdx.x;
     output [idx] = input[idx]*2;
  
     __syncthreads();
  }
  
  void printVector (const float *p, const int N) {
      for (int i=0; i&lt;N; i++)
      printf("%f\n",p[i]);
  }
  
  int main(int argc, char *argv[])
  {
      if(argc &lt; 3)
          printf("Usage: cuda &lt;layers&gt; &lt;inputs&gt;\n");
      else
      {
          L = atoi(argv[1]);
          N = atoi(argv[2]);
          I = atoi(argv[2]);
          inputs = (float*)malloc(I*sizeof(float));
          float * weights = (float*)malloc(I*sizeof(float));
  
          // and fill with some arbitrary values
          for (int i=0; i&lt;I; i++)
          {
              inputs[i] = 1;
          }
          for (int i=0; i&lt;I; i++)
          {
              weights[i] = 1.5;
          }
  
          // allocate device memory
          float * devInputs = NULL;
          float * devTemp = NULL;
          float * devWeights = NULL;
  
          cudaMalloc ( (void**)&amp;devInputs, I*sizeof(float) );
          cudaMalloc ( (void**)&amp;devTemp, I*sizeof(float) );
          cudaMalloc ( (void**)&amp;devWeights, I*sizeof(float) );
  
          // set kernel launch configuration
          dim3 threadsMul = dim3(512, 1);
          int blocksCount = floor(I / threadsMul.x) + 1;
          dim3 blocksMul  = dim3(blocksCount, 1);
  
          dim3 threadsSum = dim3(512, 1);
          blocksCount = floor(I / threadsSum.x) + 1;
          dim3 blocksSum  = dim3(blocksCount, 1);
  
          cudaMemcpy      ( devInputs, inputs, I*sizeof(float), cudaMemcpyHostToDevice );
          cudaMemcpy      ( devWeights, weights,I*sizeof(float), cudaMemcpyHostToDevice );
  
          //kernels calling in this cycle
          for(int j=0;j&lt;L;j++)
          {
              // copying data to see that's ok
            cudaMemcpy      ( inputs, devInputs, I*sizeof(float), cudaMemcpyDeviceToHost );
            cudaMemcpy      ( weights, devWeights, I*sizeof(float), cudaMemcpyDeviceToHost );
  
              // print it
            printf("inputs:\n");
            printVector (inputs, N);
            printf("weights:\n");
            printVector (weights, N);
            printf("\n");
  
              // running first kernel
            mulKernel&lt;&lt;&lt;blocksMul, threadsMul&gt;&gt;&gt;(devTemp, devInputs);//, devWeights);
  
              // copying and printing data. We can see thats array weights contains a wrong values
            cudaMemcpy      ( inputs, devInputs, I*sizeof(float), cudaMemcpyDeviceToHost );
            cudaMemcpy      ( weights, devWeights, I*sizeof(float), cudaMemcpyDeviceToHost );
  
            printf("inputs:\n");
            printVector (inputs, N);
            printf("weights:\n");
            printVector (weights, N);
            printf("\n");
  
            if(cudaDeviceSynchronize() == cudaSuccess)
              printf("threads syncronized\n");
  
            cudaMemcpy      ( inputs, devInputs, I*sizeof(float), cudaMemcpyDeviceToHost );
            cudaMemcpy      ( weights, devWeights, I*sizeof(float), cudaMemcpyDeviceToHost );
  
            printf("inputs:\n");
            printVector (inputs, N);
            printf("weights:\n");
            printVector (weights, N);
            printf("\n");
  
            sumKernel&lt;&lt;&lt;blocksSum, threadsSum&gt;&gt;&gt;(devInputs, devTemp);
  
            cudaMemcpy      ( inputs, devInputs, I*sizeof(float), cudaMemcpyDeviceToHost );
            cudaMemcpy      ( weights, devWeights, I*sizeof(float), cudaMemcpyDeviceToHost );
  
            printf("inputs:\n");
            printVector (inputs, N);
            printf("weights:\n");
            printVector (weights, N);
            printf("\n\n");
  
            if(cudaDeviceSynchronize() == cudaSuccess)
              printf("threads syncronized\n");
  
            cudaMemcpy      ( inputs, devInputs, I*sizeof(float), cudaMemcpyDeviceToHost );
            cudaMemcpy      ( weights, devWeights, I*sizeof(float), cudaMemcpyDeviceToHost );
  
            printf("inputs:\n");
            printVector (inputs, N);
            printf("weights:\n");
            printVector (weights, N);
            printf("\n\n");
          }
  
          cudaMemcpy      ( inputs, devInputs, I*sizeof(float), cudaMemcpyDeviceToHost );
  
          cudaFree         ( devInputs   );
          cudaFree         ( devTemp   );
          cudaFree         ( devWeights   );
  
          printVector (inputs, N);
  
          free(inputs);
          free(weights);
      }
      return 0;
  }
  </code></pre>
  
  <p>And look the output. After callig first kernel, the devWeights array lost its data. But it doesn't used anywhere. I just copy it to the memory, run kernels (that don't affects it) and copy back to host. And in output I see that it changed. Why? What am I doing wrong?</p>
  
  <p>In main function you can see cycle for. In it I run two kernels: sumKernel and mulKernel. Before running kernel, after it, and after synchronization threads I copy arrays to host and print it. So, I see wrong data after calling kernel. See comments in code.</p>
  
  <p>I don't see any error (only cudaSuccess).</p>