Lec 21 · GPGPU (범용 GPU 컴퓨팅)과 CUDA programming model

1.1GPGPU 의 배경: graphics 의 FLOP·bandwidth 요구와 throughput 지향 설계, 초기 GPGPU 에서 CUDA·OpenCL 로.

1.2CPU vs GPU: latency 최적화 vs massive TLP·throughput, warp 기반 SIMT, fine-grained multithreading 으로 latency hiding.

1.3CUDA 모델: kernel·grid·thread block·thread 계층, block 독립성과 atomic, memory space(register/local/shared/global/constant/texture), cudaMemcpy.

1.4마지막에 C for CUDA 문법(__global__, <<< >>>, threadIdx 등). 슬라이드 32–41(vector_add 예제, branch divergence, dynamic warp formation)은 다음 시간.

상황·문제병렬 아키텍처 두 번째 파트로 GPGPU 를 다룸.

원인GPU 는 게임·3D modeling 용이지만 연산 능력이 매우 큼.

해결GPGPU(general-purpose computing on GPUs) = 병렬성이 충분한 일반 workload 를 GPU 로 계산.

상황·문제단일 GPU 가 실시간 rendering 하는 NVIDIA 데모 영상을 재생.

원인Zoom 으로는 해상도·부드러움이 떨어짐.

해결강의 노트의 링크로 직접 보라고 안내. 화면 뒤의 대량 연산을 생각해 보라는 취지.

상황·문제full HD 화면은 약 2 million pixel, 50–100 Hz refresh, vertex 당 10–100 연산, fragment 당 더 많은 연산.

원인이 연산은 CPU core 의 add·multiply·compare 와 본질적으로 같은 숫자 계산이며, data 공급도 필요.

해결GPU 는 높은 FLOP 과 memory hierarchy 전반의 높은 bandwidth 를 갖추게 됨 → 다른 계산에도 활용 가능.

상황·문제graphics workload 는 task·data 수준의 massive parallelism 을 가짐.

원인사람은 millisecond 단위로 인지, 하드웨어는 nanosecond 단위 → 6 orders of magnitude gap.

해결개별 연산 latency 보다 ms 안에 끝내는 총 연산량(throughput)을 최대화하도록 GPU 를 설계.

상황·문제vertex 에서 triangle 을 거쳐 framebuffer(GPU DRAM)에 그림을 얻는 pipeline.

원인오랫동안 쓰인 graphics workload 의 공통 추상화.

해결초기 GPU 는 이 연산만을 위해 만들어졌고, 성능이 커지며 general purpose 용도가 검토됨.

상황·문제초기 GPU 는 graphics 연산만 이해해 general purpose 계산에 비친화적.

원인알고리즘을 graphics API 연산으로 다시 표현해야 해 번거롭고 usability 가 제한됨.

해결그래도 physics simulation 등에서 좋은 성능 → CUDA 개발로 이어짐. 1저자 Ian Buck 은 NVIDIA 에서 CUDA 개발에 참여.

상황·문제GPU 계산을 위한 programming language 가 필요.

원인CUDA 는 C/C++ 에 extension 을 더한 NVIDIA 전용 환경이라 AMD GPU 등에서 못 씀.

해결CUDA: thread 를 대량 생성해 각자 작은 계산, host(CPU)-GPU data 이동·sync·memory 관리는 programmer 가 명시. 대안으로 OpenCL(Apple 시작, Khronos 유지), 최근 AMD ROCm.

상황·문제GFLOPS 기준 고성능 GPU 가 CPU 보다 throughput 이 훨씬 큼 (2019년까지 그래프).

원인throughput 이 전부는 아니어서 CPU 에 맞는 task 와 GPU 에 맞는 task 가 따로 있음.

해결두 processor 에 workload 를 나눠 매핑. 딥러닝 성공 3요소 = data, 새 ML 기법, GPU. Summit 에서 GPU 가 peak 의 95% 이상(tensor core FP16 포함 99.7%).

상황·문제CPU 와 GPU 는 서로 다른 목적에 최적화됨.

원인CPU 는 single thread latency 최소화(out-of-order, 큰 cache, 복잡한 control). GPU 는 thread 당 latency 보다 massive TLP 로 throughput.

해결GPU 는 수많은 thread 를 hardware 로 penalty 없이 전환해 latency 를 숨김(register 를 core 에 유지). serial 부분은 CPU, parallel 부분은 GPU 로 heterogeneous 실행.

상황·문제CPU 는 ALU 를 늘릴수록 활용도를 위해 control logic 도 복잡해짐.

원인GPU 는 같은 instruction 을 여러 thread 가 실행 → control 하나로 여러 ALU 구동.

