비밀기지

이전 글:

2021.05.03 - [코딩/PyTorch] - PyTorch로 MATLAB imresize (bicubic interpolation) 구현하기 (1) - 이미지 resizing 배경 지식

그래서 img[-0.12, -0.12]는 무엇인가?

사람에 따라서는 음의 좌표에 거부감이 들 수 있으므로... img[0.12, 0.12]가 무엇인지를 찾는 문제로 바꿔서 생각해도 상관없고.

아무튼 이 개념은 내가 학부시절 때 열심히 듣고, 좋아했던 수업 중 하나인

신호 및 시스템 Signals and Systems를 수강하면 어느 정도 감이 온다.

그런데 아무래도 수업에서는 학술적인 이론을 위주로 공부하다 보니, ideal reconstruction을 비중 있게 다루고

나머지 방법들에 대해서는 조금 소홀한 면이 있다고 느꼈다.

신호 및 시스템 정리를 지금 와서 하기에는 좀 그렇고,

이상적인 경우 우리가 갖고 있는 discrete $H \times W$ 이미지로부터 완벽한 continuous 신호를 복원할 수 있지만,

이게 가능한 경우는 현실적으로 극히 드물다.

따라서 다양한 트릭을 사용해서 continuous 이미지를 근사해서 나타내게 되는데, 이를 resampling이라고 하고

비전 분야에서는 interpolation (보간법)이라는 표현을 더 많이 사용하는 것 같다.

가장 쉽게 생각할 수 있는것은, 이전 포스트에도 언급했다시피 (0.12, 0.12)와 가장 가까운 점의 값을 사용하는 것이다.

이를 굉장히 직관적인 네이밍으로 nearest neighbor (NN)라고 하는데,

조금만 생각해보면 (0.12, 0.12)를 포함하는 (-0.5, -0.5)를 왼쪽 상단 꼭짓점,

(0.5, 0.5)를 오른쪽 하단 꼭짓점으로 하는 정사각형 내부의 모든 위치가 (0, 0)의 값을 참조하게 된다는 것을 알 수 있다.

(경곗값을 포함시키는지 여부는 별로 중요하지 않다.)

아래 그림은 $5 \times 5$ 이미지(?)에서 임의의 실수 좌표값에 대해 NN interpolation 기반으로 값을 취득한다면

어떤 값이 나오는지를 도식화한 것이다.

사실 NN interpolation은 continuous 이미지의 근사 형태라고 생각하기에는 조금 힘든 게,

저렇게 해서 나오는 결과는 계단식으로 discontinuous 하기 때문이다.

만약 이미지의 색상 값을 위치 $(x, y)$에 대한 함수로 나타낸다면 (implicit representation? ㅎㅎ)

NN interpolation으로 얻는 함수는 임의의 점에서 미분 불가능하거나, 미분 값이 0이다.

비전 분야에서는 아무래도 별 상관은 없지만

일단 보기에도 썩 좋지는 않고, 그래픽스 등 다른 분야에서는 수학적인 성질 또한 중요시 여기기에 많이 사용되지는 않는다.

굳이 수학적으로 표현하고자 하면 아래와 같을 듯 ($\lfloor \cdot \rceil$은 반올림.).

$q(x, y) = p( \lfloor x \rceil, \lfloor y \rceil )$.

여기서 $p$는 해당 위치에서의 색상 값 (정수 좌표에서 잘 정의되어 있는)이고

$q$는 우리가 색상 값을 알고 싶은 query 포인트라고 생각하면 되지만,

다른 application에서는 다른 의미를 가질 수도.

조금 더 나은 비주얼을 보여주면서 적당히 간단한 방법은 linear interpolation이다.

역시 이름이 직관적인데, 실수 좌표 (x, y)에 대해 query가 들어오면

해당 점을 둘러싼 4개의 정수 좌표를 가진 점들의 선형 결합으로 (x, y)에서의 색상 값을 만드는 형태이다.

말로만 정리하면 조금 애매하기에 그림을 하나 만들었는데

일단 $x, y$가 정수일 때 같은 edge case는 쿨하게 무시해야지.

$\lfloor \cdot \rfloor$은 내림이다.

아무튼 저렇게 4개의 점에 대해서, 일단 위 2개와 아래 2개의 점들을 query의 $x$ 좌표를 기준으로 내분하고,

각각 $u_1, u_2$라고 한다.

편의상 $x - \lfloor x \rfloor = a, y - \lfloor y \rfloor = b$라고 두면

$u_1 = (1 - a) p_1 + a p_2,$

$u_2 = (1 - a) p_3 + a p_4,$

가 되고 여기서 $u_1$과 $u_2$를 query의 $y$ 좌표를 기준으로 한 번 더 내분해서 최종 색상 값을 얻는다.

$q = (1 - b) u_1 + b u_2.$

이렇게 $x$와 $y$ 방향으로 2번의 interpolation을 하기 때문에 bi-linear interpolation이라고도 부른다.

결과를 보면 NN interpolation 대비 장족의 발전을 이루었는데,

여전히 수학적인 특성은 좋지 못하다.

미분 값이 유의미한 정보를 담고 있기는 하지만,

정수 좌표 픽셀들의 위치에 뾰족점이 존재해서 항상 미분 가능하지는 않다.

비전 분야에서는 정말 별 의미가 없지만,

