[펌] GPGPU를 이용한 레이트레이싱

http://www.imaso.co.kr/?doc=bbs/gnuboard.php&bo_table=article&wr_id=43180

 

 

GPGPU를 이용한 레이트레이싱

프린트 이메일 전송

월간 마이크로소프트웨어 2013년 7월호 / Smart & Contents

GPGPU(General-purpose GPU)는 GPU의 강력한 성능을 일반적인 컴퓨팅 분야에 쓰도록 지원하는 것이다. GPGPU 활용을 위해서는 여러 가지 방법이 있는데, 현재는 OpenCL이나 CUDA를 이용하는 것이 보편화된 방법이다. 이 글에서는 CUDA를 이용해서 레이트레이싱(raytracing)을 하는 과정의 기초 부분을 살펴보기로 한다.

김혁 hybrid@nexon.co.kr｜서강대학교 대학원에서 컴퓨터 그래픽스를 전공하면서 레이트레이싱을 사용한 렌더링 기술을 연구했다. 현재는 넥슨 데브캣 스튜디오에서 신규 게임 프로젝트 K를 개발 중에 있다.

게임에서 사용되는 렌더링 방식은 래스터라이제이션(rasterization)이라고 한다. 각각의 삼각형을 스크린에 투영해 렌더링될 화면의 픽셀 색상을 결정하는 이 방식과는 달리, 레이트레이싱에서는 픽셀(pixel) 기준으로 픽셀에 칠해질 색상을 결정하는 방식으로서 계산하는 방향이 다르다.

< 그림 1> 레이트레이싱

레이트레이싱의 과정을 가장 간단하게 요약하면 다음과 같다.

1) 카메라 정보로부터 스크린의 각 픽셀에 해당하는 광선(ray) 생성
2) 각 광선 기준으로 카메라와 가장 가까운 삼각형을 찾음
3) 그 삼각형과 광선이 교차(hit)하는 지점을 중심으로 셰이딩(shading) 계산을 수행

이 과정에서 속도를 빠르게 하기 위한 알고리즘을 제대로 도입하지 않으면 제 아무리 강력한 GPU 병렬 처리의 힘을 빌린다 하더라도 매우 많은 시간이 걸린다. 빠른 처리를 위해서는 단순히 강력한 하드웨어뿐만이 아니라 적절한 알고리즘과 하드웨어에 대한 이해가 필수적이다.

레이트레이싱을 위한 가장 간단한 구현 방식은 픽셀에 해당하는 광선을 현재 렌더링할 장면(scene)에 있는 모든 삼각형과 교차 검사를 하는 것이다. 이는 굉장히 많은 시간이 걸리기 때문에 일반적으로는 특수한 공간 구조를 만들어 삼각형을 배치한다. 여기서는 BSP트리(Binary Space Partitioning Tree)의 일종인 KD트리(KD-Tree)를 기준으로 설명한다. 최근에는 BVH (Bounding Volume Hierarchy) 방식도 많이 쓰이는 추세다. BVH는 KD트리보다는 동적 장면(dynamic scene)에 적합한 특징을 가지고 있는데, 근본적인 것은 KD트리를 기반으로 하기 때문에 KD트리를 알면 쉽게 이해하고 구현할 수 있다.

< 그림 2> 일반적인 BSP와 KD트리

KD트리
만약 KD트리와 같은 공간가속구조(spatial acceleration structure)를 전혀 사용하지 않는다면 하나의 광선에 대해 (1) 모든 삼각형과 일일이 교차 계산을 하거나 (2) 삼각형을 광선 기준으로 정렬 후 가까운 곳부터 교차 계산을 해야 한다. 공간가속구조는 삼각형을 엄밀하게 정렬하기보다는 공간을 나눠 가까운 곳부터 먼저 계산할 수 있도록 해주는 장치다.

< 그림 3> 광선이 1, 2, 3번의 리프 노드(공간)를 순차적으로 탐색

