Skip to content
Search

CUDA:Memory

lastmod 2024-07-22created 2024-07-135 edits

Data Transfer

Overlap Data Transfers

About asynchronous data transfer.

If you use a stream for the cudaMemcpyAsync call, you can insert an event into the stream after the asynchronous transfer and then synchronize on that event using cudaEventSynchronize.

This guarantees that the copy has completed, but does not rely on the device being idle or the stream being empty.

APIs

cudaMemcpyAsync 

result = cudaMemcpyAsync(d_a, a, N, cudaMemcpyHostToDevice, stream1)

Getting the memory alignment value

invetigation on cudaMalloc alignment => aligned to at least 512 bytes 

#include <sys/time.h>
#include <cuda_runtime.h>
#include <stdio.h>

void test(int size)
{
    float *d1, *d2;

    cudaMalloc(&d1, size);
    cudaMalloc(&d2, size);
    printf("Alignment: %ld\n", (d2 - d1) * sizeof(float));
    cudaFree(d1);
    cudaFree(d2);
}

int main(int argc, char **argv)
{
    // set up device
    int dev = 0;
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, dev);
    printf("Using Device %d: %s\n", dev, deviceProp.name);
    cudaSetDevice(dev);

    test(1);

    return(0);
}

Result:

Using Device 0: GeForce GTX TITAN X
Alignment: 512

in AWS p2.xlarge (k80)

Using Device 0: Tesla K80
Alignment: 512

This document says as https://www.classes.cs.uchicago.edu/archive/2011/winter/32102-1/reading/CUDA_C_Best_Practices_Guide.pdf

cudaMalloc(), is guaranteed to be aligned to at least 256 bytes

But, it looks actually aligned to at least 512 bytes. It looks the document is old.

ref. https://stackoverflow.com/questions/36534599/cuda-malloc-minimum-and-typical-actual-alignment

256 on Fermi, 512 on the Kepler, Maxwell

So, it was 256 bytes for old GPU such as Fermi, but it is 512 bytes nowadays.

See also

Favorite site

References

  1. CUDA_C_and_CPP_-Streams_and_Concurrency-Steve_Rennich-_NVIDIA.pdf ↩

author Zer0 sitemap