CUDA의 구조 이해: SM, Block, Thread완전 정리

🧠 1. 전체 구조 개요

GPU는 다음과 같이 계층적으로 구성되어 있습니다.

GPU

GPU
 ┣━━ SM0
 │    ┣━━ Block0
 │    │     ┣━━ Thread0
 │    │     ┣━━ Thread1
 │    │     ┣━━ Thread2
 │    │     ┗━━ Shared Memory (Block0 전용)
 │    ┣━━ Block1
 │    │     ┣━━ Thread0
 │    │     ┣━━ Thread1
 │    │     ┣━━ Thread2
 │    │     ┗━━ Shared Memory (Block1 전용)
 │    ┗━━ Register File, L1 Cache
 │
 ┃
 ┗━━ Global Memory (VRAM, GPU 외부 DRAM)


GPU
 ┣━━ 여러 GPC
 ┃     ┣━━ 여러 SM
 ┃     │     ┣━━ Warp Scheduler / CUDA Cores / Tensor Cores / Shared Memory / Registers
 ┃     │     ┗━━ Block 단위 실행
 ┃     ┗━━ L2 Cache, Memory Controller
 ┗━━ Global Memory (VRAM)
        ┗━━ GPU 전체 데이터 저장

GPU → SM → Block → Thread
이 순서로 구성되어 있으며, 병렬 연산을 수행합니다.

 

구성요소	역할
GPC (Graphics Processing Cluster)	여러 SM을 묶은 클러스터 단위 (SM 관리)
SM (Streaming Multiprocessor)	실제 연산이 일어나는 핵심 유닛
L2 Cache	모든 SM이 공유하는 캐시 (Global Memory 접근 완화)
Memory Controller	VRAM(Global Memory)와의 데이터 송수신 관리
Global Memory (VRAM)	GPU 외부의 대용량 메모리 (수 GB, 느림)
PCIe / NVLink	CPU와 GPU 간 데이터 전송 통로

---

⚙️ 2. SM (Streaming Multiprocessor)

• GPU의 핵심 연산 단위
• CPU로 치면 “코어(Core)”와 비슷한 역할
• 내부에는 수백 개의 CUDA Core, 레지스터, shared memory, warp scheduler 등이 포함됨

SM구성

SM의 역할

• Block을 받아 실행
• Thread들을 warp 단위로 스케줄링
• shared memory를 통해 thread 간 데이터 교환 수행

구성요소	역할
Warp Scheduler	32-thread 묶음(=warp)을 스케줄링, 동시에 실행
Dispatch Unit	Warp Scheduler가 명령어를 각 유닛에 분배
CUDA Cores (FP32 ALU)	부동소수점/정수 연산 수행 (가장 기본적인 계산기)
Tensor Cores	행렬·딥러닝용 연산 전용 유닛 (FP16/BF16 등)
Load/Store Unit	메모리 읽기/쓰기 관리 (Global, Shared 등)
Register File	각 Thread가 사용하는 초고속 변수 저장 공간
Shared Memory	Block 내부 Thread들이 공유하는 데이터 공간
L1 Cache / Texture Cache	SM 내부 캐시 (자주 쓰는 데이터 빠르게 접근)
Special Function Unit (SFU)	삼각함수, 제곱근 등 특수 연산 수행
Instruction Buffer / Decode	명령어를 받아 해석하고 파이프라인에 전달

 

---

🧩 3. Block (Thread Block) - 작업팀

• 하나의 커널 함수(global) 실행 시, 연산이 여러 Block으로 나뉨
• 각 Block은 하나의 SM 위에서 실행됨
• Block은 내부적으로 여러 Thread들의 집합

Shared Memory는 SM 내부에 물리적으로 존재하지만
논리적으로는 각 Block별로 분리되어 사용됩니다.

구성 요소 역할

구성요소	역할
Thread 집합	여러 Thread가 같은 코드(kernel)을 동시에 실행
Shared Memory (SM 내부 일부)	Block 내 Thread들이 데이터 공유
Thread Index (threadIdx)	Block 내 Thread의 고유 인덱스
Block Index (blockIdx)	전체 Grid 내 Block의 위치 지정
Sync Barrier (__syncthreads)	Block 내부 Thread 간 동기화 명령

---

🔹 4. Thread (스레드) - 작업자

• 실제 연산을 수행하는 가장 작은 단위
• 각 thread는 고유한 인덱스를 가짐
• 예: 벡터 덧셈

구성 요소	역할
Register	Thread 전용 변수 저장 (초고속)
Program Counter (PC)	실행할 명령어 위치
Local Memory	Register가 부족할 때 spill되는 임시 저장공간
Thread ID	Block 내 위치(threadIdx.x/y/z)
Warp 단위 실행 (32개)	같은 명령어를 동시에 실행하는 SIMD 구조
Divergence Handling	분기문(if/else) 처리 시 warp 내 병렬제어

 

 

