10

In "CUDA C Programming Guide 5.0", p73 (also here) says "Any address of a variable residing in global memory or returned by one of the memory allocation routines from the driver or runtime API is always aligned to at least 256 bytes". I do not know the exact meaning of this sentence. Could anyone show an example for me? Many thanks.

A derivative question: So, what about allocating an one-dimensional array of basic elements (like int) or self-defined ones? The starting address of the array will be multiples of 256B, while the address of each element in the array is not necessarily multiples of 256B?

einpoklum
  • 118,144
  • 57
  • 340
  • 684
jsc0218
  • 427
  • 5
  • 14
  • 3
    Start address of any cudaMalloc will be multiple of 256. 256 is two symbols in hex, they will be zero in addrexx; so you can get address like 0x0456ad00 but not 0x0456ad80. – osgx Dec 29 '12 at 15:50

1 Answers1

14

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.

sgarizvi
  • 16,623
  • 9
  • 64
  • 98
  • 3
    Providing for the convenient binding of textures to memory allocated via cudaMalloc() without the need to resort to texture offsets is another reason besides performance for the pointer alignment provided by cudaMalloc(). – njuffa Dec 29 '12 at 19:17
  • I added a derivative question. please help me to figure it out @sgar91. – jsc0218 Dec 30 '12 at 01:45
  • 1
    Are you sure that any result of `cudaMalloc()` has alignment no less than `cudaDeviceProp::textureAlignment` ? – einpoklum Apr 12 '16 at 07:14
  • 1
    @einpoklum Yes, as far as I have tested, I am getting same behavior across different devices. Alignment is not less than `cudaDeviceProp::textureAlignment`. – sgarizvi Apr 12 '16 at 08:28