현재 SW Occlusion Culling + HW Occlusion Culling을 사용중이다. 일단 복셀지형을 놓고 볼때 어지간하면 오브젝트 개수가 5만에서 8만개씩 나오기 때문에 S/W Occlusion Culling을 먼저 걸어주지 않으면 H/W Occlusion Culling에서도 꽤 큰 손해를 본다. draw call이 엄청 증가하기 때문이다. 그런데 S/W Occlusion Culling은 또 CPU 처리에서 상당한 지연을 일으킨다. 그래서 이걸 async로 처리하고 있긴 한데 그 … More S/W Raytracing을 이용한 복셀 오브젝트 Culling
2020년 8월 25일 방송
몇일전에 CUDA의 Unified Memory System에 문제가 있다고 포스팅했었다. CUDA Unified Memory 사용시 시스템이 멈춰버리는 현상. 아..아쉽다. 정말 멋진 기능이었는데. CUDA프로그래밍 해본 사람은 알텐데 cudaMalloc()으로 GPU메모리를, cudaMallocHost()로 CPU측 메모리를 할당하고 이 두가지 다른 메모리의 내용을 수동으로 동기화시키는 일은 상당히 번거롭다. 특히 Voxel Horizon프로젝트처럼 시스템 메모리에 구축한 KD-Tree를 GPU 메모리로 옮길땐 아주 짜증나는 작업을 해야한다. GPU측, CPU측 … More Baking Light-map. CUDA vs CPU
게임 클라이언트에서 크래시 발생시 덤프파일을 생성하도록 해놨다. 데스크탑버전은 적용이 되어있는데 UWP는 적용이 안되어있어서 UWP에도 적용하려고 데스크탑버전의 덤프생성 코드를 테스트했다. 그런데 문제 발생. 어라? 근데 덤프 생성에 문제가 있다. 덤프파일은 만들어졌지만 정상적으로 디버거에서 읽히지 않는다. MiniDumpWriteDump()함수가 실패하네? 해당 에러값으로 이틀동안 구글링을 해봐도 도움되는 답이 없다. 테스트를 하다보니 64비트에서만, MiniDumpWithFullMemory플래그를 줄 경우만 발생하는걸 알게 됐다. ’64비트는 뭐가 … More CUDA사용시 MiniDumpWriteDump()실패
오늘 메모리 사용량 최적화를 끝으로 CUDA적용은 마무리 지었다. 추가적으로 CUDA 사용이 필요할 경우 CUDA기능을 모아둔 dll라이브리에 기능을 추가하면 된다. Voxel Horizon은 라이트맵(Light map)을 사용한다. 지형지물간의 그림자도 라이트맵을 구울때 ray와 삼각형들간 교차테스트를 수행해서 계산한다. Voxel Horizon은 실시간으로 지형지물이 변화한다. 따라서 라이트맵을 다시 굽는 일이 수시로 발생한다. 네트워크로 복셀 오브젝트를 스트리밍했을때, 복셀들을 추가/삭제/편집 해서 복셀들간 지형지물이 변경될때 … More Game Dev – Voxel Horizon에 CUDA적용 마무리.
맨날 테스트하는 복셀 1500만개짜리 맵에서 tree구조의 메모리를 71MB 소모했었다. 정확히는 node의 메모리는 얼마 안되고 말단 node(leaf)에서 들고있는 삼각형 데이터의 메모리가 대부분이다. 교차 테스트를 위해서 leaf마다 삼각형배열을 가지고 있는데 이게 메모리를 제법 차지한다. 외장 그래픽 카드를 장착한 데스크탑이나 게이밍 노트북에선 이 정도 메모리 소모는 별 문제가 아니다. 하지만 내 테스트머신중 하나인 Surface book 1은 GPU메모리가 1GB밖에 … More CUDA측 Tree자료구조 메모리 줄이기.
1. 코딩/디버깅환경이 거지같은 Compute Shader를 작성하기 전에 성능을 가늠해볼수 있다. 보수적으로 잡았을때 Compute Shader가 CUDA의 50% – 80%정도 나온다고 예측해보면 CUDA코드가 CPU코드보가 3-4배 빠르면 충분히 GPU코드를 추가 작성할 가치가 있다. 2. Compute Shader를 추가작성하기로 결심했더라도 CPU코드 -> Compute Shader로 바로 가기는 어렵다. 디버깅 환경이 거지같으니 코드 짜놓고도 이게 맞는건지 틀린건지 확인하기 어렵다. CPU -> Compute … More Voxel Horizon프로젝트에 CUDA를 적용하고 있는 이유
KD-Tree를 순회할땐 stack이 필요하다. stack을 구현한다는게 CPU코드에선 아무 문제도 아니지만 CUDA에선 고민이 좀 필요하다. CUDA에서 자유롭게 쓸 수 있는 메모리는 Global Memory이다. 그런데 겁나 느리다. 시스템 메모리보단 엄청 빠르지만 GPU보단 한참 느리다. stack을 global memory에 올려놓으면 아무래도 억세스할때마나 stall이 걸린다. CUDA에서 사용할 수 있는 고속 메모리는 shared memory가 있다. 여기다 stack을 올려놓으면 물론 빠르다. 그런데 … More CUDA에서의 stack 구현
엊그제 포스팅했던 CUDA를 이용한 라이트맵 베이킹의 성능향상에 대한 얘기다. 좀더 GPU를 빡시게 사용하려고 cuda Stream + 멀티스레드 까지 적용했는데 생각보단 성능이 많이 오르지 않는다. 프로파일링 해보니 Occupancy가 이론상 50% 달성할 수 있는데 실제로는 29.3%만 달성한 것으로 나온다. Occupancy가 영 만족스럽지 못하다. 그래서 이런저런 튜닝을 좀 했다. 오랫만에 Occupancy Calculator를 실행해서 SM 5.2와 6.1기준으로 숫자를 넣어 … More CUDA를 이용한 복셀 월드의 라이트맵 계산 – 성능개선
복셀 월드의 라이트맵 계산을 CUDA로 하도록 만들었다. 2010년에 이미 CUDA로 라이트맵을 구웠기 때문에 별로 대단한 작업은 아니다. 다만 이번엔 CUDA 6부터 추가된 Unified Memory 체계를 사용했다는 점이 다르다. 우선 기본적으로 라이트맵 계산을 하는 원리를 보자. 복셀 오브젝트마다 공간상의 점과 텍스처 좌표를 연결시킬 patch배열을 가지고 있다. 각 patch들로부터 라이트까지 ray를 쏴서 거리과 각도에 따라 라이트값을 구한다. … More CUDA를 이용한 복셀 월드의 라이트맵 계산
yuchi's development 
댓글을 달려면 로그인해야 합니다.