[CUDA] 01. 2D Memory Buffer 할당하기

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. Image Data의 구조

 1-1. 다룰 데이터의 형태

압축이 해제된 상태의 단일 이미지 (또는 영상 중 한 프레임) 데이터를 이야기 하고자 합니다.
JPEG, PNG등 고유의 압축 방식을 압축 해제하면, 모두 동일한 형태의 데이터가 나옵니다.
동영상 또한 연속된 사진을 고유의 압축방법으로 압축한 것이기 때문에 Decoding 시 한 Frame에 대해서는
사진과 동일한 형태의 데이터가 나옵니다.

※ 물론 Color Space는 다를 수 있습니다.
그럼에도 불구하고 이 이야기를 하는 이유는 의외로 이를 받아들이지 못하는 분들이 많기 때문입니다.
Video Processing 관련 개발을 하다보면, 현업에서조차 관련지식이 부족한 엔지니어에게서
영상인데 왜 '사진'으로 처리하냐 라는 질문을 받을때가 꽤 있습니다. 슬픈 일이죠.

(다시 본론으로 돌아와서)
따라서 이 데이터는 해상도 크기 만큼의 각 픽셀의 색상 데이터를 모두 사람이 눈으로 확인할 수 있는 형태입니다.
물론 머리로 읽고 상상할 수는 없지만, 최소한 (X,Y) 픽셀이 무슨 색일지 직관적으로 계산할 수 있습니다.

1-2. Color Space

