제 2 장: GPU 프로그래밍 모델

그래픽 처리 장치(GPU)는 고정 기능 그래픽 가속기에서 다양한 애플리케이션을 가속화할 수 있는 고도로 병렬화되고 프로그래밍 가능한 컴퓨팅 엔진으로 발전해왔습니다. GPU의 대규모 병렬 처리 기능을 효과적으로 활용할 수 있도록 NVIDIA CUDA, OpenCL, DirectCompute와 같은 다양한 병렬 프로그래밍 모델과 API가 개발되었습니다. 이러한 프로그래밍 모델은 GPU 하드웨어의 저수준 세부 사항을 숨기면서 프로그래머가 애플리케이션의 병렬 처리를 표현할 수 있도록 하는 추상화를 제공합니다.

이 장에서는 GPU용 병렬 프로그래밍 모델의 핵심 개념과 원리를 탐구할 것입니다. 실행 모델, GPU 명령어 집합 아키텍처(ISA), NVIDIA GPU ISA 및 AMD의 Graphics Core Next(GCN) ISA에 초점을 맞출 것입니다. 또한 이러한 개념이 실제 어떻게 적용되는지를 보여주는 예제도 제공할 것입니다.

실행 모델

GPU 프로그래밍 모델의 실행 모델은 커널이라는 개념을 기반으로 합니다. 커널은 GPU에서 병렬로 실행되는 많은 스레드에 의해 실행되는 함수입니다. 커널을 실행할 때 프로그래머는 생성할 스레드의 수와 이들을 그리드, 블록(또는 협력 스레드 어레이 - CTA) 및 개별 스레드의 계층 구조로 구성하는 방법을 지정합니다.

그리드는 전체 문제 공간을 나타내며 하나 이상의 블록으로 구성됩니다.
블록은 서로 협력하고 공유 메모리와 배리어를 통해 동기화할 수 있는 스레드 그룹입니다. 블록 내의 스레드는 동일한 GPU 코어(스트리밍 멀티프로세서 또는 계산 장치라고 함)에서 실행됩니다.
각 스레드는 자신의 블록과 그리드 내에서 고유한 ID를 가지며, 이 ID를 사용하여 메모리 주소를 계산하고 제어 흐름 결정을 내릴 수 있습니다.

이 계층적 구조를 통해 프로그래머는 데이터 병렬성(동일한 작업이 여러 데이터 요소에 적용됨)과 작업 병렬성(다른 작업이 병렬로 실행됨)을 모두 표현할 수 있습니다.여기는 한국어 번역입니다. 코드 부분은 주석만 번역했습니다.

그림 2.1: GPU 실행 모델의 스레드 계층 구조

            Grid
       ________________
      /   /   /   /   /
     /   /   /   /   / 
    /   /   /   /   /
   /   /   /   /   /
  /__/__/__/__/__/
    |   |   |   |
    |   | Block |
    |   |   |   |
  스레드 스레드 ...

그림 2.1: GPU 실행 모델의 스레드 계층 구조.

SIMT 실행

CUDA와 OpenCL과 같은 GPU 프로그래밍 모델은 Single-Instruction, Multiple-Thread (SIMT) 실행 모델을 따릅니다. SIMT 모델에서 스레드는 warp(NVIDIA 용어) 또는 wavefront(AMD 용어)라고 불리는 그룹으로 실행됩니다. 같은 warp 내의 모든 스레드는 동시에 같은 명령어를 실행하지만, 각 스레드는 서로 다른 데이터를 다룹니다.

그러나 모든 처리 요소가 동기화되어 실행되는 전통적인 Single-Instruction, Multiple-Data (SIMD) 모델과 달리, SIMT는 스레드가 독립적인 실행 경로를 가지고 분기 명령어에서 갈라질 수 있게 합니다. warp가 분기 명령어를 만나면, GPU 하드웨어는 warp 내의 각 스레드에 대한 분기 조건을 평가합니다. 모든 스레드가 같은 경로를 타면(수렴), warp는 정상적으로 실행을 계속합니다. 일부 스레드가 다른 경로를 타면(분기), warp는 서로 다른 경로를 따르는 두 개 이상의 subwarp로 분할됩니다. GPU 하드웨어는 분기된 경로의 실행을 직렬화하여, 각 subwarp에서 비활성 스레드를 마스크 처리합니다. 모든 경로가 완료되면, subwarp는 다시 수렴하고 동기화되어 실행을 계속합니다.

