스트리밍 멀티 프로세서, 블록 및 스레드 (CUDA)
CUDA 코어, 스트리밍 멀티 프로세서, CUDA 블록 및 스레드 모델 간의 관계는 무엇입니까?
무엇이 병렬화되고 무엇이 병렬화되며 어떻게 매핑됩니까? 더 효율적인 것은 무엇입니까? 블록 수 또는 스레드 수를 최대화합니까?
내 현재 이해는 멀티 프로세서 당 8 개의 cuda 코어가 있다는 것입니다. 모든 cuda 코어는 한 번에 하나의 cuda 블록을 실행할 수 있습니다. 해당 블록의 모든 스레드는 특정 코어에서 직렬로 실행됩니다.
이 올바른지?
스레드 / 블록 레이아웃은 CUDA 프로그래밍 가이드 에 자세히 설명되어 있습니다. 특히 4 장은 다음과 같이 설명합니다.
CUDA 아키텍처는 확장 가능한 다중 스레드 SM (Streaming Multiprocessor) 배열을 중심으로 구축됩니다. 호스트 CPU의 CUDA 프로그램이 커널 그리드를 호출하면 그리드 블록이 열거되고 사용 가능한 실행 용량이있는 멀티 프로세서에 배포됩니다. 스레드 블록의 스레드는 하나의 멀티 프로세서에서 동시에 실행되고 여러 스레드 블록은 하나의 멀티 프로세서에서 동시에 실행될 수 있습니다. 스레드 블록이 종료되면 비워진 다중 프로세서에서 새 블록이 시작됩니다.
각 SM에는 8 개의 CUDA 코어가 포함되어 있으며 한 번에 32 개 스레드의 단일 워프를 실행하므로 전체 워프에 대해 단일 명령을 실행하는 데 4 클럭 사이클이 필요합니다. 주어진 워프의 스레드가 잠금 단계에서 실행된다고 가정 할 수 있지만 워프간에 동기화하려면 __syncthreads().
GTX 970에는 각각 128 개의 Cuda 코어가있는 13 개의 스트리밍 멀티 프로세서 (SM)가 있습니다. Cuda 코어는 스트림 프로세서 (SP)라고도합니다.
블록을 GPU에 매핑하는 그리드를 정의 할 수 있습니다.
스레드를 스트림 프로세서 (SM 당 128 Cuda 코어)에 매핑하는 블록을 정의 할 수 있습니다.
하나의 워프는 항상 32 개의 스레드로 구성되며 워프의 모든 스레드는 동시에 실행됩니다.
GPU의 가능한 모든 기능을 사용하려면 SM에 SP가있는 것보다 SM 당 훨씬 더 많은 스레드가 필요합니다. 각 컴퓨팅 기능에는 한 번에 하나의 SM에 상주 할 수있는 특정 수의 스레드가 있습니다. 정의한 모든 블록은 대기열에 추가되고 SM이 리소스 (사용 가능한 SP 수)를 가질 때까지 기다린 다음로드됩니다. SM이 Warps를 실행하기 시작합니다. 하나의 워프에는 32 개의 스레드 만 있고 SM에는 예를 들어 128 개의 SP가 있으므로 SM은 주어진 시간에 4 개의 워프를 실행할 수 있습니다. 문제는 스레드가 메모리 액세스를 수행하면 스레드가 메모리 요청이 충족 될 때까지 차단된다는 것입니다. 숫자로 볼 때 : SP의 산술 계산은 18-22 사이클의 지연 시간이있는 반면 캐시되지 않은 전역 메모리 액세스는 최대 300-400 사이클이 걸릴 수 있습니다. 즉, 한 워프의 스레드가 데이터를 기다리는 경우 128 개의 SP 중 일부만 작동합니다. 따라서 스케줄러는 가능한 경우 다른 워프를 실행하도록 전환합니다. 이 워프가 차단되면 다음 작업이 실행됩니다. 이 개념을 대기 시간 숨김이라고합니다. 워프 수와 블록 크기는 점유 (SM이 실행하도록 선택할 수있는 워프 수에서)를 결정합니다. 점유율이 높으면 SP에 대한 작업이 없을 가능성이 높습니다.
각 cuda 코어가 한 번에 한 블록을 실행한다는 귀하의 진술은 잘못되었습니다. 스트리밍 멀티 프로세서에 대해 이야기하면 SM에있는 모든 스레드에서 워프를 실행할 수 있습니다. 한 블록의 크기가 256 스레드이고 GPU가 SM 당 2048 개의 스레드가 상주하도록 허용하는 경우 각 SM에는 SM이 실행할 워프를 선택할 수있는 8 개의 블록이 있습니다. 실행 된 워프의 모든 스레드는 병렬로 실행됩니다.
다양한 컴퓨팅 기능 및 GPU 아키텍처에 대한 숫자는 https://en.wikipedia.org/wiki/CUDA#Limitations에서 찾을 수 있습니다.
Nvidia 점유 계산 시트 (Nvidia 제공) 에서 점유 계산 시트를 다운로드 할 수 있습니다 .
Compute Work Distributor는 SM에 스레드 블록 (공유 메모리, 워프, 레지스터, 장벽 등)에 대한 충분한 리소스가있는 경우에만 SM에서 스레드 블록 (CTA)을 예약합니다. 공유 메모리와 같은 스레드 블록 수준 리소스가 할당됩니다. 할당은 스레드 블록의 모든 스레드에 대해 충분한 워프를 생성합니다. 리소스 관리자 는 SM 하위 파티션에 라운드 로빈 을 사용하여 워프를 할당 합니다. 각 SM 하위 파티션에는 워프 스케줄러, 레지스터 파일 및 실행 단위가 포함됩니다. 워프가 하위 파티션에 할당되면 완료 될 때까지 하위 파티션에 남아 있거나 컨텍스트 전환 (파스칼 아키텍처)에 의해 선점됩니다. 컨텍스트 스위치 복원시 워프는 동일한 SM 워프 ID로 복원됩니다.
워프의 모든 스레드가 완료되면 워프 스케줄러는 워프에서 발행 한 모든 미해결 명령이 완료 될 때까지 기다린 다음 리소스 관리자가 워프 ID 및 등록 파일을 포함하는 워프 수준 리소스를 해제합니다.
스레드 블록의 모든 워프가 완료되면 블록 수준 리소스가 해제되고 SM은 Compute Work Distributor에게 블록이 완료되었음을 알립니다.
워프가 하위 파티션에 할당되고 모든 리소스가 할당되면 워프는 활성화 된 것으로 간주되어 워프 스케줄러가 워프 상태를 적극적으로 추적하고 있음을 의미합니다. 각주기에서 워프 스케줄러는 중단 된 활성 워프와 명령을 실행할 수있는 워프를 결정합니다. 워프 스케줄러는 우선 순위가 가장 높은 워프를 선택하고 워프에서 1-2 개의 연속 명령을 발행합니다. 이중 발행에 대한 규칙은 각 아키텍처에 따라 다릅니다. 워프가 메모리로드를 발생 시키면 종속 명령어에 도달 할 때까지 독립 명령어를 계속 실행할 수 있습니다. 그런 다음 워프는로드가 완료 될 때까지 중단되었다고보고합니다. 종속 수학 명령어도 마찬가지입니다. SM 아키텍처는 워프 사이의주기 당 전환을 통해 ALU와 메모리 대기 시간을 모두 숨기도록 설계되었습니다.
이 답변은 잘못된 멘탈 모델을 도입하기 때문에 CUDA 코어라는 용어를 사용하지 않습니다. CUDA 코어는 파이프 라인 단 정밀도 부동 소수점 / 정수 실행 단위입니다. 문제 비율 및 종속성 지연 시간은 각 아키텍처에 따라 다릅니다. 각 SM 하위 파티션 및 SM에는로드 / 저장 단위, 배정 밀도 부동 소수점 단위, 반 정밀도 부동 소수점 단위, 분기 단위 등 다른 실행 단위가 있습니다.
성능을 최대화하기 위해 개발자는 블록 대 워프 대 레지스터 / 스레드의 균형을 이해해야합니다.
점유라는 용어는 SM에서 최대 워프에 대한 활성 워프의 비율입니다. Kepler-Pascal 아키텍처 (GP100 제외)에는 SM 당 4 개의 워프 스케줄러가 있습니다. SM 당 최소 워프 수는 적어도 워프 스케줄러 수와 같아야합니다. 아키텍처에 6주기 (Maxwell 및 Pascal)의 종속 실행 대기 시간이있는 경우 대기 시간을 처리하려면 SM 당 24 개 (24/64 = 37.5 % 점유) 인 스케줄러 당 최소 6 개의 워프가 필요합니다. 스레드에 명령어 수준의 병렬 처리가 있으면이 값을 줄일 수 있습니다. 거의 모든 커널은 80-1000주기가 걸릴 수있는 메모리로드와 같은 가변 대기 시간 명령을 발행합니다. 대기 시간을 숨기려면 워프 스케줄러 당 더 많은 워프가 필요합니다. 각 커널에 대해 워프 수와 공유 메모리 또는 레지스터와 같은 다른 리소스 사이에 트레이드 오프 포인트가 있으므로 다른 희생이 발생할 가능성이 있으므로 100 % 점유를 최적화하는 것은 권장되지 않습니다. CUDA 프로파일 러는 개발자가 해당 균형을 결정하는 데 도움을주기 위해 명령 발행 률, 점유 및 지연 이유를 식별하는 데 도움이 될 수 있습니다.
스레드 블록의 크기는 성능에 영향을 줄 수 있습니다. 커널에 큰 블록이 있고 동기화 장벽을 사용하는 경우 장벽 지연이 지연 원인이 될 수 있습니다. 이것은 스레드 블록 당 뒤틀림을 줄임으로써 완화 될 수 있습니다.
하나의 장치에 여러 스트리밍 멀티 프로세서가 있습니다.
SM은 여러 블록을 포함 할 수 있습니다. 각 블록에는 여러 스레드가 포함될 수 있습니다.
SM에는 여러 CUDA 코어가 있으며 (개발자로서 워프에 의해 추상화 되었기 때문에 신경 쓰지 않아도됩니다.) 스레드에서 작동합니다. SM은 항상 스레드의 워프 작업을합니다 (항상 32). 워프는 동일한 블록의 스레드에서만 작동합니다.
SM과 블록 모두 스레드 수, 레지스터 및 공유 메모리 수에 제한이 있습니다.
참고 URL : https://stackoverflow.com/questions/3519598/streaming-multiprocessors-blocks-and-threads-cuda
'code' 카테고리의 다른 글
| 명령 줄에서 WAR의 클래스를 어떻게 실행합니까? (0) | 2020.11.18 |
|---|---|
| Java 스레드를 정상적으로 중지하는 방법은 무엇입니까? (0) | 2020.11.18 |
| Redis에서 세트를 비우거나 삭제 하시겠습니까? (0) | 2020.11.18 |
| 데이터 폴더의 URI로 카메라 결과를 얻는 방법은 무엇입니까? (0) | 2020.11.18 |
| XPath :: 다음 형제 가져 오기 (0) | 2020.11.18 |