색공간(= Color Space)이라 함은 위에서 말한 각각 픽셀이 어떤 색일지 표현하는 표현방식 입니다.
흔히 알고있는 RBG Color Space라면, 색 표현은 순서대로 빨강, 파랑, 초록 순서대로 배치됩니다.
보통 픽셀당 24비트 Color Depth 가 할당되므로, 각각 8비트 (0~256)의 수치로 색을 표현합니다.
배치 순서에 따라, BGR이 되기도 하고 투명도(Alpha)를 더해서 RGBA 일때도 있습니다.
동영상 데이터였다면 색의 3원색이 아닌, 휘도와 색차 라는 별도 표현 방식인 YUV 형태를 띄며,
세부 형태에 따라 I420, NV12 등의 색공간으로 표현될 수 있습니다.
(RGB 쉬운 설명: https://ko.wikipedia.org/wiki/RGB)

1-3. Packed format과 Planar format

이 두 표현법은 꼭 알고있어야 멀티미디어 쪽 이것저것을 다루는데에 유리합니다.
Image Data뿐 아니라, Audio Data를 다룰 때도 등장하기 때문입니다.

Packed는 한 픽셀에 대한 정보가 연이어 등장하는 방식 입니다.
- RGB Color Space가 이에 해당하며, 예를 들어 첫번째 픽셀이 빨강일때 데이터를 바이트 단위로 읽어보면,
  255, 0, 0 이 순서대로 오게 됩니다.
- Audio 데이터의 경우, 첫번째 음에 대해 각 채널별 데이터 (e.g. 스테레오면 좌/우) 를 연이어 표현하게 됩니다.

 

RGB Format (Packed)의 Image Data 형태

 

Planar는 전체 데이터에 대해 한종류의 정보를 쭉 늘어놓고, 다음 정보를 쭉 늘어놓는 방식입니다.
- YUV의 I420이 이에 해당하며, 예를 들어 전체 Image에 대한 Y성분 데이터가 쭈욱 표현된 후, U 성분이 쭈욱, V 성분이 쭈욱 표현됩니다.
- Audio 데이터의 경우, 약속된 단위 사이즈 만큼의 첫 채널 데이터가 쭈욱 나오고, 연이어 채널 데이터가 일렬로 들어섭니다.

 

YUV Planar에 대한 대략적인 구조 (실제는 좀더 복잡할 수 있음)

 

이러한 기본 개념들을 알고 있어야 NPPi 든, ffmpeg이든 뭐든 다룰 때 시행착오를 줄일 수 있습니다.
저는 대부분의 예제들을 RGB Color Space에 맞춰 다뤄보겠습니다.

---

2. RGB Data의 Access 방법

Image Data는 가로 X 세로 즉, 2차원 데이터입니다. 이는 2차원 배열로 표현될 수 있으며,
프로그래밍 기초 시간에 배우듯이 2차원 배열은 실질적으로 그냥 1차원 배열로 할당되며 그대로 Access 할 수 있습니다.

다음의 두가지 Memory Access 방식은 결과적으로 동일합니다.
중요한 건 두번째 방식의 의미입니다.

///////////////// 다차원 메모리 구조 //////////////////
// 640 x 480 Size RGB Image Data
unsigned char imageData[480][640][3];

// Set 7,2 Pixel to Red
int x = 6, y = 1, width = 640;
imageData[y][x][0] = (unsigned char)255; // Red
imageData[y][x][1] = (unsigned char)0;   // Blue
imageData[y][x][2] = (unsigned char)0;   // Green
/////////////////////////////////////////////////////

/////////////// 1차원 배열로 Access 시 /////////////////
// 640 x 480 Size RGB Image Data
unsigned char imageData[480 * 640 * 3];


// Set 7,2 Pixel to Red
int x = 6, y = 1, width = 640;
int yPos = y * width * 3;
int xPos = x * 3;
imageData[yPos + xPos + 0] = (unsigned char)255; // Red
imageData[yPos + xPos + 1] = (unsigned char)0;   // Blue
imageData[yPos + xPos + 2] = (unsigned char)0;   // Green
/////////////////////////////////////////////////////

특정 Y 좌표로 이동하기 위해서는 Pointer를 'width 길이에 해당하는 만큼씩 y번' 이동해야 했고,
특히 RGB Packed 형태였기 때문에 * 3 한 만큼 (정확히는 X 3바이트) 이동합니다.
그 다음 X에 해당하는 만큼 * 3 만큼을 뒤로 이동하면 원하는 좌표 지점의 픽셀 데이터를 Access 할 수 있습니다.

그런데 여기서 만약에, Image 의 가로 길이 (width)는 640 인데,
담을 메모리는 가로 크기를 넉넉잡아 700만큼 할당하고 싶다면 어떻게 될까요?

// 640 x 480 Size RGB Image Data (But width size is 700)
unsigned char imageData[480 * 700 * 3];

// Set 7,2 Pixel to Red
int x = 6, y = 1, width = 640, widthSize = 700;
int yPos = y * widthSize * 3;
int xPos = x * 3;
imageData[yPos + xPos + 0] = (unsigned char)255; // Red
imageData[yPos + xPos + 1] = (unsigned char)0;   // Blue
imageData[yPos + xPos + 2] = (unsigned char)0;   // Green

기존에 yPos를 구하기 위해 width를 곱했는데, widthSize를 곱하게 됩니다.
세로로 한 라인을 넘어가기 위해서는 widthSize * 3 이 필요하다는 결론이 나오고,
이를 좀 더 있어보이게(?) 바꿔보면

// 640 x 480 Size RGB Image Data (But width size is 700)
unsigned char imageData[480 * 700 * 3];

// Set 7,2 Pixel to Red
int x = 6, y = 1, width = 640, widthSize = 700;
int lineSteps = widthSize * 3;
int yPos = y * lineSteps;
int xPos = x * 3;
imageData[yPos + xPos + 0] = (unsigned char)255; // Red
imageData[yPos + xPos + 1] = (unsigned char)0;   // Blue
imageData[yPos + xPos + 2] = (unsigned char)0;   // Green

이 개념은 실제로 활용되는 개념입니다.
메모리 버퍼가 실제 Width와 동일하게 할당되지 않을 때가 많고, 따라서 한 Line을 넘어가기 위한 단위 크기를 API에 입력하도록 요구하는 경우가 매우 많습니다.
흔히 이런 길이는 Steps, Line Step, Stride 등으로 불리며, Byte Size 단위로 입력받는 경우가 많습니다.
(예제처럼 RGB Color Space면 픽셀당 3바이트 라는 점을 고려해야하구요.)
API Description에 구체적으로 어떻게 입력해야하는지 (바이트 단위인지, 아니면 픽셀 수만 입력하면 되는지) 기재되어 있는 경우 그를 따르거나, 안되면 약간의 시행착오를 겪어서 알아내야합니다.

이러한 구조를 갖는 이유는 두가지입니다.
1) 색상 요소별로 데이터 크기가 다른데, 동일한 2차원 메모리 규격을 활용할 때.
2) 버퍼의 크기를 H/W Architecture가 연산하기 좋은 형태로 할당할 때.

1) 색상 요소별로 데이터 크기가 다른 경우