<그림 3>과 같이 광선은 1, 2, 3번의 리프 노드(공간)를 순차적으로 탐색하게 된다. 1번을 탐색한 후 삼각형이 없기 때문에 넘어가고, 2번 노드에서는 두 삼각형을 모두 교차 계산해 그림의 왼쪽 삼각형이 더 가까움을 알 수 있고, 이미 교차를 했기 때문에 3번의 노드는 탐색할 필요가 없다. <그림 3>에서처럼 공간이 이미 나눠져 있으면 광선 기준으로 가까운 곳부터 탐색해 빠르게 교차 계산을 수행할 수 있다.

이때 KD트리를 위한 레이트레이싱 알고리즘은 이미 많이 최적화돼 왔고, 많은 참고자료들이 존재한다. 그 내용들은 다른 자료들을 참고하길 바라면서 CUDA를 이용한 레이트레이싱에 집중해 설명하기로 한다.

CUDA의 이용
지면 관계상 여기서 제시되는 소스 코드나 레이트레이싱 구현에 관한 내용은 생략한다. 마찬가지로 CUDA에 관한 기본적인 내용 또한 가능한 간단하게 설명한다. CUDA를 이용하면 GPU를 이용해서 많은 연산을 병렬로 처리할 수 있다. 전부는 아니더라도 상당히 많은 부분의 C/C++ 코드를 사용할 수 있긴 하지만, GPU의 처리 구조와 하드웨어 구조를 이해해야만 GPU의 성능을 제대로 끌어낼 수 있다(상황에 따라서는 잘못 구현할 경우, 충분히 GPU에 적합한 응용 분야라 하더라도 멀티스레딩을 이용한 CPU 계산보다 느릴 수도 있다).

최근의 하드웨어는 많이 나아진 편이지만, GPU 연산을 위한 처리상의 제약사항은 제법 많은 편이다. 레지스터 사용 개수, 메모리 액세스, 한번에 병렬 처리하는 양(당연한 말이지만, GPU에서는 모든 GPU 내의 스레드가 동시에 수행되지 않는다)을 프로그래머가 고려해야 하고, 이런 이슈들로 인해 많은 성능 차이가 날 수 있다. 여기서 몇 가지 이슈를 다뤄본다.

먼저 각 스레드에는 독립적인 광선 하나씩이 배분된다. 교차하는 삼각형 중 가장 가까운 삼각형을 찾아야 하는데, 앞서 설명했듯이 트리를 탐색하고 리프 노드에 도달할 경우 리프 노드에 있는 삼각형들과 교차 검사를 한 후 교차하는 삼각형들이 없을 경우 다른 노드를 탐색하게 된다. 이때 다른 노드를 찾아가는 과정에 스택을 사용하게 된다. 이 스택의 액세스 빈도는 상당히 높다.

struct StackInfo;
StackInfo SharedStack[STACK_SIZE];

CUDA 코드 내에서 이 스택을 위와 같이 로컬 변수와 같이 선언할 경우 일반적으로 글로벌 메모리에 할당되게 된다. CUDA에서 글로벌 메모리는 굉장히 느려서 잦은 읽기/쓰기 처리에 적합하지 않다.

__shared__ StackInfo SharedStack[STACK_SIZE];

만약 CUDA 코드 내에서 위와 같이 선언해주면 이 메모리는 쉐어드 메모리(shared memory)에 잡히게 된다.

< 그림 4> 스레드와 불럭, 그리고 메모리

CUDA는 여러 개의 스레드가 하나의 블럭을 형성하는데, 이 블럭에는 각각 일정 양의 쉐어드 메모리를 할당할 수 있다. 블럭은 프로그래머가 지정한 개수만큼 여러 개의 스레드로 구성된다. 즉, 여러 스레드가 모여서 1개의 블럭을 형성하고, 같은 블럭 내의 여러 스레드는 쉐어드 메모리를 공유해서 사용한다(<그림 4>를 보면 알 수 있듯이 두 스레드는 각각 자신만의 레지스터와 로컬 메모리를 가지고 있고 쉐어드 메모리를 공유해서 사용한다). 

