title: CUDA: use of shared memory to return an array from device functions
body: <p>is it possible to allocate shared memory for a kernel (inside or extern) and use it in other device functions called from the kernel?
  Specially interesting for me will be, if/how i can use it as a returned parameter/array.</p>
  
  <p>It seems to be no problem to use shared memory as input parameter in device functions (at least i get no problems, errors or unexpected results.</p>
  
  <p>When I use it as a return parameter, I get several problems:</p>
  
  <ul>
  <li><p>I can run the program when it was built from debug configuration.</p></li>
  <li><p>But i can't debug it -> it crashes in the device functions when i use the shared memory</p></li>
  <li><p>Also i get errors with <code>cuda-memchecker</code> -> invalid <code>__global__</code> read
  because address is out of bound an it read from shared address space</p></li>
  </ul>
  
  <p>So is it possible to use shared memory for returning arrays from device functions to kernels?</p>
  
  <p>EDIT:</p>
  
  <p>I wrote a very simple example to exclude other errors done by me.</p>
  
  <pre><code>#define CUDA_CHECK_RETURN(value) {                                      \
      cudaError_t _m_cudaStat = (value);                                  \
      if (_m_cudaStat != cudaSuccess) {                                   \
          printf( "Error %s at line %d in file %s\n",                     \
                  cudaGetErrorString(_m_cudaStat), __LINE__, __FILE__);   \
          exit(-1);                                                       \
      } }
  
  __device__ void Function( const int *aInput, volatile int *aOutput )
  {
      for( int i = 0; i &lt; 10; i++ )
          aOutput[i] = aInput[i] * aInput[i];
  }
  
  __global__ void Kernel( int *aInOut )
  {
       __shared__ int aShared[10];
  
      for(int i=0; i&lt;10; i++)
          aShared[i] = i+1;
  
      Function( aShared, aInOut );
  }
  
  int main( int argc, char** argv )
  {
      int *hArray = NULL;
      int *dArray = NULL;
  
      hArray = ( int* )malloc( 10*sizeof(int) );
      CUDA_CHECK_RETURN( cudaMalloc( (void**)&amp;dArray, 10*sizeof(int) ) );
  
      for( int i = 0; i &lt; 10; i++ )
              hArray[i] = i+1;
  
      CUDA_CHECK_RETURN( cudaMemcpy( dArray, hArray, 10*sizeof(int), cudaMemcpyHostToDevice ) );
      cudaMemcpy( dArray, hArray, 10*sizeof(int), cudaMemcpyHostToDevice );
  
      Kernel&lt;&lt;&lt;1,1&gt;&gt;&gt;( dArray );
  
      CUDA_CHECK_RETURN( cudaMemcpy( hArray, dArray, 10*sizeof(int), cudaMemcpyDeviceToHost ) );
      cudaMemcpy( hArray, dArray, 10*sizeof(int), cudaMemcpyDeviceToHost );
  
      free( hArray );
      CUDA_CHECK_RETURN( cudaFree( dArray ) );
      cudaFree( dArray );
  
      return 0;
  }
  </code></pre>
  
  <p>I excecute the kernel by one threadblock and one thread per block. It's no problem to build the program and run it. I get the expected results.
  But if the program is testet with cuda-memchecker it terminates the kernel and following log appears.</p>
  
  <pre><code>Error unspecified launch failure at line 49 in file ../CuTest.cu
  ========= Invalid __global__ read of size 4
  =========     at 0x00000078 in /home/strautz/Develop/Software/CuTest/Debug/../CuTest.cu:14:Function(int const *, int volatile *)
  =========     by thread (0,0,0) in block (0,0,0)
  =========     Address 0x01000000 is out of bounds
  =========     Device Frame:/home/strautz/Develop/Software/CuTest/Debug/../CuTest.cu:25:Kernel(int*) (Kernel(int*) : 0xd0)
  =========     Saved host backtrace up to driver entry point at kernel launch time
  =========     Host Frame:/usr/lib/libcuda.so (cuLaunchKernel + 0x34b) [0x55d0b]
  =========     Host Frame:/usr/lib/libcudart.so.5.0 [0x8f6a]
  =========
  ========= Program hit error 4 on CUDA API call to cudaMemcpy 
  =========     Saved host backtrace up to driver entry point at error
  =========     Host Frame:/usr/lib/libcuda.so [0x24e129]
  =========     Host Frame:/usr/lib/libcudart.so.5.0 (cudaMemcpy + 0x2bc) [0x3772c]
  =========     Host Frame:[0x5400000]
  =========
  ========= ERROR SUMMARY: 2 errors
  </code></pre>
  
  <p>Does the shared memory have to be aligned, do I have to do something else or can it be ignored - don't think so?</p>