RGB는 Packed Format이기 때문에 동일할 수 밖에 없지만, Planar Format인 YUV는 그렇지 않습니다.
인간이 식별하기 좋은 Y 요소는 해상도 크기만큼 동일하게 픽셀당 1 바이트의 데이터를 갖는 반면,
인간이 둔감한 색차정보인 U,V 요소는 4픽셀 (가로 세로 2x2) 당 1 바이트의 데이터 크기를 갖습니다.
따라서 2차원 메모리에서 공간을 표시하면 대략 이런 방식이 됩니다.

 

YUV 중 IMC3 포멧의 메모리 배열

 

(출처: docs.microsoft.com/en-us/windows/win32/medfound/recommended-8-bit-yuv-formats-for-video-rendering)

메모리 배열의 가로 크기가 이미지의 Width 만큼 할당될 경우,
Y 요소에 대해 (X,Y) 좌표의 픽셀 값을 찾을 때는 정확한 메모리 크기만큼 Line Step을 활용할 수 있지만,
U 또는 V 요소에 대해 (X,Y) 좌표의 픽셀 값을 찾을 때에는 실제 차지하는 공간보다 큰 영역 만큼의 Line Step을 활용하는 셈입니다.

 

YUV 중 YV12 포멧의 메모리 배열

 

(출처: docs.microsoft.com/en-us/windows/win32/medfound/recommended-8-bit-yuv-formats-for-video-rendering)

반대로 위와 같이 여분 공간(Space 혹은 Padding 이라고 부름)이 없는 경우,
U와 V의 Line Step은 Width / 2가 됩니다. 즉, 색상 요소별로 다른 Line step을 갖게되는 거죠.
YUV Color Space에서 갖는 메모리 버퍼의 Line Step은 Stride 라는 용어가 가장 많이 쓰이는 듯 합니다.

Chromium의 libyuv API를 한번 보시죠.

 

libyuv의 API (I420을 ARGB로 컨버팅)

 

(출처: chromium.googlesource.com/libyuv/libyuv/+/refs/heads/master/include/libyuv/convert_argb.h)

Y,U,V 요소 뿐만 아니라, ARGB 에 대해서도 Stride를 요구하고 있습니다.
여기서의 특징은 I420은 Planar 형식이기 때문에 각각의 Stride가 필요했고, ARGB는 Packed여서 Stride가 하나만 필요했다는 점 입니다.

2) 버퍼의 크기를 H/W Architecture가 연산하기 좋은 형태로 할당할 때

예를 들어 RGB Color Space로 표현되는 이미지의 크기가 640x480 이라고 할 때,
한 행의 Byte 크기는 1920 Byte가 됩니다.
컴퓨터 공학을 전공했다면, 한 행의 크기가 2048 Byte 였다면 더 자연스러웠을 거라는 생각을 하게될 겁니다.
그냥 2의 제곱근이라 마음이 편안한게 아니라, Data 처리를 행 단위로 하게될 확률이 높으므로,
성능적으로도 이득을 보게 되어있습니다.

이게 왜 그런지는 짧게 설명해보겠습니다. (컴퓨터 구조나 운영체제 과목에 자세한 이야기가 나옵니다.)
간략하게 요약하면 폰 노이만 구조에서 기억장치와 처리장치가 분리되어있고,
처리장치로 작업을 가져오기 위해 Caching을 하게 되는데, 이 가져오는 Block 단위가 2의 제곱근 이기 때문에,
거기에 꼭 맞는 단위로 데이터 영역을 할당하는게 성능적으로 유리해집니다.
(안그러면 특정 행에서는 Memory Caching 동작을 더 많이 하게되는 상황이 생깁니다.)
관련 설명: (ssoonidev.tistory.com/35)

CUDA 역시 2차원 메모리 할당 시, 요청한 Width 크기 대신,
더 빠르게 연산할 수 있는 적정 크기로 Padding을 덧붙여 메모리를 할당합니다.
CUDA 메모리 할당 관련 API에서는 이 Line Step 크기를 Pitch 라고 부릅니다.

 

CUDA 의 2D 메모리 할당 (640 x 480 크기의 RGB 버퍼 예시)

 

---

3. CUDA Memory Allocation 방법

위의 개념이 모두 이해되었다면, 실제 API 사용 방법은 단순합니다.

#include <cuda_runtime.h>

//////////////////////////////////

cudaError_t ret;
void *mem = NULL;
size_t pitch = (size_t)0;
size_t width = (size_t)(640 * 3);
size_t height = (size_t)480;

