[CUDA] 02. Memory Copy 및 Crop

Nvidia CUDA 를 활용한 Image Processing 라이브러리 예제들을 차근차근 설명해볼까 합니다.

CUDA 런타임 라이브러리 기반의 NPPi 와 같은 라이브러리를 활용하다보니,
마땅한 예제가 없어 익히는데 시간이 걸렸었는데
다음 배우는 사람에게 조금이나마 도움이 되고자 하나하나 천천히 설명해나가보고자 합니다.
CUDA를 활용해 Image 처리를 하거나, Video 처리를 하거나, 기타 이것저것 해보고자 하는
초심자들에게 조금이나마 도움이 되었으면 합니다.

- 다루고자 하는 내용
2D Memory Buffer 기초, NPPi를 활용한 Image Processing, (시간이 되면) Video Codec SDK

※ 현재 저의 개발환경은 다음과 같습니다.
Ubuntu 18.04, Nvida 450.66, CUDA 11.0 (GPU: GTX-1050Ti)
제가 다룰 예제들은 Cuda 10.0 에서도 크게 문제없이 동작합니다.

1. System Memory 와 Device Memory

일반적으로 프로그램에서 변수를 할당하거나 표준 라이브러리로 메모리를 할당 시, System Memory가 할당됩니다.
이는 RAM이라는 공간일 수도 있고, 하드디스크나, CPU Cache일 수도 있지만 어찌되었건
OS가 관리하는 본래의 우리가 알고있던 메모리 공간 입니다.

그 외 CUDA에서는 자체적인 GPU 내 메모리를 할당 해제 등 관리할 수 있습니다.
따라서 CUDA를 활용한 여러 연산을 수행할 때에는 데이터를 GPU 메모리로 복제 후 연산을 수행할 수 있습니다.

CUDA Runtime API에서는 System Memory를 Host, GPU Memory를 Device로 구분합니다.
따라서 데이터를 System Memory에서 GPU Memory로 (HostToDevice),
또는 GPU Memory에서 GPU Memory로 (DeviceToDevice),
또는 GPU Memory에서 System Memory로 (DeviceToHost) 복제할 때 메모리 복제가 어디서 어디로 수행될지 상수를 구분해 사용해야 합니다.

+ 참고로 덧붙이자면, 이러한 데이터의 이동은 물리적인 Bus를 거치므로, 필연적으로 병목을 형성할 수 있습니다.
특히나 어떤 데이터를 읽어서 GPU로 처리하는 일련의 과정은 Storage -> System Memory -> Device Memory로 데이터 복제가
여러번 반복되며 성능 저하의 원인이 될 수 있습니다.
Nvidia에는 이러한 입출력 병목(I/O Bottleneck)을 해소하기 위해, GPU Direct 라는 기술을 제공하고 있습니다.
Storage에서 System Memory를 거치지 않고, GPU Memory로 바로 읽어들이는 방식입니다.

Nvidia GPU Direct 기술 개요도

