최적화 썰

RTX그래픽 카드 입수 기념 잡담방송중에 내 게임을 플레이하면서 약간 이상한점을 느꼈다.
어떤 맵들은 처음 입장후 마우스를 휘릭 움직이는 순간 2-3초 정도 화면이 멈춤다. 사실 이전에도 비슷한 증상을 느끼곤 했는데 그냥 다른 중요한 이슈들이 있어서 신경을 못쓰고 있었다.
그런데 다른 사람들이 보고 있으니 그 2-3초나 끊기는게 너무 쪽팔렸다.

방송 끝내고 조사를 시작했다.
원인은 최초 맵 로딩 후 CUDA측 메모리에 KD-Tree와 각 leaf마다 삼각형 데이터를 업데이트할때 시간이 오래 걸리는 것이었다.
처음엔 삼각형 긁어오는게 느린줄 알았다.
그래서 멀티스레드로 코드를 뜯어고치기 시작했다. 멀티스레드를 쓰려면 어차피 cudaStream을 사용해야한다. cudaStream을 사용하면 CPU<->GPU 메모리간 전송과 커널 호출을 비동기적으로 중첩해서 진행할 수 있다.

멀티 스레드를 적용하기 시작하니 코드가 몇 배씩 복잡해지기 시작한다. 이걸 어찌 수습하나.
근데 혹시나 해서 leaf에서 삼각형 데이터 가져오는 코드를 찾아보니 이건 삼각형 긁어올때 느린게 아니다. 즉 멀티스레드로 빨라질 문제가 아니다.

CUDA측 코드를 다시 살펴봤다.
흠. node 하나 업데이트할때마다 cudaMalloc()과 cudaMallocHost()를 호출한다.
cudaMalloc()은 GPU측 메모리 할당이다. 이 메모리는 맵이 떠 있는동안 동안 계속 유지되어야 한다. 일단 손대기 껄끄러우니 제외.
그리고 cudaMallocHost()는 GPU메모리로 node와 삼각형 데이터를 전송하기 위해 필요한 upload buffer를 할당하기 위해 사용한다.

‘임시로 사용할 메모리인데 이걸 매번 cudaMallocHost()로 할당할 필요가 있나? 한번에 왕창 잡아두고 앞에서부터 끊어쓰면 되지 않나…?’ 라는 정상적인 생각을 이제서야 했다.

이런 작업용 메모리는 충분히 작은 사이즈라면 스택에 로컬변수로, 좀 크면 공용으로 큰 블럭 하나 잡아두고 사용하면 된다.그런데 cudaStream을 이용해서 비동기 카피를 하게 되면 카피 끝날때까지 메모리가 유지는 되어야 한다. pinned memory임이 보장되어야 하므로 cudaMallocHost()를 사용해야겠고. 그러니 스택이나 임시 메모리를 한 블럭을 공유하는 방법은 사용할 수 없다.

나중에 좀 제대로 만들기로 하고 일단 cudaMallocHost()로 100MB짜리 큰 메모리 블럭 하나를 할당했다. cudaMallocHost()를 호출할 상황에서 이 메모리 블럭의 앞부분부터 순차적으로 끊어서 사용한다.

3만개의 node를 업데이트할 때 기존에는 3만번 cudaMallocHost()를 호출했지만, 이 코드 수정으로 cudaMallocHost()한번 호출에 메모리를 가져갈 위치를 가리키는 offset변수의 덧셈만 3만번 수행하게 된다.

여기까지 고치고 나니 전체 트리 업데이트에 걸리는 시간이 2.7초 정도에서 -> 120ms 정도로 왕창 줄었다.

이제 처음 하려던대로 cudaStream을 적용했다.
cudaMemcpy()로 GPU<->CPU메모리간 전송을 하면 이것은 PCI버스로 통신하는 I/O작업이기 때문에 데이터 사이즈에 상관없이 기본 비용(시간지연)이 상당히 발생한다.
이 시간동안 cudaMemcpy()를 호출한 스레드는 그냥 대기(blocking)하게 된다.
cudaSream을 적용해서 cudaMemcpyAsync()를 호출하면 스레드는 blocking되지 않는다.
모든 작업을 완료(정확히는 작업요청을 완료) 한 후 cudaStream핸들에 대해서 완료여부를 대기하면 된다.
싱글 스레드 상태에서도 cudaStream을 사용하면 카피요청을 해놓고 cpu작업을 계속 진행할 수 있다. 백그라운드에선 PCI버스를 타고 CPU <-> GPU메모리간 전송이 진행된다.

싱글스레드 상태에서 cudaStream 한개를 적용하고 트리 업데이트 시간은 120ms -> 70ms로 또 줄었다.

일반적으로 copy engine이 2개 이상인지라 여러개의 cudaStream을 적용했는데 전송량 자체가 많은게 아닌지라 오히려 약간 성능이 떨어졌다. cudaStream개수가 많아지면 동기화할때 대기하는 시간이 길어진다.

최종적으로 싱글스레드에 cudaStream 1개 사용하는것으로 확정.

최적화의 비법은 의외로 가까운 곳에 있나니… 사실 대부분의 경우 성능을 갉아먹는 요인은 메모리 할당/해제 쪽이다.
내 코드의 흠(불필요한 heap메모리 할당/해제)만 잘 찾아봐도 대단한 노력 없이 성능을 향상시킬 수 있다.

그런데 진짜 하고 싶은 말은 이게 아니고….
나도 내 게임 테스트 잘 안하는데 방송하면서 설렁설렁 플레이하다보면 버그가 나오곤 한다.
쪽팔리기도 하고 해서 방송 끝나자마자 엄청 집중해서 원인을 찾게 되는데 이게 꽤 도움이 된다.

결론은 역시 남에게 보여주려고 하면 버그가 잘 드러난다.


답글 남기기

댓글을 게시하려면 다음의 방법 중 하나를 사용하여 로그인 하세요:

WordPress.com 로고

WordPress.com의 계정을 사용하여 댓글을 남깁니다. 로그아웃 /  변경 )

Google photo

Google의 계정을 사용하여 댓글을 남깁니다. 로그아웃 /  변경 )

Twitter 사진

Twitter의 계정을 사용하여 댓글을 남깁니다. 로그아웃 /  변경 )

Facebook 사진

Facebook의 계정을 사용하여 댓글을 남깁니다. 로그아웃 /  변경 )

%s에 연결하는 중