ret = cudaMallocPitch(&mem, &pitch, width, height);
//////////////////////////////////

메모리를 할당 시 cudaMallocPitch API를 활용하게 되는데, 
여기에 입력하는 width 값은 픽셀 크기가 아니라 바이트 사이즈 여야 합니다. (예제는 RGB 이므로 3을 곱했습니다.)
Output으로는 Memory Pointer와 pitch 값이 나오게 되는데,
향후 모든 버퍼 관련 연산에서 Pitch 값을 반드시 함께 입력하게 되어있습니다. (그래야 계산이 되니까요.)
따라서 둘을 함께 묶어서 관리하는게 좋습니다.

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

---

4. NPPi Memory Allocation 방법

NPPi는 Nvidia의 2D Image 처리 라이브러리입니다.
이미지 변형을 GPU로 빠르게 처리하고 싶을 때 활용하면 좋습니다.
NPPi를 처음 접하면 API 개수와 네이밍에 압도되고, 그만보고 싶어지는데, 따라서 네이밍 룰을 먼저 이해하고갈 필요가 있습니다.

1) NPPi Naming Rule

(docs.nvidia.com/cuda/npp/general_conventions_lb.htmldocs.nvidia.com/cuda/npp/nppi_conventions_lb.html)

C 뒤에 숫자가 붙을 때: Packed format으로 몇채널인지 입니다. (예를 들어 RGB면 C3, RGBA면 C4)
P 뒤에 숫자가 붙을 때: Planar format으로 몇채널인지 입니다. (이 경우 각각의 Line Step을 입력하게 됨)
A가 붙어있을 때: Alpha (투명도) 채널의 영향을 무시하고 이미지를 처리
R이 붙어있을 때: 이미지 내에서 특정 ROI (Region of Interest) 영역에 대해 처리하는 API라는 의미

API 이름 형식은 실제로 <API명>_<색상요소의 자료형>_<위에 설명한 이것저것> 입니다.

Resize 함수이고 + 색상 하나당 8bit unsigned로 표현되고 + 3채널 Packed format인 경우,
nppResize_8u_C3R 이 됩니다.

이정도만 알면 추후 본인의 목적에 맞게 NPPi의 메모리도 쉽게 할당할 수 있습니다.

2) NPPi로 메모리 할당하기

위와 동일하게 3채널의 640 x 480 크기 이미지 버퍼를 생성한다고 했을 때,
(여기서는 좀 더 구체적으로 한 색상당 8bit unsigned의 자료형이라고 했을 때)

#include <nppi.h>

//////////////////////////////////////////////////////
Npp8u *mem = NULL;
int stepBytes = 0;
int width = 640;
int height = 480;

mem = nppiMalloc_8u_C3(width, height, &stepBytes);
//////////////////////////////////////////////////////

 

여기서 재미있는 점은, cudaMallocPitch보다 상위레벨 API 답게,
Width 값을 바이트로 입력하지 않고 Pixel 수로 입력했습니다.
대신 API 이름에 픽셀당 3바이트임이 명시되어 있습니다.

API Description도 보다 명쾌하게 정리되어 있습니다.

 

nppiMalloc_8u_C3

 

(출처: docs.nvidia.com/cuda/npp/group__image__memory__management.html#ga64e742b7897ffab51c455b1401164e86)

사실 이 API는 내부적으로 cudaMallocPitch와 동일한 동작을 하게 됩니다.
따라서 교차 사용이 가능하다는 특징이 있습니다.
cudaMallocPitch 로 할당한 메모리를 NPPi에 활용해도 문제되지 않고,
nppiMalloc 으로 할당한 메모리를 다른 NVENC 같은 API에 그대로 입력해도 문제되지 않습니다.

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

---

4. 정리

단순히 2차원 메모리를 할당하는 데에도, 꽤 많은 지식이 필요했습니다.
사실 1회성으로는 모르고 그냥 쓸수도 있는 부분이지만, 정확히 이해하지 못하고 활용하면
언젠가 원인모를 오동작에 해메는 일이 발생할 수 있습니다.
또 깊이 이해해야만, 처음 접하는 동종 API나 유사한 S/W Framework를 마주했을 때, 통찰력을 얻을 수 있습니다.

다음에는 이렇게 할당한 GPU상의 메모리를 Access하거나 복사하는 법을 소개하겠습니다.