컴퓨터 구조론 8장 [고성능 컴퓨터시스템 구조 / 그래픽처리유니트(GPU)]

GPU는 실시간 그래픽 처리용으로 개발되었으나, 계산 능력이 강화되어 일반적인 산술 연산 처리에도 쉽게 이용할 수 있도록 개선되면서 다양한 데이터 병렬 응용(data parallel application)들을 위한 가장 이상적인 보조 프로세서로서의 입지를 굳혀가고 있다.

초기 GPU는 몇 가지 치명적인 결점들을 가지고 있었다. 프로그래밍을 위해서는 GPU 내부 구조와 그래픽 API(application program interface)에 대한 지식이 필요하다는 것과 그래픽 프로그래밍 자체가 너무 복잡하다는 것이다.

기억장치에 대한 직접적인 읽기 및 쓰기와 같은 기본적인 연산은 과학 기술 응용들에 필수적인 부동소수점 연산이 지원되지 않았고, 그러한 문제점에도 불구하고 GPU가 널리 보급되기 시작한 것은 NVIDIA사가 두가지 주요 기술을 도입한 것이 계기가 되었다.

• 통합된 그래픽 및 계산 구조(unified graphic and compute architecture) 
  • 그래픽 응용뿐 아니라 일반적인 계산에도 적합한 통합적인 하드웨어 구조
• CUDA(Compute Unified Device Architecture) 프로그래밍 모델
  • 일반적인 프로그래밍 언어들을 이용하여 GPU를 위한 프로그램을 쉽게 작성할 수 있게 해주는 프로그래머-친화적 병렬 컴퓨팅 플랫폼으로서, NVIDIA사에 의해 개발된 프로그래밍 모델

위의 두 가지 기술들이 도입됨으로써, GPU 사용자는 C 프로그램에 CUDA 확장자들만 추가하면 대규모 병렬 프로세서를 이용할 수 있게 되었다.

• CUDA를 기반으로 한 새로운 GPU 프로그래밍 기술인 'GPU 컴퓨팅'으로 인하여 CUDA를 기반으로 한 새로운 GPU 프로그래밍 기술인 GPU 컴퓨팅으로 인항여 GPU의 응용 분야가 크게 확대된 것이다.
• 특히 그래픽 처리뿐 아니라 일반적인 과학계산 처리에도 적합한 구조를 가지게 되면서 GPU를 GPGPU(general-purpose GPU)라고 부르기도 한다.
• CUDA 프로그램을 처리하는 컴퓨팅 시스템은 일반적인 CPU 혹은 컴퓨터(PC, SMP 등)를 지칭하는 호스트와 GPU 같은 병렬프로세서를 지칭하는 디바이스로 구성된다.

전형적인 GPU의 간략화된 내부 구성도

호스트와 GPU 간의 연결 및 간략화된 내부 구조를 보여주고 있다.

• 호스트로부터 병렬 프로그램 코드들이 GPU로 전송되어 오면, 그들은 스케줄러와 디스패치 유니트에 의해 스트리밍 다중 프로세서(SM)들로 적절히 할당되어 처리된다.
• GPU 내부에는 여러 개의 SM들이 존재하며, 각 SM은 실제 프로그램 코드를 수행하는 다수의 스트리밍 프로세서(SP)들로 구성된다.
• 근본적으로 GPU는 산술 보조프로세서이다. 그런데 일반 CPU 칩에 포함되어 있는 부동소수점 보조프로세서들과의 차이점은 많은 수의 동일 연산들을 서로 다른 데이터들에 대하여 SPMD(single-program multiple-data)형으로 병렬처리 한다는 점이다.
• 즉, 다수의 SP 들이 동일한 프로그램 코드의 복사본을 받아 서로 다른 데이터에 대하여 연산을 처리하는 과정에서 호스트는 순차적인 부분을 처리하고, GPU는 계산량이 많고 병렬 처리가 가능한 부분을 담당하게 된다. 이러한 관계를 CUDA 프로그램 내부 구조를 이용하여 살펴볼 수 있다.

CUDA 프로그래밍 모델

• CUDA 프로그램은 호스트가 실행하는 호스트 코드(host Code)와 디바이스가 실행하는 디바이스 코드(device code)로 이루어진 통합 소스 코드이며, 컴파일 과정에서 분리된다. 호스트 코드는 순차적인 프로그램인 ANSI C 기반으로 작성되며, 디바이스 코드는 ANSI C 코드와 커널 함수로 이루어진다.

커널 함수(kernel function)

• 병렬함수 및 데이터 구조를 명시하는 키워드들로 작성되는 코드로서 디바이스에 의해 실행된다.
• 이 커널 함수의 한 인스턴스를 스레드라고 부르며, GPU 내부의 기본 구성요소인 GPU 프로세서 코어(CUDA코어)에 의해 실행된다.

• 큰 사각형은 Grid, 그 내부의 작은 사각형들은 블록이라고 부르며 각 블록에는 여러 개의 스레드들이 포함되어 있다.
• 하나의 병렬 커널에 의해 생성되는 스레드 전체를 그리드라고 하며, 스레드들의 수가 많을 때는 적절한 개수의 스레드들을 묶은 블록 단위로 분할되는 것이다.
• 블록은 하나의 스트리밍 프로세서(SM)로 할당된다.