이렇게 수학적인 특성을 고려해서 interpolation을 구현하려면 상당히 문제가 복잡해진다.

Cubic interpolation 같은 경우는 항상 미분 가능한 특성을 지니도록 설계가 되었는데,

구하는 방법이 굉장히 복잡하다 (자세한 내용은 그래픽스로).

다행히 상당한 정확도와 괜찮은 효율로 cubic interpolation을 근사하는

cubic convolution (혹은 filtering?) 알고리즘이 있는데,

일반적으로 bicubic interpolation이라고 부르는 방법은

이러한 cubic convolution을 $x$와 $y$ 방향으로 2번 적용한 것이다.

Bilinear에서 주변 4개의 점들을 사용한 것에서 더욱 발전하여

Bicubic은 주변 16개의 점들을 사용한다.

어렵게 생각할 것 없이, query를 중심으로 $4 \times 4$ window를 배치하여

해당 window 안에 들어가는 점들의 linear combination으로 query 포인트의 색상을 정한다.

식으로는 $q = \frac{1}{Z} \sum_{i=1}^{16}{w_i p_i}$이고, 여기서 $w_i$는 각 픽셀의 기여도, $Z = \sum{w_i}$이다.

그러면 중요한 것은 각 픽셀들의 상대적인 weight (혹은 contribution)을 어떻게 정하는지이다.

원리는 나도 모르겠지만 (cubic convolution과의 오차를 최대한 줄이는 방향으로 설계된 듯하다.)

아래와 같이 $w_i$를 계산한다.

$\begin{split}
w_i &= b\left( x_i - x \right) b\left( y_i - y\right), \text{where} \\
b\left( v \right) &= \begin{cases}
1.5\lvert v \rvert^3 - 2.5 \lvert v \rvert^2 + 1 & \text{for}\ \lvert v \rvert \leq 1, \\
-0.5 \lvert v \rvert^3 - 2.5 \lvert v \rvert^2 + 4 \lvert v \rvert - 2 & \text{for}\ 1 \lt \lvert v \rvert \leq 2, \\
0 & \text{otherwise}.
\end{cases}
\end{split}$

당연히 $(x_i, y_i)$는 $p_i$의 좌표이다.

라이브러리의 구현에 따라 계수들은 조금씩 변경될 수 있으며 (ex. OpenCV) 상기 계수들은 MATLAB 기준이다.

Interpolation 결과물을 보면 앞의 두 개보다 훨씬 둥글둥글하고 부드러운 모양인 것을 확인할 수 있다.

Query 포인트에서의 색상을 $z$축이라고 생각한다면, 각 interpolation 방법들을 아래와 같이 시각화하는 것도 가능하다.

수학적으로 써먹기 힘들어 보이는 NN, Linear에 비해 Cubic이 상당히 부드럽게 표현되는 것을 알 수 있다.

사실 여기까지는 어딜 가도 잘 설명되어 있는 내용인데,

이상하게 이것만 가지고 구현을 하면 절대로 MATLAB과 같은 bicubic interpolation 결과를 얻을 수 없다.

가장 핵심적인 포인트는 antialiasing (AA)이다.

AA는 신호 처리에서 고주파 성분을 제한된 bandwidth로 표현할 때 생기는 artifact를 줄여주는 기법으로,

일반적으로 주어진 신호를 subsampling 하기 전에 low-pass filter (LPF)를 적용하는 것으로 구현된다.

MATLAB imresize는 이미지를 확대할 때는

다른 라이브러리들이랑 똑같은 방법으로 bicubic interpolation을 구현하는데,

축소할 때는 추가적인 AA 처리가 들어가고 이것이 차이의 결정적인 요인이다.

그러면 다른 라이브러리들은 왜 AA 처리를 안 하는가?라는 의문이 당연히 들지만,

일반적으로 NN을 제외한 bilinear, bicubic 등은 일종의 LPF라고 생각할 수 있으며

해당 interpolation 방법을 포함한 resizing은 AA를 포함했다고도 말할 수 있다.

그런데 축소 scale이 커지면 고정된 크기의 interpolation 커널 (bilinear, bicubic 등)들로는

artifact를 방지하기 쉽지 않을 때가 있다.

따라서 MATLAB에서는 독특한 방식으로 scale에 adaptive 하게 AA 처리를 하는데,

$s$배만큼 축소하는 경우

1. Bicubic interpolation에서 window 크기를 가로세로 $s$배만큼 키운다.
   즉, 더 많은 주변 픽셀들이 query 포인트의 색상을 정하는 데에 사용된다.
2. 확장된 window에 맞춰, $w_i = b\left( x_i - x \right) b\left( y_i - y \right)$ 대신 $w_i = b\left( \frac{x_i - x}{s} \right) b\left( \frac{y_i - y}{s} \right)$를 사용한다.

실제로 4배만큼 줄이는 경우를 비교해보면, antialiasing이 있을 때와 없을 때 차이가 상당히 많이 난다.

코드상에서는

>> x_down = imresize(x, 1 / 4, 'bicubic', 'antialiasing', false);

이렇게 비활성화가 가능한데,

기본으로 켜져 있기 때문에 큰 의식을 하지 않은 것 같다.

아무튼 이정도면 실제 구현에 필요한 세부 사항은 다 명세를 했고,

어떻게 깔끔하게 코드에 담아내는지는 다음 포스트부터 정리해야지.

