allocating shared memory

by

CUDA supports dynamic shared memory allocation. If you define the kernel like this:

__global__ void Kernel(const int count)
{
    extern __shared__ int a[];
}

and then pass the number of bytes required as the the third argument of the kernel launch

Kernel<<< gridDim, blockDim, a_size >>>(count)

then it can be sized at run time. Be aware that the runtime only supports a single dynamically declared allocation per block. If you need more, you will need to use pointers to offsets within that single allocation. Also be aware when using pointers that shared memory uses 32 bit words, and all allocations must be 32 bit word aligned, irrespective of the type of the shared memory allocation.

More Related Contents:

  1. What is option -O3 for g++ and nvcc?
  2. cudaMemcpy segmentation fault
  3. How do I start a new CUDA project in Visual Studio 2008?
  4. CUDA, Using 2D and 3D Arrays
  5. How do I start a CUDA app in Visual Studio 2010?
  6. Why does cudaMalloc() use pointer to pointer?
  7. Structure of Arrays vs Array of Structures
  8. CUDA how to get grid, block, thread size and parallalize non square matrix calculation
  9. Trouble launching CUDA kernels from static initialization code
  10. Cuda kernel returning vectors
  11. CUDA and Classes
  12. Compiling Cuda code in Qt Creator on Windows
  13. How performing multiple matrix multiplications in CUDA?
  14. Why does nvcc fails to compile a CUDA file with boost::spirit?
  15. Should I unify two similar kernels with an ‘if’ statement, risking performance loss?
  16. Why is NVIDIA Pascal GPUs slow on running CUDA Kernels when using cudaMallocManaged
  17. Reset Cuda Context after exception
  18. Why is CUDA pinned memory so fast?
  19. LNK2038: mismatch detected for ‘RuntimeLibrary’: value ‘MT_StaticRelease’ doesn’t match value ‘MD_DynamicRelease’ in file.obj
  20. Compiling code containing dynamic parallelism fails
  21. Wierd pointer construct
  22. What is happening in this C++ code in C++ (Exponent Function)
  23. uint8_t can’t be printed with cout
  24. C++ Convert string (or char*) to wstring (or wchar_t*)
  25. How to generate a random number in C++?
  26. Detecting TCP Client Disconnect
  27. How to make C++ cout not use scientific notation
  28. Calling R Function from C++
  29. c++ undefined references with static library
  30. How does “using std::swap” enable ADL?

Leave a Comment