422. 스레드 블록 (Thread Block)과 워프 (Warp)

핵심 인사이트 (3줄 요약)

1. 본질: 스레드 블록 (Thread Block)과 워프 (Warp)는 GPU (Graphics Processing Unit)가 수많은 스레드를 "개별 작업자"가 아니라 "정해진 단위의 실행 묶음"으로 다루게 만드는 계층 구조다.
2. 가치: 워프는 실행 스케줄링의 최소 단위, 블록은 협업과 동기화의 최소 단위라서, 이 둘을 구분해야 병렬성·메모리 재사용·지연 은닉 (Latency Hiding)을 함께 얻을 수 있다.
3. 판단 포인트: 블록 크기는 단순히 크게 잡는다고 좋은 것이 아니라, 워프 배수·공유 메모리 (Shared Memory) 사용량·레지스터 압박·분기 발산을 함께 고려해 정해야 한다.

Ⅰ. 개요 및 필요성

스레드 블록과 워프는 CUDA (Compute Unified Device Architecture) 기반 GPU 실행 모델에서 스레드를 조직하는 기본 단위다. 워프는 하드웨어가 실제로 한 번에 실행하는 최소 묶음이고, 블록은 같은 스트리밍 멀티프로세서 (Streaming Multiprocessor, SM) 안에서 함께 배치되어 협력할 수 있는 묶음이다.

이 구조가 필요한 이유는 GPU가 CPU (Central Processing Unit)처럼 스레드 하나하나를 무겁게 관리하지 않기 때문이다. GPU는 수천~수만 개의 스레드를 동시에 걸어 두고, 메모리 지연이 생기면 다른 워프로 즉시 교체하면서 연산기를 쉬지 않게 만든다. 만약 스레드가 아무 질서 없이 흩어져 있다면, 하드웨어는 누가 함께 실행되어야 하는지, 누가 메모리를 공유할 수 있는지, 어디까지를 동기화할 수 있는지 판단하기 어렵다.

즉 워프는 "실행 효율"을 위한 규격이고, 블록은 "협업 범위"를 위한 규격이다. 이 둘이 분리되어 있기 때문에 GPU는 한편으로는 같은 명령을 대량 병렬 실행하고, 다른 한편으로는 블록 내부에서 데이터 재사용과 장벽 동기화까지 제공할 수 있다.

이 그림은 스레드가 어떻게 계층적으로 묶여 실행과 협업의 단위로 바뀌는지 보여준다.

┌──────────────────────────────────────────────────────────────┐
│        GPU 스레드 조직: 실행 단위와 협업 단위의 분리         │
├──────────────────────────────────────────────────────────────┤
│ Grid                                                        │
│  ├─ Block 0  ─┬─ Warp 0  ─ Thread 0  ... Thread 31          │
│  │            ├─ Warp 1  ─ Thread 32 ... Thread 63          │
│  │            └─ Shared Memory + __syncthreads() 가능       │
│  ├─ Block 1  ─┬─ Warp 0  ─ Thread 0  ... Thread 31          │
│  │            └─ 다른 SM에 배치될 수 있으며 Block 0과 분리   │
│  └─ ...                                                    │
├──────────────────────────────────────────────────────────────┤
│ Warp  = 실행 스케줄링 최소 단위                             │
│ Block = 협업·동기화·자원 할당 최소 단위                     │
└──────────────────────────────────────────────────────────────┘

핵심은 블록과 워프가 같은 계층을 다른 말로 부르는 것이 아니라는 점이다. 워프는 "같이 달리는 묶음"이고, 블록은 "같은 작업실을 쓰는 묶음"이다. 따라서 CUDA 커널을 설계할 때는 항상 "이 코드는 워프 관점에서 잘 실행되는가"와 "이 데이터는 블록 관점에서 잘 협업되는가"를 동시에 봐야 한다.

• 📢 섹션 요약 비유: 워프는 32명이 동시에 출발하는 달리기 레인이고, 블록은 같은 공구함을 쓰는 작업반이다. 같이 뛰는 규칙과 같이 일하는 범위가 다르기 때문에 둘을 섞어 이해하면 GPU 성능이 꼬인다.

Ⅱ. 아키텍처 및 핵심 원리

워프와 블록의 차이를 이해하려면 먼저 하드웨어 자원 배치 방식을 봐야 한다. 하나의 블록은 반드시 하나의 SM 안에 통째로 올라가며, 그 블록 안의 스레드들은 여러 개의 워프로 쪼개져 워프 스케줄러의 선택을 받는다. 예를 들어 블록 크기가 256이면, 이는 256개의 독립 실행 단위가 아니라 8개의 워프로 해석된다.

이 그림은 블록이 SM 안에 배치될 때 실제로 어떤 자원 제약을 받는지 압축해서 보여준다.