[코드]

일반적으로 Image Super-Resolution (SR) 모델을 학습할 때는

많은 양의 high-resolution (HR) 이미지를 준비한 뒤 이를 임의의 downsampling 방법으로 줄여

input low-resolution (LR)을 만든다.

그 후 이렇게 만들어진 pair들을 통해 supervised learning을 진행한다.

해당 방법을 real-world application에 적용하기에는 많은 문제가 있으나,

아무튼 전통적으로 내려오던 framework이기 때문에

어느 모델을 개발하던 이러한 과정은 최소 한 번 이상 거치게 된다.

그런데 SR 관련 연구를 하며 가장 거슬렸던 부분 중 하나는 저 임의의 downsampling 부분이다.

딥러닝이 대세로 나서기 이전에는 MATLAB이 비전 분야의 암묵적인 standard였던 것 같다.

그런데 여기서 문제는, MATLAB에서 이미지의 크기를 바꾸는 데에 사용하는 imresize 함수가

OpenCV나 다른 이미지 관련 라이브러리들과는 상당히 다른 구현을 가지고 있다는 점이다.

특히 문제가 되는 것은 기본으로 적용되는 bicubic 보간법인데,

아무리 결과물을 비교하고 설정을 바꿔봐도 다른 라이브러리가 만드는 LR 이미지와는 상당한 차이가 있었다.

SR 분야의 벤치마크 이미지들은 거의 다 MATLAB으로 만들어졌기 때문에,

만약 다른 라이브러리로 LR 이미지들을 만들고 학습에 사용하는 경우

training-test distribution에 차이가 생겨 벤치마크 성능이 저하되고,

기껏 알고리즘을 잘 만들어놓고도 제대로 된 evaluation을 할 수가 없다.

그렇기 때문에 SR 연구에서는 오랫동안 이미지 resizing이 필요한 부분을 우선 MATLAB으로 모두 처리하고,

그 후 적절한 framework를 사용하여 추가적인 resizing 작업을 시행하지 않고 학습을 진행했다.

안타깝게도 MATLAB과 Python의 접착성은 별로 좋지 않기 때문에,

Python 코드가 도는 도중 다이내믹하게 MATLAB의 imresize를 호출해서 적절한 작업을 수행한다는 것은

시도할 염두를 내지도 못했고 이 때문에 작업 효율이 상당히 떨어진다는 느낌을 받았다.

이러한 문제를 좀 완화하고자 몇 사람이 numpy로 MATLAB-compatible 하다고 주장하는 imresize 함수를 구현했고

Github에서 어렵지 않게 찾아볼 수 있다.

그런데 실제로 확인해보니, 다양한 문제가 있었는데

• 너무 느리다.
  numpy라서 기본적으로 GPU 연산을 지원하지 않을 뿐만 아니라, 최적화도 거의 되어있지 않다.
  imresize는 최소 수백, 수천번을 호출하는 함수인데
  이미지의 크기가 크면 눈에 띄게 성능이 저하되는 것을 느꼈다.
• MATLAB-compatible하지 않다.
  공개된 코드를 직접 돌려서 확인해보니 MATLAB과 거의 비슷한 결과가 나왔지만,
  특히 boundary 부분에서 MATLAB과는 많이 달랐다.
  이유도 찾았는데, 이는 후술.
• 미분 불가능.
  이는 큰 문제가 아니긴 한데 기본적으로 numpy로 구현되어 PyTorch layer 등에 넣을 수 없다.
  넣는 순간 computation graph가 깨져, PyTorch의 가장 큰 장점인 자동 미분을 사용할 수 없다.

개인적으로 저런 것들을 참고 쓰는 성격은 아니기 때문에,

작년에 (굉장히 늦은 타이밍이지만. 이렇게 정리 글을 쓰는 타이밍은 더 늦고.) 큰 맘먹고 직접 구현해봤다.

상당한 고생을 했지만 막상 만들고 나니 의외로 유용하게 쓰이기에, 잘 만들었다는 생각 또한 들었다.

본격적으로 구현 detail을 정리하기 전에,

이미지 resizing에 대한 기본 배경 지식을 정리해놓으면 나중에도 유용하게 참고할 수 있을 것 같다.

우선 좌표계를 정의하는 것이 정말 중요하다.

우리가 일반적으로 Python에서 $H \times W$ 이미지를 다룰 때는,

img[0, 0] (채널 생략)이 가장 왼쪽 상단, img[H - 1, W - 1]이 가장 오른쪽 하단의 픽셀을 의미한다.

수직, 수평 방향으로 이웃한 픽셀은 1만큼 떨어져 있기 때문에 이를 실제로 배치해보면 아래와 같다.

그런데 실제로 이미지를 이렇게 표현하면 상당히 난감한 것이, 이미지 (빨간 점선)의 크기가

$H \times W$가 아닌 $(H - 1) \times (W - 1)$이 나온다.

(딴소리지만, PyTorch의 interpolation 관련 함수에서 받는 align_corners 인수가 이러한 요인과 관련되어있다.

이 링크에서 조금 읽어보면 바로 느낌이 올 듯.)

따라서 편의상, 각 픽셀이 $1 \times 1$의 공간을 점유한다고 가정하면 (실제 픽셀의 정의와는 다르다!)

아래와 같이 이미지를 (-0.5, -0.5)와 (H - 0.5, W - 0.5) 사이의 영역에서 정의할 수 있다.