정리

• 호스트 코드 : ANSI C 기반의 순차적 프로그램
• 디바이스 코드 : ANSI C 코드와 커널 함수(kernel function)로 구성
• 커널 함수 : 병렬 함수 및 데이터 구조를 명시하는 키워드들을 확장한 형태로 작성된 코드로서, 디바이스가 실행하는 부분 (병렬 처리가 가능한 부분)
• 스레드 (thread) : 커널 함수의 한 인스턴스 (instance)로서, GPU의 SP(CUDA 코어)에 의해 실행
• 그리드 (grid) : 하나의 병렬 커널에 의해 실행되는 스레드 전체
• 블록 (block) : 그리드 내의 스레드들을 적절한 수의 스레드들로 분할한 단위 -> 하나의 스트리밍 프로세서 (SM)에게 할당

그리드, 블록 및 스레드 간의 관계

• 응용 프로그램의 커널 함수에 의해 생성되는 전체 스레드들을 포함하는 그리드는 여러 개의 블록들을 포함하며, 그들은 2차원으로 배열될 수 있다. 각 블록에 포함된 스레드들은 그 수에 따라 2차원 혹은 3차원 배열로 구성되며, 별도의 스레드ID(thread identifier)가 주어진다. 
• 스레드 ID는 커널 함수 내에서 내장 변수인 threadIdx로 액세스되며, 각 스레드가 CUDA 코어에 의해 실행될 때 자신에 할당된 데이터들의 기억장치 주소를 계산하는데도 활용된다.
• 이 모든 스레드들은 동일한 프로그램 코드로 이루어지지만, 서로 다른 데이터를 처리하기 때문에 CUDA 프로그램 실행 모델을 SPMD로 분류하는 것이다.

CUDA 프로그램이 실행되는 과정

1. 호스트가 순차적 코드를 실행한다.
2. 커널 함수가 호출되면, 병렬 커널 코드가 디바이스로 보내진다.
3. 디바이스에서 커널이 그리드(다수의 블록 및 스레드들 포함)를 생성하여 SM들에게 할당하며, SM 내의 SP들이 한 스레드씩을 담당하여 실행한다.
4. 모든 스레드들의 실행이 완료되면 그리드가 종료되고, 결과값들이 호스트로 전송된다.
5. 호스트는 다음 순차적 코드를 실행하며PU는 실시간 그래픽 처리용으로 개발되었으나, 계산 능력이 강화되어 일반적인 산술 연산 처리에도 쉽게 이용할 수 있도록 개선되면서 다양한 데이터 병렬 응용(data parallel application)들을 위한 가장 이상적인 보조 프로세서로서의 입지를 굳혀가고 있다.

CUDA 프로그램을 실행하기 위한 디바이스 측의 기억장치 모델을 보면 각 스레드는 별도의 기억장치(local memory 레지스터 세트로 구현)를 가지고 있는데, 입력되는 데이터와 결과 값을 일시적으로 저장할 수 있을 정도의 적은 용량이면 된다.

• 같은 블록에 포함된 스레드들은 캐시와 공유 기억장치(shared memory)를 공동으로 사용한다. 
• 기억장치의 마지막 계층으로는 호스트가 처리할 모든 응용 프로그램들이 생성한 그리드들이 함께 사용할 수 있는 전역 기억장치(global memory)가 있다.
• 마지막으로, CUDA 프로그램 모델에서 호스트와 디바이스 간의 정보 전송 방식을 보면, 병렬 커널을 실행하기 위해서는 호스트가 디바이스로 프로그램 코드와 처리할 데이터들을 보내주어야 하며, 결과 데이터들을 전송받아야 한다.
• 이러한 정보의 전송은 전역 기억장치를 통해 이루어진다.

CUDA 실행을 위한 기억장치 구조

• 호스트가 디바이스로 보내주는 정보를 전역 기억장치의 지정된 위치에 저장하고, 디바이스는 각 블록이 필요한 정보를 공유 기억장치의 지정된 위치에 저장하고, 디바이스는 각 블록이 필요한 정보를 공유 기억장치로 이동하거나 직접 액세스 한다.
• SP들이 최종적으로 처리한 결과들을 다시 전역 기억장치에 저장하면, 호스트가 그들을 인출해 가는 것이다.
• CUDA 코어에 할당된 각 스레드는 지역 기억장치로 사용될 수 있는 별도의 레지스터 세트를 가지며, 다른 스레드의 레지스터 세트는 액세스할 수 없다.
• 같은 SM에 할당된 스레드들이 데이터 공유를 통하여 서로 협ㄹ력할 수 있는 유일한 방법은 공유 기억장치를 이용하는 것이다.
• 레지스터 세트와 공유 기억장치에 저장된 데이터들은 스레드 블록의 처리가 종료되면 더 이상 사용할 수 없게 된다.
• 또한 같은 응용 프로그램으로부터 생성된 스레드들도 일단 서로 다른 블록에 소속되면 공유 기억장치를 통해서도 협력할 수 없으며, 필요한 경우에는 액세스 속도가 느린 전역 기억장치를 이용할 수밖에 없다.

