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

안녕하세요.

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 가 될 것입니다.

감사합니다.