7장: GPU 설계에서의 스트리밍 멀티프로세서 설계

스트리밍 멀티프로세서(SM)는 NVIDIA GPU 아키텍처의 기본 구성요소입니다. 각 SM에는 SIMT(Single Instruction, Multiple Thread) 방식으로 명령어를 실행하는 CUDA 코어 집합이 포함되어 있습니다. SM은 워프 관리 및 스케줄링, 분기 발산 처리, 공유 메모리 및 캐시에 대한 빠른 접근을 담당합니다. 이 장에서는 SM의 파이프라인, 워프 스케줄링 메커니즘, 레지스터 파일 설계, 공유 메모리 및 L1 캐시 조직 등 SM의 마이크로아키텍처를 살펴볼 것입니다.

SM 마이크로아키텍처와 파이프라인

SM은 동시에 수백 개의 스레드를 효율적으로 실행하도록 설계된 고도로 병렬화되고 파이프라인화된 프로세서입니다. 그림 7.1은 NVIDIA Volta 아키텍처의 SM에 대한 간소화된 블록 다이어그램을 보여줍니다.

                                 명령어 캐시
                                         |
                                         v
                                    워프 스케줄러
                                         |
                                         v
                               디스패치 유닛(4개 워프)
                                 |   |   |   |
                                 v   v   v   v
                               CUDA 코어(FP64/FP32/INT)
                               CUDA 코어(FP64/FP32/INT)
                               CUDA 코어(FP64/FP32/INT)
                               ...
                               텐서 코어
                               텐서 코어
                               ...
                               로드/스토어 유닛
                               로드/스토어 유닛
                               ...
                               특수 기능 유닛
                                         ^
                                         |
                                레지스터 파일(64 KB)
                                         ^공유 메모리 / L1 캐시 (96 KB)

그림 7.1: NVIDIA Volta 아키텍처의 SM (Streaming Multiprocessor)의 단순화된 블록 다이어그램.

SM의 주요 구성 요소는 다음과 같습니다:

명령어 캐시: 지연 시간을 줄이고 처리량을 높이기 위해 자주 액세스되는 명령어를 저장합니다.
워프 스케줄러: 실행을 준비한 워프를 선택하고 사용 가능한 실행 장치에 배포합니다.
디스패치 유닛: 초당 최대 4개의 워프에 대한 명령어를 가져오고 해석하며 적절한 실행 장치에 배포합니다.
CUDA 코어: 다양한 정수 및 부동 소수점 연산을 지원하는 프로그래밍 가능한 실행 장치. Volta에는 각 SM당 64개의 CUDA 코어가 있습니다.
Tensor 코어: 딥 러닝 및 AI 워크로드 가속을 위해 설계된 전용 실행 장치. Volta에는 각 SM당 8개의 Tensor 코어가 있습니다.
로드/저장 유닛: 전역 메모리, 공유 메모리 및 캐시에 대한 메모리 작업(로드 및 저장)을 처리합니다.
특수 기능 유닛: 초월 함수 및 기타 복잡한 수학 연산을 실행합니다.
레지스터 파일: 스레드 전용 레지스터에 빠르게 액세스할 수 있습니다. Volta에는 각 SM당 64KB 레지스터 파일이 있습니다.
공유 메모리 / L1 캐시: 소프트웨어 관리 캐시(공유 메모리) 또는 하드웨어 관리 L1 데이터 캐시로 사용할 수 있는 구성 가능한 메모리 공간입니다.

SM 파이프라인은 여러 워프를 동시에 실행하고 메모리 지연을 숨길 수 있도록 설계되어 처리량을 최대화합니다. 그림 7.2는 SM 파이프라인의 단순화된 모습을 보여줍니다.

    명령어 가져오기
            |
            v
    명령어 디코딩
            |
            v
    피연산자 수집
            |
            v
    실행(CUDA 코어, Tensor 코어, 로드/저장 유닛, 특수 기능 유닛)
            |
            v
    결과 기록

그림 7.2: 단순화된 SM 파이프라인.

파이프라인 단계는 다음과 같습니다:

