복셀 월드의 라이트맵 계산을 CUDA로 하도록 만들었다.
2010년에 이미 CUDA로 라이트맵을 구웠기 때문에 별로 대단한 작업은 아니다.
다만 이번엔 CUDA 6부터 추가된 Unified Memory 체계를 사용했다는 점이 다르다.
우선 기본적으로 라이트맵 계산을 하는 원리를 보자.
- 복셀 오브젝트마다 공간상의 점과 텍스처 좌표를 연결시킬 patch배열을 가지고 있다.
- 각 patch들로부터 라이트까지 ray를 쏴서 거리과 각도에 따라 라이트값을 구한다. 기본 광도를 구한 다음에 radiosity를 적용해서 연쇄적으로 난반사를 적용해도 되고 그냥 이대로 사용할수도 있다. 어쨌든 기본적으로 이런식으로 구현한다.
- ray가 라이트에 도달하기 전에 지형지물을 만나면 라이트 적용을 받지 못한다. 이 경우 그림자가 드리워진것과 같다.
- 그러기 위해선 월드상의 지형지물에 대해서 ray충돌 테스트를 해야한다. 오브젝트 한개당 1400-1500개 정도의 patch를 들고 있고 오브젝트 개수는 5만개 이상, 삼각형 개수는 100만개 이상이다. 따라서 (모든 patch x 모든 삼각형)으로 검사하면 면 엄청나게 느리다.
- 월드는 KD-Tree로 공간분할 되어있다. 복셀 오브젝트들은 KD-Tree에 들어있다. 그리고 각 복셀 오브젝트들은 자기 자신의 최적화된 삼각형 데이터를 들고 있다.
- 따라서 KD-Tree 순회를 수행하여, N x M 검사보다 압도적으로 빠르게 충돌하는 노드를 찾을 수 있다.
- 노드가 리프일 경우 리프 안에 복셀 오브젝트들에 대해서 ray충돌검사를 한다. 각 오브젝트들은 AABB를 가지고 있으므로 ray vs AABB체크를 먼저 한다.
- 오브젝트와 충돌할 경우 오브젝트 안의 삼각형들에 대해서 ray 충돌 검사를 한다.
- 최종적으로 충돌할 경우 이 패치는 라이트를 받지 않는다.
- 충돌하지 않으면 라이트값을 계산해서(N dot L등) 패치에 저장한다.
- 라이트 개수가 추가되면 각각의 라이트에 대해서 이 과정을 반복하면서 패치에 저장된 광도값을 더해간다.
CPU로 계산할 경우
- 논리코어 개수만큼의 스레드를 만들어서 오브젝트 1개 – 1스레드로 위에서 언급한 KD-Tree Traversal을 수행하며 패치의 값을 채운다.
[GPU로 계산할 경우]
- 시스템 메모리 상에 존재하는 KD-Tree를 GPU메모리로 옮겨야한다.
- 이게 예전엔 꽤 성가신 작업이었다. cudaMallocHost()로 시스템 메모리를 할당받고, cudaMalloc()으로 GPU메모리를 할당 받아서 cudaMemcpy()를 이용해서 수동으로 내용을 맞춰줘야했다. 코드는 짜도 일단 CPU상에서 디버깅을 할땐 GPU메모리에 올린 트리가 제대로 만들어졌는지 확인할 방법이 없다.
- 최근의 CUDA는 Unified Memory 체계를 제공한다. 따라서 cudaMallocManaged()로 할당한 메모리는 cpu측에서도 억세스 가능, GPU측에서도 억세스 가능하다. 따라서 CUDA 함수들을 작성할때 __host__ __device__ 지정자를 모두 적용해 주면 한벌의 코드로 CPU/GPU 모두 사용 가능하다. CPU 디버깅으로 충분히 테스트하고 검증할 수 있다.
- 오브젝트 -> CUDA의 Block으로 맵핑한다. Block한개는 최대 1024스레드를 수행할 수 있는데 오브젝트 한개는 대개 1400-1500개 정도의 패치를 가지므로 Block한개로는 오브젝트 한개를 처리하지 못한다. 따라서 오브젝트 한개의 패치들을 여러개의 Block에 나눠담는다.
- Block에 맵핑될 구조체에는 패치의 메모리 포인터, 패치 개수, 오브젝트의 위치와 AABB를 담고 있다. CUDA커널 함수 안에서 Block Index로 이 구조체의 배열을 억세스한다.
- CUDA안에서 KD-Tree를 순회할때 Stack이 필요하다. 이것은 로컬 메모리 배열을 잡고 Stack자료구조에 맵핑하는 식으로 구현하면 된다. Stack을 shared memory에 잡으면 빠르다. 문제는 이 경우 Shared Memory용량이 부족해서 많은 수의 Block들이 맵핑되지 못한다. 그래서 되려 느려질수 있다. 이번 코드에선 최대한 많은 Block을 맵핑하기 위해 Shared Memory는 아예 사용하지 않았다.
[결과]
맵 구성 : 예의 1500만개 복셀, 오브젝트 5만4천개의 월드에서 방향성 라이트 1개
하드웨어 : GTX970, i7 8700K
-no 최적화
처음 만들었을땐 GPU점유율이 20%를 넘기지 못했다. 시간도 20초 넘게 걸렸다.
– 커널 함수 일부 최적화 , 메모리 할당 빈도 줄임
GPU점유율 35% ,16초 정도
– cudaStream적용
cudaStream을 이용해서 async방식으로 고치고 GPU점유율 75% – 80%까지 올렸다. 약 3.4초 소요.
목표가 2초인데 남은 20%를 다 쓸 수 있으면 2.8초 정도까진 가능할것 같다.
현재 8개의 cudaStream을 사용하고 있다. stream개수를 조절해보면 좀더 GPU 점유율을 높일 수 있을지 모르겠다.
현재 싱글스레드로 돌아가는데 이걸 멀티스레드로 바꾸면 남은 20%더 쓸 수 있을것 같긴하다. 하지만 코드 복잡해지니 생각좀 해봐야겠다.
[추가 – CPU코드와의 성능 비교]
현재 CPU코드는 라이트맵을 일괄적으로 한방에 굽지 않도록 되어있다. 게임 플레이중에 스트리밍한 복셀 오브젝트들에 대해 실시간으로 라이트 계산을 하는게 목적이었기 때문이다.
그래서 정확한 비교가 불가능해서 어제는 CPU코드와의 비교를 하지 않았다. 하지만 역시 이 부분은 정확히 체크할 필요가 있어서 동일 조건으로 한방에 몰아서 계산하는 코드를 만들어 넣었다.
CUDA코드에선 stream개수를 16개로 늘렸다. 성능이 약간 더 향상되었다.
결과는 다음과 같다.
테스트 맵 : 복셀오브젝트 54056개, 데이터상의 복셀 수 1500만여개.
CPU : i7 8700K ,6 core , 6 threads ( HT로 12스레드를 지원하지만 성능차가 나는 경우도 대개 10%이내이므로 6스레드만 사용했다.)
– 8831.12 ms
GPU : GTX970
– 4360.28 ms
CUDA코드가 2배 정도 빠르다.
[추가]
KD-Tree를 GPU측에 구성하는 시간이 포함되어있다.
실제 라이트 계산만 감안하면 약 2.9초 걸린다. 2배보다 더 빠르다.
yuchi's development 