그림 2.2는 분기된 제어 흐름에서의 SIMT 실행을 보여줍니다.

         Warp
   ________________
  /   /   /   /   /
 /   /   /   /   / 
/   /   /   /   /
   |   |   |
   | 분기   |
   |   |   |
  / \ / \ / \
 /   X   \   \
/   / \   \   \
   /   \   \
  /     \   \
 /       \   \
/         \   \
           \
            \
             \
   수렴

그림 2.2: 분기된 제어 흐름에서의 SIMT 실행.

이러한 분기 처리 메커니즘을 통해 SIMT는 더 유연한 제어 흐름을 지원할 수 있습니다.여기는 한국어 번역본입니다:

메모리 계층

GPU는 병렬 작업 부하의 높은 대역폭과 낮은 지연 시간 요구 사항을 지원하기 위해 복잡한 메모리 계층 구조를 가지고 있습니다. 메모리 계층은 일반적으로 다음과 같습니다:

글로벌 메모리: 가장 크지만 가장 느린 메모리 공간으로, 커널의 모든 스레드에서 접근할 수 있습니다. 글로벌 메모리는 일반적으로 고대역폭 GDDR 또는 HBM 메모리를 사용하여 구현됩니다.
공유 메모리: 블록 내의 모든 스레드가 공유하는 빠른 온칩 메모리 공간. 공유 메모리는 스레드 간 통신과 블록 내 데이터 공유에 사용됩니다.
상수 메모리: 모든 스레드에 읽기 전용 데이터를 브로드캐스트하는 데 사용되는 읽기 전용 메모리 공간.
텍스처 메모리: 공간적 지역성에 최적화되고 텍스처 캐시를 통해 액세스되는 읽기 전용 메모리 공간. 텍스처 메모리는 주로 그래픽 작업 부하에서 사용됩니다.
로컬 메모리: 각 스레드의 전용 메모리 공간으로, 레지스터 스필링과 대용량 데이터 구조에 사용됩니다. 로컬 메모리는 일반적으로 글로벌 메모리에 매핑됩니다.

GPU에서 높은 성능을 달성하기 위해서는 메모리 계층의 효과적인 활용이 중요합니다. 프로그래머는 공유 메모리의 사용을 극대화하고 글로벌 메모리 액세스를 최소화하여 메모리 지연 시간과 대역폭 병목 현상을 줄여야 합니다.

그림 2.3은 GPU 메모리 계층을 보여줍니다.

      ____________
     |            |
     |   Global   |
     |   Memory   |
      ____________
           |
      ____________
     |            |
     |  Constant  |
     |   Memory   |
      ____________
           |
      ____________
     |            |
     |  Texture   |
     |   Memory   |
      ____________
           |
           |
      ____________
     |            |
     |   Shared   |
     |   Memory   |
      ____________
           |
      ____________ 
     |            |
     |   Local    |
     |   Memory   |
      ____________

그림여기는 한국어 번역입니다:

GPU 명령어 세트 아키텍처

GPU 명령어 세트 아키텍처(ISA)는 소프트웨어와 하드웨어 간의 저수준 인터페이스를 정의합니다. 이는 GPU에서 지원되는 명령어, 레지스터, 메모리 주소 지정 모드를 지정합니다. GPU ISA에 대한 이해는 효율적인 GPU 코드 개발 및 성능 최적화에 필수적입니다.

이 섹션에서는 두 주요 GPU 벤더, NVIDIA와 AMD의 ISA를 탐색할 것입니다. NVIDIA의 Parallel Thread Execution(PTX)과 SASS ISA, 그리고 AMD의 Graphics Core Next(GCN) ISA에 중점을 둘 것입니다.

NVIDIA GPU ISA