명령어 가져오기: 워프 스케줄러가 실행할 준비가 된 워프를 선택합니다.Here is the Korean translation of the provided markdown file, with the code comments translated, while the code itself remains unchanged:
Instruction Fetch: 워프 단위로 다음 명령어를 명령어 캐시에서 가져옵니다.
Instruction Decode: 가져온 명령어를 디코딩하여 연산 유형, 피연산자, 대상 레지스터를 결정합니다.
Operand Collection: 명령어 실행에 필요한 피연산자를 레지스터 파일 또는 공유 메모리에서 수집합니다.
Execution: 명령어를 적절한 실행 장치(CUDA 코어, 텐서 코어, 로드/저장 장치, 특수 기능 장치)에서 실행합니다.
Writeback: 실행 결과를 레지스터 파일 또는 공유 메모리에 기록합니다.

높은 성능을 달성하기 위해 SM은 리소스 활용도를 최대화하고 지연 시간을 숨기기 위한 여러 가지 기술을 사용합니다:

이중 발행: SM은 한 사이클에 두 개의 독립적인 명령어를 워프에 발행할 수 있어, 명령어 수준 병렬성이 향상됩니다.
파이프라인화된 실행 장치: 실행 장치가 파이프라인화되어 있어, SM이 이전 작업이 완료되기 전에 새로운 작업을 시작할 수 있습니다.
지연 숨김: SM은 사이클 단위로 워프 간에 전환할 수 있어, 메모리 액세스와 지연 시간이 긴 작업의 지연 시간을 다른 워프의 명령어 실행으로 숨길 수 있습니다.

예제 7.1은 두 벡터의 요소 단위 덧셈을 수행하는 단순한 CUDA 커널을 보여줍니다.

__global__ void vectorAdd(int *a, int *b, int *c, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        c[tid] = a[tid] + b[tid];
    }
}

예제 7.1: 벡터 덧셈을 수행하는 CUDA 커널.

이 예에서, 커널의 각 스레드는 입력 벡터 a와 b의 해당 요소를 더하여 출력 벡터 c에 결과를 저장합니다. SM은 각 스레드를 CUDA 코어에 할당하고 사용 가능한 코어에 워프 단위로 스레드를 스케줄링하여 이 커널을 실행합니다. 로드/저장 장치는 전역 메모리에서 입력 데이터를 가져오고 결과를 다시 기록하는 데 사용됩니다.

워프 스케줄링 및 분기 처리여기는 효율적인 워프 스케줄링이 SM(Streaming Multiprocessor) 성능 극대화에 매우 중요하다는 내용입니다. 워프 스케줄러는 실행 준비가 된 워프를 선택하고 사용 가능한 실행 유닛으로 전송하는 역할을 합니다. 워프 스케줄러의 주된 목표는 항상 실행 준비가 된 워프가 존재하도록 하여 실행 유닛을 계속 사용할 수 있도록 하는 것입니다.

SM은 두 단계의 워프 스케줄링 메커니즘을 사용합니다:

워프 스케줄링: 워프 스케줄러는 라운드-로빈 또는 가장 오래된 순서 등의 스케줄링 정책에 따라 실행 준비가 된 워프를 선택합니다. 그런 다음 선택된 워프를 사용 가능한 실행 유닛으로 전송합니다.
명령 스케줄링: 각 워프 내에서, SM은 명령어 간 의존성과 실행 유닛의 가용성에 따라 명령어를 스케줄링합니다. SM은 단일 사이클에서 동일한 워프로부터 여러 개의 독립적인 명령어를 발급하여 명령어 수준 병렬성을 극대화할 수 있습니다.

그림 7.3은 두 단계의 워프 스케줄링 메커니즘을 보여줍니다.

    워프 풀
    워프 1 (실행 준비)
    워프 2 (대기 중)
    워프 3 (실행 준비)
    ...
    워프 N (실행 준비)
        |
        v
    워프 스케줄러
        |
        v
    전송 유닛
        |
        v
    실행 유닛

그림 7.3: 두 단계의 워프 스케줄링 메커니즘.

워프 스케줄링의 핵심 과제 중 하나는 분기 발산(branch divergence)을 처리하는 것입니다. SIMT 실행 모델에서, 동일한 워프의 모든 스레드는 동기화되어 같은 명령어를 실행합니다. 하지만 워프가 분기 명령어(예: if-else 문)를 만나면, 일부 스레드는 if 경로를 다른 스레드는 else 경로를 실행하게 됩니다. 이러한 상황을 분기 발산이라고 합니다.

