32-2의 warp의 early terminate에 대해 질문드립니다.

Question

안녕하세요! 정말 재미있게 강의를 보고 있습니다. 그런데 강의 32-2를 보다 궁금증이 생겨 질문 드립니다! warp의 early terminate가 정확히 어떤 동작인지 궁금한데요. 제가 생각해 볼 때 __syncthreads() 함수를 호출하게 되면, 블록 내 모든 스레드가 해당 함수에 도달하기까지 대기해야 하고 그러면 뒤에 더 이상의 동작이 없는 Warp들이라 하더라고 해당 함수에서 대기하고 있는 상태여야 할거 같습니다. 따라서 그러면 해당 warp는 다른 warp들이 동작을 마칠때까지 기다리게 되어 리소스를 반환하거나 하는 동작이 없을거 같은데 왜이렇게 성능이 향상되는지 궁금합니다.

onemoresipofcoffee · Answer

안녕하세요. CUDA 프로그래밍 강의를 수강해 주셔서 감사합니다.   질문에서 말씀하셨던 내용이 모두 맞습니다. __syncthreads() 함수가 적용되므로, 더 이상 동작이 없는 warp 라도 대기 상태로 있는 것도 맞습니다.   다만, 생각보다 문제의 크기가 크다는 점이 차이를 일으키고 있습니다.   32-2 강의를 본다면, 8번 슬라이드에서 수행하는 것이나, 12번 슬라이드에서 수행하는 것이나, 현재의 thread 갯수가 작을 때는, 별 차이가 없습니다. 그게 맞습니다. 다만, 8번 슬라이드나, 12번 슬라이드가 설명을 위해서 작은 싸이즈로 그린 것이고, 실제 thread 의 총 갯수는, 17번 슬라이드를 보시면, 예제로 돌릴 때, thread 갯수를 16M 개로 설정하고, thread block size 는 1024개로 설정했습니다. (즉, 16M / 1K = 16K 개의 thread block 이 돌게 됩니다.)   __syncthreads() 함수가 thread block 단위로 적용된다는 점도 고려하면, 이제 전체적인 efficiency는 thread 1개가 아니라, thread block 단위로 고려해 보시면 될 겁니다. 실행되면, 16M / 1K = 16K 개의 thread block 이 돌기 시작하는데, 8번 슬라이드의 방식이라면, 대부분의 thread block 이 끝까지 살아 남아 있을 겁니다. 반면, 12번 슬라이드 방식에서는, 시작하자마자, 절반 정도 되는 8K개의 thread block 이 모두 early terminate 되어 버립니다. 다음 단계에 다시 4K 개가 early terminate 되어 버릴 것이고요. CUDA 전체로 보아서는 thread block 의 수가 급격히 감소하면서, 효율성이 급격히 좋아지게 됩니다. 그래서, 전체 속도가 매우 빨라지게 되는 것입니다.   생각보다 thread block 의 갯수가 많은, 꽤 큰 사이즈로 예제를 돌렸다는 것이 key point 가 될 것입니다. 감사합니다.