32-2의 warp의 early terminate에 대해 질문드립니다.
안녕하세요!
정말 재미있게 강의를 보고 있습니다.
그런데 강의 32-2를 보다 궁금증이 생겨 질문 드립니다!
warp의 early terminate가 정확히 어떤 동작인지 궁금한데요.
제가 생각해 볼 때 __syncthreads() 함수를 호출하게 되면, 블록 내 모든 스레드가 해당 함수에 도달하기까지 대기해야 하고 그러면 뒤에 더 이상의 동작이 없는 Warp들이라 하더라고 해당 함수에서 대기하고 있는 상태여야 할거 같습니다.
따라서 그러면 해당 warp는 다른 warp들이 동작을 마칠때까지 기다리게 되어 리소스를 반환하거나 하는 동작이 없을거 같은데 왜이렇게 성능이 향상되는지 궁금합니다.
Answer 1
1
안녕하세요.
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 가 될 것입니다.
감사합니다.
BFS, DFS 활용이 되는 상황에서의 방향성
0
3
1
코딩살구클럽 승인
0
10
1
코딩살구클럽승인
0
9
1
리눅스로 진행해도 될까요?
0
20
2
코딩살구클럽 승인
0
43
2
3-D 관련 질문
0
33
2
코살구 회원가입 문의
0
38
2
코살구 로그인 문제
0
58
2
26년 1회차 기출 강의
0
54
2
3-A 문제 풀이 관련 질문
0
51
3
2-O 질문 있습니다
0
38
2
2-T 문제에 관한 질문
0
38
2
코딩 살구 클럽 접속 및 사용방법 문의
0
56
2
기출문제
0
55
1
안녕하세요~. 현재 코살코딩클럽 사이트가 접속이 안됩니다~
0
64
2
코딩살구클럽 로그인문제
0
74
3
정보처리기사 실기 zip
0
44
2
데이터베이스 노션
0
38
2
본 강의와는 상관없는 내용입니다만..
0
35
1
수업노선자료파일 다운
0
34
2
코딩 살구 클럽 로그인 문제
0
79
2
[30-4] hist-shared.cu 에서 질문있습니다.
0
153
2
선생님 gpu->cpu 속도 개선에 대해서 질문드려요
0
616
1
28-4 shuffle shared - shared memory 강의에서 질문이 있습니다.
0
429
1