분기 발산을 처리하기 위해, SM은 프레디케이션(predication) 기술을 사용합니다. 워프가 발산된 분기를 만나면, SM은 두 경로를 순차적으로 실행하면서 각 경로에 해당하지 않는 스레드를 마스킹합니다. 그런 다음 프레디케이트 레지스터를 사용하여 결과를 결합하여 각 스레드가 올바른 결과를 받도록 합니다.

예제 7.2는 발산된 분기가 있는 CUDA 커널을 보여줍니다.이 Markdown 파일의 한국어 번역은 다음과 같습니다. 코드의 경우 코드 자체는 번역하지 않고 주석만 번역했습니다.

__global__ void 분기가_발생하는_커널(int *data, int *result) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (data[tid] > 0) { // 데이터가 0보다 크면
        result[tid] = data[tid] * 2; // 데이터의 2배를 result에 저장
    } else { // 데이터가 0 이하면
        result[tid] = data[tid] * 3; // 데이터의 3배를 result에 저장
    }
}

예제 7.2: 분기가 발생하는 CUDA 커널.

이 예제에서 분기 조건 data[tid] > 0으로 인해 한 워프 내의 일부 스레드는 if 경로를, 나머지는 else 경로를 실행하게 됩니다. SM(Streaming Multiprocessor)은 이 분기 상황을 처리하기 위해 두 경로를 순차적으로 실행하고 각 경로에서 비활성 스레드를 마스킹합니다.

그림 7.4는 분기가 발생한 워프의 프레디케이션 과정을 보여줍니다.

    워프(32개 스레드)
    스레드 1: data[1] = 5, result[1] = 10
    스레드 2: data[2] = -3, result[2] = -9
    ...
    스레드 32: data[32] = 7, result[32] = 14

    분기:
    if (data[tid] > 0) {
        result[tid] = data[tid] * 2;
    } else {
        result[tid] = data[tid] * 3;
    }

    프레디케이션:
    1단계: if 경로 실행(마스킹 적용)
        스레드 1: result[1] = 10
        스레드 2: (비활성화)
        ...
        스레드 32: result[32] = 14

    2단계: else 경로 실행(마스킹 적용)
        스레드 1: (비활성화)
        스레드 2: result[2] = -9
        ...
        스레드 32: (비활성화)

    최종 결과:
    스레드 1: result[1] = 10
    스레드 2: result[2] = -9
    ...
    스레드 32: result[32] = 14

그림 7.4: 분기가 발생한 워프의 프레디케이션 과정.

프레디케이션을 사용하여 SM은 명시적인 분기 명령어 또는 제어 흐름 분기 없이 분기 상황을 처리할 수 있습니다. 그러나 분기가 발생하면 SM이 두 경로를 순차적으로 실행해야 하므로 실제 병렬 수행 능력이 감소됩니다.

레지스터 파일 및 오퍼랜드 수집기

레지스터 파일은 SM의 핵심 구성 요소로, 스레드 전용 레지스터에 대한 빠른 접근을 제공합니다. 각 SM은 많은 활성 스레드를 지원하고 워프 간 효율적인 문맥 전환을 가능하게 하기 위해 큰 레지스터 파일을 가지고 있습니다.Here is the Korean translation of the provided markdown file, with the code comments translated:

NVIDIA Volta 아키텍처에서, 각 SM은 2 KB씩 32개의 뱅크로 구성된 64 KB 레지스터 파일을 가지고 있습니다. 이 레지스터 파일은 많은 동시 스레드를 지원하기 위해 높은 대역폭과 낮은 대기 시간을 제공하도록 설계되었습니다.

성능 향상과 뱅크 충돌 최소화를 위해, SM은 피연산자 수집이라는 기술을 사용합니다. 피연산자 수집기는 레지스터 파일 뱅크에서 피연산자를 모아 실행 장치로 전달하는 특수한 장치입니다. 피연산자 수집기를 사용함으로써 SM은 뱅크 충돌의 영향을 줄이고 실행 장치의 활용도를 높일 수 있습니다.

