Pinned Memory

References

• Professional CUDA C Programming

Contents

• Pinned Memory

---

Pinned Memory

할당된 Host 메모리는 기본적으로 pageable합니다. 즉, OS에 의해 host의 가상 메모리에서 다른 물리 메모리로 데이터를 이동하는 페이지 폴트(page fault) 동작이 일어날 수 있습니다. L1 캐시가 물리적으로 사용 가능한 메모리보다 훨씬 더 많은 on-chip 메모리를 제공하는 것처럼 가상 메모리는 물리적으로 사용 가능한 것보다 훨씬 더 많은 메모리를 제공합니다.

 

GPU에서는 host OS가 데이터를 물리적으로 이동하는 시점을 제어할 수 없기 때문에 pageable한 host 메모리에 있는 데이터를 안전하게 액세스할 수 없습니다. pageable한 host 메모리에서 device 메모리로 데이터를 전송할 때, CUDA 드라이버는 먼저 임시 page-locked 또는 pinned host 메모리를 할당하고, host 데이터를 pinned memory로 복사한 다음에 아래 그림의 왼쪽과 같이 pinned memory에서 device memory로 데이터를 전송합니다.

CUDA 런타임은 다음의 API를 통해 직접 pinned host memory를 할당할 수 있습니다.

cudaError_t cudaMallocHost(void **devPtr, size_t count);

이 함수는 page-locked하고, device에서 액세스할 수 있는 count bytes만큼의 host 메모리를 할당합니다. device에서 직접 pinned memory에 액세스할 수 있기 때문에, 이는 pageable memory보다 더 높은 read/write bandwidth를 갖습니다. 그러나 과도한 양의 pinned memory를 할당하면 가상 메모리 데이터를 저장하는데 사용하는 host 시스템의 사용가능한 pageable memory 양을 감소시키므로, host 시스템의 성능이 저하될 수 있습니다.

 

아래의 코드 일부는 error 핸들링을 포함하여 pinned memory를 할당하는 방법을 보여줍니다.

할당된 pinned host memory는 다음의 함수로 메모리를 해제할 수 있습니다.

cudaError_t cudaFreeHost(void *ptr);

 

pageable host memory가 아닌 pinned host memory를 사용하면 성능의 향상이 얼마나 되는지 직접 확인해보겠습니다.

먼저, 다음의 코드는 기본적인 pageable host memory를 사용하여 간단하게 메모리 전송을 수행하는 코드입니다.

#include <stdio.h>
#include <cuda_runtime.h>

int main(int argc, char** argv)
{
    // setup device
    int dev = 0;
    cudaSetDevice(dev);

    // memory size
    unsigned int iSize = 1 << 22;
    unsigned int nBytes = iSize * sizeof(float);

    // get deivce information
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, dev);
    printf("device %d: %s memory size %d nbyte %.2fMB\n", dev, deviceProp.name,
            iSize, nBytes/(1024.f*1024.f));
    
    // allocate the host memory
    float *h_a = (float*)malloc(nBytes);
    // allocate the device memory
    float *d_a;
    cudaMalloc((void**)&d_a, nBytes);

    // initialize the host memory
    for (auto i = 0; i < iSize; i++)
        h_a[i] = 0.5f;

    // transfer data from the host to the device
    cudaMemcpy(d_a, h_a, nBytes, cudaMemcpyHostToDevice);

    // transfer data from the device to the host
    cudaMemcpy(h_a, d_a, nBytes, cudaMemcpyDeviceToHost);
    
    // free memory
    cudaFree(d_a);
    free(h_a);

    // reset device
    cudaDeviceReset();
    return 0;
}

컴파일 후에, nvprof로 프로파일링을 수행해보겠습니다.

약 5.68ms의 시간이 걸렸습니다.

 

다음의 코드는 pageable host memory를 pinned host memory로 바꾸어서 데이터 전송을 수행하는 코드입니다.

#include <stdio.h>
#include <cuda_runtime.h>

int main(int argc, char** argv)
{
    // setup device
    int dev = 0;
    cudaSetDevice(dev);

    // memory size
    unsigned int iSize = 1 << 22;
    unsigned int nBytes = iSize * sizeof(float);

    // get deivce information
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, dev);
    if (!deviceProp.canMapHostMemory) {
        printf("Device %d does not support mapping CPU host memory!\n", dev);
        cudaDeviceReset();
        return 0;
    }
    printf("device %d: %s memory size %d nbyte %.2fMB canMap %d\n", dev, deviceProp.name,
            iSize, nBytes/(1024.f*1024.f), deviceProp.canMapHostMemory);
    
    // allocate the host memory
    float *h_a;
    cudaMallocHost((void**)&h_a, nBytes);
    // allocate the device memory
    float *d_a;
    cudaMalloc((void**)&d_a, nBytes);

    // initialize the host memory
    for (auto i = 0; i < iSize; i++)
        h_a[i] = 0.5f;

    // transfer data from the host to the device
    cudaMemcpy(d_a, h_a, nBytes, cudaMemcpyHostToDevice);

    // transfer data from the device to the host
    cudaMemcpy(h_a, d_a, nBytes, cudaMemcpyDeviceToHost);
    
    // free memory
    cudaFree(d_a);
    cudaFreeHost(h_a);

    // reset device
    cudaDeviceReset();
    return 0;
}

마찬가지로 컴파일 후, 실행하면 다음의 결과를 확인할 수 있습니다.

5.68ms가 걸리던 것에 비해서 약 5.13ms로 0.5ms의 속도 향상이 있습니다.

 

 

Pinned Memory는 Pageable Memory보다 메모리 할당 및 해제에 cost가 더 많이 소모되지만, 대용량 데이터 전송에서 더 높은 처리량을 제공합니다. Pageable Memory가 아닌 Pinned Memory를 사용함으로써 얻을 수 있는 속도 향상은 compute capability에 따라서 다릅니다. 예를 들어, Fermi 디바이스에서는 일반적으로 10MB 이상의 데이터를 전송할 때 Pinned Memory를 사용하는 것이 유리합니다.