NVIDIA GPU는 두 수준의 ISA를 지원합니다: PTX(Parallel Thread Execution)와 SASS(Streaming ASSembler). PTX는 CUDA 컴파일러의 안정적인 대상을 제공하는 가상 ISA이며, SASS는 NVIDIA GPU의 네이티브 ISA입니다.

PTX (Parallel Thread Execution)

PTX는 NVIDIA GPU용 저수준 가상 ISA입니다. LLVM IR 또는 Java 바이트코드와 유사하게, 컴파일러의 안정적이고 아키텍처 독립적인 대상을 제공합니다. CUDA 프로그램은 일반적으로 PTX 코드로 컴파일되며, 이는 다시 NVIDIA GPU 드라이버에 의해 네이티브 SASS 명령어로 변환됩니다.

PTX는 산술, 메모리, 제어 흐름 명령어의 광범위한 범위를 지원합니다. 무제한 개수의 가상 레지스터와 술어화를 지원하여 제어 흐름의 효율적인 구현을 가능하게 합니다. PTX는 또한 스레드 동기화, 원자 연산, 텍스처 샘플링을 위한 특수 명령어를 제공합니다.

다음은 간단한 벡터 덧셈 커널의 PTX 코드 예제입니다:

.version 7.0
.target sm_70
.address_size 64

.visible .entry vecAdd(
    .param .u64 vecAdd_param_0,
    .param .u64 vecAdd_param_1,
    .param .u64 vecAdd_param_2,
    .param .u32 vecAdd_param_3
)
{
    .reg .b32 %r<4>;
    .reg .b64 %rd<6>;

    ld.param.u64 %rd1, [vecAdd_param_0];
    ld.param.u64 %rd2, [vecAdd_param_1];
    ld.param.u64 %rd3, [vecAdd_param_2];
    ld.param.u32 %r1, [vecAdd_param_3];
    cvta.to.global.u64 %rd4, %rd1;
    cvta
}
```이 마크다운 파일의 한국어 번역은 다음과 같습니다. 코드에 대해서는 주석만 번역했습니다.

```to.global.u64 %rd5, %rd2;
    mov.u32 %r2, %tid.x;
    mul.wide.u32 %rd6, %r2, 4;
    add.s64 %rd7, %rd4, %rd6;
    add.s64 %rd8, %rd5, %rd6;

    // 입력 벡터 요소를 전역 메모리에서 로드합니다.
    ld.global.f32 %f1, [%rd7];
    ld.global.f32 %f2, [%rd8];
    add.f32 %f3, %f1, %f2;

    cvta.to.global.u64 %rd9, %rd3;
    add.s64 %rd10, %rd9, %rd6;
    // 결과 벡터 요소를 전역 메모리에 저장합니다.
    st.global.f32 [%rd10], %f3;

    ret;
}

이 PTX 코드는 vecAdd 커널 함수를 정의합니다. 이 함수는 입력 및 출력 벡터에 대한 포인터와 벡터의 크기를 매개변수로 받습니다. 이 커널은 글로벌 스레드 ID를 계산하고, 해당 요소를 입력 벡터에서 읽어와 더한 다음, 결과를 출력 벡터에 기록합니다.

SASS (Streaming ASSembler)

SASS는 NVIDIA GPU의 네이티브 ISA (Instruction Set Architecture)입니다. SASS는 저수준 기계 특정 ISA로, GPU 하드웨어에 직접 매핑됩니다. SASS 명령어는 NVIDIA GPU 드라이버에 의해 PTX 코드에서 생성되며, 일반적으로 프로그래머에게 노출되지 않습니다.

SASS 명령어는 메모리 대역폭과 명령어 캐시 크기를 줄이기 위해 compact한 형식으로 인코딩됩니다. SASS는 레지스터, 즉시 값, 다양한 메모리 액세스 모드와 같은 폭넓은 연산자 유형을 지원합니다.

다음은 벡터 덧셈 커널의 SASS 코드 예입니다:

code_version_number 90
                     @P0 LDG.E R2, [R8];
                     @P0 LDG.E R0, [R10];
                     @P0 FADD R0, R2, R0;
                     @P0 STG.E [R12], R0;
                         EXIT;

