[CUDA] 03. Nvidia NPPi의 Remap 함수로 이미지 변형

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. Remap 기본 개념

Remap이란 특정 Mapping Table을 활용해 이미지를 완전히 자유롭게 반전시키는 방법입니다.

매핑 테이블의 자유도가 높아서, Crop, Resize, Rotate, Flip, 그 외 자유 변형에 활용할 수 있습니다.
OpenCV의 Remap을 활용하시던 분들이라면 쉽게 이해하고 활용할 수 있습니다.
그러나, 저는 이 개념이 익숙하지 않아서 NPPi의 Remap을 이해하는 데에 꽤 많은 시간이 쓰였습니다.

Remap을 위해서는 X Map과 Y Map이 필요합니다.
이 Map들은 결과물로 나올 이미지와 동일한 2차원 크기의 버퍼로 구성되며,
색상 값이 아닌 원본 이미지의 좌표 값을 Float 형태로 담습니다.

말로만 표현하기는 너무 복잡하니, 이미지를 1/4 크기로 Resize 하는 Remap 예제를 가져 와보겠습니다.

1/4 Resizing을 위한 Remap 동작 흐름도

원본 이미지가 Remap 함수를 통과하며 1/4 (가로 1/2, 세로 1/2 크기) 로 작아지는 것을 볼 수 있습니다.

이 때, 원본 이미지를 "가로 10 pixel", "세로 8 pixel"의 크기로 가정하고,
결과 이미지를 "가로 5 pixel", "세로 4 pixel"의 크기로 가정해 보겠습니다.
이 경우, X Map과 Y Map은 결과 이미지의 크기와 동일한 크기의 2차원 배열이 필요합니다.

이 X Map과 Y Map을 조합해 실제 결과 이미지의 각 픽셀마다 원본 이미지의 어떤 RGB 값을 가져와야할 지
우측 하단과 같은 Mapping Table이 생성됩니다.
(논리적으로 개념을 이해하기 위해 생성된다는 표현했습니다. 실제 연산에서는 굳이 별도로 필요하지 않습니다.)
이 Mapping Table의 좌표를 통해, 보라색 참조 선과 같이 원본의 픽셀 좌표를 엑세스 해 RGB 값을 가져옵니다.
가로 세로 각각 1/2 씩 줄였으므로, 한 칸씩 띄엄띄엄 좌표를 읽어오는 걸 알 수 있습니다.

본 예시에서는 데이터가 모두 정수로 똑 떨어졌는데, 위에 명시했다시피 Map의 데이터는 Float형 입니다.
이 말인 즉슨 딱 떨어지는 좌표가 아닌 어정쩡한 중간 좌표도 입력이 가능하다는 뜻입니다.
따라서 Remap 함수에는 반드시 Interpolation을 어떻게 할 지가 상수로 입력됩니다.

NPPi의 Interpolation 상수는 현재 Cuda 11 기준으로 다음과 같습니다.

typedef enum 
{
    NPPI_INTER_UNDEFINED         = 0,
    NPPI_INTER_NN                = 1,        /**<  Nearest neighbor filtering. */
    NPPI_INTER_LINEAR            = 2,        /**<  Linear interpolation. */
    NPPI_INTER_CUBIC             = 4,        /**<  Cubic interpolation. */
    NPPI_INTER_CUBIC2P_BSPLINE,              /**<  Two-parameter cubic filter (B=1, C=0) */
    NPPI_INTER_CUBIC2P_CATMULLROM,           /**<  Two-parameter cubic filter (B=0, C=1/2) */
    NPPI_INTER_CUBIC2P_B05C03,               /**<  Two-parameter cubic filter (B=1/2, C=3/10) */
    NPPI_INTER_SUPER             = 8,        /**<  Super sampling. */
    NPPI_INTER_LANCZOS           = 16,       /**<  Lanczos filtering. */
    NPPI_INTER_LANCZOS3_ADVANCED = 17,       /**<  Generic Lanczos filtering with order 3. */
    NPPI_SMOOTH_EDGE             = (int)0x8000000 /**<  Smooth edge filtering. */
} NppiInterpolationMode; 