🔸 5. Warp (워프)

• 한 SM은 Thread들을 32개 단위로 묶어서 동시에 실행
• 동일한 명령어를 동시에 실행하는 SIMD 방식
• 조건문 등으로 분기되면 warp divergence(성능 저하) 발생

---

💾 6. 메모리 계층 (Memory Hierarchy)

메모리 종류 접근 범위 속도 설명

Register	Thread 전용	🔥 매우 빠름	Thread 내부 변수 저장
Shared Memory	Block 내 공유	⚡ 빠름	Thread 간 데이터 교환
Global Memory	모든 Thread	🐢 느림	GPU의 DRAM
Constant / Texture	모든 Thread	⚙️ 중간	읽기 전용 캐시

🧭 7. 실제 연산 흐름 예시 (벡터 덧셈)

1. CPU → GPU로 데이터 복사
2. kernel<<<gridDim, blockDim>>>(A, B, C) 호출
3. SM이 여러 Block을 받아 처리
4. 각 Block 안의 Thread들이 동시에 A[i] + B[i] 계산
5. 결과를 Global Memory에 저장
6. CPU가 결과를 가져옴

[데이터 흐름]
CPU (Host RAM)
   ↓  (PCIe / NVLink)
GPU VRAM (Global Memory)
   ↓
Memory Controller (메모리 접근 제어)
   ↓
L2 Cache (GPU 전체 공용, on-die)
   ↓
Streaming Multiprocessor (SM)
   ┣━━ L1 Cache
   ┣━━ Shared Memory (Block 단위 공유)
   ┗━━ Warp Scheduler / CUDA Cores / Tensor Cores
        ↓
     Register (Thread 전용 초고속 메모리)
        ↓
     ALU (산술논리연산기, 실제 트랜지스터 단위 연산)


[물리적구조]
트랜지스터 (Transistor)
   ↓
ALU (Arithmetic Logic Unit)
   ↓
CUDA Core (FP32 연산 단위)
   ↓
SM (Streaming Multiprocessor)
   ↓
GPC (Graphics Processing Cluster)
   ↓
GPU 다이 (Die, 실리콘 웨이퍼 상의 회로)
   ↓
GPU 칩 (패키징된 반도체)
   ↓
GPU 제품 (그래픽카드 완제품)

 

[참고]

CUDA명칭	물리적 위치	공유범위	속도	실제부품
Register	SM 내부	Thread 단위	🔥 매우 빠름	칩 내부 SRAM
Shared Memory	SM 내부	Block 단위	⚡ 빠름	칩 내부 SRAM
L1 / L2 Cache	칩 내부	SM 및 GPU 전체	⚙️ 중간	칩 내부 SRAM
Global Memory	GPU 칩 외부	GPU 전체	🐢 느림	VRAM (GDDR6 등)
Constant / Texture	칩 내부 캐시 + VRAM	GPU 전체	⚙️ 중간	VRAM + 전용 캐시

 

용어	제조	위치
GPU 다이 (GPU Chip)	NVIDIA / TSMC 웨이퍼	연산, SM/CUDA 포함
VRAM 칩 (GDDR6 등)	삼성 / 마이크론 웨이퍼	데이터 저장

 

 

0. (사전) CPU 쪽 준비

1. CPU에서 데이터 준비CPU는 연산에 쓸 배열 A, B, C를 RAM에 올려둠.
2.  
   float *hA, *hB, *hC; // host 메모리(CPU)
3. GPU 쪽 메모리 공간 예약→ 여기서 할당되는 건 Global Memory(VRAM) 영역.
4.  
   float *dA, *dB, *dC; // device 메모리(GPU) cudaMalloc(&dA, N*sizeof(float)); cudaMalloc(&dB, N*sizeof(float)); cudaMalloc(&dC, N*sizeof(float));

---

1. CPU → GPU 데이터 전송

1. CPU → GPU로 복사
   • 이 순간 데이터는
     • CPU RAM → (PCIe/NVLink) → GPU Global Memory(VRAM) 으로 감
   • 아직 연산은 시작 안 했고, 데이터만 GPU 쪽에 올라간 상태.
2.  
   cudaMemcpy(dA, hA, N*sizeof(float), cudaMemcpyHostToDevice); cudaMemcpy(dB, hB, N*sizeof(float), cudaMemcpyHostToDevice);

---

2. 커널 호출 (작업 설명서 주기)

1. 커널 실행 요청여기서 일어나는 일:
   • CPU는 “이런 커널을, 이런 그리드/블록 크기로, 이 데이터에 대해 실행해”라고 명령만 보냄.
   • 실제 연산 스케줄링은 이제 GPU 내부 하드웨어(SM, 스케줄러) 가 맡음.
   • gridSize 개의 Block이 생성된다고 보면 됨.