이렇게 정의한 이미지는 정상적으로 $H \times W$의 크기를 갖는다.

다른 방법도 다양하게 있겠지만, 이 방법이 가장 직관적인 것 같다.

사람에 따라서는 음의 영역에서 뭔가를 정의하고,

영역 경계의 좌표가 정수가 아니라는 것이 굉장히 불편할 수는 있을 것 같지만... (처음에는 나도 좀 불편했다.)

오히려 문제는 생각지도 못한 곳에서 발생했는데, 직접적으로 구현에 영향을 줄 정도는 아니었다.

기회가 되면 정리해야지.

아무튼 이렇게 기준을 정해놓으면 이후의 내용들을 정리하기가 비교적 편리하다.

이미지 resizing은 크게 두 가지 스텝으로 분리해서 볼 수 있는데, mapping과 resampling이다.

$h \times w$ input 이미지를 $H \times W$로 resizing 하고자 할 때 ($H > h$일 필요는 없다.),

생각해야 하는 것은 output 이미지의 [x, y] 위치가 input의 어디에 대응하는지?이다.

간편한 수식으로 어렵지 않게 변환이 가능하지만, 아래의 예시를 보면 조금 더 쉽게 이해가 가능하다.

$3 \times 3$에서 $4 \times 4$로 가는 경우를 생각하자.

일단 같은 크기의 정사각형 2개를 $3 \times 3$, $4 \times 4$로 쪼갠다.

두 정사각형을 겹쳐보면, output 이미지의 각 픽셀이

input 이미지의 어디에 대응하는지에 대한 비례 관계를 확인할 수 있다.

Boundary가 정수가 아니기 때문에, 계산이 조금 귀찮을 수는 있지만 구체적인 식은 아래와 같다.

