Zero-Copy Memory & Unified Virtual Addressing

References

• Professional CUDA C Programming

Contents

• Zero-Copy Memory
• Unified Virtual Addressing

---

Zero-Copy Memory

일반적으로 host는 device 변수에 직접 액세스할 수 없고, device는 host 변수에 직접 액세스할 수 없습니다. 하지만 여기에 한 가지 예외인 zero-copy 메모리가 있습니다. host와 device는 모두 zero-copy 메모리에 액세스할 수 있습니다.

 

GPU 스레드들은 직접 zero-copy 메모리에 액세스할 수 있습니다. zero-copy 메모리를 커널 내에서 사용하면 다음의 몇 가지 이점을 얻을 수 있습니다.

• Leveragin host memory when there is insufficient device memory
• Avoiding explicit data trasfer between the host and device
• Improving PCIe transfer rates

Zero-copy 메모리를 사용하여 host와 device 간의 데이터를 공유할 때, host와 device의 액세스 사이에서 반드시 메모리를 동기화해주어야 합니다. 동시에 host와 device에서 zero-copy 메모리의 데이터를 수정하는 것은 정의되지 않은 동작을 유발합니다.

 

Zero-copy 메모리는 device 메모리 공간에 매핑되는 pinned (non-pageable) memory 입니다.

다음의 함수를 통해 매핑되는 pinned memory 공간을 생성할 수 있습니다.

cudaError_t cudaHostAlloc(void **pHost, size_t count, unsigned int flags);

이 함수는 page-locked되고 device에서 액세스할 수 있는 count bytes만큼의 host 메모리를 할당합니다. 이 함수에 의해서 할당된 메모리는 반드시 cudaFreeHost를 통해 해제되어야 합니다. 

flags 파라미터는 할당된 메모리의 특별한 속성을 추가로 구성할 수 있습니다.

• cudaHostAllocDefault
• cudaHostAllocPortable
• cudaHostAllocWriteCombined
• cudaHostAllocMapped

cudaHostAllocDefault는 cudaHostAlloc의 동작을 cudaMallocHost와 동일하게 합니다.

cudaHostAllocPortable을 설정하면 할당된 메모리뿐만 아니라 모든 CUDA 컨텍스트에서 사용할 수 있는 pinned memory를 반환합니다.

cudaHostAllocWriteCombined 플래그는 write-combined 메모리를 반환하며, 이는 일부 시스템 구성에서 PCI Express를 통해 더 빠르게 전송될 수 있지만, 대부분의 host에서 효율적으로 read할 수는 없습니다. 그러므로 write-combined 메모리는 host에 의해서 write되고 device에 의해서 read되는 버퍼를 사용하기에 좋은 옵션입니다.

Zero-copy 메모리와 가장 관련이 있는 플래그는 cudaHostAllocMapped인데, 이는 device address 공간에 매핑된 host memory를 반환합니다.

 

다음의 함수를 사용하여 매핑된 pinned memory에 대한 device pointer를 얻을 수 있습니다.

cudaError_t cudaHostGetDevicePointer(void **pDevice, void *pHost, unsigned int flags);

이 함수는 매핑된 pinned host memory에 액세스하기 위해 device에서 참조할 수 있는 pDevice에 device pointer를 리턴합니다. 만약 device가 mapped, pinned memory를 지원하지 않는다면, 이 함수는 실패합니다.

flags는 현재 0으로만 설정할 수 있습니다.

 

read/write 작업이 빈번한 device memory에 zero-copy memory를 사용하면 성능이 매우 저하될 수 있습니다. 이는 매핑된 메모리로의 모든 memory transcation은 PCIe 버스를 통과해야하기 때문에 Global Memory와 비교해도 상당한 latency가 추가됩니다.

 

간단하게 두 배열의 합을 구하는 커널을 사용하여 zero-copy memory를 사용하는 방법과 일반적인 device global memory 방법의 차이를 비교해보도록 하겠습니다. 배열을 더하기 위해서 사용되는 커널 함수는 동일하고, 이름만 다릅니다.

__global__
void sumArrays(float* A, float* B, float* C, const int N)
{
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
    
    if (idx < N)
        C[idx] = A[idx] + B[idx];
}

__global__
void sumArraysZeroCopy(float* A, float* B, float* C, const int N)
{
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
    
    if (idx < N)
        C[idx] = A[idx] + B[idx];
}

전체 코드는 아래 링크를 참조해주세요.

https://github.com/junstar92/parallel_programming_study/blob/master/CUDA/vectorAdd/sumArrayZerocopy.cu

GitHub - junstar92/parallel_programming_study: Study parallel programming - CUDA, OpenMP, MPI, Pthread

 

위의 코드를 아래의 커맨드로 컴파일하고,

nvcc -O3 -o sumArrayZerocopy sumArrayZerocopy.cu -I..

nvprof로 프로파일링해보도록 하겠습니다. 커맨드 argument는 입력하지 않았기 때문에 배열의 크기는 \(2^10\)으로 설정됩니다.

