스레드 블록(CUDA 프로그래밍)

Thread block (CUDA programming)

스레드 블록은 직렬 또는 병렬로 실행될 수 있는 스레드 그룹을 나타내는 프로그래밍 추상화입니다.더 나은 프로세스 및 데이터 매핑을 위해 스레드는 스레드 블록으로 그룹화됩니다.이전에는 아키텍처에서 스레드 블록의 스레드 수가 블록당 총 512개의 스레드로 제한되었지만, 2010년 3월 현재 컴퓨팅 기능 2.x 이상의 경우 블록에 최대 1024개의 스레드가 포함될 수 있습니다.동일한 스레드 블록의 스레드는 동일한 스트림 프로세서에서 실행됩니다.[1]동일한 블록의 스레드는 공유 메모리, 장벽 동기화 또는 원자 작동과 같은 다른 동기화 프리미티브를 통해 서로 통신할 수 있습니다.

여러 개의 블록이 결합되어 격자를 형성합니다.동일한 그리드에 있는 모든 블록은 동일한 수의 스레드를 포함합니다.블록의 스레드 수는 제한적이지만 병렬로 작동하기 위해 많은 스레드 블록이 필요하고 사용 가능한 모든 멀티프로세서를 사용하기 위해 그리드를 사용할 수 있습니다.

CUDA는 병렬성을 활용하기 위해 상위 수준의 언어가 사용할 수 있는 병렬 컴퓨팅 플랫폼 및 프로그래밍 모델입니다.CUDA에서 커널은 스레드의 도움을 받아 실행됩니다.스레드는 커널의 실행을 나타내는 추상적인 엔티티입니다.커널은 특별한 장치에서 실행되도록 컴파일하는 기능입니다.다중 쓰레드 응용프로그램은 병렬 계산을 구성하기 위해 동시에 실행되는 많은 쓰레드를 사용합니다.모든 스레드에는 인덱스가 있으며, 이 인덱스는 메모리 주소 위치를 계산하고 제어 결정을 내리는 데 사용됩니다.

치수

CUDA는 호스트 디바이스 응용 프로그램을 실행하는 데 사용되는 이기종 프로그래밍 모델을 기반으로 작동합니다.OpenCL과 유사한 실행 모델을 가지고 있습니다.이 모델에서는 일반적으로 CPU 코어인 호스트 디바이스에서 애플리케이션을 실행하기 시작합니다.장치는 처리량 지향 장치, 즉 병렬 연산을 수행하는 GPU 코어입니다.커널 함수는 이러한 병렬 실행을 수행하는 데 사용됩니다.이러한 커널 기능이 실행되면 컨트롤은 직렬 실행을 재개하는 호스트 장치로 다시 전달됩니다.

많은 병렬 응용 프로그램에는 다차원 데이터가 포함되므로 스레드 블록을 1D, 2D 또는 3D 스레드 배열로 구성하는 것이 편리합니다.그리드의 블록 간 통신이나 협력이 불가능하므로 그리드의 블록을 독립적으로 실행할 수 있어야 합니다.커널이 실행되면 스레드 블록당 스레드 수가 지정되고 스레드 블록 수가 지정되면 이는 실행된 총 CUDA 스레드 수를 정의합니다.[2]블록의 최대 x, y, z 차원은 1024, 1024 및 64이며 블록당 최대 스레드 수인 x × y × z ≤ 1024가 되도록 할당해야 합니다.[3]블록은 x, y 및 z 차원에서 각각 최대 2-131, 65,535 및 65,535 블록의 1, 2 또는 3차원 그리드로 구성할 수 있습니다.[3]블록당 최대 스레드와 달리 최대 그리드 치수와 구별되는 그리드당 블록이 없습니다.

인덱싱

1D 인덱싱

CUDA의 모든 스레드는 배열의 메모리 위치를 계산하고 액세스할 수 있도록 특정 인덱스와 연결됩니다.