2.  
   int blockSize = 256; // block당 thread 수 int gridSize = (N + 255)/256; // block 개수 vecAdd<<<gridSize, blockSize>>>(dA, dB, dC, N);

---

3. GPU 안에서의 스케줄링 (SM 레벨)

1. GPU가 Block들을 여러 SM에 배치
   • GPU 안에는 예를 들어 SM이 80개 있다고 하면,
   • 생성된 Block 수가 80개보다 많을 확률이 높음 → 대기열처럼 관리
   • 각 SM은 여유가 생길 때마다 Block을 하나씩 가져가서 실행
   • 이때, Block 하나는 무조건 하나의 SM 안에서만 돈다 (분산 안됨)

---

4. SM 안에서의 Block 실행

1. SM이 Block을 받으면:
   • 그 Block에 대해 Shared Memory 공간을 분할/할당
     (커널이 __shared__ 선언했으면 여기서 잡힘)
   • Block 안의 Thread들이 사용할 Register도 Register File에서 할당
   • 그 다음 Block 안의 Thread들을 32개씩 잘라서 Warp로 만듦

---

5. Warp 단위 연산 (Thread 실제 실행)

1. Warp Scheduler 동작
   • SM 안에는 보통 여러 Warp가 동시에 “준비 완료” 상태
   • Warp Scheduler가 “이번 사이클에는 Warp #3 실행!” 이런 식으로 고름
   • 그러면 그 Warp에 속한 Thread 32개가 같은 명령어를 동시에 실행
     → 이게 우리가 보는 C[i] = A[i] + B[i]; 이런 한 줄
2. 각 Thread의 인덱스 계산
   커널 안에서는 보통 이렇게 되어 있음:
   • blockIdx.x : 이 Block이 Grid에서 몇 번째 Block인지
   • blockDim.x : 하나의 Block 안에 Thread가 몇 개 있는지
   • threadIdx.x: Block 안에서 내가 몇 번째 Thread인지
   • 이 셋을 조합해서 전역 인덱스(global index) 하나가 만들어짐 → idx
   • 이 덕분에 모든 Thread가 같은 코드를 실행해도 서로 다른 데이터를 처리하게 됨
3.  
   __global__ void vecAdd(float *A, float *B, float *C, int N) { int idx = blockIdx.x * blockDim.x + threadIdx.x; if (idx < N) { C[idx] = A[idx] + B[idx]; } }

---

6. 메모리 접근 경로

1. Thread가 A[idx], B[idx] 읽을 때
   • 우선 SM의 L1/Shared에서 찾음
   • 없으면 → L2 Cache → 그래도 없으면 → Global Memory(VRAM) 에서 가져옴
   • Global Memory는 느리니까, 보통은 coalesced access 되게 인덱스를 짜거나,
     Shared Memory에 한 번 올려놓고 여러 Thread가 같이 쓰도록 함
2. 덧셈 연산
   • 실제 연산은 CUDA Core(FP32 ALU)에서 수행
   • C[idx] = A[idx] + B[idx]; 이 한 줄이
     • Load A
     • Load B
     • FADD
     • Store C
       이런 식으로 내부에서는 여러 마이크로옵으로 실행
3. 결과 저장
   • 계산된 값은 다시 Global Memory(dC[idx]) 에 저장됨
   • 이때도 Store 경로는 L1/L2를 타고 내려가 VRAM에 반영됨

---

7. 모든 Block 끝날 때까지 반복

1. SM이 Block 하나를 끝내면
   → 대기 중인 다음 Block을 가져와 또 실행
   → 이걸 GPU 전체 SM들이 동시에 하고 있으니 병렬성이 큰 것

---

8. CPU가 결과 다시 가져오기

1. CPU로 다시 복사
   • GPU Global Memory → (PCIe) → CPU RAM
   • 이 시점에서야 CPU는 결과를 볼 수 있음
2.  
   cudaMemcpy(hC, dC, N*sizeof(float), cudaMemcpyDeviceToHost);
3. GPU 자원 해제
4.  
   cudaFree(dA); cudaFree(dB); cudaFree(dC);

---

📌 한 장으로 다시 적어보면

1. Host 준비 → 데이터 CPU RAM에 있음
2. Device 메모리 할당 → GPU Global Memory에 공간 만듦
3. Host→Device 복사 → PCIe 통해 VRAM으로
4. 커널 호출 → grid, block 크기 전달
5. GPU가 Block을 SM들에 분산
6. SM이 Block을 받아 warp로 쪼개서 실행
7. 각 Thread가 Global Memory에서 자기 데이터 읽음
8. ALU(CUDA Core)에서 연산
9. 결과를 다시 Global Memory에 씀
10. 모든 Block 끝나면 CPU가 Device→Host로 결과 복사
11. CPU가 최종 결과 사용