Using maximum shared memory in Cuda

by

from here:

Compute capability 7.x devices allow a single thread block to address the full capacity of shared memory: 96 KB on Volta, 64 KB on Turing. Kernels relying on shared memory allocations over 48 KB per block are architecture-specific, as such they must use dynamic shared memory (rather than statically sized arrays) and require an explicit opt-in using cudaFuncSetAttribute() as follows:

cudaFuncSetAttribute(my_kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, 98304);

When I add that line to the code you have shown, the invalid value error goes away. For a Turing device, you would want to change that number from 98304 to 65536. And of course 65536 would be sufficient for your example as well, although not sufficient to use the maximum available on volta, as stated in the question title.

In a similar fashion kernels on Ampere devices should be able to use up to 160KB of shared memory (cc 8.0) or 100KB (cc 8.6), dynamically allocated, using the above opt-in mechanism, with the number 98304 changed to 163840 (for cc 8.0, for example) or 102400 (for cc 8.6).

Note that the above covers the Volta (7.0) Turing (7.5) and Ampere (8.x) cases. GPUs with compute capability prior to 7.x have no ability to address more than 48KB per threadblock. In some cases, these GPUs may have more shared memory per multiprocessor, but this is provided to allow for greater occupancy in certain threadblock configurations. The programmer has no ability to use more than 48KB per threadblock.

Although it doesn’t pertain to the code presented here (which is already using a dynamic shared memory allocation), note from the excerpted documentation quote that using more than 48KB of shared memory on devices that support it requires 2 things:

  1. The opt-in mechanism already described above
  2. A dynamic rather than static shared memory allocation in the kernel code.

example of dynamic:

extern __shared__ int shared_mem[];

example of static:

__shared__ int shared_mem[1024];

Dynamically allocated shared memory also requires a size to be passed in the kernel launch configuration parameters (an example is given in the question).

The assumption here is also that you intend to allocate/use all the shared memory via dynamic allocation. The proper limit, however, is that the sum of your dynamic request and static request cannot exceed the device maximum. So if you allocate 4kbytes statically, then the total dynamic allocation request cannot exceed (device limit)-4kbytes.

More Related Contents:

  1. What is the canonical way to check for errors using the CUDA runtime API?
  2. Unspecified launch failure on Memcpy
  3. Different CUDA versions shown by nvcc and NVIDIA-smi
  4. Thrust inside user written kernels
  5. GPU Emulator for CUDA programming without the hardware [closed]
  6. How is CUDA memory managed?
  7. Using GPU from a docker container?
  8. How to measure the inner kernel time in NVIDIA CUDA?
  9. How do CUDA blocks/warps/threads map onto CUDA cores?
  10. Understanding CUDA grid dimensions, block dimensions and threads organization (simple explanation) [closed]
  11. What can I do against ‘CUDA driver version is insufficient for CUDA runtime version’?
  12. How to get the CUDA version?
  13. Can anyone provide sample code demonstrating the use of 16 bit floating point in cuda?
  14. What kind of variables consume registers in CUDA?
  15. In CUDA, what is memory coalescing, and how is it achieved?
  16. CUDA: How many concurrent threads in total?
  17. CUDA apps time out & fail after several seconds – how to work around this?
  18. multi-GPU basic usage
  19. How to use 2D Arrays in CUDA?
  20. How to implement device side CUDA virtual functions?
  21. How to create a CUDA context?
  22. CUBLAS: Incorrect inversion for matrix with zero pivot
  23. Why does my CUDA kernel crash (unspecified launch failure) with a different dataset size?
  24. Cuda Mutex, why deadlock?
  25. CUDA determining threads per block, blocks per grid
  26. How does CUDA assign device IDs to GPUs?
  27. How to disable a specific nvcc compiler warnings
  28. printf() in my CUDA kernel doesn’t result produce any output
  29. Multiply Rectangular Matrices in CUDA
  30. “invalid configuration argument ” error for the call of CUDA kernel?

Leave a Comment