이 Algorithm들이 어떻게 동작할 지는 다른 외부 검색을 통해 확인하시면 좋습니다.
(여기에도 일부 있네요. https://matplotlib.org/mpl_examples/images_contours_and_fields/interpolation_methods.hires.png)
그리고 Remap에서 모든 상수를 지원하지는 않습니다.
NPPI_INTER_NN NPPI_INTER_LINEAR NPPI_INTER_CUBIC NPPI_INTER_CUBIC2P_BSPLINE NPPI_INTER_CUBIC2P_CATMULLROM NPPI_INTER_CUBIC2P_B05C03 NPPI_INTER_LANCZOS 만 지원한다고 나와있네요.
(https://docs.nvidia.com/cuda/npp/group__image__remap.html Detailed Description 참고)

1/4 Resizing Sample을 코드 단위로 설명해 보겠습니다.

(전체 샘플 코드는 여기에 있습니다. https://github.com/z-wony/CudaPractice/blob/master/src/Practice_05_Nppi_Remap/main.c)

Source Image를 읽어오고 (https://z-wony.tistory.com/21)
NPPi로 이미지 버퍼 할당하는 (https://z-wony.tistory.com/20)
기초적인 동작은 생략하겠습니다.

이를 생략하면 (1) X Map 생성, (2) Y Map 생성, (3) Remap 수행. 3단계로 정리됩니다.

(1) X Map 생성

void _createXMap(int width, int height, Npp32f **xMap, int *steps)
{
    printf("====== Step 5-1. Alloc CPU memory ======\n"); 
    float *cpuMem = (float *)malloc(sizeof(float) * width * height);
    printf("====================================================\n\n");

    printf("====== Step 5-2. Set X Map data ======\n");
    for (int y = 0; y < height; y++) {
        for (int x = 0; x < width; x++) {
            float xCoordMap = x * 2.0f;  // Gathering data from twice scale position (e.g. 3,0 from 6,0)
            cpuMem[y * width + x] = xCoordMap;
        }
    }
    printf("====================================================\n\n");

    

우선 코드를 반으로 나눴습니다.
일반적인 Malloc으로 CPU Memory를 할당합니다.
이미지의 크기 만큼의 Float으로 구성된 2차원 배열이 필요합니다.
(편의상 1차원으로 Array로 만들었습니다. 이게 똑같다는 것 역시 지난 예제에서 설명했습니다.)

이후 이중 for문을 통해, X Map의 실질 참조 픽셀 데이터를 입력해줍니다.
X 좌표에 대한 값만 입력하는 것이며, "결과 이미지의 (X,Y)위치에서 참조할 원본 이미지의 X 좌표를 입력한다" 는 점만 명확히 이해하시면 됩니다.
X 표 기준으로는 1/2 축소를 진행하므로, 작아진 결과 이미지 기준으로 본인 X 좌표의 2배수로 원본 이미지의 X 좌표를 참조합니다.

    printf("====== Step 5-3. Alloc GPU memory ======\n");
    Npp32f *map = nppiMalloc_32f_C1(width, height, steps);
    printf("====================================================\n\n");

    printf("====== Step 5-4. cudaMemcpy2D (Host To Device) ======\n");
    cudaMemcpy2D(map,           // dst Buffer (Device)
        (size_t)*steps,         // Pitch size of dst (devPtr)
        (void *)cpuMem,         // src Buffer (Host)
        (size_t)(width * sizeof(float)), // Pitch size of src (imgBuffer)
        (size_t)(width * sizeof(float)), // Width size of src (imgBuffer)
        (size_t)height,
        cudaMemcpyHostToDevice); // Direction of copy (IMPORTANT)
    printf("====================================================\n\n");

    *xMap = map;
    free(cpuMem);
}

NPPi Remap의 Map 데이터는 GPU 메모리(정확히는 Npp32f * Type) 여야 합니다.
위에서 CPU 메모리를 활용한 이유는 데이터 셋업을 위해서였고, 실제 Map으로 활용하기 위한 메모리를
nppiMalloc_32f_C1() 으로 할당해 줍니다.

그런 뒤, cudaMemcpy2D로 CPU Memory의 데이터를 NPPi의 GPU Memory로 복사해 줍니다.

(2) Y Map 생성

void _createYMap(int width, int height, Npp32f **yMap, int *steps)
{
    printf("====== Step 6-1. Alloc CPU memory ======\n");
    float *cpuMem = (float *)malloc(sizeof(float) * width * height);
    printf("====================================================\n\n");

    printf("====== Step 6-2. Set Y Map data ======\n");
    for (int y = 0; y < height; y++) {
        for (int x = 0; x < width; x++) {
            float yCoordMap = y * 2.0f;  // Gathering data from twice scale position (e.g. 3,0 from 6,0)
            cpuMem[y * width + x] = yCoordMap;
        }
    }
    printf("====================================================\n\n");

    printf("====== Step 6-3. Alloc GPU memory ======\n");
    Npp32f *map = nppiMalloc_32f_C1(width, height, steps);
    printf("====================================================\n\n");

    printf("====== Step 6-4. cudaMemcpy2D (Host To Device) ======\n");
    cudaMemcpy2D(map,           // dst Buffer (Device)
        (size_t)*steps,         // Pitch size of dst (devPtr)
        (void *)cpuMem,         // src Buffer (Host)
        (size_t)(width * sizeof(float)), // Pitch size of src (imgBuffer)
        (size_t)(width * sizeof(float)), // Width size of src (imgBuffer)
        (size_t)height,
        cudaMemcpyHostToDevice); // Direction of copy (IMPORTANT)
    printf("====================================================\n\n");

    *yMap = map;
    free(cpuMem);
}

Y Map은 X Map과 동일합니다.
X Map과 동일하게, "결과 이미지의 (X,Y)위치에서 참조할 원본 이미지의 Y 좌표를 입력한다" 라고 이해하면 됩니다.
yCoordMap에 입력되는 데이터가 무엇을 의미하는지만 명확히 이해하기 바랍니다.

(3) Remap 수행

NPPi Remap중 8bit Unsigned 3채널 이미지를 위한 nppiRemap_8u_3CR 함수 프로토타입 입니다.

nppiRemap 함수 Prototype

위에 서술한 X Map, Y Map, Interpolation 외에도, srcSize, srcROI, dstSizeROI을 입력해주어야 합니다.

    NppiSize srcSize = { srcWidth, srcHeight }; 
    NppiRect srcRoi = { 0, 0, srcWidth, srcHeight };
    NppiSize dstSize = { dstWidth, dstHeight };

    printf("====== Step 7. Try Remap (nppiRemap_8u_C3R) ======\n");
    NppStatus ret = nppiRemap_8u_C3R(srcMem->memory,
                     srcSize,
                     (int)srcMem->pitch,
                     srcRoi,
                     xMap,
                     xMapSteps,
                     yMap,
                     yMapSteps,
                     dstMem->memory,
                     (int)dstMem->pitch,
                     dstSize,
                     NPPI_INTER_LINEAR);
    if (ret)
        printf("Remap error (%d)\n", (int)ret);
    printf("====================================================\n\n");
    nppiFree(xMap);
    nppiFree(yMap);

Map 만 잘 준비되면, 다른 파라미터들은 부수적인 것들입니다.

Interpolation 상수는 우선은 Linear로 지정하였습니다.
이 상수는 본인의 실제 실행환경에 맞춰 다양하게 조정해보시기 바랍니다.

본 예제는 좌표 매칭도 매우 단순하지만, 현업에서 소수점 단위로 정교한 좌표매칭을 하는 경우,
Interpolation 알고리즘에 따라 결과 이미지의 품질도 다양하게 달라지며,

밀리초, 나노초 단위 성능에 민감한 S/W의 경우
알고리즘별 수행시간도 다르므로 프로파일링을 통해 고민해봐야할 문제입니다.

샘플을 수행하면 output.jpg가 1/4 사이즈로 작아진 것을 보실 수 있습니다.

 

Resize만 되는게 아니라는 걸 보여드리기 위해 몇가지 예제를 더 넣었습니다.

2. Flipping 예제

Horizontal Flip

Hozirontal Flip을 위한 Remap 동작 개념도

대략적인 흐름은 이렇습니다.
위의 Resize 예시와 유사한데, Y Map은 원본과 동일한 배치를 갖고, X Map만 X 좌표를 반전시킨 것을 볼 수 있습니다.
이렇게 되면 이미지의 X 좌표만 거꾸로 참조하게 되므로 뒤집힌 이미지를 생성하게 됩니다.

전체 Sample 코드는 여기에 있습니다.
(https://github.com/z-wony/CudaPractice/blob/master/src/Practice_06_Nppi_Remap_flip_horizontal/main.c)

(1) X Map 생성

void _createXMap(int width, int height, Npp32f **xMap, int *steps)
{
    printf("====== Step 5-1. Alloc CPU memory ======\n");
    float *cpuMem = (float *)malloc(sizeof(float) * width * height);
    printf("====================================================\n\n");
    
    printf("====== Step 5-2. Set X Map data ======\n");
    for (int y = 0; y < height; y++) {
        for (int x = 0; x < width; x++) {
            float xCoordMap = (float)(width - x - 1);  // Invert data using inversion of X coordinate
            cpuMem[y * width + x] = xCoordMap;
        }
    }   
    printf("====================================================\n\n");

X Map에서 (width - x - 1)로 참조 좌표를 뒤집어주게 됩니다.

(2) Y Map 생성

void _createYMap(int width, int height, Npp32f **yMap, int *steps)
{
    printf("====== Step 6-1. Alloc CPU memory ======\n");
    float *cpuMem = (float *)malloc(sizeof(float) * width * height);
    printf("====================================================\n\n");

    printf("====== Step 6-2. Set Y Map data ======\n");
    for (int y = 0; y < height; y++) {
        for (int x = 0; x < width; x++) {
            float yCoordMap = (float)y;     // Nothing to do
            cpuMem[y * width + x] = yCoordMap;
        }
    }
    printf("====================================================\n\n");

Y Map은 정상적으로 y 값을 그대로 입력해 준 것을 볼 수 있습니다.

 

Vertical Flip

Vertical Flip을 위한 Remap 동작 개념도

Vertical Flip은 위의 Horizontal Flip과 반대로, "X Map은 x 좌표 그대로", "Y Map은 반전 좌표"를 주게되면 됩니다.

샘플 코드는 다음과 같습니다.
(https://github.com/z-wony/CudaPractice/blob/master/src/Practice_07_Nppi_Remap_flip_vertical/main.c)

3. 보너스 응용

NPPi Remap의 자유도를 보여드리기 위해, 다이아몬드로 변형하는 샘플을 작성해 보았습니다.

예제 코드: https://github.com/z-wony/CudaPractice/blob/master/src/Practice_08_Nppi_Remap_Diamond/main.c

결과물은 다음과 같습니다.

Diamond Remap 결과물

이 예제 코드는 그리 깔끔하게 짜여지지는 않았습니다.

Remap을 이해하신다면 이런류의 변형이 자유롭다는 걸 보여드리기 위한 예제이며,
각자의 방식으로 자유롭게 활용하시기 바랍니다.