이번에 카메라 피드로 받아오는 rgb pixel들을 실시간으로 연산해야 하는 작업을 처리해야 했는데, 어떻게 하면 빠르게 할 수 있을까 고민을 하다가 Metal을 이용하게 되었습니다. 이번 기회에 Metal을 이용해보면서 알게된 사항들을 정리해보았습니다.
Metal은 GPU를 이용해서 그래픽 연산이 많은 종류의 앱, 대표적으로 게임 혹은 3D 이미지 처리 그리고 머신러닝에도 이용될 수 있는 프레임워크입니다. GPU 연산의 핵심은 병렬처리 입니다. 여러개의 thread가 동시에 일을 처리하는 것입니다. 일꾼을 여러 명 사용해서 작업을 처리하기 때문에 처리해야 할 양이 많은 연산을 할 때 더욱 빠르게 작업을 완료할 수 있게 되는 것입니다.
For loop 연산 대체하기
처음 사용해보는 프레임워크였지만, Apple에서 Metal을 이용한 기본 연산 예제를 제공해 주어 참고하면서 전반적인 사용법을 익힐 수 있었습니다. 해당 예제에서는 두 개의 배열에서 동일한 인덱스에 위치한 값들을 더하는 연산을 보여줍니다.
크기가 n인 두 개의 배열의 동일한 인덱스 값들을 더하는 연산을 하려면, 기본적으로 for loop을 이용해 처리할 수 있을 것입니다. 그러면 0부터 n-1까지 차례로 연산을 하게 되는데, 그러다보니 데이터 양이 10만개, 100만개 늘어날 수록 연산을 완료하는 데 시간에 지연이 느껴지게 될 것입니다.
이 때 GPU를 활용하면 어떨까요? 여러 개의 thread가 연산을 동시에 처리하기 때문에 동일한 시간에 처리할 수 있는 연산의 양이 훨씬 많아지고, 결과적으로 모든 연산을 완료하는 데 소요되는 총 시간이 줄어들 게 됩니다. 간단하게 index 0의 값들은 thread 0이 담당하고 index n의 값들은 thread n이 담당하고 이런 식으로 이해해 볼 수 있겠습니다. 아쉽게도 일꾼 thread의 수가 무한대는 아니기 때문에 얼만큼 사용할 지 설정해야 되기는 합니다.
한 번 예제 코드를 보도록 하겠습니다. 참고로 Metal은 MSL 언어를 사용합니다.
#include <metal_stdlib>
using namespace metal;
/// This is a Metal Shading Language (MSL) function equivalent to the add_arrays() C function, used to perform the calculation on a GPU.
kernel void add_arrays(device const float* inA,
device const float* inB,
device float* result,
uint index [[thread_position_in_grid]])
{
// the for-loop is replaced with a collection of threads, each of which
// calls this function.
result[index] = inA[index] + inB[index];
}
보시면, add_arrays라는 함수에 여러 파라미터가 들어가는 것을 확인할 수 있습니다. 하나 씩 살펴보겠습니다.
- inA는 첫번째로 주어지는 리스트(float 값이 들어있는)를 가르키는 device memory의 pointer를 의미합니다.
1. device memory?
2. pointer?
- inB는 inA과 동일하죠? 다만 두번째로 주어지는 리스트에 대한 정보입니다.
- result는 연산한 값을 저장할 리스트를 가르키는 pointer 입니다. inA, inB와의 차이점은 const 키워드가 붙어있지 않다는 점인데요, 값을 읽기만 하는 inA, inB와는 다르게 write도 해야하기 때문입니다.
- 이제 핵심인 index! 이름과 uint 타입에서 느껴지 듯이 접근 할 리스트의 인덱스를 의미합니다. 그런데 값이 특이해 보이지 않나요? [[thread_position_in_grid]]는 MSL에서 기본적으로 제공하는 값입니다. 이름 그대로 GPU grid 상 thread의 position을 이야기합니다.
이게 왜 필요할까요? GPU에서는 여러 thread가 동시에 작동을 한다고 했죠? 그러다 보니, 어떤 thread가 어떤 값을 이용할 지 알려주는 것이 필요하고, 그것을 위치 즉, Index로 명시해주는 것입니다. 요새 핫한 Nvidia에서 한 번 CPU와 GPU의 작동 방식을 시연으로 보여 준 적이 있었는데, 그 영상을 한 번 보시면 바로 이해가 되실 겁니다.
https://www.youtube.com/shorts/vGWoV-8lteA
GPU 부분에서 보시면 아시겠지만, 여러 개의 잉크 노즐이 동시에 작동해서 한 번에 모나리자 그림을 찍어내는데요. 그렇게 동작을 시키려면 각각의 노즐은 자신이 그림의 어떤 부분을 그려내야 되는 지 알고 있어야 겠죠? 그것을 알 수 있는 정보가 바로 [[thread_position_in_grid]] 입니다.
Grid에 대해 상세하게 설명을 해주는 애플 문서가 있으니, 꼭 참고해보시면 좋겠습니다. ThreadGroup이라는 개념도 있는데, 같이 알아두면 좋을 것 같습니다.
https://developer.apple.com/documentation/metal/creating-threads-and-threadgroups
Creating Threads and Threadgroups | Apple Developer Documentation
Learn how Metal organizes compute-processing workloads.
developer.apple.com
다시 예제에 대한 설명으로 돌아가 보면, 여기에서는 grid 내 thread의 position을 리스트 값에 접근할 index로 이용하고 있습니다. 예를 들어서, grid가 10 * 1 크기라면, thread의 position은 0부터 9의 값을 갖게 될 것입니다. 리스트의 크기가 10보다 작지 않다는 가정 하에, 해당 position 값을 이용해서 리스트 내부 값에 접근할 수 있겠죠?
이렇게 파라미터에 대해서 간단하게 알아보았습니다. 하나 하나 자세히 이해하기 위해서는 시간이 필요하지만, function으로 제공해야할 파라미터는 어떤 것들이 있어야 하는 지 대략적으로 알 수 있습니다.
1. 데이터가 들어있을 리스트들
2. 결과 값을 담을 리스트
3. 접근할 인덱스 값
요렇게 제공하면 GPU가 알아서 값들을 사용해주나 봅니다.
이중 for loop 연산은 어떻게 할 수 있을까?
저는 이중 for loop 연산을 했어야 했기 때문에 예제 만큼 간단하지는 않았지만, 또 많이 어렵지도 않았습니다. 이 때도 역시나 thread와 grid에 대한 개념에 대한 이해가 필요했습니다! 그렇다면 2차원 배열 연산은 어떻게 할 수 있는 지 한 번 알아볼까요? 저와 비슷한 고민을 하던 분들께 도움이 되었으면 좋겠네요.
먼저, 코드는 아래와 같습니다. 위 예제에서 보았던 것과 굉장히 유사하죠?
#include <metal_stdlib>
using namespace metal;
kernel void add_arrays(device const float* firstArray [[buffer(0)]],
device const float* secondArray [[buffer(1)]],
device float3* output [[buffer(2)]],
uint2 grid_size [[threads_per_grid]],
uint2 gid [[thread_position_in_grid]])
{
uint index = gid.y * grid_size.x + gid.x;
output[index] = firstArray[gid.x] + secondArray[gid.y];
}
firstArray는 이중 for loop에서 바깥 loop에 해당되는 값을 줄 것이고, secondArray는 안쪽 loop에 해당하는 값을 줄 것입니다.
그리고 이중 for loop 연산이기 때문에 uint가 아닌 uint2를 이용했고, index를 알기 위해 약간의 계산이 필요했기 때문에 [[threads_per_grid]] 라는 정보가 더 필요했습니다.
우선 uint2는 무엇이냐면, uint로 구성되어 있는 vector 값입니다. vector라고 해서 어렵게 느낄 필요 없이, 2d 좌표계 값이라고 생각하면 편합니다. 아래에 빨간색으로 표현된 (9, 10) 과 같은 값이 uint2 입니다. 각각의 값은 x, y로 접근하면 됩니다.
[[thread_position_in_grid]]는 앞에서 알아 본 것 처럼 grid에서 thread의 위치 정보입니다. 이전 예제에서는 1D 였기 때문에 uint 값이었지만, 2D 인 상황에서는 (9, 10), (0, 0) 처럼 (x, y) 값으로 주어집니다.
[[threads_per_grid]]는 이름에서 알 수 있다시피 Grid 내에 몇 개의 thread가 있는지 알려주는 uint2 값 입니다. 만약에 grid가 4 x 4 크기라면 (4, 4) 라는 값이 주어지게 됩니다.
그러면 이제 output buffer에 값을 저장할 때 이용될 index를 구하기 위해 사용한 식을 이해해 보겠습니다.
uint index = gid.y * grid_size.x + gid.x;
function의 파라미터들 중에 계산값을 저장할 output이 있었죠? 이 값은 'buffer'입니다(firstArray, secondArray도 마찬가지!). 그리고 buffer는 1D입니다. 하지만 지금 thread grid는 2D를 사용하고 있죠? 그렇기 때문에 vector로 받게된 grid의 위치가 1D 버퍼에서는 어디에 위치할 지 값을 계산해주어야 하는 것입니다. (짧게 말해서 flatten 한다고 하면 될까요?)
조금 복잡해 보이지만 이미지를 통해서 한 번 설명해드리겠습니다. 아래와 같은 4 X 4 grid에서 만약 'L'의 위치 즉, (3, 2)는 1D array에서 몇 번째 index에 위치할까요?
이 grid에 적혀있는 알파벳들을 1d로 나열하면 아래처럼 되겠습니다.
[A, B, C, D, E, F, G, H, I, J, K, L, M, N, O, P]
위 배열에서의 각각의 알파벳의 위치를 찾는 것이 위 식이 설명하는 것이 되는 것입니다.
'L'을 보면, 11번째 index에 위치하고 있네요. 이 위치를 계산하려면, (2d grid의 폭 * L의 y 값)을 구해주고 (L의 x값)을 더해주면 됩니다.
계산해보면, (4 * 2) + 3 = 11 로 정확한 값이 나왔음을 알 수 있습니다.
그래서 firstArray의 값, secondArray의 값을 gid.x (0부터 n-1까지), gid.y (0부터 n-1까지) 하나씩 다 더해준다는 것이 아래 식이 됩니다. 위에서 구해 준 index 값을 이용해서 output에 저장해주면 되는 것입니다.
output[index] = firstArray[gid.x] + secondArray[gid.y];
index 만드는 식에 조금 시간이 걸릴 수는 있지만, 생각보다 이중 for loop을 metal로 처리하는 방법도 생각보다 쉽지 않나요? (ㅎㅎ..)
이렇게 metal을 이용해서 1d와 2d 연산을 하는 방법을 알아보았습니다. 저와 비슷한 연산이 필요한 분들께 도움이 되었으면 좋겠네요.
이어서는 열심히 작성한 metal method를 swift에서 사용하는 방법에 대해서 알아보고, 저희의 관심사인 CPU를 사용했을 때와 비교해서 ✨시간이 얼마나 빨라졌는지✨도 알아보도록 하겠습니다!
감사합니다 🐥
---
혹시 두 개의 array에 덧셈 연산을 하는 예제가 궁금하신 분들은 Apple에서 제공하는 예제를 참고하시면 됩니다. 간단한 예제이지만 기본적인 사용법을 익히기 좋습니다!
https://developer.apple.com/documentation/metal/performing-calculations-on-a-gpu