호스트와 디바이스 간의 정보 전송을 지원하기 위한 몇 가지 함수

cudaMalloc((void**)&Md, size);

• 호스트가 정보를 전송하기 전에 미리 기억장치 영역을 할당받기 위한 것으로서, Md는 디바이스 측의 전역 기억장치를 가리킨다. 이 함수가 수행되면 할당된 영역에 대한 포인터의 주소와 그 영역의 크기가 리턴된다.

cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice);

cudaMemcpy(Md, M, size, cudaMemcpyDeviceToHost);

• 두 번째 함수와 세 번째 함수는 각각 호스트로부터 디바이스로, 그리고 디바이스로부터 호스트로 정보를 보내기 위한 것이다.

cudaFree(Md);

• 마지막 함수는 더 이상 사용되지 않는 기억장치 영역을 해제하기 위함

GPU의 내부 구조

• Fermi 구조는 16개의 SM들로 이루어져 있으며, 각 SM은 32개의 CUDA 코어들을 포함하고 있다.
• Fermi GPU는 전체적으로 (16개 SM) X (32개 CUDA 코어) = 512개의 CUDA 코어들을 가지고 있다.
• SM들은 공유 L2 캐시를 중심으로 상하에 8개씩 위치한다.

• Giga Thread 전역 스케줄러는 호스트로부터 입력되는 스레드 블록들을 SM 내부 스케줄링을 담당하는 왑 스케줄러(warp scheduler)들에게 분배해주는 역할을 수행한다.

Fermi GPU의 주요 구성 요소

• GPU 프로세서 코어들(전체적으로 32개의 CUDA 코어들)
• 2개씩의 왑 스케줄러 및 디스패치 포트
• 16개의 적재/저장 유니트
• 4개의 SFU(Special Function Unit)
• 32K개의 32-비트 레지스터들
• 공유 기억장치 및 L1 캐시(전체 64KB)

단일 SM의 내부 구조

• SM에 포함된 32개의 코어들은 디스패치 포트를 통해 들어오는 오퍼랜드의 임시 저장을 위한 레지스터인 오퍼랜드 수집기, 정수 산술논리유니트(ALU)와 부동소수점 유니트(FPU), 그리고 출력 큐 레지스터로 구성된다.
• ALU와 FPU는 독립적인 파이프라인 구조이며, 별도의 데이터 통로를 가지고 있지만, 동시에 처리할 수는 없다.

왑(Warp)

• 동시에 처리될 수 있는 32개의 스레드 묶음

SM이 한번에 수행할 수 있는 32개의 스레드 묶음을 왑이라고 하는데, 각 스레드에게는 순차적인 ID가 부여된다.

• 왑 스케줄러는 스레드에게 별도의 명령어 주소 카운터와 레지스터를 지정하여 CUDA 코어로 할당함으로써, 각 스레드의 독립적인 분기와실행이 가능하게 해준다. GPU에서는 가능한 많은 스레드들을 처리해여 코어들의 이용률을 극대화시킬 수 있기 때문에, Fermi에서는 SM당 두 개의 왑 스케줄러와 디스패치 유니트가 한 번에 두 개씩의 왑들을 코어에 할당한다.
• 각 SM에는 16개의 적재/저장 유니트들이 있어서 매 클록마다 각 스레드를 위한 16개의 오퍼랜드 주소들을 계싼할 수 있다.
• 그 주소들은 모두 캐시나 DRAM에 대한 액세스를 위한 것이다.
• 네 개의 SFU들은 sine, cosine, square root와 같은 초월 함수 계산을 32-비트 혹은 64-비트 부동소수점 수에 대하여 처리한다.
•  Fermi 구조에서 각 SM은 32K(32,768)개의 32-비트 레지스터들을 가지고 있으며, 각 스레드는 최대 64개씩의 레지스터들을 할당 받을 수 있다.
• 레지스터들은 나노초(ns) 단위의 빠른 액세스 시간을 가지며, 캐시 및 공유기억 장치와 함께 기억장치 계층 구조를 이루고 있따. 1차(L1) 캐시는 명령어 캐시와 데이터 캐시로 분할되어 있는데, L1 데이터 캐시는 필요에 따라 공유 기억장치를 적절히 분할하여 사용할 수 있다.
• L2 캐시는 SM들의 중간에 위치하며 데이터와 명령어를 모두 저장하는 통합 캐시로서 모든 SM들에 의해 공유된다. 

---

Reference

컴퓨터 구조론 개정 5판

https://junstar92.tistory.com/246

CUDA Programming Model

https://butter-shower.tistory.com/41

그래픽 처리 유니트 (GPU) - GPU의 개념과 CUDA 프로그래밍

https://www.researchgate.net/figure/NVIDIA-Fermi-GPU-architecture_fig2_272199765

https://www.researchgate.net/figure/Nvidia-Fermi-GPU-architecture-17_fig1_350054481