이 SASS 코드는 앞서 보여준 PTX 코드에 해당합니다. 입력 벡터 요소를 전역 메모리에서 로드하고(LDG.E), 덧셈 연산을 수행하며(FADD), 결과를 다시 전역 메모리에 저장합니다(STG.E). 마지막으로 커널을 종료합니다.

AMD Graphics Core Next ISA

AMD GPU는 Graphics Core Next (GCN) 아키텍처와 ISA를 사용합니다. GCN은 그래픽과 계산 작업을 모두 지원하는 RISC 기반 ISA입니다. GCN은 고성능, 확장성, 전력 효율성을 목표로 설계되었습니다.

GCN은 다음과 같은 주요 기능을 도입했습니다:

스칼라 명령어 세트
벡터 명령어
고급 메모리 액세스 모드
동적 스케줄링
멀티 스레딩 지원

이러한 기능들은 AMD GPU의 성능과 효율성을 향상시키는 데 기여합니다.여기는 효율적인 스칼라 연산 및 흐름 제어를 위한 ALU에 대한 정보입니다.

데이터 병렬 연산의 병렬 실행을 위한 벡터 ALU.
원자 연산과 공유 메모리에 대한 저지연 액세스를 지원하는 고대역폭 메모리 시스템.
기본+오프셋 및 스칼라+벡터 주소 지정을 지원하는 유연한 메모리 연산 주소 모드.

다음은 벡터 추가 커널에 대한 GCN ISA 코드의 예입니다:

.text
.globl vecAdd
.p2align 2
 
.type vecAdd,@function
vecAdd:
    .set DPTR, 0
 
    # 커널 인수를 메모리에서 로드합니다.
    s_load_dwordx4 s[0:3], s[4:5], 0x0
    s_load_dword s4, s[4:5], 0x10
    s_waitcnt lgkmcnt(0)
 
    # 입력 벡터 요소를 메모리에서 로드합니다.
    v_lshlrev_b32 v0, 2, v0
    v_add_u32 v1, vcc, s1, v0
    v_mov_b32 v3, s3
    v_addc_u32 v2, vcc, s2, v3, vcc
    flat_load_dword v1, v[1:2]
 
    v_add_u32 v3, vcc, s0, v0
    v_mov_b32 v5, s3
    v_addc_u32 v4, vcc, s2, v5, vcc
    flat_load_dword v0, v[3:4]
 
    # 벡터 요소를 더하고 결과를 메모리에 저장합니다.
    v_add_f32 v0, v0, v1
    flat_store_dword v[3:4], v0
    s_endpgm

이 GCN 코드는 flat_load_dword를 사용하여 입력 벡터 요소를 로드하고, v_add_f32를 사용하여 더한 다음, flat_store_dword를 사용하여 결과를 메모리에 저장합니다. s_load_dwordx4 및 s_load_dword 명령어는 커널 인수를 메모리에서 로드하는 데 사용됩니다.

GPU 아키텍처에 알고리즘 매핑하기

GPU 아키텍처에 효율적으로 알고리즘을 매핑하는 것은 높은 성능을 달성하는 데 crucial합니다. 주요 고려 사항은 다음과 같습니다:

충분한 병렬성 노출

알고리즘은 GPU의 병렬 처리 기능을 완전히 활용할 수 있도록 많은 세분화된 스레드로 분해되어야 합니다. 이는 종종 서로 독립적으로 실행될 수 있는 데이터 병렬 부분을 식별하는 것을 포함합니다.

브랜치 분기 최소화

워프/웨이브프론트 내의 분기형 제어 흐름은 직렬화와 SIMD 효율성 저하를 초래할 수 있습니다. 알고리즘은 가능한 한 브랜치 분기를 최소화하도록 구조화되어야 합니다. 이는 데이터 종속적인 분기 사용을 줄임으로써 달성할 수 있습니다.메모리 계층 구조 활용하기

전역 메모리에 접근하는 것은 비용이 많이 듭니다. 알고리즘은 공유 메모리와 레지스터의 사용을 최대화하여 전역 메모리 접근을 줄여야 합니다. 또한 데이터는 워프 내의 스레드가 연속적인 메모리 위치에 접근할 수 있도록 메모리에 배치되어야 합니다. 메모리 계층 구조를 효과적으로 사용하면 메모리 지연 시간과 대역폭 병목 현상을 크게 줄일 수 있습니다.

