CUDA - Host Memory와 Device Memory의 사이의 데이터 전송

Mapped Memory란?
Mapped Memory는 CUDA 2.2부터 추가된 새로운 기능입니다. 지금까지 CUDA 프로그래밍 구조에서는 Host Memory(호스트머신 위의 메인 메모리)와 Device Memory(GPU상의 비디오메모리)는 완전히 별개의 주소공간 위의 메모리로서 명확하게 나눌 수 있었습니다. 그러나 Mapped Memory를 사용하여 이러한 차이를 명확하게 구별하지 않고 CUDA 프로그래밍을 실시할 수 있게 됩니다.

무엇보다, Mapped Memory를 사용하는 본래의 목적은 위와 같은 편리성보다 실제적인 성능 향상에 있습니다. Host Memory와 Device Memory의 사이의 데이터 전송은 레지스터나 공유 메모리의 액세스에 비하면 훨씬 더 저속 PCI 버스를 통해 실행되기 때문에 많은 경우 이 부분이 보틀 넥이 됩니다.

이 통신을 고속화하는 것은 매우 중요한 과제가 되지만 Mapped Memory에 의한 전송과 CUDA의 통상적인 메모리 API를 이용한 전송에서 대략적으로 1.5배에서 2배 정도의 성능 향상이 있습니다. 하지만, Mapped Memory의 몇가지 특성을 파악한 다음 사용하지 않으면 반대로 성능이 저하되기도 합니다. 다음은 Mapped Memory의 간단한 사용법과 실제로 사용했을 때의 벤치마크와 그 결과에 대해 말합니다.

기본적인 사용법
Mapped Memory 사용 순서는 다음과 같습니다.

1. 디바이스에 Mapped Memory를 사용하기 위한 플래그를 세트 한다
2. Host Memory를 Mapped Memory 전용으로 할당한다
3. 2로 확보한 Host Memory를 Device Memory에 맵하여 디바이스로부터 Mapped Memory에 액세스하기 위한 포인터를 동기획득 한다
4. 3으로 동기획득 한 포인터를 사용하여 Kernel 내부로부터 Mapped Memory에 액세스하여 적당한 크기의 데이터를 로드, 스토어 한다
5. 계산이 종료되면 Mapped Memory를 해제한다

다음 예제는 일차원 배열의 각 요소를 인크리먼트(increment)하는 CUDA 프로그램입니다.

// includes, system
#include <stdio.h>

// includes, project
#include <cuda.h>

// some defines
#define VECTOR_SIZE 1024
#define BLOCK_SIZE 256

///////////////////////////////////////////////////////////////////////////////
// Kernel
////////////////////////////////////////////////////////////////////////////////
__global__ void vectorInc(unsigned int* A)
{
    unsigned int index = blockDim.x * blockIdx.x + threadIdx.x;
    unsigned int a;

    a = A[index];

    a++;

    A[index] = a;

    return;
}

////////////////////////////////////////////////////////////////////////////////
// Program Main
////////////////////////////////////////////////////////////////////////////////
int main(int argc, char** argv)
{
    // 디바이스에 플래그를 세트
    cudaSetDeviceFlags(cudaDeviceMapHost);
 
    // Host Memory를 Mapped Memory로서 할당
    unsigned int* h_A;
    cudaHostAlloc((void**) &h_A, VECTOR_SIZE*sizeof(unsigned int), cudaHostAllocMapped);
 
    // 배열의 초기화
    for( unsigned int i = 0; i < VECTOR_SIZE; i++) {
 h_A[i] = (unsigned int) i;
    }

    // Host Memory를 Device Memory에 맵 한다
    unsigned int* d_A;
    cudaHostGetDevicePointer((void **) &d_A, (void *) h_A, 0);
 
    // Kernel를 실행한다
    vectorInc<<<VECTOR_SIZE/BLOCK_SIZE, BLOCK_SIZE>>>(d_A);
 
    // 결과의 표시
    for( unsigned int i = 0; i < VECTOR_SIZE; i++) {
 printf("h_A[%d]:%d\n", i, h_A[i]);
    }
    
    // 메모리의 해제
    cudaFreeHost(h_A);
 
    cudaThreadExit();
 
    return 0;

}