┌──────────────────────────────────────────────────────────────┐
│              Block 배치와 Warp 실행의 실제 모습              │
├──────────────────────────────────────────────────────────────┤
│ SM                                                          │
│  ├─ 자원 한도: Registers / Shared Memory / Max Threads      │
│  ├─ Block A (256 threads)                                   │
│  │   ├─ Warp 0                                               │
│  │   ├─ Warp 1                                               │
│  │   ├─ ...                                                  │
│  │   └─ Warp 7                                               │
│  ├─ Block B (256 threads)                                   │
│  │   └─ 동일 방식으로 분해                                   │
│  └─ Warp Scheduler: 준비된 Warp를 번갈아 실행                │
├──────────────────────────────────────────────────────────────┤
│ Block이 너무 크면  : 동시에 올라갈 Block 수 감소             │
│ Block이 너무 작으면: Warp 수 부족 또는 자원 단편화 발생      │
└──────────────────────────────────────────────────────────────┘

여기서 중요한 정량 감각은 다음과 같다. 블록 크기가 48이면 하드웨어는 이를 1.5 워프로 처리할 수 없으므로 실제로는 2개 워프로 본다. 즉 48개 중 32개는 첫 번째 워프를 채우지만, 나머지 16개는 두 번째 워프에 들어가고 16개 슬롯은 비어 있는 셈이다. 그래서 블록 크기를 32의 배수로 잡는 것이 기본 원칙이 된다.

하지만 32의 배수라고 항상 최적은 아니다. 블록을 1024처럼 크게 잡으면 워프 수는 많아지지만, 블록 하나가 레지스터와 공유 메모리를 많이 점유해 같은 SM에 동시에 올라갈 수 있는 블록 수가 줄 수 있다. 그러면 메모리 접근을 기다리는 동안 교체할 워프 풀이 부족해져 지연 은닉 효과가 약해진다. 결국 좋은 블록 크기는 "워프 낭비가 적고, 동시에 충분한 워프를 유지하며, 블록 내부 협업 비용도 감당 가능한 지점"이다.

• 📢 섹션 요약 비유: 블록 크기 설정은 창고에 박스를 쌓는 일과 같다. 너무 작은 박스만 쌓으면 공간이 자투리로 남고, 너무 큰 박스를 넣으면 몇 개 못 들어가서 회전율이 떨어진다.

Ⅲ. 비교 및 연결

스레드 블록과 워프는 자주 같이 언급되지만, 설계 질문이 다르다. 워프는 "한 번에 어떤 스레드들이 같은 명령 흐름으로 묶이는가"를 묻고, 블록은 "어디까지를 같은 협업 구역으로 볼 것인가"를 묻는다. 이 차이를 구분해야 워프 발산 (Warp Divergence) 문제와 블록 동기화 문제를 분리해서 볼 수 있다.

예를 들어 if 조건이 워프 내부에서 갈라지면 같은 워프에 속한 스레드들이 번갈아 실행되어 효율이 떨어진다. 반면 블록 내부에서 __syncthreads()를 너무 자주 호출하면, 각 워프가 모두 장벽에 도착할 때까지 기다려야 하므로 협업 비용이 커진다. 즉 워프는 제어 흐름의 문제를, 블록은 협업 구조의 문제를 더 직접적으로 드러낸다.

이 개념은 다른 병렬 컴퓨팅 모델과도 연결된다. CPU 멀티스레딩은 스레드마다 제어 흐름 자율성이 크지만 문맥 교환 비용이 무겁다. 반대로 GPU 워프는 자율성을 줄이는 대신 제어를 단순화해 대량 병렬 처리량을 얻는다. 또한 블록은 분산 시스템의 "한 노드 안 협업"과 비슷하고, 블록 간 관계는 "노드 간 독립 작업"에 가깝다. 그래서 GPU 커널 설계는 미세하게 보면 하드웨어 스케줄링 문제이고, 크게 보면 병렬 작업 분할 문제이기도 하다.

• 📢 섹션 요약 비유: 워프는 같은 버스에 탄 승객들이고, 블록은 같은 캠프에 배정된 팀이다. 버스 안에서는 모두 같은 길로 움직여야 하고, 캠프 안에서는 장비를 나눠 쓰며 작업 순서를 맞춰야 한다.

Ⅳ. 실무 적용 및 기술사 판단

실무에서 블록과 워프를 다룰 때 가장 흔한 질문은 "블록 크기를 몇으로 잡을 것인가"다. 정답은 고정값이 아니라, 커널의 성격에 따라 기준이 달라진다. 단순 벡터 연산처럼 공유 메모리 사용이 적고 제어 흐름이 단순한 커널은 128~256 스레드 블록이 안정적인 출발점이 되곤 한다. 반면 타일링 기반 행렬 곱처럼 블록 내부 데이터 재사용이 중요하면 16×16, 32×8처럼 메모리 배치와 맞물린 2차원 블록 구성이 더 적절할 수 있다.

실무 판단 체크리스트