공용 메모리는 무척 빠르지만 용량이 매우 제한돼 있고, 더욱이 공용 메모리를 많이 사용할수록 L1 캐시 메모리의 사용이 제한된다. 실제 우리가 캐시 메모리를 제어하진 않지만, 로컬 메모리나 글로벌 메모리를 액세스할 때는 캐시 메모리를 통해 접근하게 되는데, 공용 메모리를 사용할 경우 그만큼 L1 캐시 메모리의 사용량이 줄어든다. 반면 글로벌 메모리는 많은 공간을 가지고 있지만, 상대적으로 속도가 많이 느리기 때문에 프로그래머는 최대한 글로벌 메모리의 액세스를 줄여야 한다는 것을 염두에 둬야 한다.

CUDA 아키텍처에서는 이처럼 여러 메모리가 각각의 장단점을 가지고 있다. 또한 GPU 하드웨어에 따라 성능이나 제약에도 차이가 있기 때문에 사용하는 하드웨어의 특성에 대한 이해도 필요하다. 

다시 레이트레이싱의 스택으로 돌아와서 스택을 살펴보자. 사실 레이트레이싱에서 스택은 제법 많은 메모리를 필요로 한다. 렌더링할 장면에 따라 다르지만 스택이 수백 개가 필요한 경우도 많다.

< 그림 5> 자동차의 앞부분을 간략하게 렌더링한 장면(스택 크기를 200개로 잡아서 스택이 부족한 부분은 검은색으로 깨져서 표현됨)

앞서 용량 제약을 말했다시피 일반적인 방법으로는 이렇게 큰 스택을 쉐어드 메모리에 넣을 수 없다. 그리고 용량의 제약이 없는 글로벌 메모리나 로컬 메모리에 할당하기에는 너무 잦은 액세스 때문에 성능 저하에 원인이 될 수 있다. 이러한 제약 속(그리고 딱히 이런 제약과는 무관하게)에서 여러 연구가 진행됐는데, 여기서는 그 중 숏스택(short stack) 방법을 소개하겠다.

__shared__ StackInfo SharedStack[SHORT_STACK_SIZE*BLOCK _SIZE_X*BLOCK_SIZE_Y];
const int stackIndex = SHORT_STACK_SIZE*(tx+ty* BLOCK_ SIZE_X);
StackInfo *Stack = & SharedStack [stackIndex];

이 코드에서 BLOCK_SIZE_X, Y는 CUDA의 블럭 내의 스레드 개수를 2차원 배열로 지정한 형태다. 다음과 같이 CUDA의 커널(ShootRayWithShortStack 함수)을 실행할 때 지정해준다.

dim3 dimGrid( RayCount/BLOCK_SIZE_X/BLOCK_SIZE_Y, 1, 1 );
dim3 dimBlock( BLOCK_SIZE_X, BLOCK_SIZE_Y, 1 );
ShootRayWithShortStack<<<dimGrid, dimBlock>>>(…);

이 과정에서 stackIndex에 해당하는 부분이 각 스레드 자신에게 할당 받은 스택의 시작 주소가 된다. 따라서 &SharedStack [stackIndex]로 스레드 자신의 스택을 사용한다.

단, 여기서는 이 SHORT_STACK의 크기를 크게 잡지 않는다. 1~5개 정도로 매우 작은 양을 할당한다. 그러면 하나의 블럭은 SHORT_STACK*BLOCK_SIZE_X* BLOCK_SIZE_Y만큼의 스택을 쉐어드 메모리에 가지고 있게 되고, 각 스레드는 SHORT _STACK만큼의 자신만의 스택을 얻게 된다. 이런 작은 크기의 스택에서는 금방 스택 오버플로우가 발생하게 된다. 이때의 처리는 간단하게(원형 큐의 방식과 마찬가지로) 가장 아래에 있는 오래된 스택 정보를 삭제하고 새 스택 정보를 위에 넣는다.

