몇일전에 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()실패
맨날 테스트하는 복셀 1500만개짜리 맵에서 tree구조의 메모리를 71MB 소모했었다. 정확히는 node의 메모리는 얼마 안되고 말단 node(leaf)에서 들고있는 삼각형 데이터의 메모리가 대부분이다. 교차 테스트를 위해서 leaf마다 삼각형배열을 가지고 있는데 이게 메모리를 제법 차지한다. 외장 그래픽 카드를 장착한 데스크탑이나 게이밍 노트북에선 이 정도 메모리 소모는 별 문제가 아니다. 하지만 내 테스트머신중 하나인 Surface book 1은 GPU메모리가 1GB밖에 … More CUDA측 Tree자료구조 메모리 줄이기.
코룸 온라인 개발하던 시절에… 베타테스트 직전에 운영팀에서 pc방을 돌면서 게임을 테스트했다. 그리고 버그 리포트라는게 전달이 됐는데 그중에 RIVA TNT에서 크래시한다는 내용이 있었다. 당시 RIVA TNT에서 압축 테스처를 지원하지 않기 때문에 생긴 문제였다. 아마 2003년 즈음이었을텐데 그때 많이 쓰던 그래픽 카드는 GeForce 2/4 MX, 고급군으로는 GeForce 3/4 TI 였다. 코룸온라인은 최소사양이 RIVA TNT였다. 나한텐 안물어보고 정했던것 … More 기능 지원 안되는 디바이스 지원에 대한 잡설
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를 이용한 복셀 월드의 라이트맵 계산 – 성능개선
