How is CUDA memory managed?

by

The device memory available to your code at runtime is basically calculated as

Free memory =   total memory 
              - display driver reservations 
              - CUDA driver reservations
              - CUDA context static allocations (local memory, constant memory, device code)
              - CUDA context runtime heap (in kernel allocations, recursive call stack, printf buffer, only on Fermi and newer GPUs)
              - CUDA context user allocations (global memory, textures)

if you are getting an out of memory message, then it is likely that one or more of the first three items is consuming most of the GPU memory before your user code ever tries to get memory in the GPU. If, as you have indicated, you are not running on a display GPU, then the context static allocations are the most likely source of your problem. CUDA works by pre-allocating all the memory a context requires at the time the context is established on the device. There are a lot of things which get allocated to support a context, but the single biggest consumer in a context is local memory. The runtime must reserve the maximum amount of local memory which any kernel in a context will consume for the maximum number of threads which each multiprocessor can run simultaneously, for each multiprocess on the device. This can run into hundreds of Mb of memory if a local memory heavy kernel is loaded on a device with a lot of multiprocessors.

The best way to see what might be going on is to write a host program with no device code which establishes a context and calls cudaMemGetInfo. That will show you how much memory the device has with the minimal context overhead on it. Then run you problematic code, adding the same cudaMemGetInfo call before the first cudaMalloc call that will then give you the amount of memory your context is using. That might let you get a handle of where the memory is going. It is very unlikely that fragmentation is the problem if you are getting failure on the first cudaMalloc call.

More Related Contents:

  1. How to measure the inner kernel time in NVIDIA CUDA?
  2. nvidia-smi Volatile GPU-Utilization explanation?
  3. How do I use Nvidia Multi-process Service (MPS) to run multiple non-MPI CUDA applications?
  4. How does CUDA assign device IDs to GPUs?
  5. How do I choose grid and block dimensions for CUDA kernels?
  6. GPU Emulator for CUDA programming without the hardware [closed]
  7. Modifying registry to increase GPU timeout, windows 7
  8. Passing Host Function as a function pointer in __global__ OR __device__ function in 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. What is a bank conflict? (Doing Cuda/OpenCL programming)
  13. Any particular function to initialize GPU other than the first cudaMalloc call?
  14. CUDA apps time out & fail after several seconds – how to work around this?
  15. How do I select which GPU to run a job on?
  16. multi-GPU basic usage
  17. How can I make tensorflow run on a GPU with capability 2.x?
  18. why do we need cudaDeviceSynchronize(); in kernels with device-printf?
  19. CUDA determining threads per block, blocks per grid
  20. What is the correct version of CUDA for my nvidia driver?
  21. Using GPU from a docker container?
  22. Setting up Visual Studio Intellisense for CUDA kernel calls
  23. How to get the CUDA version?
  24. How to generalize fast matrix multiplication on GPU using numba
  25. In CUDA, what is memory coalescing, and how is it achieved?
  26. How do I check if PyTorch is using the GPU?
  27. CUBLAS: Incorrect inversion for matrix with zero pivot
  28. Using maximum shared memory in Cuda
  29. Cuda Mutex, why deadlock?
  30. CUDA allocation alignment is 256 bytes – seriously?

Leave a Comment