< 그림 6> 숏스택에서 오버플로우 발생 시 가장 오래된 스택 정보를 삭제

오버플로우가 발생하지 않는 환경에서는 이 스택을 평소처럼 쓰면 된다. 하지만 작은 스택 크기 때문에 스택은 금방 비게 된다. 스택이 비어 있어서 더 이상 POP을 해야 할 수 없는 상황에서는 리스타트 알고리즘을 사용한다.

리스타트 알고리즘은 사실 매우 간단한데, 부족한 스택 정보를 어떻게 하기보다는 이미 계산한 부분에 해당하는 광선을 잘라내고 트리 탐색을 처음부터 다시 시작하는 것이다.

< 그림 7> 리스타트 알고리즘의 예(이미 탐색한 부분의 광선을 자르고 트리를 처음부터 다시 탐색한다. 이 경우 스택이 필요하지 않음)

<그림 7>에서 윗 그림은 현재 광선이 1에 해당하는 부분을 탐색 완료한 시점이다. 이때 원래는 상단의 인터널 노드에 해당하는 정보를 스택으로부터 얻어 와야 하는데, 스택의 정보가 없기 때문에 <그림 7>의 아래 그림처럼 광선의 꼬리 부분을 잘라낸다. 그리고 트리를 처음부터 다시 탐색하면 1번에 해당하는 리프 노드는 다시 탐색하지 않게 된다.

이 방법을 쓰면 당연히 노드를 중복해서 탐색하게 되는데, 현실적으로 봤을 때 예상외로 중복해서 탐색하게 되는 노드의 비율이 그렇게 많지 않다(<표 1> 및 참고자료 6번 참조).

< 표 1> Daniel R. Horn 등이 실험한 중복 노드의 탐색 비율(Robot 장면)

이처럼 숏스택 방식은 빠른 액세스가 가능한 적은 개수의 스택을 유지하면서도 리스타트 알고리즘을 통해 렌더링에 문제가 없게 하는 특징이 있다. 아무리 수백 개 크기의 스택이 필요하다고 하더라도 그 중에 실제로 잦은 액세스를 하는 것들은 적기 때문에 이러한 방법을 통해 속도 향상을 얻을 수 있다.

그 외의 성능 이슈들
실제로 고려해야 하지만 자세히 설명하지 못한 성능 이슈들은 많이 있다. 삼각형과 광선의 교차 계산은 제법 많은 연산을 필요로 한다. 이 과정을 최소화하기 위해서 삼각형을 미리 Barycentric 좌표계로 변환시켜서 실제 연산 시에 빠르게 수행하도록 하는 방법이 있다. 각 노드에 들어가는 정보도 8바이트로 압축할 수 있다. 

또한 사실 최근 GPU 하드웨어의 속도가 매우 향상됐기 때문에, 어느 정도 레이트레이서의 성능을 올려놓으면, 성능의 병목 지점은 레이트레이싱 처리가 아닌 CPU 메모리(host memory)와 GPU 메모리(device memory) 간의 데이터 전송이 된다. 이를 고려해 최대한 메모리 전송이 일어나지 않게 전반적인 레이트레이서 설계를 하는 것도 중요하다.

지금까지 KD트리가 만들어진 것으로 가정하고 설명했지만, 사실 KD트리는 어떻게 만드느냐에 따라서도 트리를 만드는 속도와 렌더링 속도가 영향을 미친다. 그리고 상대적으로 KD트리는 렌더링 속도가 빠르지만, 새로 트리를 만드는 것은 느린 편이다. 이에 최근에는 KD트리와 상당 부분 흡사한 BVH(Bounding Volume Hierarchy) 방법도 사용된다. BVH는 렌더링 속도는 KD트리보다 느리지만, 트리 업데이트 기능이 있어서(이는 KD트리에서는 불가능) 어느 정도의 동적인 장면에서도 사용할 수 있다.