(출처: https://developer.nvidia.com/blog/gpudirect-storage/)

2. Host Memory를 Device로 Copy (System Memory -> GPU Memory)

링크의 실제 예제 코드를 활용해서 설명하겠습니다.
(예제 코드: https://github.com/z-wony/CudaPractice/blob/master/src/Practice_03_cudaMemcpy2D/main.c)

예제의 앞부분에서 libjpeg 으로 Image를 Decode해서 System Memory에 올려놓은 상태 입니다.

    cudaMemcpy2D(devPtr,        // dst Buffer (Device)
        pitch,                  // Pitch size of dst (devPtr)
        imgBuffer,              // src Buffer (Host)
        (size_t)(imgWidth * 3), // Pitch size of src (imgBuffer)
        (size_t)(imgWidth * 3), // Width size of src (imgBuffer)
        (size_t)imgHeight,
        cudaMemcpyHostToDevice); // Direction of copy (IMPORTANT)

첫번째와 두번째 인자가 복사를 받을 Device Memory 포인터와 Pitch값 입니다. (cudaMallocPitch로 할당)
세번째 인자인 imgBuffer는 System Memory 배열의 시작 주소이고,
그 다음은 Pitch인데 libjpeg에서 특별히 Padding 영역을 할당하지 않았으므로, width에 해당하는 바이트 크기가 그대로 쓰입니다.
실제 width를 입력하는 필드는 다섯번째 필드인데, CUDA 스럽게 Pixel 길이가 아닌 바이트 크기를 입력받고 있습니다.
그 다음 세로 길이을 입력하고, 마지막이 메모리 복제 방향입니다. (API 설명에서는 kind 라는 이름을 썼네요)

cudaMemcpy2D 함수 API Description

(출처: https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY.html#group__CUDART__MEMORY_1g3a58270f6775efe56c65ac47843e7cee)

 

3. GPU간 메모리 복제 (Device to Device)

이미지나 비디오 처리 과정에서 일련의 파이프라인으로, 몇단계의 연산을 수행할 때
필요에 의해서 파이프라인 1, 파이프라인 2 가 각각의 메모리 공간을 할당해서 사용하고 있고
데이터만 전달해야할 상황이 생길 수 있습니다.
이 때 System Memory를 활용해서 복제하면, H/W간 입출력 동작을 하게 되므로 굉장한 성능 병목을 만듭니다.
따라서 GPU Memory 내에서 Memcpy가 일어나는게 성능적으로 유리할 때가 많습니다.

위와 동일한 예제를 계속 활용중인데, 메모리 주소와 Pitch 값 등을 묶어서 관리하기 위해 구조체를 사용하고 있습니다.

typedef struct _gpuMemory {
    void *memory;
    size_t pitch;
    size_t widthByte;
    size_t height;
} gpuMemory;

////////////////////////////////////////////////////////////////////
void imageCopyBetweenGPU(gpuMemory *srcMem, gpuMemory *dstMem)
{
    ////// Skip
    ////// ...

    cudaMemcpy2D(dstMem->memory,    // dst Buffer (Device)
        dstMem->pitch,              // Pitch size of dst (dstMem->memory)
        srcMem->memory,             // src Buffer (Device)
        srcMem->pitch,              // Pitch size of src
        srcMem->widthByte,          // Width size of src
        srcMem->height,
        cudaMemcpyDeviceToDevice); // Direction of copy (IMPORTANT)
}

동일한 API이기 때문에 크게 다르지는 않으며, 가장 중요한 건 마지막 파라미터 입니다.
cudaMemcpyDeviceToDevice 를 입력해, GPU 메모리 내에서 메모리 복제가 일어나고 있습니다.

4. GPU에서 System으로 메모리 복제 (Device to Host)

화면에 직접 Rendering 하는 등의 경우 불필요할 수도 있지만, 처리된 데이터를 저장하거나 전송하는 등의 경우에서
GPU의 메모리를 다시 System Memory로 복제해야할 일이 생깁니다.
이 때에는 2의 예제와 거의 유사하되 src와 dst 파라미터의 위치만 바꿔서 입력합니다.

cudaMemcpy2D(hostMem,           // dst Buffer (Host)
        width * 3,                  // Pitch size of dst (hostMem)
        gpuMem->memory,             // src Buffer (Host)
        gpuMem->pitch,              // Pitch size of src
        gpuMem->widthByte,          // Width size of src
        gpuMem->height,
        cudaMemcpyDeviceToHost); // Direction of copy (IMPORTANT)

 

예제를 실행해보면 output.jpg에 Image가 잘 복사되어 저장된 걸 볼 수 있습니다.

2,3,4의 예제 코드에 대한 동작 개요도

 

5. Memcpy2D를 활용해 원하는 위치로 Crop 하기

링크의 실제 예제 코드를 활용해서 설명하겠습니다.
(예제 코드: https://github.com/z-wony/CudaPractice/blob/master/src/Practice_04_cudaMemcpy2D_crop/main.c)

나머지 모든 코드는 위의 예제와 동일합니다.
다만 imageCropUsingMemcpy() 함수의 Memcpy 동작을 유심히 보시죠.

void imageCropUsingMemcpy(gpuMemory *srcMem, int srcRoiX, int srcRoiY, int srcRoiWidth, int srcRoiHeight,
                          gpuMemory *dstMem, int dstX, int dstY)
{
    /////////// Skip
    /////////// ...

    // Calculate starting point address using X, Y coordinates
    unsigned char *srcPtr = (unsigned char *)srcMem->memory + (srcRoiY * srcMem->pitch) + (srcRoiX * 3);
    unsigned char *dstPtr = (unsigned char *)dstMem->memory + (dstY * dstMem->pitch) + (dstX * 3);

    // Memory Copy
    cudaMemcpy2D((void *)dstPtr,    // dst Buffer (Device)
        dstMem->pitch,              // Pitch size of dst (dstMem->memory)
        (void *)srcPtr,             // src Buffer (Device)
        srcMem->pitch,              // Pitch size of src
        (size_t)srcRoiWidth * 3,       // Width size of src
        (size_t)srcRoiHeight,
        cudaMemcpyDeviceToDevice);  // Direction of copy (IMPORTANT)

우선 복제 전에, src의 메모리 주소 포인터 값을 원하는 (srcRoiX, srcRoiY) 좌표로 이동하게 하기 위해,
지난 장에 계속 했던 2D 메모리 주소 계산을 Pitch값을 활용해서 수행합니다.
여기서 유의할 점은 void *로 캐스팅하지 않고, 8bit 크기인 unsigned char 의  포인터 형으로 놓고,
좌표를 계산했습니다. 이유는 void *로 놓을 시, 제 컴퓨터가 64bit이므로 포인터에 숫자를 더해서 주소를 계산하는 연산에서
숫자 하나당 64bit 로 취급하기 때문에 원하는 주소가 나오지 않습니다.
dstPtr 역시 동일하게 메모리 주소를 바꾸어 계산합니다.

또한 Memcpy 자체에서 복제할 width, height 크기를 필요한 만큼만 지정합니다.
결국 Memcpy 동작 내부에서는 srcMem->pitch - (srcRoiWidth * 3) 만큼을 몽땅 Padding 영역일 것으로 생각하고
메모리 복제시 참조하지 않게 됩니다.
Memcpy의 내부 동작을 알고 이용해서 Crop을 구현하는 방법입니다.

5의 예제코드에 대한 동작 개요도