1. 블록 크기가 32의 배수인가?
2. 블록당 공유 메모리와 레지스터 사용량 때문에 SM 동시 상주 블록 수가 과도하게 줄지 않는가?
3. 같은 워프에 들어가는 스레드들이 비슷한 분기 경로를 타는가?
4. 워프 내 메모리 접근이 연속적이어서 코얼레싱이 잘 일어나는가?
5. 블록 내부 동기화가 실제 데이터 재사용 이득보다 더 비싸지 않은가?

대표 안티패턴

• 워프 크기를 무시한 40, 48, 96 같은 블록 설정을 습관적으로 사용하는 경우
• 블록 내부 협업이 거의 없는데도 큰 블록과 잦은 장벽 동기화를 넣는 경우
• 분기 많은 데이터를 같은 워프에 섞어 넣어 발산 비용을 키우는 경우
• 공유 메모리를 과하게 써서 오히려 점유율을 무너뜨리는 경우

기술사 관점에서 중요한 판단 문장은 이것이다. 블록은 크게 잡을수록 좋은 것이 아니라, 워프 효율·점유율·협업 이득의 균형이 맞을 때만 좋은 선택이다. 따라서 튜닝은 "최대 스레드 수"가 아니라 "병목이 어디서 생기는가"를 기준으로 해야 한다. 메모리 병목이면 코얼레싱과 타일링을 먼저 보고, 제어 병목이면 워프 발산을 먼저 줄이는 식으로 접근해야 한다.

• 📢 섹션 요약 비유: 요리 주방에서 한 조를 너무 크게 만들면 재료는 많이 나르지만 조리대가 비좁아지고, 너무 작게 만들면 손은 남아도는데 교대 인원이 부족해진다. 적정 인원 배치가 곧 성능이다.

Ⅴ. 기대효과 및 결론

스레드 블록과 워프를 제대로 이해하면 GPU 병렬성의 본질이 보인다. GPU 성능은 "코어 수가 많다"보다 "워프가 끊기지 않고 흐르며, 블록 내부 재사용이 잘 일어난다"에 더 가깝다. 즉 계산량만 많은 커널이 아니라, 하드웨어가 좋아하는 형태로 조직된 커널이 빠르다.

이 구조의 기대효과는 분명하다. 워프 단위 스케줄링은 메모리 지연을 숨기고, 블록 단위 협업은 글로벌 메모리 왕복을 줄이며, 둘의 조합은 대규모 행렬 연산·이미지 처리·딥러닝 추론 같은 작업에서 높은 처리량을 만든다. 반면 한계도 분명하다. 워프 내부 분기 불균형이 심하거나, 블록 자원 사용이 지나치게 크거나, 데이터 배치가 하드웨어와 맞지 않으면 이론적 병렬성은 실제 성능으로 이어지지 않는다.

앞으로는 스레드 블록 클러스터 (Thread Block Cluster)처럼 블록보다 큰 협업 단위가 늘어나더라도, 기본 철학은 바뀌지 않는다. 실행은 작은 규격으로 묶고, 협업은 통제 가능한 범위로 제한하며, 데이터 배치는 하드웨어 친화적으로 설계한다는 원칙이다. 따라서 이 개념은 "GPU의 스레드 묶음"이 아니라 실행 단위와 협업 단위를 분리해 병렬성을 관리하는 방법으로 기억하는 것이 가장 정확하다.

• 📢 섹션 요약 비유: 잘 훈련된 공장은 작업자를 그냥 많이 넣는 곳이 아니라, 같이 움직일 팀과 같이 도구를 쓸 팀을 다르게 설계하는 곳이다. 워프와 블록은 GPU 공장의 그 설계도다.

📌 관련 개념 맵

📈 관련 키워드 및 발전 흐름도

CUDA (Compute Unified Device Architecture)
    │
    ▼
SM (Streaming Multiprocessor) 기반 실행 모델
    │
    ├─▶ Warp : SIMT (Single Instruction Multiple Threads) 실행
    │          │
    │          └─▶ Warp Divergence · Memory Coalescing
    │
    └─▶ Thread Block : Shared Memory · Barrier 동기화
               │
               └─▶ Occupancy 최적화 · Thread Block Cluster 확장

이 흐름은 CUDA 실행 모델이 SM을 중심으로 갈라져, 한쪽은 워프 기반 실행 효율로, 다른 한쪽은 블록 기반 협업 효율로 발전하는 구조를 보여준다.

👶 어린이를 위한 3줄 비유 설명

1. GPU 공장에서는 사람을 한 명씩 부르는 대신 32명씩 한 줄로 세워 한꺼번에 움직여요. 그 한 줄이 워프예요.
2. 그리고 여러 줄을 같은 방에 넣어 공구함을 같이 쓰게 만든 묶음이 블록이에요.
3. 줄 서는 방법과 방 배치가 잘 맞아야 공장이 쉬지 않고 빠르게 돌아가요.