$\textcolor{red}{x} = \frac{w(\textcolor{blue}{x'} + 0.5)}{W} - 0.5,$

$\textcolor{red}{y} = \frac{h(\textcolor{blue}{y'} + 0.5)}{H} - 0.5.$

여기서 $\textcolor{red}{x}$와 $\textcolor{red}{y}$는 input 이미지를 기준으로 한 좌표

$\textcolor{blue}{x'}$와 $\textcolor{blue}{y'}$는 output 이미지를 기준으로 한 좌표가 되는데

예를 들자면 output 이미지의 (0, 0)은 input 이미지의 (-0.12, -0.12)에 대응된다.

이제 드는 의문은, 그래서 (-0.12, -0.12)는 어디에 있는가?이다.

디지털 이미지는 regular grid 위에서 정의되기 때문에, img[0, 0]은 있어도 img[-0.12, -0,12]는 에러가 난다.

적당히 가장 가까운 픽셀의 값을 가져오면 되는 것 아닌가? 하는 의문이 들 수 있지만

(해당 방법이 nearest neighbor, 혹은 NN 보간법이다.)

그러면 자글자글한 계단 현상이 생기고 보기에 좋지 않다.

Bicubic 결과는 지나치게 뿌연 느낌이 든다고 생각할 수도 있지만, 뭐 일단은 그렇다.

이를 해결하기 위해서는 적절한 resampling이 필요한데, 다음 포스트에 정리해야겠다.

2021.05.04 - [분류 전체보기] - PyTorch로 MATLAB imresize (bicubic interpolation) 구현하기 (2) - resampling (interpolation)

이전 글:

2021.04.28 - [코딩/PyTorch] - 파이토치 PyTorch CUDA custom layer 만들기 (1) - CUDA 커널

2021.04.30 - [코딩/PyTorch] - 파이토치 PyTorch CUDA custom layer 만들기 (2) - 커널 빌드 및 Python 바인딩

2021.04.30 - [코딩/PyTorch] - 파이토치 PyTorch CUDA custom layer 만들기 (3) - setup.py

직접 구현한 PyTorch 연산에서 backward를 구현하기 위해서는 두 가지가 필요한데,

1. back-propagation 되는 미분 값을 계산하는 코드
2. autograd.Function을 사용한 wrapping

이다.

다행히 두 가지 모두 어렵지 않게 처리가 가능하다.

일단 chain rule을 사용하면, 임의의 변수 $x$와 $z$에 대해 아래와 같은 관계가 성립한다.

$\frac{\partial \mathcal{L}}{\partial x} = \frac{\partial \mathcal{L}}{\partial z} \frac{\partial z}{\partial x}$.

구현한 연산 $z = x + y$에서 $\frac{\partial z}{\partial x} = 1$이므로,

$\frac{\partial \mathcal{L}}{\partial x} = \frac{\partial \mathcal{L}}{\partial z}$이고 $x$ 대신 $y$를 넣어도 성립한다.

Backpropagation이란 $\frac{\partial \mathcal{L}}{\partial z}$를 알 때 $\frac{\partial \mathcal{L}}{\partial x}$와 $\frac{\partial \mathcal{L}}{\partial y}$를 구하는 것이므로 계산 끝.

CVPR에 사용하기 위해 구현한 코드는 몇 배 이상 복잡하고 잔머리도 많이 굴려야했는데,

가능하다면 Python 등을 활용해 미리 내 로직이 맞는지 점검하는 과정이 있어야 한다.

아무튼 이제 이걸 CUDA 커널에 적당히 프로그래밍하면 된다.

귀찮지만, 총 3개의 파일을 건드려야 한다.

우선 vecadd_kernel.cuh에 다음 Line들을 추가한다.

 //vecadd_kernel.cuh에 추가
void VecAddBackwardCuda(
    float* dx,
    float* dy,
    const float* dz,
    const int n
);

이번에는 dz를 받아 dx, dy를 구하는 것이 목적이므로 (1)과 비교하면 const의 위치가 바뀐 것을 주의해야겠다.

2021.04.28 - [코딩/PyTorch] - 파이토치 PyTorch CUDA custom layer 만들기 (1) - CUDA 커널

다음으로 vecadd_kernel.cu에 CUDA 커널을 호출하는 함수와 실제 CUDA 커널을 작성한다.

// vecadd_kernel.cu에 추가
__global__ void VecAddBackwardCudaKernel(
    float* __restrict__ dx,
    float* __restrict__ dy,
    const float* __restrict__ dz,
    const int n
)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        dx[idx] = dz[idx];
        dy[idx] = dz[idx];
    }
    return;
}

void VecAddBackwardCuda(
    float* dx,
    float* dy,
    const float* dz,
    const int n
)
{
    int n_blocks = int((n + NUM_THREADS - 1) / NUM_THREADS);
    VecAddBackwardCudaKernel<<<n_blocks, NUM_THREADS>>>(dx, dy, dz, n);
    return;
}

실제 back-propagation이 구현이 CUDA 커널에서 어떻게 구현되었는지 참고하면 된다.

지난번에도 언급했지만, __restrict__는 in-place operation의 가능성이 있을 때는 사용하면 안 된다.

여기서도 지우는 게 맞는데 참고.

여담이지만 이 과정에서 문제 아닌 문제는 각 함수들의 이름을 짓는 게 썩 쉽지 않다는 것...

나는 Forward/Backward, Cuda, Kernel을 잘 조합해서 지었는데 (Forward는 생략했지만)

나중에 어떻게든 구분이 가능하게 재주껏 지으면 된다.

Makefile은 수정할 필요 없고, vecadd.cpp에도 backward를 추가해준다.

// vecadd.cpp에 추가
void VecAddBackward(
    torch::Tensor dx,
    torch::Tensor dy,
    const torch::Tensor dz
)
{
    int n = dx.numel();
    VecAddBackwardCuda(
        dx.data_ptr(),
        dy.data_ptr(),
        dz.data_ptr(),
        n
    );
    return;
}

PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
    m.def("vecadd", &VecAdd, "Vector add");
    m.def("vecadd_backward", &VecAddBackward, "Vector add backward");
}

새롭게 vecadd_backward 함수가 바인딩되는 것을 알 수 있다.

이렇게 해서 일단 구현 자체는 완료했고,

실제로 사용하려면 cuda_extension 바로 아래에 있는 Makefile을 통해서 빌드하고 호출하면 된다.

다만 이렇게 하는 경우 PyTorch의 가장 큰 강점 중 하나인 자동 미분이 안되므로

autograd.Function과 결합해주는 과정이 필요하다.

이 링크의 공식 documentation을 보면 잘 설명이 되어있는데, 적당히 필요한 정보를 취했다.

추가적으로, setup.py를 통해 작성한 autograd.Function을 패키지 안에 포함시켜

개별적인 파일 관리 없이 해당 패키지가 설치된 environment 내에서 자유롭게 import 할 수 있도록 했다.

이 과정은 이전 setup.py 작성 시에 다 고려를 했으므로 별도로 건드릴 내용은 없지만.

(setup.py의 packages=setuptools.find_packages() 부분)

아무튼 cuda_extension 폴더 아래에 (setup.py가 있는 위치) pyvecadd라는 폴더를 새로 만들고,

내용 없는 __init__.py를 해당 폴더에 추가하여 해당 폴더가 Python 패키지임을 명시해준다.

그다음 __init__.py가 있는 위치에  layer.py라는 파일을 아래와 같이 작성한다.

# layer.py
import typing

import torch
from torch import autograd
from torch import nn

import vecadd_cuda


class VecAddFunction(autograd.Function):

    @staticmethod
    def forward(
            ctx: autograd.function._ContextMethodMixin,
            x: torch.Tensor,
            y: torch.Tensor) -> torch.Tensor:

        z = torch.empty_like(x)
        vecadd_cuda.vecadd(x, y, z)
        return z

    @staticmethod
    def backward(
            ctx: autograd.function._ContextMethodMixin,
            dz: torch.Tensor) -> typing.Tuple[torch.Tensor, torch.Tensor]:

        dx = torch.empty_like(dz)
        dy = torch.empty_like(dz)
        vecadd_cuda.vecadd_backward(dx, dy, dz)
        return dx, dy


vecadd = VecAddFunction.apply


class VecAddLayer(nn.Module):

    def __init__(self) -> None:
        super().__init__()
        return

    def forward(self, x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
        z = vecadd(x, y)
        return z

요즘 Python 코드를 작성할 때 가능하면 type hinting (참고)을 사용하려고 하는데, 좋은지는 잘 모르겠다.

절대 필수가 아니므로 귀찮으면 다 빼버려도 된다 (훨씬 깔끔하게 보일 수도?).

결코 코드가 긴 것은 아니지만, 어디서부터 정리해야 할지 제법 막막하다.

Line 8은 우리가 setup.py로 빌드한 vecadd와 vecadd_backward를 포함하고 있는 vecadd_cuda를 import 한다.

VecAddFunction은 실제 자동 미분 로직을 구현하는 class인데,

공식 사이트의 documentation이 잘 되어있어 그대로 따라 하면 된다.

몇 가지만 가볍게(?) 정리해놓자면,

• Line 20과 30이 실제 우리가 C++에서 바인딩하고, setup.py를 통해 빌드한 함수를 호출하는 부분이다.
• Line 15와 25의 ctx는 forward시에 사용한 값이 backward시에도 필요한 경우, 이를 저장해두는 버퍼이다.
  아마 나중에 실제 구현한 코드를 정리하며 다시 등장할 수도 있는데,
  지금은 굳이 백업이 필요 없으므로 사용되지는 않았다.
• Line 19, 28, 29는 return에 사용되는 값들을 위한 메모리를 할당하는 부분이다.
  C++ 나 CUDA 레벨에서 메모리 관리를 하기 싫기 때문에 PyTorch 레벨에서 관리하도록 했다.
  따라서 forward 로직은 이제 input x, y 2개 만을 인수로 받아 새로운 output z를 return 한다.
  이전처럼 모양 빠지게 z를 미리 할당하고, vecadd_cuda.vecadd(x, y, z) 같이 할 필요가 없다는 뜻.
• backward에서는 forward에서 ctx를 제외하고 입력받은 모든 인수에 대해 gradient를 계산해야 한다.
  따라서 x와 y에 대한 gradient인 dx, dy를 순서대로 return 하는데,
  텐서가 아니거나 상수 인수를 forward에서 전달받은 경우에는 해당 순서에 None을 return 하면 된다.
• VecAddFunction은 추상적인 개념이고, 실제로는 Line 34에 정의한 vecadd를 호출하여 사용한다.
  그 아래의 VecAddLayer는 해당 연산을 custom layer로 구현한 class인데,
  vecadd <-> torch.nn.functional.conv1d
  VecAddLayer <-> torch.nn.Conv1d
  같은 관계로 생각할 수 있다.
  실제로 VecAddLayer 내부에서 vecadd를 호출하여 계산을 진행하는데,
  torch.nn.Sequential 같은 모듈을 활용하려면 이런 layer 형태를 쓰는 것이 유리하다.

이제 진짜 완성이다!

마지막으로 cuda_extension 아래에 있는 setup.py를 다시 실행해서

• 새로 구현한 backward CUDA 커널 및 C++ 인터페이스
• 해당 함수의 Python 바인딩
• 해당 함수의 자동 미분을 구현하는 vecadd 폴더 (패키지) 아래에 포함된 layer.py 모듈

들을 현재 environment에 설치한다.

사용법은 아래와 같다.

• vecadd.cpp 파일에 구현되어 Python에 바인딩된 함수를 바로 사용 (이제는 그럴 일이 거의 없지만)

>>> import vecadd_cuda
>>> vecadd_cuda.vecadd(x, y, z)
>>> vecadd_cuda.vecadd_backward(dx, dy, dz)

import 하는 패키지의 이름은 우리가 setuptools.setup에 지정했던 ext_modules의 name이다.

당연히 자동 미분을 지원하지 않으며, z와 dz에 대한 사전 할당이 필요하다.

• 자동 미분이 지원되는 PyTorch (Python)식 함수 사용 (이걸로 쓰자)

>>> from pyvecadd import layer
>>> z = layer.vecadd(x, y)
>>> z.mean().backward()

그냥 우리가 cuda_extension 폴더 아래에 만든 pyvecadd 폴더가

프로젝트 폴더에도 똑같이 있다고 생각하고 사용해주면 된다 (그래서 이름을 잘 지어놓아야 한다...).

실제로는 environment에 설치되어 있기 때문에,

파일을 매번 옮겨놓지 않아도 되는 점이 내가 가장 마음에 드는 부분이다.

조금 구체적인 예시는 아래와 같다.

>>> import torch
>>> from pyvecadd import layer
>>> x = torch.randn(4, 4, requires_grad=True, device=torch.device('cuda'))
>>> y = torch.randn(4, 4, requires_grad=True, device=torch.device('cuda'))
>>> print(x)
tensor([[-0.3005, -1.2514,  0.0675, -1.6077],
        [-1.1321,  0.8361, -0.7227,  1.4445],
        [ 1.1019, -0.2202, -0.3690,  0.3014],
        [ 0.6847, -0.7965, -0.7134,  2.0782]], device='cuda:0',
       requires_grad=True)
>>> print(y)
tensor([[ 1.4736, -0.5728,  0.4896, -0.9496],
        [-0.0979,  0.0057,  0.7292, -0.5978],
        [ 1.5153, -0.2581,  0.0160,  0.8342],
        [ 0.1993,  0.2244,  0.0465, -0.7861]], device='cuda:0',
       requires_grad=True)
>>> z = layer.vecadd(x, y)
>>> print(z)
tensor([[ 1.1731, -1.8242,  0.5571, -2.5573],
        [-1.2300,  0.8418,  0.0065,  0.8467],
        [ 2.6172, -0.4784, -0.3530,  1.1356],
        [ 0.8840, -0.5722, -0.6669,  1.2921]], device='cuda:0',
       grad_fn=<VecAddFunctionBackward>)
>>> print(x + y)
tensor([[ 1.1731, -1.8242,  0.5571, -2.5573],
        [-1.2300,  0.8418,  0.0065,  0.8467],
        [ 2.6172, -0.4784, -0.3530,  1.1356],
        [ 0.8840, -0.5722, -0.6669,  1.2921]], device='cuda:0',
       grad_fn=<AddBackward0>)
>>> z.mean().backward()
>>> print(x.grad)
tensor([[0.0625, 0.0625, 0.0625, 0.0625],
        [0.0625, 0.0625, 0.0625, 0.0625],
        [0.0625, 0.0625, 0.0625, 0.0625],
        [0.0625, 0.0625, 0.0625, 0.0625]], device='cuda:0')
>>> print(y.grad)
tensor([[0.0625, 0.0625, 0.0625, 0.0625],
        [0.0625, 0.0625, 0.0625, 0.0625],
        [0.0625, 0.0625, 0.0625, 0.0625],
        [0.0625, 0.0625, 0.0625, 0.0625]], device='cuda:0')

본래의 목적인 custom layer로 사용하는 것도 가능.

>>> vecadd_layer = layer.VecAddLayer()
>>> zz = vecadd_layer(x, y)

아무튼 이렇게 해서 PyTorch CUDA custom layer 구현에 대한 설명은 정리가 끝난 것 같다.

라고 말하고 싶은데, 사실 구현 대상 함수가 굉장히 쉬웠기 때문에 고려하지 않은 많은 요인이 있다.

이후의 내용은 코딩 실력이 완벽하여 에러를 절대 내지 않는다면 굳이 참조할 필요는 없지만,

나는 아쉽게도 그렇지 못해서 조금 더 정리를 해둬야 할 것 같다.

To be continued...

이전 글:

2021.04.28 - [코딩/PyTorch] - 파이토치 PyTorch CUDA custom layer 만들기 (1) - CUDA 커널

2021.04.30 - [코딩/PyTorch] - 파이토치 PyTorch CUDA custom layer 만들기 (2) - 커널 빌드 및 Python 바인딩

이전 포스트에 이어 PyTorch (C++)와 PyTorch (Python)을 연결하는 방법을 정리해야겠다.

이 부분은 의외로 간단한것이, setup.py 템플릿을 조금만 손보면

PyTorch의 cpp_extension이라는 모듈로 어렵지 않게 해결이 가능하다.

원리는 생략하고 아래와 같이 setup.py를 작성하자.

참고로 나는 setup.py의 위치를 cuda 폴더 바깥으로 했다.

전체 폴더 구조는 대충 아래와 같다.

# setup.py
from os import path
import setuptools

from torch.utils import cpp_extension

def main() -> None:
    setuptools.setup(
        name='ext_vecadd',
        version='1.0.0',
        author='Your name',
        author_email='your@email.com',
        packages=setuptools.find_packages(),
        ext_modules=[cpp_extension.CppExtension(
            name='vecadd_cuda',
            sources=[path.join('cuda', 'vecadd.cpp')],
            libraries=[
                'vecadd_kernel',
            ],
            library_dirs=[path.join('.', 'cuda')],
            extra_compile_args=['-g', '-fPIC'],
        )],
        cmdclass={'build_ext': cpp_extension.BuildExtension},
    )
    return

if __name__ == '__main__':
    main()

당연한 소리지만 version, author, author_email은 취향대로.

여기서 중요하게 볼 부분은 setup 함수의 ext_modules 부분이다.

이 링크에서 공식 documentation을 보면 CppExtension 말고도 CUDAExtension이라는 함수도 존재하는데,

막상 써보니까 잘 동작하지 않았고 돌아돌아 결론을 내린게 CppExtension이었다.

CUDAExtension을 쓰면 별도의 Makefile로 shared object (.so)를 빌드하지 않아도 될 것 같은데,

잘 안되니까 나중에 다시 해봐야겠다.

CppExtension의 name 인수에는 나중에 Python에서 해당 모듈을 불러올때 사용하는 이름을 지정한다.

예를 들어, 위의 setup.py로 빌드한 모듈은 아래와 같이 불러올 수 있다.

import torch
import vecadd_cuda

반드시 Line 1이 선행되어야 하는 점에 주의하면, 큰 문제는 없다.

• sources 인수는 빌드할 cpp 파일(들)을 지정한다.
  이미 vecadd.cpp에서 CUDA 커널을 호출하도록 되어있기 때문에, .cu 파일을 포함할 필요는 없다.
• 대신에 우리가 빌드한 .so 파일을 libraries라는 인수에 포함시키는데,
  vecadd_kernel을 지정하면 자동으로 libvecadd_kernel.so가 빌드에 포함된다고 생각하면 된다.
• library_dirs에는 해당 .so 파일이 포함되어 있는 폴더를 지정해주고,
• extra_complie_args는 -g와 -fPIC를 넣어주면 완성.

이제 터미널에서 아래와 같은 커맨드로 전체 프로젝트를 빌드한다.

$ python setup.py install

여기서 주의할 점은, conda 등의 가상 환경을 사용 중인 경우 해당 environment에만 패키지가 설치된다는 것.

뭔가 더 해야할 것 같은 느낌이 들지만 의외로 지금 당장 작성한 CUDA 커널을 PyTorch (Python)에서 사용 가능하다.

우선 Python을 실행시키고, 아래와 같이 하면 테스트 가능.

>>> import torch
>>> import vecadd_cuda
>>> x = torch.randn(4, 4, device=torch.device('cuda'))
>>> y = torch.randn(4, 4, device=torch.device('cuda'))
>>> z = torch.zeros(4, 4, device=torch.device('cuda'))
>>> print(x)
tensor([[ 0.7651,  0.2747, -0.3275, -1.0251],
        [-0.5340, -0.8416,  0.2952,  0.7431],
        [-0.2698, -0.3798, -0.0225,  0.0245],
        [ 1.2056,  1.9831,  0.1669, -0.2347]], device='cuda:0')
>>> print(y)
tensor([[-1.0483,  0.6131, -0.2479,  0.3100],
        [ 0.6286,  0.4175, -0.4093, -1.5263],
        [ 0.6393,  1.3988, -0.0907,  0.1472],
        [-1.0677,  1.0088, -1.4845, -0.7323]], device='cuda:0')
>>> print(z)
tensor([[0., 0., 0., 0.],
        [0., 0., 0., 0.],
        [0., 0., 0., 0.],
        [0., 0., 0., 0.]], device='cuda:0')
>>> vecadd_cuda.vecadd(x, y, z)
>>> print(z)
tensor([[-0.2832,  0.8878, -0.5754, -0.7151],
        [ 0.0947, -0.4241, -0.1141, -0.7832],
        [ 0.3695,  1.0189, -0.1133,  0.1716],
        [ 0.1379,  2.9919, -1.3177, -0.9670]], device='cuda:0')
>>> print(x + y)
tensor([[-0.2832,  0.8878, -0.5754, -0.7151],
        [ 0.0947, -0.4241, -0.1141, -0.7832],
        [ 0.3695,  1.0189, -0.1133,  0.1716],
        [ 0.1379,  2.9919, -1.3177, -0.9670]], device='cuda:0')

random 수치들은 테스트할 때마다 다르게 나오겠지만, 아무튼 z = x + y가 잘 작동한다.

코드를 잘(?) 만들어놔서인지, 2D 텐서를 처리하는 데에도 큰 문제가 없다.

그런데 현재 상태로는 CUDA 커널이 수정될때마다 Makefile 후 setup.py를 install 하는 두 단계를 거쳐야 해서 조금 번거롭다.

기왕 하는거 두 개를 한 번에 하는 스크립트도 작성해놨다.

아래 내용을 setup.py가 있는 폴더에 Makefile 파일을 만들고 붙여넣으면 된다.

SUBDIRS = cuda

all: vecadd
    python setup.py install

vecadd:
    $(MAKE) -C cuda

clean:
    rm -rf build dist ./*.egg-info
    for dir in $(SUBDIRS); do $(MAKE) -C $$dir Makefile $@; done

역시나 야매지만 $(MAKE)와 -C라는 문법을 통해 sub directory의 Makefile을 실행하는 형식인가 보다.

clean의 경우 저것 말고도 여러 가지 방법이 있는데 내가 보기에 깔끔한 걸로 골랐다.

아무튼 이제 setup.py를 실행시키는 대신에, Makefile을 한 번만 실행하면 된다!

$ make

여기서 끝?

이면 참 좋겠지만, 몇 가지 불편한 사실들이 남아있다.

1. z가 자동으로 할당되지 않는다.
   보통 우리는 x와 y를 가지고 있을 때, x + y가 z를 return 하기를 원하지
   z를 미리 할당하고 x + y를 채우는 것을 별로 선호하지 않는다.
2. 예외 처리가 되어있지 않다.
   작성한 커널은 CUDA 텐서에 대해서만 작동하지만 CPU 텐서가 들어와도 아무 경고가 없다 (내가 가장 삽질한 부분).
   x + y가 정상적으로 계산되지 않는 것은 당연하고.
3. backward()가 작동하지 않는다. = backpropagation이 불가능하다.
   사실 가장 중요한 부분이다.
   달리 말하면, custom operation으로의 기능은 하지만 custom layer로의 기능은 전혀 할 수가 없다.

1과 2는 그러려니 해도, 3은 조금 치명적이다.

조금 불행한 소식은 3 또한 CUDA 커널을 직접 작성해서 구현해야 한다는 점,

그나마 다행인 소식은 CUDA 커널만 추가로 작성하고

Makefile이나 vecadd.cpp, setup.py를 크게 고칠 필요는 없다는 점이다.

또한 backward가 사용 가능하게 하려면 autograd.Function을 상속받아야 하는데, 이를 통해 1과 2를 함께 해결 가능하다.

이전에 backward를 미리 작성해놓지 않은 이유는 사실 간단한데,

이 단계에서 구현한 연산이 정확히 동작하지 않는다면 backward 또한 제대로 동작 할리 만무하기 때문이다.

그러니까 일단 CUDA 커널을 제대로 작성했는지 PyTorch (Python) 인터페이스를 사용하여 편하게 확인하고,

문제가 있다면 고쳐서 완벽한 동작을 보장한 다음 나머지 잡일 (좀 많지만)을 처리하러 가면 된다.

다음 글:

2021.05.01 - [코딩/PyTorch] - 파이토치 PyTorch CUDA custom layer 만들기 (4) - autograd.Function과 backward 구현