[NVIDIA GPU] Pinned Host Memory (CUDA)

연구실에서 기업과제로 CUDA를 사용하여 병렬화 코딩을 진행하고 있다. Multi-GPUs 관련 코딩을 하면서 Pinned Host Memory라는 부분을 알게 되어 짧게 정리하였다. 다양한 최적화 기법을 시도해 보았다. 그 중 Pinned Host Memory를 사용하여 상당한 성능 향상을 얻을 수 있었다.

보통 Host (CPU)에서 메모리 allocation(할당)을 하면 Pageable Memory이다. Pageable Memory의 경우 GPU 메모리로 바로 복사가 불가능하다고 한다. 우선 Pageable Memory와 Non-Pageable Memory (Page-Locked)에 대해서 정리하였다.

  • Pageable Memory: Memory 내용(contents)이 DRAM에서 하드디스크 (Secondary Storage Device)로 page out 되거나 반대로 하드디스크에서 DRAM으로 page in이 가능한 메모리를 의미한다. Page in/Page out을 하기 위해서는 CPU (Host)의 도움이 필요하다고 한다. (OS가 어떻게 page in/page out을 하는지는 다음에 시간 되면 더 상세히 정리할 계획이다) 보통 OS에서 User Memory Space의 경우 Pageable Memory이다.
  • Non-Pageable Memory: Pageable Memory와 반대로 page in/page out이 불가능한 메모리를 Non-Pageable Memory라 한다. 결과적으로 하드디스크로 데이터를 page out/page in 하는 작업이 필요없다. OS에서 Kernel Memory Space는 보통 Non-Pageable Memory라고 한다.
  • 추가: Non-Pageable Memory를 너무 많이 할당하면 OS의 안전성이 떨어진다고 한다. “여기서 말하는 너무 많이라 하는 메모리 크기”는 실제 시스템 마다 차이가 있고 정확한 수치는 알려져 있지 않다.

Pageable Memory를 GPU로 복사하기 위해서 GPU driver가 Pinned Memory (Temporary Page-Locked Host Array 또는 Non-Pageable Memory) 를 할당해야 한다. Pinned Memory를 할당 후 Pageable Memory의 데이터를 Pinned Memory로 복사 후 GPU 메모리로 복사할 수 있다. 그림 1은 Pageable Data Transfer의 예를 보여준다.

그림 1: Pageable Data Transfer의 예

반면 host에서 데이터를 할당할 때 Pinned Memory를 사용하면 Pageable Memory에서 Pinned Memory로 복사하는 과정을 제거할 수 있다. 그림 2는 Pinned Memory 할당 후 바로 GPU Memory로 복사하는 예를 보여준다.

그림 2: Pinned Data Transfer의 예

CUDA 프로그램 시 Pinned Memory를 사용하기 위해서는 Host memory 을 할당할 때 cudaMallocHost() 또는 cudaHostAlloc() (deallocation시 cudaFreeHost())를 사용하면 된다. 그림 3은 cudaMallocHost() 의 예를 보여준다. 그림 4는 (출처 1)에서 가져온 성능 결과이다. 결과에서 볼 수 있 듯 성능이 최대 2배 가까이 차이 나는 것을 확인할 수 있다.

그림 3: cudaMallocHost() 사용 예

그림 4: Pinned Memory 성능

연구실 기업과제를 할 때 cudaMallocHost을 사용하면 더 좋은 성능을 얻는 것을 확인하였다. 특히 Kernel 계산 시간보다 memory 복사 시간이 더 오래 걸리는 경우 cudaMallocHost로 실행 시간을 많이 단축할 수 있다.

출처

  1. https://devblogs.nvidia.com/parallelforall/how-optimize-data-transfers-cuda-cc/
  2. https://en.wikipedia.org/wiki/CUDA_Pinned_memory
  3. http://www.orangeowlsolutions.com/archives/443
  4. https://www.quora.com/What-is-the-difference-between-pageable-and-non-pageable-memory-Is-user-memory-space-pageable-or-non-pageable

3 thoughts on “[NVIDIA GPU] Pinned Host Memory (CUDA)”

  1. Pinned Memory 개념이 커널 메모리 영역쪽 버퍼라고 생각하면 될까요??
    Pageable memory 없이 Pinned Memory 사용하는건 Memory mapped IO 개념이구요??

    Reply
    • 안녕하세요. 답변이 많이 늦었습니다.

      CPU 메인 메모리에서 GPU 메모리로 데이터를 복사할 때 Pinned Memory를 통해서 GPU 메모리로 데이터가 복사됩니다.
      처음 코딩을 하실 때부터 고정된 Pinned Memory를 사용하고 CPU의 데이터를 Pin Memory에 저장하시면 중간에 Pinned Memory 저장하는 과정을 제거할 수 있습니다.

      Pageable Memory와 Pin 메모리는 모두 메인 메리로 공간을 의미합니다.
      감사합니다.

      Reply
  2. OpenCL 공부중에 유사한 궁금증이 있어 방문하게 되었으며, 이해에 많은 도움이 되었습니다.
    감사합니다!

    Reply

Leave a Reply to SUNG-JIN KANG Cancel reply