CUDA allocation alignment is 256 bytes – seriously?

by

The pointers which are allocated by using any of the CUDA Runtime’s device memory allocation functions e.g cudaMalloc or cudaMallocPitch are guaranteed to be 256 byte aligned, i.e. the address is a multiple of 256.

Consider the following example:

char *ptr1, *ptr2;

int bytes = 1;

cudaMalloc((void**)&ptr1,bytes);
cudaMalloc((void**)&ptr2,bytes);

Suppose the address returned in ptr1 is some multiple of 256, then the address returned in ptr2 will be atleast (ptr1 + 256).

This is a restriction imposed by the device on which the memory is being allocated. Mostly, pointers are aligned due to performance purposes. (Some NVIDIA guy should be able to tell if there is some other reason also).

Important:

Pointer alignment is not always 256. On my device (GTX460M), it is 512. You can get the device pointer alignment by the cudaDeviceProp::textureAlignment field.

Alignment of pointers is also a requirement for binding the pointer to textures.

More Related Contents:

  1. Different CUDA versions shown by nvcc and NVIDIA-smi
  2. Cuda atomics change flag
  3. Thrust inside user written kernels
  4. Modifying registry to increase GPU timeout, windows 7
  5. How can I implement a custom atomic function involving several variables?
  6. Passing Host Function as a function pointer in __global__ OR __device__ function in CUDA
  7. What is the purpose of using multiple “arch” flags in Nvidia’s NVCC compiler?
  8. How are 2D / 3D CUDA blocks divided into warps?
  9. What can I do against ‘CUDA driver version is insufficient for CUDA runtime version’?
  10. What is a bank conflict? (Doing Cuda/OpenCL programming)
  11. nvidia-smi Volatile GPU-Utilization explanation?
  12. Can anyone provide sample code demonstrating the use of 16 bit floating point in cuda?
  13. Any particular function to initialize GPU other than the first cudaMalloc call?
  14. CUDA and nvcc: using the preprocessor to choose between float or double
  15. Polymorphism and derived classes in CUDA / CUDA Thrust
  16. CUDA: How many concurrent threads in total?
  17. How do I select which GPU to run a job on?
  18. multi-GPU basic usage
  19. device function pointers
  20. How can I make tensorflow run on a GPU with capability 2.x?
  21. cudaMemset() – does it set bytes or integers?
  22. How to implement device side CUDA virtual functions?
  23. CUDA determining threads per block, blocks per grid
  24. How does CUDA assign device IDs to GPUs?
  25. How to disable a specific nvcc compiler warnings
  26. printf() in my CUDA kernel doesn’t result produce any output
  27. CUDA function pointers
  28. Multiply Rectangular Matrices in CUDA
  29. Using an array of device function pointers
  30. “invalid configuration argument ” error for the call of CUDA kernel?

Leave a Comment