그림 7.5는 SM의 레지스터 파일과 피연산자 수집기의 간략한 다이어그램을 보여줍니다.

    레지스터 파일 (64 KB)
    뱅크 1 (2 KB)
    뱅크 2 (2 KB)
    ...
    뱅크 32 (2 KB)
        |
        v
    피연산자 수집기
        |
        v
    실행 장치

그림 7.5: SM의 레지스터 파일과 피연산자 수집기.

피연산자 수집기는 여러 명령어와 여러 워프의 피연산자를 모아서, SM이 한 사이클에 다른 워프의 명령어를 실행 장치로 발행할 수 있게 합니다. 이를 통해 레지스터 파일 접근 지연 시간을 숨길 수 있어 SM의 전체 처리량이 향상됩니다.

예제 7.3은 두 벡터의 내적을 계산하는 CUDA 커널을 보여줍니다.

__global__ void dotProduct(float *a, float *b, float *result, int n) {
    // 각 스레드가 계산한 부분합을 저장하는 공유 메모리 배열
    __shared__ float partialSum[256];
    int tid = threadIdx.x;
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    // 각 스레드가 맡은 부분합을 0으로 초기화
    partialSum[tid] = 0;
 
    // 각 스레드가 자신의 범위에서 내적 계산
    while (i < n) {
        partialSum[tid] += a[i] * b[i];
        i += blockDim.x * gridDim.x;
    }
 
    // 모든 스레드가 내적 계산을 완료할 때까지 기다림
    __syncthreads();
 
    // 병렬 리덕션을 통해 부분합을 합침
    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid < s) {
            partialSum[tid] += partialSum[tid + s];
        }
        __syncthreads();
    }
 
    // 최종 결과를 결과 배열에 저장
    if (tid == 0) {
        result[blockIdx.x] = partialSum[0];
    }
}

이 예제에서 각 스레드는 자신에게 할당된 부분의 내적을 계산하여 부분합을 구합니다.다음은 제공된 마크다운 파일의 한국어 번역입니다. 코드 부분은 번역하지 않았고, 주석 부분만 번역했습니다.

입력 벡터의 요소들로부터 부분 합을 계산합니다. 부분 합은 공유 메모리 배열 partialSum에 저장됩니다. 모든 스레드가 자신의 부분 합을 계산한 후, 병렬 감소 연산을 수행하여 최종 내적 결과를 얻습니다.

오퍼랜드 수집기는 이 예제에서 핵심적인 역할을 합니다. 효율적으로 공유 메모리 액세스와 산술 연산을 위한 오퍼랜드를 수집함으로써, 뱅크 충돌을 방지하고 실행 유닛의 활용도를 높입니다.

결론

스트리밍 멀티프로세서는 현대 GPU 아키텍처의 핵심 연산 단위입니다. 그 설계는 세부적인 다중 스레딩, SIMT 실행, 그리고 효율적인 오퍼랜드 수집을 통해 처리량을 최대화하고 메모리 지연 시간을 숨기는 데 초점을 맞추고 있습니다.

SM의 주요 구성 요소에는 실행할 워프를 선택하는 워프 스케줄러, 분기 발산과 수렴을 처리하는 SIMT 스택, 스레드 전용 레지스터에 빠르게 액세스할 수 있는 레지스터 파일과 오퍼랜드 수집기, 그리고 데이터 공유와 재사용을 지원하는 공유 메모리와 L1 캐시 등이 있습니다.

GPU 아키텍처가 계속 발전함에 따라, 분기 발산 처리, 워프 스케줄링, 레지스터 파일 설계 등의 분야에 대한 연구가 미래 GPU의 성능과 효율성을 향상시키는 데 필수적일 것입니다. 동적 워프 형성, 스레드 블록 압축, 오퍼랜드 재사용 캐시 등의 새로운 기술은 SM의 기능을 크게 향상시키고 병렬 컴퓨팅 워크로드에서 새로운 수준의 성능을 실현할 수 있는 잠재력을 가지고 있습니다.

Chapter 6 Gpu Performance Metrics and Analysis Chapter 8 Interconnect and on Chip Networks