배열의 초기화나 결과 출력의 코드가 들어가 있있지만 main 함수는 다음 순서대로 구현되고 있는 것을 알 수 있습니다.

통상적인의 Device Memory와 비교
비교를 위하여 통상적인 메모리를 사용하고 전송을 명시적으로 실시했을 경우의 프로그램의 main 함수를 봅시다.

int main(int argc, char** argv)
{

    // Host Memory를 확보
    unsigned int* h_A;
    cudaHostAlloc((void**) &h_A, VECTOR_SIZE*sizeof(unsigned int), cudaHostAllocMapped);
 
    // 배열의 초기화
    for( unsigned int i = 0; i < VECTOR_SIZE; i++) {
 h_A[i] = (unsigned int) i;
    }

    // Device Memory를 할당
    unsigned int* d_A;
    cudaMalloc((void**) &d_A, VECTOR_SIZE*sizeof(unsigned int));
 
    // Host Memory로부터 Device Memory에 데이터를 전송 한다
    cudaMemcpy(d_A, h_A, VECTOR_SIZE*sizeof(unsigned int), cudaMemcpyHostToDevice);

    // Kernel를 실행한다
    vectorInc<<<VECTOR_SIZE/BLOCK_SIZE, BLOCK_SIZE>>>(d_A);

    // Kernel의 종료를 기다린다
    cudaThreadSynchronize();

    // Device Memory로부터 Host Memory에 데이터를 전송 한다
    cudaMemcpy(h_A, d_A, VECTOR_SIZE*sizeof(unsigned int), cudaMemcpyDeviceToHost);

 
    // 결과의 표시
    for( unsigned int i = 0; i < VECTOR_SIZE; i++) {
 printf("h_A[%d]:%d\n", i, h_A[i]);
    }
    
    // 메모리의 해제
    cudaFree(d_A);
    cudaFreeHost(h_A);
 
    cudaThreadExit();
 
    return 0;

}

코드를 살펴보면 알수 있겠지만 통상적인 Device Memory와 Mapped Memory를 사용할 때 다른 부분은 다음과 같습니다.

1. 플래그를 설정 유무
2. cudaMemcpy에 의한 전송 실시 유무
3. Device Memory의 메모리 할당과 해제 실시 유무

포인트는 Kernel측의 변경은 일절 실시하지 않고 사용할 수 있는 것입니다. Kernel 내부의 액세스 전에 통상적인 Global Memory인 경우는 Device Memory의 액세스로 처리됩니다. Mapped Memory인 경우에는 자동적으로 PCI를 개입시켜 데이터 전송이 실행됩니다. Kernel 내부의 계산과 PCI 버스를 개입시킨 데이터 전송을 비동기로 실행할 수 있게 됩니다. 이것이 Mapped Memory가 효율적으로 데이터 전송을 실시할 수 있는 이유입니다.

벤치마크
Mapped Memory를 사용해, 실제의 성능측정을 실시한 결과가 다음의 그래프입니다. 비교 대상은 통상적인 메모리 영역(malloc()로 할당한 것)와 페이지 락을 실시한 메모리 영역(cudaHostAlloc()로 할당한 것)와 디바이스 메모리간의 전송을 측정하였습니다.

Mapped Memory의 전송 성능이 그래프에 나타난 것처럼 Mapped Memory의 전송 성능이 cudaMemcpy에 의해 전송을 실시하는 것보다 고속입니다. 이것은 Kernel 실행 중에 비동기화 메모리 전송을 실시할 수 있어서, cudaMemcpy()에 의한 데이터 전송이 PCI 버스를 한쪽방향 통신 밖에 사용할 수 없는 것에 비해 Mapped Memory에 의한 데이터 전송은 양방향 통신이 가능한 것이 고속데이터 전송이 가능한 이유라 할 수 있습니다.

다만 여기서 나타낸 예는 Global Memory에 대해서 시퀀셜 액세스 했을 경우의 벤치마크입니다. 직접접근의 경우는 PCI의 통신이 적은 데이터로 많은 회수를 사용하면 반대로 전송 성능이 낮아질 수 있습니다. Mapped Memory를 사용할 때, Global Memory의 액세스 패턴의 최적화의 중요성은 더욱 커지기 때문에 이 부분을 주의하여 프로그램 설계할 필요가 있습니다.

출처 - http://goparallel.egloos.com/1968095