해결warp(예: 32 thread)=hardware 가 만든 사실상의 SIMD 연산. die area 를 compute 에 투자. A100 SM 에 tensor core(FMA 전용, 256 FMA/cycle).

상황·문제warp = 같은 instruction 을 다른 data 에 실행하는 thread 집합, common PC 공유.

원인hardware scheduler 가 준비된 warp 중 1~2개를 SIMD pipeline 에 보냄. programmer 는 vector 가 아니라 thread 로 작성(SPMD).

해결SIMT: thread 를 hardware 가 warp 로 묶어 실행. SIMD 대비 thread 별 다른 처리(branch)가 쉬움, warp 를 유연하게 구성.

상황·문제cache miss 난 warp 는 data 가 올 때까지 block 됨. GPU 는 cache hit rate 가 낮고 GDDR·on-chip 경로 latency 가 김.

원인모든 thread 의 register 가 register file 에 있어 warp 전환에 비용이 없음(fine-grained multithreading).

해결warp 가 충분하면 다른 warp 실행으로 memory latency 를 숨겨 compute 활용도 유지. 대신 GPU 는 bandwidth 집약적이고 per-thread cache 는 작음.

상황·문제많은 thread 를 동시에 실행하려면 실행 계층이 필요.

원인warp 는 hardware 개념이라 programmer 에게 노출되지 않음.

해결kernel(GPU 함수) → grid → thread block → thread. block 은 SM 하나에, thread 는 SP 하나에 매핑. launch 때 grid·block 차원 지정, threadIdx·blockIdx 로 data 를 나눔.

상황·문제block·thread 의 실행 순서는 알 수 없음.

원인block 간 synchronization(공유 lock)은 deadlock 을 쉽게 유발.

해결block 간 coordination 은 shared variable + atomic operation(shared queue pointer 는 OK)으로. atomic 남용은 serialization 병목.

상황·문제GPU 는 register 와 memory 뿐인 CPU 보다 memory space 가 복잡.

원인register·local(thread 전용), shared(block 내, programmer 관리 on-chip scratchpad), global(모든 thread), constant·texture(GPU 는 read only).

해결host 는 global·constant·texture 만 접근. CPU↔GPU data 는 cudaMemcpy 로 명시적으로 복사(GPU 0/1 지정 가능).

상황·문제질문: thread block 하나를 SM 두 개에 나눠 실행할 수 있나 → 불가, block 하나는 SM 하나에만 매핑. 이후 출석 확인.

상황·문제CUDA 는 C/C++ 에 최소한의 extension 만 추가.

원인host(CPU)와 device(GPU) 코드, memory 위치, launch 구성을 구분해야 함.

해결__global__(CPU 가 호출, GPU 실행), __device__(GPU 내부 호출), __constant__/__shared__, <<<grid_dim, block_dim>>> 로 launch, __syncthreads() 는 block 내 sync 만.

graphics 용 GPU 로 병렬성이 충분한 일반 계산을 수행.

NVIDIA 의 C/C++ 기반 GPU programming 언어·환경. NVIDIA 전용.

Apple 이 시작, Khronos 가 유지하는 royalty-free 병렬 표준. CUDA 와 모델 유사.

GPU 는 개별 연산 latency 보다 ms 안에 끝내는 총 연산량을 최대화.

serial 부분은 CPU, parallel 부분은 GPU 로 나눠 실행.

GPU core 묶음. thread block 하나는 SM 하나에만 매핑.

matrix FMA 전용 unit. A100 에서 256 FMA/cycle, SM 당 4개.

같은 instruction 을 실행하는 thread 집합(32), common PC. OpenCL 은 wavefront.

thread 로 programming(SPMD), hardware 가 warp 로 묶어 SIMD 처럼 실행.

register 를 register file 에 유지해 warp 를 cycle 단위로 바꿔 latency 를 숨김.

programmer 가 지정하는 thread 묶음과 그 묶음. block 은 서로 독립이어야 함.

block 간 shared variable 갱신용 (atomicAdd 등). 과용 시 serialization.

block 내 thread 가 공유하는 on-chip scratchpad, programmer 가 직접 관리.

host(CPU) memory 와 device(GPU) memory 사이 명시적 복사 함수.

4.1GPU 가 개별 연산 latency 를 덜 중시해도 되는 이유는?

4.2GPU 의 context switch 가 사실상 무료인 이유는?

4.3warp 와 thread block 의 차이는?

4.4block 간 lock 을 쓰면 안 되는 이유와 대안은?

4.5host 가 접근할 수 있는 GPU memory space 는?

4.6dim3 grid_dim(100, 50), block_dim(4, 8, 8) 의 block·thread 수는?