sumArrays 커널과 sumArraysZeroCopy 커널을 비교해보면, 1024개의 요소를 처리하는데 sumArraysZeroCopy가 조금 더 느립니다. 두 커널에서 모두 device에서 host로 계산 결과를 전달해야하기 때문에 이에 대한 소요 시간도 포함해야합니다. 하지만, 처음 host에서 device로의 입력 배열 데이터 전달은 sumArrays에서만 수행되므로, 이걸 고려하면 sumArraysZeroCopy도 충분히 빠르다고 판단할 수 있습니다.

 

배열의 크기를 변경하여 다시 비교해보도록 하겠습니다.

- 16KB

- 64KB

- 256KB

- 1MB

- 16MB

- 64MB

 

실행 결과로부터 host와 device 사이에서 작은 양의 데이터를 공유하는 경우에 zero-copy 메모리 사용이 괜찮다는 것을 알 수 있습니다. PCIe Bus를 통해 연결된 개별 GPU가 있는 대규모 데이터셋의 경우에 zero-copy 메모리는 좋지 않으며 상당한 성능의 저하를 유발합니다.

 

---

Unified Virtual Addressing

compute capability 2.0 이상의 device에서는 Unified Virtual Addressing(UVA)라는 특별한 addressing mode를 지원합니다. CUDA 4.0에서 도입된 UVA는 64비트 리눅스 시스템에서 지원됩니다. compute capability 3.0 이상(CUDA 6.0)부터는 64비트 윈도우 시스템에서 Unified Memory라는 이름으로 지원됩니다.

 

UVA에서, host 메모리와 device 메모리는 하나의 virtual address space를 공유합니다.

 

UVA가 도입되기 전에는 host 메모리를 참조하는 포인터와 device memory를 참조하는 포인터를 관리해야했습니다.

하지만 UVA를 사용하면, 포인터에 의해서 참조되는 메모리 공간이 application 코드에서 명확해집니다. UVA에서 cudaHostAlloc으로 할당된 pinned host memory는 동일한 host 및 device 포인터를 갖습니다. 

따라서, 반환된 포인터를 커널 함수에 직접 전달할 수 있습니다.

 

위에서 살펴본 zero-copy 예제 코드 일부를 다시 살펴보겠습니다. 

int main(int argc, char** argv)
{
	...
    
    // part 2: using zero-copy memory for array A and B
    // allocate zero-copy memory
    CUDA_CHECK(cudaHostAlloc((void**)&h_A, nBytes, cudaHostAllocMapped));
    CUDA_CHECK(cudaHostAlloc((void**)&h_B, nBytes, cudaHostAllocMapped));

    // initialize data at host side
    initialData(h_A, nElem);
    initialData(h_B, nElem);
    memset(hostRef, 0, nBytes);
    memset(gpuRef, 0, nBytes);

    // pass the pointer to device
    CUDA_CHECK(cudaHostGetDevicePointer((void**)&d_A, (void*)h_A, 0));
    CUDA_CHECK(cudaHostGetDevicePointer((void**)&d_B, (void*)h_B, 0));
    
    // add vector at host side for result checks
    sumArraysOnHost(h_A, h_B, hostRef, nElem);

    // execute kernel with zero copy memory
    sumArraysZeroCopy<<<grids, blocks>>>(d_A, d_B, d_C, nElem);
    
    ...
}

Zero-copy 메모리를 사용하기 위해서 cudaHostAlloc으로 메모리를 할당하고, 할당된 메모리의 주소를 cudaHostGetDevicePointer를 통해서 얻었습니다. 하지만, UVA를 지원하는 device에서는 cudaHostGetDevicePointer로 주소를 따로 얻을 필요없이, 그냥 h_A, h_B가 가리키는 주소를 커널로 직접 전달할 수 있습니다.

 

따라서 위의 코드에서 cudaHostGetDevicePointer 부분을 주석처리하고, 컴파일 후 실행해도 정상적으로 동작합니다.

int main(int argc, char** argv)
{
	...
    
    // part 2: using zero-copy memory for array A and B
    // allocate zero-copy memory
    CUDA_CHECK(cudaHostAlloc((void**)&h_A, nBytes, cudaHostAllocMapped));
    CUDA_CHECK(cudaHostAlloc((void**)&h_B, nBytes, cudaHostAllocMapped));

    // initialize data at host side
    initialData(h_A, nElem);
    initialData(h_B, nElem);
    memset(hostRef, 0, nBytes);
    memset(gpuRef, 0, nBytes);

    // pass the pointer to device
    //CUDA_CHECK(cudaHostGetDevicePointer((void**)&d_A, (void*)h_A, 0));
    //CUDA_CHECK(cudaHostGetDevicePointer((void**)&d_B, (void*)h_B, 0));
    
    // add vector at host side for result checks
    sumArraysOnHost(h_A, h_B, hostRef, nElem);

    // execute kernel with zero copy memory
    sumArraysZeroCopy<<<grids, blocks>>>(h_A, h_B, d_C, nElem);
    
    ...
}

cudaHostGetDevicePointer 함수 호출 부분을 주석 처리하고, 커널에 직접 h_A와 h_B를 전달해줍니다.

 

정상적으로 컴파일되며,

정상적으로 실행됩니다.