title: is it possible to define a Global memory array at CUDA device?
body: <p>I tr to allocate a cuda global memory array. I have the kernel of summation as:</p>
  
  <pre><code>__device__ float R_d = 0;
  __global__ void perform_summation(float* A, int N){
      int idx = blockDim.x*blockIdx.x+threadIdx.x;
      extern __shared__ float sharedArray [];
      float result[]; //THIS IS THE THING i TRIED TO CREATE
      if(idx &lt; N){
          sharedArray[threadIdx.x] = A[idx];
          //  }else{
          //      sharedArray[threadIdx.x] = 0 ;
          //  }
  
          for (unsigned int stride = 1; stride &lt; blockDim.x; stride *= 2) {
              __syncthreads();
              if(threadIdx.x % (2*stride) == 0){
                  sharedArray[threadIdx.x]+=sharedArray[threadIdx.x+stride];
              }
          }
      }
      if(idx % blockDim.x == 0){
  //      R_d += sharedArray[threadIdx.x];
          result[blockIdx.x] = sharedArray[threadIdx.x];
      }
      for (int i = 0; i &lt; gridDim.x; ++i) {
          R_d += result[i];
      }
  }
  </code></pre>
  
  <p>As summary y kernel takes an array and finds the summation of the elements by map reduce approach. Each blocks take the relative elements into shared memory and sums all the data inside than put the result to the global array that I try to create. At the final, I will sum all the numbers of the global array to find the last answer.</p>
  
  <p>As the first method I did not use global array to collect the results of each block and I just sum the results of the blocks into the variable <code>R_d</code> but it does not work and shows only the value comes from the last block as the result. I guess since I do not have sync. between blocks last block overwrites all the values at the end. Here is the what I've done at first attempt at the end of the kernel</p>
  
  <pre><code>f(idx &lt; N){
          sharedArray[threadIdx.x] = A[idx];
          //  }else{
          //      sharedArray[threadIdx.x] = 0 ;
          //  }
  
          for (unsigned int stride = 1; stride &lt; blockDim.x; stride *= 2) {
              __syncthreads();
              if(threadIdx.x % (2*stride) == 0){
                  sharedArray[threadIdx.x]+=sharedArray[threadIdx.x+stride];
              }
          }
  
          if(threadIdx.x == 0){
              R_d += sharedArray[threadIdx.x];
          }
      }
  </code></pre>
  
  <p>So I have actually two questions. How to define a global memory array at the device memory for the first solution I propose and is there any solution for the second solution that just uses the <code>R_d</code> variable?</p>