참고로 나는 고등학교 1학년 때 쿠다에 빠져버렸다. 당시에 한참 영상 편집에 빠져 있는 때였는데,
혹시라도 경험이 있다면 알 거다. NLE(흔히 말하는 편집 프로그램)에서 렌더링 돌리는 순간 굉장히 긴 시간을 기다려야 한다는 것을.
지금이야 폰에도 하드웨어 영상 인코더가 달려나와서 10분짜리 영상이야 10분 이내로 끝나는 시절이지만
그 당시만 해도 10분짜리 영상 렌더링하면 1시간이 걸리던 시절이었다.
그런데 어느 날, 어도비 프리미어에 머큐리 플레이백이라는 영상 인코더가 달리고 나서 놀라운 일이 벌어졌다. 10분짜리 영상 렌더링이 8분에 끝나 버린 것이다(어도비 프리미어는 대학생인 친척을 통해서 3년치 새뱃돈 모은 걸로 샀으니 걱정 무).
그 때 해당 연상 렌더링 엔진이 쿠다로 만들어졌단 걸 알게 됐고, GPGPU와 처음 만나게 됐다.
아무튼 그 덕에 수능 망치고 재수를 했고, 대학에 와서 OpenCL이라는 규격도 있다는 걸 알게 됐다.
써 봤고, 두 가지를 확신했다.
1. GPGPU는 앞으로 컴퓨팅의 대세가 될 것(당시에 모든 대학 동기들한테 홍보하고 다녔다)
2. 쿠다가 원탑을 먹을 것
둘 다 써본 입장에서 쿠다가 원탑을 먹을 거란 걸 확신할 수 있었다. 이유는 다음과 같다.
1. 쿠다가 훨씬 추상화가 잘 되어 있다.
2. 실제로 응용을 개발할 수 있는 라이브러리들이 잘 만들어져 있다.
3. 파편화가 적다
이제 하나하나 살펴 보자
1. 쿠다가 훨씬 추상화가 잘 되어 있다
기본적으로 쿠다는 Heterogenous, thread based SIMD supercomputing 모델이다. Heterogenous란 말은 이종이란 뜻이고, 서로 다른 연산 특성을 갖는 튜링 머신이 하나의 목적을 위해 동작하기 때문에 '이종 간' 을 뜻하는 Heterogenous를 쓰고, 서로 다른 쓰레드가 연산 유닛들에서 동일한 연산을 여러 개의 데이터에 대해 수행하는 데에 최적화된 모델이라 thread based, SIMD(플린 연산 분류에서 Single Instruction Multiple Data를 뜻함)이다. 이와 같은 환경에서 가장 중요한 것 중 하나는 연산을 수행할 유닛의 수와 그 분배 구조를 결정하는 것인데, 쿠다는 이를 C++에 덧붙인 아주 간단한 연산자를 통해서 해결한다. 다음은 쿠다에서 Hello world를 실행하는 방법이다:
몇 가지 일반적인 C++에서 보이지 않는 예약어들이 보일 것이다. 첫 번째는 __global__이라는 지시자, 두 번쨰는 생소한 <<<>>> 연산자이다. 쿠다는 컴퓨팅 환경을 두 가지 host와 device로 분류하는데, host는 일반적인 CPU-시스템 메모리 환경, device는 GPU-비디오 메모리 환경을 뜻한다. 또한 GPU에서 실행되는 모든 함수들을 커널이라고 부르는데, __global__ 지시자는 커널이므로 device에서 실행되나, host에서 호출 가능함을 의미한다. __device__라는 지시자도 있는데, 이름에서 알 수 있듯 device에서 실행되며, host에는 노출되지 않는 커널을 뜻한다.
<<<>>>는 쿠다에서 Triple brackets, Triple chevron 연산자라고 부르는데, 쿠다에서 실행되는 가장 작은 연산유닛의 추상화된 유닛이 쓰레드, 쓰레드가 여러개 모여서 블록, 블록 여러개가 모여서 그리드를 구성한다(그냥 그렇다)이 때 몇 개의 쓰레드, 블록, 그리드에서 실행할지를 host가 global 커널에 지정하는 것이 바로 이 연산자다. 많은 개념들이 나왔지만, 개념을 이해한다는 가정 하에서는 매우 짧고 직관적으로 표현 가능함을 알 수 있다. C++의 템플릿과도 유사한 점이 있어서 그런지 Chevron을 사용한 연산자를 정의한 것도 무척 마음에 든다. 그럼 동일한 코드를 OpenCL로 작성하면 어떨까?
놀랍게도 다음과 같다.
실제로는 위와는 좀 다른 예제긴 하다. 쿠다의 예제는 단순히 printf를 하는 예제였고 아래의 예제는 실제로는 행렬합을 하는 예제이기 때문. 중요한 것은 쿠다의 경우에 생략되는 실행가능 여부를 점검하는 것이 OpenCL에서는 일일이 다 해야 하고, GPU에서 실행할 명령을 미리 적재해 두는 등의 주 목표와 상관없는 복잡한 코드들이 들어간다는 것, 그리고 함수를 단순히 호출하는게 아니라, String으로 작성된 GPU에서 실행될 코드를 런타임에서 컴파일하고(?) 실행할 함수의 인자를 설정하는 함수를 실행하고(??) 쿠다에서는 간단히 Triple bracket으로 설정한 함수의 연산 유닛 수를 결정하는 부분이 enqueNDRangeKernel(해석하면 N-Dimension 범위 커널 큐에 적재 정도가 되겠다)라는 함수를 통해서 설정한다는 것. 쿠다로 한줄로 끝날 작업을 몇 줄에 거쳐서 반복하고 있는 것이다
벌써 여기에서 게임 끝인것 같지만 추가로 더 알아보도록 하자.
2. 실제로 응용을 개발할 수 있는 라이브러리들이 잘 만들어져 있다
결국 우리는 GPU가 너무 사랑스럽거나 성욕을 느껴서 GPGPU를 하는게 아니다, 엔지니어링은 언제나 도구, GPGPU는 GPU를 써서 무언가를 행하기 위한 도구일 뿐이다. 쿠다는 초기 출시 당시에는 영상 렌더링 위주로 개발되었는데, 이 당시에도 행렬 연산 등을 쉽게 수행하기 위한 퍼스트파티 라이브러리들을 매우 많이 개발했다. 이후 대규모 정보 분석과 기계 학습이 뜨면서 아예 cuDNN(쿠다 딥 뉴럴 네트워크) 등의 라이브러리들을 본인들이 직접 만들어서 배포했다. 텐서플로우(뭐... 아무튼 처음 나왔을땐 이것도 감지덕지했다), 토치 등 유명한 라이브러리들도 결국 이런 모태 위에서 탄생한 것이다.
또한 기계학습 분야는 컴퓨터 사이언스 중에서도, 컴퓨터보다는 순수 수학에 유독 접점이 높은 분야다. 이런 분야 사람들일수록 개발 과정 자체를 최대한 쉽게 해줘야 그 위에서 사용 가능한 코드베이스 에코시스템의 규모가 커지고, 이것이 쿠다를 승리로 이끌었다.
3. 파편화가 적다.
Cuda는 엔비디아가 적극적으로 미는 주력 상품이자 플랫폼이다. 따라서 엔비디아에서 자사 하드웨어들을 쿠다 실행에 용이하게 통제하며, 이들을 실행하는 소프트웨어 환경인 드라이버도 책임감 있게 관리한다. 때문에 현재 주력 세대 GPU들이라면 모두 쿠다 정상 실행을 보장한다.
반면 OpenCL은 좀 복잡하다. OpenCL 표준 자체는 애플이 인텔 맥으로 넘어가며 부족해진 병렬 연산력을 보강하기 위해 GPGPU를 수행할 목적으로 만들었으나, 후에 크로노스 재단으로 넘어가게 된다. 이 때 기업들의 후원금에 의존하는 재단답게 프로젝트의 규모가 어마어마하게 커지면서 Heterogenous Supercomputing 전반을 위한 표준으로 확고하게 거듭나게 되었는데(실질적 의미와 관계 없이 설계 철학에만 따르면 NPU, FPGA 등도 커버할 목적이었다는 뜻이다. 쿠다보다 압도적으로 세밀한 설정이 필요한 이유다. 이 모든 디바이스들을 추상화하는 것이 불가능하기 때문), 크로노스가 자신들이 관리하는 표준을 해당 얼라이언스의 모든 기업들이 제대로 드라이버와 개발 도구들을 구현하도록 통제하는데 실패하면서 결국 이도 저도 아닌 플랫폼이 되어버렸다.
게다가 실질적으로는 연산 유닛의 종류별로 최적화 전략과 설계 자체가 달라져서, 모든 플랫폼을 커버하는 단 하나의 코드가 아니라, 모든 플랫폼을 지원하려면 서로 다른 플랫폼을 위한 코드를 일일이 작성해야 하는 프레임워크가 되어버렸다. 이와 같은 이유로 OpenCL은 쿠다와의 경쟁에서 완전히 밀려버렸다.
글이 길었는데 요약하면 이렇다.
쿠다가 이긴 이유
1. OpenCL보다 압도적으로 개발하기 쉬움
2. OpenCL로 작성하면 정상 동작 보장이 안 됨
3. OpenCL의 유일한 장점인 넓은 연산환경 커버리지는 유명무실함
