여러가지 하면서 느낀건데 GPU 에 대해서는 언젠가 아주 자세하게 알아야 할 필요가 있다고 뼈저리게 느꼈다. 따라서 이번 기회를 통해 GPGPU 에 대해서 깊게 이해할 뿐만 아니라 , 직접 GPGPU 코드도 작성해보면서 이를 더 깊게 이해해보려고 한다. ( 문법이 꽤나 C++ / C 와 다른 점이 많다. )
AD102 를 기준으로 GPU 구조를 설명한다
AD102 한 장 요약 (칩 구조)
- 12 GPC (Graphics Processing Clusters)
- GPC당 6 TPC → 총 72 TPC
- TPC당 2 SM → 총 144 SM
- 메모리 인터페이스: 384-bit (12×32-bit 컨트롤러)
보기 쉽게 정리하면 다음과 같이 정리된다.
GPU (AD102 전체)
└── GPC × 12 개 (Graphics Processing Cluster)
└── TPC × 6 개 (Texture Processing Cluster)
└── SM × 2 개 (Streaming Multiprocessor)
└── Processing Block × 4 개
└── CUDA Core (FP32) × 32 개
- GPU = 공장 전체
- GPC = 공장 동
- TPC = 작업 라인 묶음
- SM = 실제 작업장
- warp = 32명으로 묶인 작업팀
- register = 각 작업자의 개인 책상
- shared memory = 팀이 같이 쓰는 화이트보드
- L2 cache = 공장 전체가 같이 쓰는 중앙 창고
먼저 하나씩 대략적으로 본 뒤 wrap + SM 쪽 구조를 중요하게 확인해야한다 ( 젤 중요 )
GPC
GPC 란 ?
- GPU 전체를 여러 구역으로 나눈 최상단 단위 이다. ( AD102 에는 총 12개가 존재한다 )
- 각 GPC 는 독립적으로 돌아간다.
GPC 안에는 무엇이 있나 ?
- 6 TPC : 각 TPC 마다 ( SM x 2 / PolyMorph Engine x 1 )
- 1개의 Raster Engine : 3D 그래픽스에서 삼각형을 픽셀로 변환하는 고정 기능 유닛 ( 그래픽스에서 쓰는 )
- 2개의 rop 파티션 ( 각 파이션에 8개 ROP 유닛 ) : 최종 픽셀을 메모리(프레임 버퍼) 에 기록
왜 이런 묶음 (GPC) 를 구성했는가 ?
그래픽 파이프라인은 연산(SM) 만 빠르다고 끝이 아니라
- 기하 처리(PolyMorph),
- 래스터( Raster ),
- 픽셀 출력(ROP)
이 세 단계가 같이 돌아야 프레임이 뽑힌다.
그래서 " SM 묶음 + 그래픽 전용 블록 " 을 한 덩어리 (GPC) 로 묶어서
-> 여러 GPC 가 병렬로 프레임 처리하도록 한다.
( GPC 는 그래픽 파이프라인 스케일링 단위로 보면 정확 )
TPC
GPC 안에 더 작은 단위
- SM 2 개
- 1 개의 PolyMorph Engine ( 3D 그래픽스에서 vertex 처리 담당 )
그리고 가장 중요한 SM 구조에 대해서는 매우 자세히 살펴보도록 하겠다.
SM ( Streaming Multiprocessor )
SM 은 GPU 에서 실제 연산이 일어나는 핵심 단위이다.
( SM = “연산 유닛 몇 개”가 아니라, 스케줄러 + 레지스터 + 캐시 + load/store + 특수 연산 유닛까지 다 갖춘 작은 병렬 실행 엔진 )
하나의 SM 에는 다음과 같은 것이 담겨져 있다.
- 4개의 processing block
- 128KB L1 Cache / Shared Memory
- 4 Texture Units ( 그래픽스에서 이미지 샘플링 빠르게 하려고 만든 것 )
- 1개 3rd Gen RT core ( Ray Tracking 전용 가속기 )
- 2개 FP64 units ( 64bit float 연산용 유닛 )
( 교수님이 중요하다고 짚어주거나 언급한 부분은 따로 또 부연설명을 하겠다 )
L1 Cache / Shared Memory ( 128 kb )
- SRAM 하나의 풀을 L1 Cache / Shared Memory 로 나눠서 사용하는 형태이다.
- 둘다 SM 의 on-chip Memory 라서 빠르다.
( SM 기준으로 VRAM 은 off-chip 이다. )
L1 Cache
- 하드웨어가 알아서 관리
Shared Memory
- __shared__ 로 명시적으로 배열 만들고 / thread 들이 global 에서 가져온 데이터들을 shared 에 올려놓는다.
- 같은 block 의 thread 끼리 공유해서 재사용 하는 테크닉.
cf .그럼 왜이렇게 설계했는가 ?
둘 다 sm 근처의 빠른 SRAM 을 필요로 하는데 워크로드에 따라서 요구가 다르다.
- 행렬곱/컨볼루션/타일링 -> shared 크게 주면 이득 ( 딥러닝에서 자주 쓰이는. )
Processing block
Processing block = “warp를 뽑아서(issue) 연산/메모리를 처리하는 미니-SM”
하나의 SM 을 4개의 Processing Block 으로 나눌 수 있다. 각 블록 역시 독립적으로 명령어 스케줄링하고 실행할 수 있다.
또한 Processing Block 안에는 다음과 같은 것들이 존재한다.
- 16 FP32 CUDA cores + 16 FP32/INT32 CUDA cores ( 128 CUDA cores per SM )
- L0 instruction cache
- 1 warp scheduler + 1 dispatch unit
- 64KB register file (16,384 × 32-bit)
- 1 SFU ( 나눗셈 / 사인&코사인 계산이 원클락에 안되니까 그걸 잘 하기 위한 추가 special function )
- 4 Load/Store units
- 1 Ada 4th Gen Tensor Core
(A) L0 cache
- instruction ( 명령어 ) 미리 가져다 놓는 캐시
- 매번 메모리에서 가져오면 느리니까 Processing block 에 직접 달아버림.
- SM 내부에 있다.
(B) CUDA core ( 16 FP32 CUDA cores + 16 FP32/INT32 CUDA cores ( 128 CUDA cores per SM ) )
CUDA core : FP 32 부동소수점 연산 1개를 1클럭에 처리하는 ALU 1개
- ALU : 산술 연산 하는 하드웨어 회로
- CPU 안에도 ALU 가 있고 , GPU 의 CUDA core 도 ALU 다. 다만 GPU 용으로 단순화 되어 있는 대신에 양이 매우매우 많다. ( 16384 개있음 ) ( = 32개 X 4 ( processing block ) x 2 ( Sm ) x 6 (TPC) x 12 ( GPC ))
- 32 x 4 여서 128 개 있다고 하는거다 ( SM 안에 )
두종류의 CUDA core 가 존재한다.
- INT 32 도 지원하는 연산 회로는 FP 32 만 하는 회로보다 트랜지스터가 많이 필요하다.
fp32 전용 core -> 단순회로 , 면적 작음
fp32/int 32 -> 복잡회로 , 면적 큼
여기서 32 라는 숫자가 매우 중요하다. (어찌나 여러번 말하시던지 외워버림)
뒤에서 더 설명하겠지만 warp 라는 개념이 나오는데 이것과 관련이 있다.
- 일단 GPU 는 warp ( 32 threads) 를 기본 실행 단위로 SIMT 방식으로 처리한다.
- thread 1개 단위가 아니라 warp 단위로 스케줄/실행된다. ( 중요 )
1. warp1개 = 32 threads
2. 32 thread 가 SM 안의 32 개 CUDA Core 에 1:1 mapping 되어서 작동한다.
3. 동일한 명령어를 32개가 동시에 1클럭에 실행한다 ( SIMT )
( warp 는 같은 instruction 을 같이 실행하는 단위이다 - 다른 instruction 을 수행하는게 아님. )
( 뭐 조건이 복잡하면 1클럭은 아닐 수 있지만 편의상 그렇게 설명한 듯 싶다. )
문제
" 이렇게 core(레지스터) 가 많아도 부족하다 "
( 이건 뭐 차차 알아가도록 하고.. )
(C) Tensor Core
- 오른쪽에 맨 초록색 영역 보이는가 ? 거기가 텐서 코어다.
- CUDA core (ALU) 와 별개로, GPU 안에 있는 행렬곱 전용 유닛이다.
- D = A x B + C 기본형태는 이렇다 ( 딥러닝에서 매우매우 자주 쓰이는 모양 )
cf.
RTX 4090 성능 비교:
CUDA Core (FP32): 82.6 TFLOPS
Tensor Core (FP16): 330 TFLOPS ← 4배
Tensor Core (FP8): 1,321 TOPS ← 16배
(D) WRAP
가장 중요한 파트라고 생각해서 기존 설명뿐만 아니라 조금 더 내용을 추가해 두려고 한다.
1. warp 란 무엇인가 ?
- warp = 32 개의 thread 를 묶은 실행 단위
- SM 이 다루는 기본 실행 단위이다. ( 즉 , 직접 다루는 단위가 thread 하나가 아니라 warp 이다. )
- GPU 는 내부적으로 32 개의 Thread 를 한 덩어리로 처리하겠다는 전제로 설계
SM은 4개의 Processing block 으로 나뉘고 , 각 processing block 에는 warp scheduler 가 한개씩있다
( SM 안에는 총 4개의 warp sceheduler )
-> SM 하나가 통쨰로 한번에 warp 하나만 처리하는게 아니라 . 여러 warp 를 resident 상태로 들고 있으면서 , 각 scheuler 가 매 순간 실행 가능한 warp 를 골라 instruction 을 issue 한다.
SM
├─ Processing Block 0 ─ Warp Scheduler 0 ──> 어떤 warp 하나 선택
├─ Processing Block 1 ─ Warp Scheduler 1 ──> 어떤 warp 하나 선택
├─ Processing Block 2 ─ Warp Scheduler 2 ──> 어떤 warp 하나 선택
└─ Processing Block 3 ─ Warp Scheduler 3 ──> 어떤 warp 하나 선택issue 시점마다 ( 매 cycle ? ) 각 scheduler 는 " 지금 당장 다음 instruction을 실행할 수 있는 warp가 누구냐? " 를 보고 골라서 보낸다.
2. resident / active /ready warp
- 이부분이 강의를 들으면서 좀 헷갈렸던 부분이다. 아마 편의상 그냥 warp 로 설명하시고 넘어 가신 것 같은데 듣는 입장에서는 약간 의아했던 부분이 있어서 정리를 좀 해보려고 한다 (내 집중력이 안좋았을수도 있고..) . 일단 먼저 warp 를 3가지로 구분하면 이해가 쉽다.
A) resident warp
- 현재 SM 안에 올라와 있는 warp
- block 이 SM 에 배정되면 그 block 의 warp 들이 SM 의 자원
- register
- shared memory
- thread slots
- warp slots
을 차지하면서 resident 상태가 된다.
B) active warp
보통 resident 이면서 아직 실행이 끝나지 않은 warp 를 말한다.
( 사실상 resident == active 비슷하다고 보면 된다 )
C) ready warp
바로 다음 instruction 을 실행할 수 있는 warp
어떤 warp 가 global memory load 결과를 기다리는 중이면 resident 이긴 하지만 ready 는 아닐 수 있다.
(맞나 이게?)
그니까 1개의 SM 에 최대 48개의 warp 를 가질 수 있고 실행할 수 있다는건
보통 생각하는 것 처럼 그냥 하나의 warp 가 cuda core 에 붙어서 연산이 끝날때 까지 쭉 계산하는 형태가 아니라
스위칭을 통해서 엄청 자주 바뀌면서 돌아가는 것 같다 ( 마치 시프에서 배웠던 context switiching 을 통해서 하면 쭉 연산하는 것 처럼 보이게 하는 착각을 일으키는 ? )
위 내용이 맞는 것 같아서 다시 정리해보면.
GPU에서 warp execution은 CPU처럼 “한 스레드가 코어를 오래 점유”하는 느낌보다
여러 warp가 SM 안에 대기하고 있고, scheduler가 그때그때 ready한 warp를 빠르게 골라 instruction을 발행하는 구조
즉 어떤 warp가:
- arithmetic instruction 몇 개 수행
- global memory load 만나서 stall
- 잠시 대기
- 그동안 다른 warp 수행
- 나중에 다시 돌아와 이어서 수행
처럼 돌아간다고 한다.
또 , SM 에 최대 48warps 가 있다고 이야기 한건 정확히는
"한 SM 이 최대 48개의 resident / active warp 를 동시에 보유할 수 있음" 이다.
이 말이 곧 48개가 모두 동시에 한 사이클에 실행된다는 뜻이 아니라. SM 안에 최대 48warp 의 문맥을 돌려두고
그중 ready 한 warp 들을 스케줄러들이 선택해서 issue 하는 느낌이다.
( 그니까 보유 capacity 와 issue 된 것들과는 다른 것이고 , 이를 다 합쳐서 48개라고 하는 뜻. )
하지만 엄연히 CPU context swithcing 과는 다른데, 이 이유를 zero-cost context swithcing 에서 알아보도록 하겠다.
2. Warp processing
- SM 에서는 여러 warp 가 독립적으로 처리된다.
- SM 은 이 warp 들을 여러개 동시에 올려놓고 관리한다 ( 즉 , sm 이 한번에 warp 하나만 들고 있는 것이 아니라 , 여러 warp 를 resident 상태를 유지하면서 그중 실행 가능한 warp 를 골라서 처리한다. )
execution context 란 ?
각 warp의 32 개 thread 에 대해 다음과 같은 실행 상태가 SM 내부 ( on-chip ) 에 유지된다.
- ( program counters, registers, and so on ) 이런 정보를 통틀어서 execution context 라고 볼 수 있다.
어떤 warp 가 도중에 멈춰도 warp 상태가 이미 SM 내부에 저장되어 있어서 나중에 다시 이어서 실행할 수 있다.
* 추가설명 ( ai 와 질문하면서 필요한 부분 그대로 발췌 )
SM 이 warp 를 들고 있는 형식이 정확히 어떤 의미인지 이해가 잘 안갔다. ( 아래 zero-cost context switching 에서도 또 나오는 개념이므로 )
그니까 뭐 당연히 어떠한 정보들을 유지하는 형식일 것 같기는한데 더 깊게 알아보고 싶다.
1. SM 이 warp 를 들고있다.
= SM 안에 resident한 warp들의 실행 상태(execution context) 를 유지할 수 있는 하드웨어 자원들이 있다는 뜻이다.
상태에는 대표적으로:
- 다음에 실행할 instruction 위치
- 각 thread의 register 값들
- warp의 active mask / 실행 상태
- 스케줄링 대상 여부
같은 정보가 포함된다. 강의자료와 CUDA Programming Guide도 warp의 execution context로 program counters, registers, etc. 를 들고, 이것이 warp lifetime 동안 on-chip에 유지된다고 설명한다.
즉,
- warp = 32개 thread의 실행 단위
- SM = 그 warp들을 실행시키는 하드웨어
- SM이 warp를 들고 있다 = 그 32개 thread가 지금까지 어디까지 실행했고 어떤 값을 갖고 있는지 SM 내부 자원에 저장해 둔 상태
라는 뜻이다.
2. 그럼 SM은 뭘 실제로 저장하나?
이걸 감각적으로 보면, resident warp 하나당 SM은 대충 이런 정보를 관리한다고 생각하면 된다.
warp 단위로 필요한 것
- 이 warp가 다음에 실행할 instruction 위치
- 이 warp가 현재 ready인지 / stall인지
- 현재 instruction에 참여하는 active thread mask
- 스케줄러가 참고할 warp 상태
thread 단위로 필요한 것
warp 안에는 32개 thread가 있으니까, 각 thread마다:
- 자기 register 값들
- 자기 thread-local execution state
- (Volta 이후에는 더 독립적인 per-thread state)
같은 것이 필요하다.
강의자료는 “every warp currently resident in an SM”에 대해 각 32 threads의 execution context 가 on-chip에 유지된다고 하고 CUDA Programming Guide는 SM의 32-bit registers가 warps 사이에 partition 되며, shared memory는 thread blocks 사이에 partition된다고 설명한다.
즉 물리적으로는:
- register file
- warp scheduling state
- PC / mask / control state
- shared memory allocation info
- 기타 pipeline bookkeeping
같은 하드웨어 자원들에 resident warp들의 상태가 배치되어 있다고 보면 된다.
register file이 핵심이다
이 부분이 가장 중요하다.
강의자료 1에서 Ada SM의 각 processing block에는 64 KB register file: 16,384 x 32-bit 가 있다고 나온다.
또 CUDA Programming Guide는 register가 SM에 위치하며 thread-local storage로 쓰인다고 설명한다.
이 말의 뜻은:
warp 안 32개 thread 각각이 자기 register들을 갖고 있고, 그 register 값들이 SM 내부 register file에 잡혀 있다는 뜻이다.
예를 들어 어떤 thread가 kernel 실행 중에:
float x = ...
float y = x * 2;
이런 중간값들을 만들면, 이런 값들 중 많은 부분은 register에 들어간다.
warp 하나에는 32 threads가 있으니,
실제로는 32개 thread 각각의 register 묶음이 필요하다.
즉 resident warp가 많아질수록:
- 많은 thread의 register 상태를 동시에 들고 있어야 하므로
- register file 용량이 occupancy를 제한하게 된다.
그래서 CUDA PG도 resident block/warp 수가 kernel의 register 사용량과 shared memory 사용량에 따라 달라진다고 설명한다.
“그럼 warp는 어디 있다?”를 더 물리적으로 말하면
이건 아주 거칠게 말하면 다음처럼 보면 된다.
warp 그 자체
- thread ID 32개 묶음
- scheduler가 한 단위로 다루는 엔트리
warp의 상태는 어디?
- register file 일부
- warp state table / scheduler bookkeeping
- PC / mask / control state
- shared memory allocation과 연관된 block 상태
즉 “warp가 저장되는 하나의 통 박스” 가 있다기보다,
SM 내부 여러 자원에 그 warp를 재개하는 데 필요한 상태가 분산되어 있다고 이해하는 쪽이 가깝다.
질문 1
warp를 그러면 뭐 어떻게 들고있는데?
답:
resident warp를 다시 실행할 수 있게 해 주는 상태 정보들을 SM 내부 하드웨어 자원에 유지한다는 뜻이다.
특히:
- program counter
- active mask / warp state
- 32개 thread 각각의 register 값
- scheduling 상태
등이다.
질문 2
진짜로 thread들은 32개가 있을꺼고.. warp 에 대한 정보를 뭐 어떻게 들고있다는건데 SM 이 ?
답:
warp는 32개 thread의 묶음이고,
SM은 그 thread들의 register state 와 warp의 control state 를 들고 있다.
특히 register 값은 SM의 register file 에 있고,
warp scheduler가 참고하는 상태 정보도 on-chip에 있다. Ada SM의 processing block마다 register file과 warp scheduler가 있다고 강의자료가 설명한다.
즉:
- thread 값들 → register file
- warp가 어디까지 왔는지 / 지금 실행 가능한지 → warp control/scheduling state
- block 간 공유 데이터 → shared memory
이런 식으로 역할이 나뉜다고 보면 된다.
3. Zero-cost context switching
앞에서 언급한 execution context 와 이어지는 내용이다.
이해하기 쉽게 CPU 에서는 thread 를 바꿀때는 보통
- 현재 상태 저장 ( PC 저장 / Register 저장 / stack pointer 저장 etc ... )
- 다른 thread 상태 복원 ( pc 불러오고 / register 불러오고 ... )
이런식으로 저장/복원 오버헤드가 상대적으로 크다.
하지만 GPU 에서는 warp 들의 execution context 가 이미 SM on-chip 에 유지되고 있기 때문에
어떤 warp 가 stall 되면 다른 warp 로 바꾸는데 별도의 무거운 과정이 필요 하지 않는다.
( 이미 정보들이 SM 위에 on chip 으로 있기 때문에 switching cost 가 거의 없다. )
이로 인해 GPU는 warp를 번갈아 실행하면서 memory latency를 숨기고 throughput을 높일 수 있다.
-> 이걸 zero-cost context switching 이라고 부른다.
GPU는 메모리 latency 를 어떻게 hiding 하는가 ?
한 warp 가 memory access 떄문에 멈추면 , SM 이 다른 ready warp 를 바로 실행해서 기다리는 시간을 가린다.
GPU 는 latency 를 없애는 것이 아니라 다른 warp 일을 시켜서 기다리는 시간이 안보이게 만드는 것이고
이를 " latency hiding " 이라고 한다.
( 예시 ) GPU의 latency hiding
- 어떤 warp가 global memory load를 수행하면, 데이터가 올 때까지 오래 기다릴 수 있다.
- 이때 그 warp는 stall 되어 당장 다음 instruction을 실행할 수 없다.
- 하지만 SM 안에는 여러 resident warp가 이미 올라와 있다.
- warp scheduler는 stall된 warp 대신 다른 ready warp를 골라 instruction을 issue한다.
- 따라서 메모리 응답을 기다리는 동안에도 ALU는 다른 warp의 연산을 계속 수행할 수 있다.
- 결과적으로 긴 memory latency가 다른 warp 실행에 의해 가려진다(hide 된다).
많은 warp 를 띄워두는게 좋다.
경우 A: warp가 별로 없음
예를 들어 SM 안에 ready warp가 2개밖에 없다고 하자.
- W0 실행
- W0 memory access → stall
- W1 실행
- W1도 memory access → stall
- 이제 실행할 ready warp가 없음
- 그러면 SM의 연산 유닛이 놀게 됨
즉 latency를 숨길 수 없다.
경우 B: warp가 많음
예를 들어 SM 안에 warp가 32개, 48개 가까이 있다고 하자.
- W0 stall
- 그럼 W1
- W1 stall
- 그럼 W2
- W2 stall
- 그럼 W3
- ...
- 이런 식으로 ready warp가 계속 있으면 연산 유닛이 쉬지 않음
즉 memory 기다리는 시간을 다른 warp들로 메꿀 수 있다.
- 이상 첫번째 GPGPU 수업 정리를 마친다.
일단 이정도만 알아도 성공적인 첫주인듯.
'GPGPU' 카테고리의 다른 글
| GPGPU 총정리 - (3) (0) | 2026.03.25 |
|---|---|
| GPGPU 총정리 - (2) (0) | 2026.03.21 |