512개 요소의 배열이 있는 예를 생각해 보십시오.조직 구조 중 하나는 512개의 스레드를 가진 단일 블록으로 그리드를 취하는 것입니다.각각 512 요소인 두 개의 배열 A와 B의 요소별 곱셈으로 구성된 512 요소의 배열 C가 있다고 생각합니다.모든 스레드는 인덱스 i를 가지며, 그것은 A와th B의 i 요소의 곱을 수행한 후 결과를th C의 i 요소에 저장합니다. i는 blockIdx(블록이 하나밖에 없으므로 이 경우 0), blockDim(블록이 512개의 요소를 가지므로 이 경우 512개) 및 각 블록마다 0에서 511까지 변화하는 threadIdx를 이용하여 계산됩니다.

CUDA 프로그래밍에서의[4] 스레드 계층구조

스레드 인덱스 i는 다음 공식으로 계산됩니다.

blockIdx.x는 x차원 블록 식별자입니다.

blockDim.x는 블록 차원의 x 차원입니다.

threadIdx.x는 스레드 식별자의 x 차원입니다.

따라서 'i'는 전체 배열을 포함하는 0부터 511까지의 값을 가질 것입니다.

1024보다 큰 배열에 대한 계산을 고려하려면 각각 1024개의 스레드가 있는 여러 개의 블록을 가질 수 있습니다.2048개의 배열 요소가 있는 예를 생각해 보십시오.이 경우 우리는 각각 1024개의 스레드를 가진 2개의 스레드 블록을 가지고 있습니다.따라서 스레드 식별자의 값은 0에서 1023까지, 블록 식별자는 0에서 1까지, 블록 치수는 1024까지 다양합니다.따라서 첫 번째 블록은 0에서 1023 사이의 인덱스 값을 얻고 마지막 블록은 1024에서 2047 사이의 인덱스 값을 갖습니다.

따라서 각 스레드는 먼저 액세스해야 하는 메모리 인덱스를 계산한 다음 계산을 진행합니다.스레드를 사용하여 배열 A와 B의 요소를 병렬로 추가하고 그 결과를 배열 C에 저장하는 예를 생각해 보십시오.스레드의 해당 코드는 다음과 같습니다.[5]

__global__ 공허한 vecAdd커널 (흘러가다 *A , 흘러가다 *B , 흘러가다 * C , 인트 n) {     인트 색인을 보다 = blockIdx.x * 블록Dim.x + threadIdx.x;     한다면 (색인을 보다 < n)     {         C[색인을 보다] = A[색인을 보다] + B[색인을 보다] ;     } }

2D 인덱싱

특히 복잡한 그리드에서도 동일한 방식으로, sreadId 뿐만 아니라 blockId는 그리드의 기하학적 구조에 따라 각 스레드에 의해 계산되어야 합니다.2차원 블록이 있는 2차원 그리드를 생각해 보십시오.threadId와 blockId는 다음 공식으로 계산됩니다.

[6]

하드웨어적 관점

스레드의 계층을 명시했지만 스레드, 스레드 블록 및 그리드는 본질적으로 프로그래머의 관점이라는 점에 유의해야 합니다.스레드 블록의 완전한 요지를 얻기 위해서는 하드웨어 관점에서 그것을 아는 것이 중요합니다.하드웨어는 동일한 명령을 실행하는 스레드를 와프(warp)로 그룹화합니다.여러 개의 와프가 스레드 블록을 구성합니다.여러 개의 스레드 블록이 SM(Streaming Multiprocessor)에 할당되며, 여러 개의 SM이 전체 GPU 유닛(전체 커널 그리드를 실행함)을 구성합니다.[citation needed]

GPU에서 스레드 블록의 하드웨어 관점 대비 프로그래머 관점의 그림 상관관계.[7]

스트리밍 멀티프로세서

GPU의 각 아키텍처(예: 케플러 또는 페르미)는 여러 개의 SM 또는 스트리밍 멀티프로세서로 구성됩니다.이것들은 낮은 클록 레이트 타겟과 작은 캐시를 가진 범용 프로세서들입니다.SM은 여러 스레드 블록을 병렬로 실행할 수 있습니다.스레드 블록 중 하나가 실행을 완료하면 바로 다음 스레드 블록을 사용합니다.일반적으로 SM은 명령어 수준 병렬화를 지원하지만 분기 예측은 지원하지 않습니다.[8]

