KAIST CS311 전산기조직 (Spring 2023) 교재: Computer Organization and Design: The Hardware/Software Interface (Patterson & Hennessy, MIPS Edition)
프로세서 관련 용어 정리
GPU를 이해하기 전에, 먼저 Processor와 관련된 핵심 용어들을 정확히 구분해두자. CPU와 GPU를 혼동 없이 비교하려면 이 용어들의 차이를 명확히 아는 것이 중요하다.
Processor, CPU, GPU
- Processor: 컴퓨터에서 데이터를 처리하고 Instruction을 실행하는 하드웨어 구성 요소의 총칭이다. CPU와 GPU가 대표적이다.
- CPU (Central Processing Unit): Core(s) 내부에서 작업(Thread)을 처리하는 프로세서이다. 각 Core가 독립적으로 Fetch, Decode, Execute를 수행한다.
- GPU (Graphics Processing Unit): 대량의 병렬 연산에 특화된 프로세서이다. GPU의 “Core"는 CPU Core와 달리 독립적인 Fetch, Decode를 수행하지 못하며, Execution Unit에 가깝다.
CPU 내부 구조
CPU 내부의 핵심 구성 요소를 정리하면 다음과 같다.
- Core: CPU 내부에서 실질적으로 작업(Thread)을 처리하는 하드웨어 단위이다.
- 각 Core마다 ALU, Register File, Controller가 독립적으로 할당된다
- Core 하나는 한 Clock Cycle 동안 한 Thread의 한 Instruction을 수행한다
- CPU Core는 독립적인 Fetch → Decode → Execute → Store 파이프라인을 갖는다
- ALU (Arithmetic Logic Unit): Core 내부에서 실질적으로 “계산"을 수행하는 Execution Unit이다.
- Multicore Processor ≡ Chip Multi-Processor (CMP): 내부에 여러 개의 Core가 존재하는 단일 프로세서 칩이다.
Multi-Processor vs Multicore
이 두 용어는 자주 혼동되므로 정확히 구분하자.
- Multicore Processor: 하나의 프로세서 칩 안에 여러 Core가 있는 것
- Multi-Processor: 여러 개의 프로세서(CPU)를 가진 컴퓨터 시스템을 지칭하는 용어
따라서 “이 CPU는 Multi-Processor이다"는 틀린 표현이고, “이 컴퓨터 시스템은 Multi-Processor System이다"가 올바른 표현이다.
┌─── Multi-Processor System ───────────────────────────────┐
│ │
│ ┌── CPU (Processor) ──┐ ┌── CPU (Processor) ──┐ │
│ │ ┌──────┐ ┌──────┐ │ │ ┌──────┐ ┌──────┐ │ │
│ │ │Core 0│ │Core 1│ │ │ │Core 0│ │Core 1│ │ │
│ │ │ ALU │ │ ALU │ │ │ │ ALU │ │ ALU │ │ │
│ │ └──────┘ └──────┘ │ │ └──────┘ └──────┘ │ │
│ │ ← Multicore → │ │ ← Multicore → │ │
│ └──────────────────────┘ └──────────────────────┘ │
│ │
│ ← Multi-Processor System → │
└──────────────────────────────────────────────────────────┘
Core, Process, Thread
Process와 Thread
- Process: 운영체제에서 실행 중인 프로그램의 인스턴스이다. 1개 이상의 Thread로 구성된다.
- Thread: Process 내에서 실행되는 작업의 기본 단위이다.
- 한 Process가 하나의 Thread만 처리할 수도 있고(Single-Threaded), 여러 Thread로 나누어 병렬 처리할 수도 있다(Multi-Threaded)
Multi-Thread Programming
기본적으로 대부분의 프로그램은 Single Thread로 동작한다. 하지만 프로그래머가 의도적으로 Process를 여러 Thread로 나누어 실행하도록 작성할 수 있으며, 이를 **멀티스레드 프로그래밍(Multi-Thread Programming)**이라 한다.
멀티스레드를 활용하면 여러 Core에 Thread를 분배하여 병렬로 실행할 수 있다. 예를 들어 4-Core CPU에서 4개의 Thread를 각 Core에 할당하면 이론적으로 4배의 처리량을 얻을 수 있다.
Hardware Multi-Threading (SMT)
개념
Software 수준에서 Thread를 여러 개 만드는 것과 별개로, 하드웨어 수준에서 한 Core가 여러 Thread의 실행을 동시에 관리하는 기술이 있다. 이를 Hardware Multi-Threading 또는 **Simultaneous Multi-Threading (SMT)**이라 한다.
Intel에서는 이를 Hyper-Threading이라는 브랜드명으로 부른다.
동작 원리
Single Threading 환경에서는 한 Core가 한 번에 하나의 Thread만 관리한다. SMT에서는 Core의 **State Hardware를 복제(duplicate)**하여 여러 Thread의 컨텍스트를 동시에 저장할 수 있게 한다.
State Hardware (Thread별로 복제되는 것):
- Register File
- Program Counter (PC)
- Instruction Buffer
- Store Buffer
- Page Table Base Register (PTBR)
공유 자원 (Thread 끼리 공유하는 것):
- Cache
- TLB (Translation Lookaside Buffer)
- Reorder Buffer
같은 Core 내의 Thread들은 동일한 메모리 시스템을 사용하므로, Cache와 TLB를 공유할 수 있다.
┌─── SMT Core (Hyper-Threading) ────────────────────┐
│ │
│ ┌─ Thread 0 State ─┐ ┌─ Thread 1 State ─┐ │
│ │ Register File │ │ Register File │ │
│ │ PC │ │ PC │ │
│ │ Inst Buffer │ │ Inst Buffer │ │
│ │ Store Buffer │ │ Store Buffer │ │
│ │ PTBR │ │ PTBR │ │
│ └───────────────────┘ └───────────────────┘ │
│ │ │ │
│ └──────┬───────────────┘ │
│ ▼ │
│ ┌─ Shared Resources ──────────────────────┐ │
│ │ ALU / Execution Units │ │
│ │ L1 Cache / L2 Cache │ │
│ │ TLB │ │
│ │ Reorder Buffer │ │
│ └─────────────────────────────────────────┘ │
└────────────────────────────────────────────────────┘
Context Switch와의 비교
Context Switch에는 두 가지 수준이 있다.
Process Context Switch: Process 간 전환으로, Memory Mapping, I/O Status, File Status 등 많은 상태 데이터를 교체해야 한다. 비용이 크다.
Thread Context Switch: 같은 Process 내의 Thread 간 전환이다. Register 값, PC, Stack 등은 교체하지만, Memory Mapping 같은 정보는 공유하므로 교체하지 않는다. Process 전환보다 가볍다.
SMT에서의 Thread 전환: 메모리를 참조하지 않고 단순히 다른 State Hardware를 선택하면 되므로, Thread Context Switch가 즉시(instantly) 수행된다. 이것이 SMT의 핵심 이점이다.
GPU 아키텍처
CPU vs GPU 구조 비교
CPU와 GPU는 설계 철학이 근본적으로 다르다.
- CPU: 커다란 고성능 Core를 소수(4~16개 정도) 보유. 복잡한 제어 로직, 큰 캐시, 분기 예측 등으로 단일 스레드의 성능을 극대화한다.
- GPU: 작은 Core를 매우 많이(수천 개) 보유. 단순한 제어 로직으로 대량 병렬 연산(throughput)을 극대화한다.
┌─── CPU ──────────────────┐ ┌─── GPU ──────────────────────────┐
│ │ │ │
│ ┌────────┐ ┌────────┐ │ │ ┌──┐┌──┐┌──┐┌──┐┌──┐┌──┐┌──┐ │
│ │ Control│ │ Control│ │ │ │SP││SP││SP││SP││SP││SP││SP│ │
│ │ │ │ │ │ │ └──┘└──┘└──┘└──┘└──┘└──┘└──┘ │
│ │ ALU │ │ ALU │ │ │ ┌──┐┌──┐┌──┐┌──┐┌──┐┌──┐┌──┐ │
│ │ ALU │ │ ALU │ │ │ │SP││SP││SP││SP││SP││SP││SP│ │
│ │ ALU │ │ ALU │ │ │ └──┘└──┘└──┘└──┘└──┘└──┘└──┘ │
│ │ ALU │ │ ALU │ │ │ ┌──┐┌──┐┌──┐┌──┐┌──┐┌──┐┌──┐ │
│ └────────┘ └────────┘ │ │ │SP││SP││SP││SP││SP││SP││SP│ │
│ │ │ └──┘└──┘└──┘└──┘└──┘└──┘└──┘ │
│ ┌────────────────────┐ │ │ ┌──┐┌──┐┌──┐┌──┐┌──┐┌──┐┌──┐ │
│ │ Cache │ │ │ │SP││SP││SP││SP││SP││SP││SP│ │
│ └────────────────────┘ │ │ └──┘└──┘└──┘└──┘└──┘└──┘└──┘ │
│ ┌────────────────────┐ │ │ │
│ │ DRAM │ │ │ (수천 개의 SP) │
│ └────────────────────┘ │ └──────────────────────────────────┘
└──────────────────────────┘
소수의 고성능 Core 다수의 단순 Core
순차 처리에 강함 병렬 처리에 강함
대표적인 성격 차이를 정리하면 다음과 같다.
| 특성 | CPU | GPU |
|---|---|---|
| Core 수 | 소수 (4~16) | 다수 (수천) |
| Core 성능 | 고성능, 복잡한 제어 | 단순, Execution Unit 수준 |
| 처리 방식 | 순차/소수 병렬 | 대량 병렬 |
| 설계 목표 | Latency 최소화 | Throughput 최대화 |
SP와 SM
GPU의 내부 구조를 이해하려면 SP와 SM 두 가지 핵심 단위를 알아야 한다.
- Streaming Processor (SP): 일반적인 의미에서 “GPU Core"라고 불리는 단위이다. 하지만 CPU Core와 달리 독립적인 Fetch, Decode, Store 능력이 없다. 실질적으로는 Execution Unit에 가깝다.
- Streaming Multiprocessor (SM): CPU에서의 Core 역할을 수행하는 단위이다. SM은 GPU 내에서 독립적으로 Fetch, Decode, Execute를 진행하기 때문에 CPU의 Core와 유사하다. SM 내부에 다수의 SP(Execution Unit)가 존재한다.
┌─── GPU ──────────────────────────────────────────────┐
│ │
│ ┌── SM 0 ──────────────┐ ┌── SM 1 ──────────────┐ │
│ │ ┌────┐┌────┐┌────┐ │ │ ┌────┐┌────┐┌────┐ │ │
│ │ │ SP ││ SP ││ SP │ │ │ │ SP ││ SP ││ SP │ │ │
│ │ └────┘└────┘└────┘ │ │ └────┘└────┘└────┘ │ │
│ │ ┌────┐┌────┐┌────┐ │ │ ┌────┐┌────┐┌────┐ │ │
│ │ │ SP ││ SP ││ SP │ │ │ │ SP ││ SP ││ SP │ │ │
│ │ └────┘└────┘└────┘ │ │ └────┘└────┘└────┘ │ │
│ │ ┌────────────────┐ │ │ ┌────────────────┐ │ │
│ │ │ Shared Memory │ │ │ │ Shared Memory │ │ │
│ │ └────────────────┘ │ │ └────────────────┘ │ │
│ │ Fetch / Decode Unit │ │ Fetch / Decode Unit │ │
│ └──────────────────────┘ └──────────────────────┘ │
│ │
│ ┌── SM 2 ──────────────┐ ┌── SM 3 ──────────────┐ │
│ │ ... │ │ ... │ │
│ └──────────────────────┘ └──────────────────────┘ │
│ │
│ ┌──────────────────────────────────────────────┐ │
│ │ Global Memory (DRAM) │ │
│ └──────────────────────────────────────────────┘ │
└──────────────────────────────────────────────────────┘
정리하면:
- GPU는 여러 개의 SM을 가진다
- 각 SM 내부에 다수의 SP(Execution Unit)가 존재한다
- SM은 독립적인 Fetch/Decode가 가능하므로 CPU의 Core에 대응된다
- SP는 ALU에 가까운 Execution Unit이다
CUDA 프로그래밍
CUDA(Compute Unified Device Architecture)는 NVIDIA GPU에서 범용 병렬 연산(GPGPU)을 수행하기 위한 프로그래밍 모델이다.
기본 용어
CUDA에서 사용하는 핵심 용어는 다음과 같다.
- Host = CPU
- Device = GPU
- Kernel = GPU에서 실행되는 함수(Function)
- CPU(Host)에서 호출(initiate)하고, GPU(Device)에서 실행(execute)된다
- 하나의 Kernel은 하나의 Grid에 의해 실행된다
Thread 계층 구조
CUDA의 Thread는 계층적으로 구성된다.
┌─── Grid (= 1 Kernel 실행 단위) ──────────────────────────────┐
│ │
│ ┌─ Block(0,0) ──┐ ┌─ Block(1,0) ──┐ ┌─ Block(2,0) ──┐ │
│ │ T(0,0) T(1,0) │ │ T(0,0) T(1,0) │ │ T(0,0) T(1,0) │ │
│ │ T(0,1) T(1,1) │ │ T(0,1) T(1,1) │ │ T(0,1) T(1,1) │ │
│ └────────────────┘ └────────────────┘ └────────────────┘ │
│ │
│ ┌─ Block(0,1) ──┐ ┌─ Block(1,1) ──┐ ┌─ Block(2,1) ──┐ │
│ │ T(0,0) T(1,0) │ │ T(0,0) T(1,0) │ │ T(0,0) T(1,0) │ │
│ │ T(0,1) T(1,1) │ │ T(0,1) T(1,1) │ │ T(0,1) T(1,1) │ │
│ └────────────────┘ └────────────────┘ └────────────────┘ │
│ │
│ Grid 내의 모든 Thread가 끝나야 Kernel 실행이 완료된다 │
└──────────────────────────────────────────────────────────────┘
Thread Block
Thread Block은 Thread의 묶음이다. Thread Block에는 다음과 같은 중요한 특성이 있다.
- SM과의 매핑: 하나의 Thread Block은 오직 하나의 SM에서 실행된다.
- SM은 여러 Thread Block을 동시에 관리할 수 있다
- 하지만 Thread Block이 여러 SM에 나뉘는 것은 불가능하다
- Shared Memory 공유: SM에 존재하는 Cache Memory(Shared Memory)를 Thread Block 내부의 Thread들끼리 공유할 수 있다.
- 동기화(Synchronize): Thread Block 내부의 Thread들 끼리는 서로 동기화가 가능하다.
Grid
Grid는 Thread Block의 묶음이다.
- 하나의 Kernel은 하나의 Grid에 의해 실행된다
- Kernel이 끝났다는 것은 Grid에 존재하는 모든 Thread가 끝났다는 것을 의미한다
CUDA 메모리 계층
CUDA에서 사용할 수 있는 메모리는 계층적으로 구성되어 있다.
┌─ Thread ─────────────────┐
│ Registers (Fastest) │ ← Thread 전용
│ Local Memory │ ← Thread 전용
└──────────────────────────┘
│
┌─ Thread Block ───────────┐
│ Shared Memory │ ← Block 내 Thread들 공유
│ (SW-managed Cache) │ SM 내에 위치
└──────────────────────────┘
│
┌─ Grid (All Threads) ────┐
│ Global Memory (DRAM) │ ← 모든 Thread 접근 가능
│ Constant Memory │ 가장 느리지만 용량 큼
│ Texture Memory │
└──────────────────────────┘
핵심은 Shared Memory이다. Shared Memory는 SM 내부에 위치하는 소프트웨어 관리 캐시(Software-managed Cache)로, Thread Block 안의 Thread들이 공유할 수 있다. Global Memory(DRAM)에 비해 접근 속도가 훨씬 빠르므로, 자주 접근하는 데이터를 Shared Memory에 올려두면 성능이 크게 향상된다.
Matrix Multiplication — CUDA 프로그래밍
행렬 곱셈(Matrix Multiplication)은 GPU 병렬 연산의 대표적인 예제이다. Host에서 Matrix Multiplication 커널을 호출하면, 하나의 Grid에 의해 전체 연산이 수행된다. Grid에서 얼마나 많은 Block으로 문제를 나누느냐에 따라 구현 방식이 달라진다.
방법 1: Only One Block
가장 단순한 방법은 한 개의 Block으로 전체 행렬 곱셈을 해결하는 것이다.
- 각 Thread가 결과 행렬의 한 원소(point) 계산을 맡는다
- Thread Block이 M과 N 행렬을 한 줄씩 읽어서 각각을 곱하고 더한 값을 P 행렬의 해당 point에 저장한다
- 각 Thread마다
for문 1개를 돌며 M의 한 행과 N의 한 열을 내적(dot product)한다
N (Width x Width)
┌───┬───┬───┬───┐
│ │ │ * │ │
│ │ │ * │ │
│ │ │ * │ │
│ │ │ * │ │
└───┴───┴───┴───┘
↓
M ┌───────┐ P
┌───┬───┬───┬───┐ │ ┌───┬───┬───┬───┐
│ │ │ │ │ │ │ │ │ │ │
│ * │ * │ * │ * │ ───┼──→ │ │ │ ● │ │ ← Thread(2,1) 계산
│ │ │ │ │ │ │ │ │ │ │
│ │ │ │ │ │ │ │ │ │ │
└───┴───┴───┴───┘ │ └───┴───┴───┴───┘
└───────┘
Row 1의 원소들과 P[1][2] = Σ M[1][k] * N[k][2]
Col 2의 원소들을 곱하여 합산
하지만 이 방법은 하나의 SM만 사용하므로 GPU의 병렬성을 충분히 활용하지 못한다.
방법 2: 여러 개의 Block 사용
행렬을 여러 Tile로 나누고, 한 Tile을 한 Block이 전담하는 방식이다.
예를 들어 16×16 행렬을 TILE_WIDTH = 2로 Tiling하면:
- 8×8 = 64개의 Tile Matrix가 생긴다
- 이러면 Tile 당 4개의 point를 맡으므로, Block당 필요한 Thread는 4개
- 총 Tile이 8×8개 필요하므로, Grid가 가져야할 Block의 수는 64개
여러 SM에 Block을 분배할 수 있으므로, GPU의 병렬성을 본격적으로 활용할 수 있게 된다.
방법 3: Tile 단위 Shared Memory 활용
행렬 곱셈에서는 같은 원소(Element)를 반복적으로 매우 많이 접근한다. 한 원소를 총 Matrix의 Width만큼 반복해서 접근하게 되므로, 외부 Memory(Global Memory) 접근 비용이 엄청나다.
여기서 Thread Block의 이점이 빛난다: Thread Block 내부의 Thread들은 SM의 Shared Memory를 공유할 수 있다!
핵심 아이디어: Tile 내에서 계산할 때, 필요한 TILE_WIDTH²개의 원소를 미리 읽어서 Shared Memory에 저장한다. 이렇게 하면 외부에서 Matrix의 Element를 읽는 것은 오직 한 번이고, 이후 Shared Memory를 TILE_WIDTH만큼 반복해서 접근하면 된다.
결과적으로 같은 원소가 한 줄의 Tile 수만큼만 접근되므로, 원래 Width번 접근하던 것이 Width / TILE_WIDTH번으로 줄어든다.
┌─── Tiled Matrix Multiplication ─────────────────────────────┐
│ │
│ Global Memory에서 Tile 크기만큼만 읽어옴 │
│ ↓ │
│ ┌─── Shared Memory ───┐ │
│ │ subTileM[][] │ ← M의 Tile 한 조각 │
│ │ subTileN[][] │ ← N의 Tile 한 조각 │
│ └─────────────────────┘ │
│ ↓ │
│ 각 Thread가 Shared Memory에서 빠르게 읽어 부분 합 계산 │
│ ↓ │
│ 다음 Tile 조각을 Global Memory에서 읽어 Shared Memory 갱신 │
│ ↓ │
│ (Width / TILE_WIDTH)번 반복 후 최종 결과를 Global Memory에 │
│ 기록 │
└─────────────────────────────────────────────────────────────┘
Barrier (동기화)
__syncthreads()
Thread Block 내부의 Thread들끼리는 __syncthreads() 함수를 통해 **동기화(sync)**할 수 있다.
Shared Memory 문제 상황
Tiled Matrix Multiplication에서 Shared Memory를 사용할 때 동기화 문제가 발생한다.
문제 상황: 한 Tile에서 계산할 때, 모든 Thread의 Element를 External Memory에서 읽어온 다음에 계산을 시작해야 한다. 각 Thread마다 한 point씩 할당되어 있으므로, 만약 다른 Thread가 아직 Shared Memory에 값을 쓰지 않았는데 다른 Thread가 그 값을 읽으려고 하면 문제가 생긴다.
해결: 모든 Thread가 Shared Memory에 값을 다 쓸 때까지 기다린 후 계산을 시작한다. 이를 __syncthreads()로 구현한다.
동기화가 필요한 두 지점:
- Shared Memory 로딩 후: 모든 Thread가 Shared Memory에 값을 다 쓸 때까지 대기
- 계산 완료 후, 다음 Tile 로딩 전: 다른 Thread가 아직 현재 Tile 계산을 안 끝냈는데 다음 iteration으로 넘어가서 Shared Memory의 값을 Overwrite하면 안 되므로 대기
Tiled Matrix Multiplication Kernel 코드
위의 개념을 종합한 CUDA Kernel 코드이다.
__global__ void MatrixMulKernel(float* M, float* N, float* P, int Width) {
// Shared Memory 선언 — Thread Block 내에서 공유
__shared__ float subTileM[TILE_WIDTH][TILE_WIDTH];
__shared__ float subTileN[TILE_WIDTH][TILE_WIDTH];
// Block, Thread 인덱스 추출
int bx = blockIdx.x; int by = blockIdx.y;
int tx = threadIdx.x; int ty = threadIdx.y;
// 이 Thread가 담당하는 결과 행렬 P의 Row, Col
int Row = by * TILE_WIDTH + ty;
int Col = bx * TILE_WIDTH + tx;
float Pvalue = 0;
// M과 N의 Tile을 순회하며 부분 합 계산
for (int m = 0; m < Width / TILE_WIDTH; ++m) {
// Global Memory → Shared Memory로 Tile 로딩
subTileM[ty][tx] = M[Row * Width + m * TILE_WIDTH + tx];
subTileN[ty][tx] = N[(m * TILE_WIDTH + ty) * Width + Col];
// 동기화: 모든 Thread가 Shared Memory에 값을 다 쓸 때까지 대기
__syncthreads();
// Shared Memory에서 읽어 부분 합 계산
for (int k = 0; k < TILE_WIDTH; ++k)
Pvalue += subTileM[ty][k] * subTileN[k][tx];
// 동기화: 다음 Tile을 로딩하기 전에 모든 Thread의 계산 완료 대기
__syncthreads();
}
// 최종 결과를 Global Memory에 기록
P[Row * Width + Col] = Pvalue;
}
코드의 핵심 포인트를 정리하면:
__shared__키워드: Shared Memory에 배열을 선언한다. Thread Block 안의 모든 Thread가 이 배열을 공유한다.blockIdx,threadIdx: 각 Thread가 Grid 내에서의 자신의 위치를 식별한다.- Tile 순회:
for (int m = 0; m < Width/TILE_WIDTH; ++m)루프를 통해 M과 N을 Tile 단위로 순회한다. - 두 번의
__syncthreads():- 첫 번째: Shared Memory 로딩 완료 보장
- 두 번째: 계산 완료 후 다음 Tile Overwrite 방지
정리
| 개념 | 설명 |
|---|---|
| CPU Core | 독립적인 Fetch/Decode/Execute, 고성능 |
| GPU SP | Execution Unit, 독립 Fetch 불가 |
| GPU SM | CPU Core에 대응, 독립 Fetch/Decode 가능, 내부에 다수 SP |
| SMT | 한 Core가 여러 Thread의 State Hardware를 보유 |
| CUDA Host/Device | CPU / GPU |
| Kernel | GPU에서 실행되는 함수, 1 Grid에 의해 실행 |
| Thread Block | SM에 매핑, Shared Memory 공유, 동기화 가능 |
| Grid | Thread Block의 묶음, Kernel 실행 단위 |
| Shared Memory | SM 내부 SW-managed Cache, Block 내 Thread 공유 |
__syncthreads() | Thread Block 내 동기화 Barrier |
| Tiled MatMul | Shared Memory 활용으로 Global Memory 접근 Width/TILE_WIDTH배 감소 |