레이트레이싱을 넘어서
기본적인 레이트레이싱을 수행하는 레이트레이서가 빠르게 구현되면, 여기서 다양한 기법들로 발전시킬 수 있다. 여기서는 광선 기준으로 가장 가까이 교차하는 삼각형을 얻는 것을 설명했는데, 이렇게 교차된 삼각형을 기준으로는 셰이딩, 라이팅을 해야 사실적인 렌더링 결과를 얻을 수 있다. 이 과정에는 기본적으로 다음과 같은 광선의 종류들이 존재한다.

- 주 광선(primary)
- 2차 광선(secondary ray)
- 그림자 광선(shadow ray)

주 광선은 카메라로부터 나오는 광선이다. 2차 광선은 반사각/입사각을 기준으로 해서 반사(혹은 굴절)되는 광선이다. 이 광선을 통해 거울과 같은 물체를 사실적으로 렌더링할 수 있다. 그림자 광선은 조명 방향으로 광선을 쏘아 조명과 셰이딩할 지점 사이에 가리는 물체가 없는지를 검사하는 것이다. 물체가 가린다면 이 부분은 그림자가 지는 부분이 된다. 이런 특징을 본다면 픽셀에 해당하는 주 광선 1개에 대해 여러 개의 2차 광선(보통 2~3개 정도의 단계를 밟는다)을 쏘고, 조명 개수만큼의 그림자 광선을 쏘기 때문에 레이트레이싱에 필요한 교차 검사는 엄청나게 많아진다.

뿐만 아니다. 포톤 매핑(Photon Mapping)을 사용하게 되면 픽셀과 상관없는 조명 기준의 많은 광선(포톤의 경로)을 사용하게 되고, 패스트레이싱(Pathtracing, 경로추적법) 등의 방법을 사용하게 되면 한 방향으로의 반사가 아니라 주변광을 모아야 하기 때문에 훨씬 많은 광선이 필요하게 된다. 이런 상황일수록 레이트레이서의 작은 성능 차이는 큰 결과의 차이를 가져온다.

결론
이 글은 레이트레이싱과 CUDA 모두 익숙하지 않은 독자에게는 혼란스러운 내용이 될 수 있지만, CUDA 사용에 있어서 메모리 사용의 제약이나 제약된 구조가 큰 걸림돌이므로 이를 위해 상황에 맞는 알고리즘들이 연구되고 개발된다는 점은 이해하면 좋겠다. CUDA는 매우 강력하게 병렬 처리를 도와주지만, 하드웨어와 구조에 대한 이해, 그리고 올바른 알고리즘의 도입 없이는 그 강력한 성능을 제대로 활용하지 못하거나 오히려 CPU보다 안 좋은 결과를 낼 수 있음을 기억하자.

참고자료
1. Ingo Wald, Realtime Ray Tracing and Interactive Global Illumination (PhD thesis), Saarland University, 2004
2. Daniel Reiter Horn et al., Interactive k-D Tree GPU Raytracing, Proceedings of the 2007 symposium on Interactive 3D graphics and games, 2007
3. Tim Foley and Jeremy Sugerman, KD-Tree Acceleration Structures for a GPU Raytracer, Proceedings of the ACM SIGGRAPH/EUROGRAPHICS conference on Graphics hardware, 2005
4. Michael Hapala et al., Efficient Stack-less BVH Traversal for Ray Tracing, Proceedings 27th Spring Conference of Computer Graphics (SCCG) 2011, 2011
5. Philip Dutre et al., Advanced Global Illumination, 2/E, ISBN-13: 978-1568813073, 2006
6. Daniel Reiter Horn et al., Kd-Tree Acceleration Structures for a GPU Raytracer, http://graphics.stanford.edu/papers/i3dkdtree/i3dkdtreeonline.ppt