몇일전에 CUDA의 Unified Memory System에 문제가 있다고 포스팅했었다.
CUDA Unified Memory 사용시 시스템이 멈춰버리는 현상.
아..아쉽다. 정말 멋진 기능이었는데.
CUDA프로그래밍 해본 사람은 알텐데 cudaMalloc()으로 GPU메모리를, cudaMallocHost()로 CPU측 메모리를 할당하고 이 두가지 다른 메모리의 내용을 수동으로 동기화시키는 일은 상당히 번거롭다.
특히 Voxel Horizon프로젝트처럼 시스템 메모리에 구축한 KD-Tree를 GPU 메모리로 옮길땐 아주 짜증나는 작업을 해야한다. GPU측, CPU측 메모리 구분 없이 cudaMallocManaged()하나로 GPU와 CPU 모두 억세스 가능한 통합 메모리 어드레스를 제공한다는 기능은 정말 달콤하고 매력적일수밖에 없다. 실제로 이번 라이트맵 베이킹 기능은 그렇게 쉽게 만들었다.
하지만 이런 망할. 앞서 언급했듯 속편하게 이 기능을 사용할 수 없다. 현재 테스트해본 대부분의 드라이버에서 Unified Memory System은 문제를 드러내고 있다. 속도가 엄청나게 느리던가, 아니면 시스템이 멎어버리던가.
하여간 어쩔수 없다. 더 이상 스트레스 받고싶지 않다. 문제가 발생하지 않는 드라이버를 기다리는것도 지쳤다.
그래서 직접 코드를 고쳐서 피해가기로 했다. 물론 내가 드라이버를 고칠순 없다. Unified Memory System이 문제가 되므로 이걸 사용하지 않도록 수정하기로 했다.
과거에 처음으로 CUDA로 라이트맵 굽는 코드를 작성했을때와 마찬가지로 GPU측 메모리와 CPU측 메모리의 포인터를 분리했다. cudaMallocManaged()로 할당한 코드들을 찾아서 명시적으로 cudaMalloc()과 cudaMallocHost()를 사용하도록 고쳤다. 상황 봐서 시스템 메모리를 캐시처럼 사용하거나 GPU 메모리 할당할때 맞춰서 같은 사이즈로 할당하거나 했다.
CUDA 커널 함수 호출하기 전에 cudaMemcpyAsync()로 시스템 메모리 -> GPU메모리 전송, 커널호출후 cudaMemcpyAsync()로 GPU메모리 -> 시스템메모리로 전송.
꽤 짜증나는 작업이었는데 어쨌든 이틀에 걸친 수정끝에 cudaMallocManaged()를 싹 제거했다.
이후 테스트를 해봤다.
일단 시스템이 멈추는 일은 현재까진 없다. 하지만 이건 좀 더 테스트를 해봐야할것 같다.
재밌는건 실제 구동 속도도 확연히 빨라졌다는 사실.
매번 테스트에 사용하는 복셀 오브젝트 5만여개의 맵 기준으로 방향성 라이트 하나 넣고 라이트맵을 구우면 집 데스크탑에서 2.xx초가 걸렸는데 cudaMallocManaged() -> cudaMalloc()/cudaMallocHost() 로 바꾸고나서 1.2초로 줄었다. 2배 빨라졌다.
nvidia 문서에는 cudaMallocManaged()를 사용할 경우 기존 방식의 90%이상의 성능은 보장한다고 적혀있었다. 지랄. 개뻥이다.
시스템이 멈추는 문제를 피하고 10배쯤 느려지는 문제를 피해도, 최상의 상태에서 돌려도 Unified Memory System은 GPU/CPU어드레스를 분리해서 처리하는 방식보다 2배는 느리다.
테스트하는 김에 cudaMallocManaged()를 제거한 버전으로 간단한 성능 비교를 해봤다.
라이트맵 베이킹 – CPU vs GPU 테스트다.
CPU 12스레드(i7 8700k)로 돌렸을때 8.7초,
CUDA(GTX970)로 돌렸을때 1.2초다. 7배 이상 빠르다.
코드 고치는게 좀 짜증나긴 했지만 빌어먹을 시스템 크래시 현상을 겪어서 차라리 다행이다. 성능을 올릴 수 있는 방법을 찾아서 다행이고 망할 Unified Memory System을 쓰면 안된다는 교훈을 얻어서 다행이다.
yuchi's development 