title: pycuda shared memory error "pycuda._driver.LogicError: cuLaunchKernel failed: invalid value"
body: <p>I have a strange problem which origin I cannot determine:</p>
  
  <p>I have a working Kernel for some special Matrix-Vector-multiplication, which I want to speed up. Basically the big matrix (10^6 times 10^6) is constructed from few small matrices. So I want to put that data in shared memory. However when I try to add the shared memory, I only get the error:</p>
  
  <p>pycuda._driver.LogicError: cuLaunchKernel failed: invalid value</p>
  
  <p>So my working kernel is:</p>
  
  <pre><code>#define FIELD_SIZE {field}
  #define BLOCK_SIZE {block}
  
  __global__ void MatrixMulKernel(double *gpu_matrix, double *gpu_b, double *gpu_y)
  {
      int tx = ... + threadIdx.x;
  
      if(tx &lt; FIELD_SIZE*FIELD_SIZE*BLOCK_SIZE) 
      { ... multiplication ... }
  }
  </code></pre>
  
  <p>And if I try to add the shared memory part it looks like</p>
  
  <pre><code>#define FIELD_SIZE {field}
  #define BLOCK_SIZE {block}
  
  __global__ void MatrixMulKernel(double *gpu_matrix_ptr, double *gpu_b, double *gpu_y)
  {
      __shared__ double gpu_matrix[BLOCK_SIZE*BLOCK_SIZE*13];
  
      int tx = ... + threadIdx.x;
      if(tx &lt; BLOCK_SIZE*BLOCK_SIZE*13) {  gpu_matrix[tx] = gpu_matrix_ptr[tx];  }
      __syncthreads();
  
      if(tx &lt; FIELD_SIZE*FIELD_SIZE*BLOCK_SIZE) 
      { ... multiplication ... }
  }
  </code></pre>
  
  <p>This is the only part I changed, so basically it has to be the gpu_matrix[tx] = gpu_matrix_ptr[tx] statement, hasnt it? But I fail to see how that should be. I basically tried to copy the tiled matrix-multiplication example from the pycuda examples. <a href="http://wiki.tiker.net/PyCuda/Examples/MatrixmulTiled" rel="nofollow">http://wiki.tiker.net/PyCuda/Examples/MatrixmulTiled</a></p>
  
  <p>The invocation is:</p>
  
  <pre><code>self.kernel.prepare([np.intp, np.intp, np.intp])
  self.kernel.prepared_call(grid_shape,
                block_shape,
                self.matrix_gpu.gpudata,
                b_gpu.gpudata,
                y_gpu.gpudata)
  </code></pre>
  
  <p>where matrix_gpu, b_gpu and y_gpu are pycuda.gpuarray instances. </p>
  
  <p>Hope that you can clear up some of my confusion...</p>