연산과 메모리 접근의 균형 잡기

알고리즘은 메모리 지연 시간을 숨기고 높은 계산 처리량을 달성하기 위해 산술 연산과 메모리 연산의 비율이 높아야 합니다. 이는 데이터 재사용 최대화, 데이터 프리페칭, 연산과 메모리 접근의 중첩을 통해 달성할 수 있습니다.

호스트-디바이스 데이터 전송 최소화하기

호스트(CPU)와 디바이스(GPU) 메모리 간 데이터 전송은 느립니다. 알고리즘은 GPU에서 최대한 많은 계산을 수행하여 이러한 전송을 최소화해야 합니다. 데이터는 대량으로 GPU로 전송되어야 하며, 전송 오버헤드를 상쇄하기 위해 가능한 한 오래 GPU에 유지되어야 합니다.

GPU 커널 개발 시 일반적으로 다음과 같은 병렬 알고리즘 설계 패턴이 사용됩니다:

Map: 각 스레드가 다른 데이터 요소에 대해 동일한 작업을 수행하여 대용량 데이터셋의 간단한 병렬 처리를 가능하게 합니다.
Reduce: 병렬 감소를 사용하여 대용량 입력 데이터에서 단일 값(예: 합계, 최대값)을 효율적으로 계산합니다. 스레드가 지역 감소를 수행하고, 이를 결합하여 최종 결과를 생성합니다.
Scan: 또한 접두사 합(prefix sum)으로 알려진 스캔은 배열 요소의 누적 합을 계산하는 데 사용됩니다. 효율적인 병렬 스캔 알고리즘은 많은 GPU 가속 애플리케이션의 핵심 구성 요소입니다.
Stencil: 각 스레드는 인접한 데이터 요소를 기반으로 값을 계산합니다. 스텐실 연산은 과학 시뮬레이션과 이미지 처리 애플리케이션에서 일반적입니다.여기는 한국어 번역입니다:
모으기/흩뿌리기: 스레드는 글로벌 메모리의 임의 위치에서 데이터를 읽어오거나(모으기) 쓰거나(흩뿌리기)합니다. 효율성을 위해서는 데이터 레이아웃과 접근 패턴을 주의깊게 구성해야 합니다.

그림 3.20은 맵 패턴의 예를 보여줍니다. 각 스레드가 입력 배열의 다른 요소에 함수(예: 제곱근)를 적용합니다.

입력 배열:  
               |  |   |   |   |   |   |   |
               v  v   v   v   v   v   v   v
              ______________________________
스레드:     |    |    |    |    |    |    |    |
             |____|____|____|____|____|____|____|
                |    |    |    |    |    |    |
                v    v    v    v    v    v    v
출력 배열:

그림 3.20: GPU 프로그래밍에서의 맵 패턴의 예.

결론

CUDA와 OpenCL과 같은 GPU 프로그래밍 모델은 현대 GPU의 병렬 처리 기능을 개발자에게 제공하여 다양한 애플리케이션을 가속화할 수 있게 합니다. 이러한 프로그래밍 모델은 세부적인 병렬 작업을 GPU 하드웨어에 효율적으로 매핑할 수 있는 추상화를 제공합니다.

GPU 코드의 고성능을 작성하기 위해서는 실행 모델, 메모리 계층 구조, 동기화 기본 요소에 대한 이해가 필수적입니다. 개발자는 스레드 구성, 분기 발산, 메모리 접근 패턴, 알고리즘 설계 등의 요인을 신중히 고려해야 합니다.

GPU 아키텍처가 계속 발전함에 따라 프로그래밍 모델과 도구도 발전해야 합니다. 새로운 하드웨어 기능과 기능을 효과적으로 활용할 수 있도록 하기 위해서는 프로그래밍 언어 설계, 컴파일러 최적화, 자동 조정 등의 연구가 중요할 것입니다.

Chapter 1 Introduction to Gpu Chip Design Chapter 3 Parallel Programming Models