스트리밍 멀티프로세서 및 그 자원의 예시.[9]

이러한 목적을 달성하기 위해 SM에는 다음과 같은 내용이 포함되어 있습니다.[8]

  • 실행 코어.(단일 정밀 부동소수점 장치, 이중 정밀 부동소수점 장치, 특수 기능 장치(SFU)).
  • 캐시:
  1. L1 캐시. (메모리 액세스 지연 시간을 줄이기 위한)
  2. 공유 메모리. (쓰레드 간 공유 데이터의 경우)
  3. 일정한 캐시(읽기 전용 메모리에서 읽기를 브로드캐스트하는 경우).
  4. 텍스쳐 캐시. (텍스쳐 메모리에서 대역폭을 집계하기 위한 것입니다.
  • 워프 스케줄러.(이는 특정 스케줄링 정책에 따라 와프에게 명령을 발행하기 위한 것입니다.)
  • 레지스터 수가 상당히 많습니다. (SM에서 한 번에 많은 수의 활성 스레드를 실행할 수 있으므로 수천 개의 레지스터를 보유해야 합니다.)

하드웨어는 SM에 스레드 블록을 예약합니다.일반적으로 SM은 여러 스레드 블록을 동시에 처리할 수 있습니다.SM에는 총 8개의 스레드 블록이 포함될 수 있습니다.스레드 ID는 해당 SM에 의해 스레드에 할당됩니다.

SM에서 스레드 블록을 실행할 때마다 스레드 블록 내부의 모든 스레드가 동시에 실행됩니다.따라서 SM 내부의 스레드 블록에 대한 메모리를 확보하기 위해서는 블록 내의 전체 스레드 세트가 실행을 완료하는 것이 중요합니다.각 스레드 블록은 워프(warp)로 알려진 예약 단위로 분할됩니다.이에 대해서는 다음 절에서 자세히 설명합니다.

NVidia의 페르미 마이크로아키텍처에 구현된 이중 [10]워프 스케줄러의 예시

SM의 워프 스케줄러는 명령어 발행 시 워프 중 우선순위를 결정합니다.[11]워프 우선순위 정책의 일부에 대해서도 다음 절에서 논의했습니다.

워프

하드웨어 측면에서 스레드 블록은 '와프'로 구성되어 있습니다.워프는 워프에 있는 모든 스레드가 동일한 명령을 실행하도록 스레드 블록 내에 있는 32개의 스레드 집합입니다.이러한 스레드는 SM에 의해 연속적으로 선택됩니다.[12]

멀티프로세서(SM)에서 스레드 블록이 실행되면 실행이 완료될 때까지 모든 와프가 상주합니다.따라서 새 블록은 새 블록의 모든 와프에 대한 충분한 수의 여유 레지스터가 있을 때까지 SM에서 실행되지 않으며 새 블록에 대한 충분한 여유 공유 메모리가 있을 때까지 실행되지 않습니다.

명령을 실행하는 32개의 나사산으로 이루어진 워프를 생각해보세요.피연산자 중 하나 또는 둘 다 준비되지 않은 경우(예: 글로벌 메모리에서 아직 가져올 수 없음), 제어권을 다른 워프로 이전하는 '컨텍스트 전환' 프로세스가 발생합니다.[13]특정 워프에서 전환할 때 해당 워프의 모든 데이터가 레지스터 파일에 남아 피연산자가 준비되면 빠르게 다시 시작할 수 있습니다.명령어에 현저한 데이터 의존성이 없는 경우, 즉 명령어의 피연산자가 모두 준비되어 있는 경우, 각 워프는 실행 준비가 된 것으로 간주됩니다.둘 이상의 워프가 실행 대상인 경우, 부모 SM은 워프 스케줄링 정책을 사용하여 다음 페치 명령을 얻을 것을 결정합니다.

실행 대상인 와프 스케줄링을 위한 다양한 정책이 아래에 설명되어 있습니다.[14]

  1. 라운드 로빈(RR) - 명령은 라운드 로빈 방식으로 가져옵니다.RR은 SM이 비지 상태로 유지되고 메모리 지연 시간에 클럭 사이클이 낭비되지 않도록 합니다.
  2. LRF(Last Recently Fetched) - 이 정책에서는 명령을 가장 오랫동안 가져오지 않은 워프가 명령을 가져올 때 우선 순위를 가집니다.
  3. 공정(FAIR)[14] - 이 정책에서 스케줄러는 모든 와프에게 해당 와프에 대해 가져온 지침의 수만큼 '공정'한 기회가 주어지도록 합니다.명령어를 최소 개수의 명령어가 불러온 워프로 가져옵니다.
  4. 스레드 블록 기반 CAWS[15](중요도 인식 워프 스케줄링) - 이 스케줄링 정책의 중점은 스레드 블록의 실행 시간을 향상시키는 것입니다.가장 오랜 시간이 걸리는 워프에 더 많은 시간 자원을 할당했습니다.이 정책은 가장 중요한 워프에 우선 순위를 부여함으로써 스레드 블록이 더 빨리 완료되어 리소스를 더 빨리 사용할 수 있습니다.

기존 CPU 스레드 컨텍스트 "스위칭"은 할당된 레지스터 값과 프로그램 카운터를 오프칩 메모리(또는 캐시)에 저장하고 복원해야 하므로 워프 컨텍스트 스위칭보다 훨씬 더 비중이 큰 작업입니다.워프의 모든 레지스터 값(프로그램 카운터 포함)은 레지스터 파일에 남아 있으며, 스레드 블록의 모든 워프 간에 공유되므로 공유 메모리(및 캐시)도 그대로 유지됩니다.

워프 아키텍처를 활용하려면 프로그래밍 언어와 개발자가 메모리 액세스를 통합하는 방법과 제어 흐름 발산을 관리하는 방법을 이해해야 합니다.워프의 각 스레드가 서로 다른 실행 경로를 선택하거나 각 스레드가 현저하게 다른 메모리에 액세스하는 경우 워프 아키텍처의 이점이 상실되고 성능이 크게 저하됩니다.

참고 항목

참고문헌

  1. ^ "Chapter 4. Hardware Implementation, The threads of a thread block execute concurrently on one multiprocessor, and multiple thread blocks can execute concurrently on one multiprocessor".
  2. ^ "CUDA Thread Model". www.olcf.ornl.gov. Archived from the original on 2016-09-23. Retrieved 2016-09-21.
  3. ^ a b "CUDA Toolkit Documentation: Features and Technical Specifications". docs.nvidia.com. Retrieved 2022-05-24.
  4. ^ "Thread Hierarchy in CUDA Programming". Retrieved 2016-09-21.
  5. ^ Kirk, David; Hwu, Wen-mei W (January 28, 2010). Programming Massively Parallel Processors: A Hands-on Approach.
  6. ^ "Thread Indexing Cheatsheet" (PDF). Retrieved 2016-09-21.
  7. ^ "Thread Optimizations (University of Mayland)" (PDF).
  8. ^ a b Wilt, Nicholas (2013). The CUDA Handbook: A Comprehensive Guide to GPU Programming.
  9. ^ "Thread Optimizations (University of Mayland)" (PDF).
  10. ^ "Thread Optimizations (University of Mayland)" (PDF).
  11. ^ "GPU Computing with CUDA Lecture 2 - CUDA Memories" (PDF).
  12. ^ "Using CUDA Warp-Level Primitives". Nvidia. 2018-01-15. Retrieved 2020-04-08. NVIDIA GPUs execute groups of threads known as warps in SIMT (Single Instruction, Multiple Thread) fashion
  13. ^ "Memory Issues in CUDA and Execution Scheduling in CUDA" (PDF).
  14. ^ a b "Effect of Instruction Fetch and Memory Scheduling on GPU Performance" (PDF).
  15. ^ "CAWS: Criticality-Aware Warp Scheduling for GPGPU Workloads" (PDF).