title: Using maximum shared memory in Cuda
body: <p>I am unable to use more than 48K of shared memory (on V100, Cuda 10.2)</p>
  <p>I call</p>
  <pre><code>cudaFuncSetAttribute(my_kernel,
                       cudaFuncAttributePreferredSharedMemoryCarveout,
                       cudaSharedmemCarveoutMaxShared);
  </code></pre>
  <p>before launching <code>my_kernel</code> first time.</p>
  <p>I use launch bounds
  and dynamic shared memory inside <code>my_kernel</code>:</p>
  <pre><code>__global__
  void __launch_bounds__(768, 1)
  my_kernel(...)
  {
      extern __shared__ float2 sh[];
      ...
  }
  </code></pre>
  <p>Kernel is called like this:</p>
  <pre><code>dim3 blk(32, 24); // 768 threads as in launch_bounds.
  
  my_kernel&lt;&lt;&lt;grd, blk, 64 * 1024, my_stream&gt;&gt;&gt;( ... );
  </code></pre>
  <p><code>cudaGetLastError()</code> after kernel call returns <code>cudaErrorInvalidValue</code>.</p>
  <p>If I use &lt;= 48 K of shared memory (e.g., <code>my_kernel&lt;&lt;&lt;grd, blk, 48 * 1024, my_stream&gt;&gt;&gt;</code>), it works.</p>
  <p>Compilation flags are:</p>
  <p><code>nvcc -std=c++14 -gencode arch=compute_70,code=sm_70 -Xptxas -v,-dlcm=cg</code></p>
  <p>What am I missing?</p>