[Multi-Core & GPU Programming] Intro to CUDA

"Intro to CUDA"

[CPU-GPU]

> CPU

· ~ 100 Threads (Core 수↓)

· 큰 Cache

· 제어 중심

→ Control-Heavy 작업에 적합 (if 구문 많은 Code)

> GPU

· ~ 10000 Threads (Core 수↑)

· 작은 Cache

· 병렬 처리 특화

→ Data-Parallel 작업에 적합 (배열 계산)

· CPU와 GPU는 메모리를 따로 가짐 ∴ Pointer를 넘긴다고 되는 것이 아니라, 데이터를 CPU 메모리에서 GPU 메모리로 복사 필요

· 병렬성이 큰 특정 Loop 또는 계산 부분만 Kernel로 떼어내는 방식

· GPU 계산이 아무리 빨라도 데이터를 자꾸 CPU↔GPU로 옮기면 전체 속도 저하

∴ CUDA 최적화의 핵심 = "연산량"보다 "데이터 이동량" 줄이기

① CPU가 Host Data 준비

② GPU의 Device Memory에 복사

③ Kernel 실행

④ 결과 필요 시 다시 Host로 복사

→ CUDA는 CPU 코드 안에 GPU 호출이 들어가는 구조

[CUDA]

※ CUDA : NVDIA GPU용 Programming Model (C/C++ 기반)

▣ Vector Addition

<C[i] = A[i] + B[i]>

· 각 원소 계산이 서로 독립이므로, Thread 하나가 원소 하나를 맡아 계산

· A_h / B_h / C_h : CPU 메모리 Pointer

· A_d / B_d / C_d : GPU 메모리 Pointer

· cudaMalloc : Device 메모리 할당

· cudaMemcpy( … HostToDevice) : 입력 A, B를 CPU에서 GPU로 복사

· vecAddKernel<<< … >>>( … ) : GPU에서 실제 연산을 수행할 함수(Kernel) 실행

· cudaMemcpy( … DeviceToHost) : 결과 C를 GPU에서 CPU로 복사

** Host Pointer / Device Pointer를 구분 → 입력 복사 → Kernel 실행 → 출력 복사

> cudaMalloc

· 일반 malloc과 비슷하지만, Pointer 자체를 수정해야 해서 Double Pointer를 받음

· 반환 값은 할당된 주소가 아니라 Error Code → 실제 결과는 인자로 전달하고 함수 반환 값은 성공/실패 상태

∴ cudaError_t cudaMalloc(void ** devPtr, size_t size)

· cudaError_t : 함수의 실행 결과를 알려주는 상태 Code (메모리 할당 성공 = 0 / 문제 발생으로 실패 = Error Code 반환)

· devPtr : 새롭게 할당된 GPU 메모리의 시작 주소를 담아올 Pointer

> cudaMemcpy

· 일반 memcpy와 다르게 복사 방향 명시 필요 (HostToHost / HostToDevice / DeviceToHost / DeviceToDevice)

∵ CPU 메모리와 GPU 메모리가 서로 다른 주소 공간

> Kernel_name<<< # of Thread Blocks, # of Threads per block >>>( )

· Kernel을 한 번 호출하는데, 실제로는 수많은 Thread가 동시에 같은 Kernel 함수를 실행

· SPMD(Single Program, Multiple Data) : 같은 Program을 여러 실행 주체가 각자 다른 데이터에 적용

** CUDA Thread는 CPU Thread와 같이 무겁고 비싼 객체가 아니라 훨씬 가벼움 → 수천 ~ 수만 개를 한 번에 띄우기 가능

▣ Kernel 함수

· __global__ : GPU에서 실행되는 Kernel 함수로, 해당 함수는 CPU(main)에서 호출하고, 실행은 GPU에서 수행

· int i = threadIdx.x + blockDim.x * blockIdx.x : 각 Thread가 전체 데이터 중 몇 번째 원소를 맡아야 하는지를 계산

· threadIdx.x : Block 안에서 해당 Thread의 번호

· blockIdx.x : Grid에서 현재 Block 번호

· blockDim.x : Block 하나에 들어있는 Thread 수

· .x : CUDA의 Index가 1차원만 있지 않고, 3차원(x, y, z)을 지원

· if (i < n) : Thread 수가 데이터 개수보다 많은 경우, 유효한 Thread만 계산하고 남는 Thread는 아무 일도 하지 않도록 막음

Ex. Block 크기 = 256

> threadIdx.x

· 1 Thread : threadIdx.x = 0

· 2 Thread : threadIdx.x = 1

· …

· Last Thread : threadIdx.x = 255

> blockIdx.x

· 1 Block : blockIdx.x = 0

· 2 Block : blockIdx.x = 1

· 3 Block : blockIdx.x = 2

> blockDim.x

· blockDim.x = 256 → Block 마다 Thread가 256개 존재

> i

· Block 0 : i = threadIdx.x + 256 x 0 → Thread 0 : i = 0 / Thread 1 : i = 1 / Thread 255 : i = 255

· Block 1 : i = threadIdx.x + 256 x 1 → Thread 0 : i = 256 / Thread 1 : i = 257 / Thread 255 : i = 511

· Block 2 : i = threadIdx.x + 256 x 2 → Thread 0 : i = 512 / Thread 1 : i = 513 / Thread 255 : i = 767

→ Block Offset + Block 내부 Offset으로 전역 Index 생성

** CUDA는 tid를 사용하는 병렬 Programming과 달리 threadidx와 blockidx를 이용해 i(Thread ID)를 직접 계산

→ tid와 i는 동일한 역할 (방법만 다름)

· Thread ⊂ Thread Block ⊂ Grid
· Thread Block의 크기는 일반적으로 HW 실행 단위인 "Warp"와 관련이 있어 32의 배수로 설정하는 것이 좋음

· 하나의 Thread Block이 가질 수 있는 최대 Thread 수는 1024개로 제한

· Grid의 최대 차원 크기 : x축 = 2³¹ - 1 / y축 & z축 = 2¹⁶ - 1 = 65,535
· Block의 최대 차원 크기 : x축 & y축 = 1024 / z축 = 64

▣ 다차원 구조체 활용

· dim3 구조체 : 다차원 Grid와 Block의 크기 지정 가능

· dim3 DimBlock(256, 1, 1) : x차원만 256인 1차원 Block을 의미

· Grid와 Block의 y, z차원을 1로 설정하여 사실상 x차원만 사용 → 앞서 ".x"로 정의했던 1차원용 vecAddKernel을 똑같이 사용 가능

▣ 1D Grid & 2D Block

▣ 2D Grid & 3D Block

· 차원이 높아질수록 Index 계산식 복잡

▣ CUDA Function Declarations

· __device__ : GPU에서 실행되며, 반드시 Device 내부에서만 호출 가능

· __global__ : Kernel 함수를 정의할 때 사용하며, GPU에서 실행되지만 호출은 CPU에서 수행 (** 반드시 반환형은 void)
· __host__ : 일반적인 C/C++ 함수처럼 CPU에서 실행하고 CPU에서 호출

** 필요에 따라 "__device__"와 "__host__"를 하나의 함수에 동시에 사용하여, 양쪽 모두에서 호출 가능한 코드를 생성 가능

▣ Image 처리

> Grayscale

<Color Image → Grayscale Image>

· 밝기(Luminance) = 0.21 x R + 0.71 x G + 0.07 B (Green 비중이 큰 이유는 사람이 밝기를 느낄 때 녹색 성분에 더 민감)

→ 단순 평균이 아니라 사람이 보는 밝기에 가까운 가중합 사용

· 각 Pixel은 서로 독립적으로 계산 가능

· 입력의 각 RGB Pixel에 동일한 Weight를 곱해서 출력 Gray Pixel 하나 생성 → Thread 하나당 Pixel 하나로 Mapping하기 좋음

<Covering 76x62 Image with 16x16 Block>

· 이미지 크기가 Block 크기의 배수가 아니므로 오른쪽/아래쪽에 남는 영역 발생

∴ if (Col < Width && Row < Height) 같은 경계 확인 필요

<Grayscale Kernel Code>

① Col, Row 계산

② 범위 Check

③ Grayscale Image의 1D Offset 계산

④ RGB Image에서 같은 Pixel의 시작 위치 계산

⑤ RGB 읽어서 가중합

⑥ 결과 저장

** 2D 좌표를 1D 배열 Index로 바꾸기 중요

· Gray는 Pixel당 1 Byte : Row x Width + Col

· RGB는 Pixel당 3 Channel : Gray Offset x Channels
→ 데이터 Layout을 이해해야 Kernel 사용 가능

> Blurring

<Image Blur>

· 출력 Pixel 하나를 만들기 위해 주변 여러 Pixel을 Read → 메모리 접근 패턴과 경계 처리가 "Grayscale"보다 더 중요

<Blurring>

· Blur Size = 1 → 중심 Pixel 주변 3x3 영역 평균을 구함 = 어떤 Pixel의 출력 값은 자기 자신뿐 아니라 주변 값까지 참조

<Blur Kernel Code>

① Col, Row 계산

② 주변 Box를 이중 for 구문으로 순회

③ 유효한 좌표만 누적

④ 평균 내서 출력

** 경계에선 참조 가능한 Pixel 수가 감소 ∴ Pixels를 세면서 평균을 구함

<Boundary Handling>

· 모서리나 가장자리에서는 주변 윈도우가 이미지 밖으로 나가므로 범위 안에 있는 Pixel만 평균내는 방식으로 구현

→ if (curRow > -1 && curRow < h && curCol > -1 && curCol < w)

▣ SM(Streaming Multi-Processor)

<NVIDIA GTX 1080>

※ SM : GPU 안에서 Thread Block과 Warp를 실제로 실행하는 핵심 단위

** Scheduling, Occupancy, Shared Memory 제한 등은 대부분 SM 기준으로 결정)

※ Warp : GPU HW가 Thread를 묶어서 관리하고 실행하는 기본 Scheduling 단위

· 32개의 Thread가 1 Warp

· 동일한 명령어 동시 실행 : 한 Warp 안에 속한 32개의 Thread는 제어 흐름을 공유하며, 모두 정확히 똑같은 명령어를 동시에 실행

· 각 Thread는 서로 다른 데이터에 접근하여 연산

→ SIMD(Single Instruction Multiple Data) / SIMT(Single Instruction Multiple Threads) 방식

> Instruction Scheduling & Control

· Warp Selector & Fetch/Decode : 하나의 SM 안에는 4개의 Warp Selector와 Fetch/Decode Unit 존재

· GPU는 Thread를 32개씩 묶어 Warp 단위로 실행

· SM은 Warp 0~63까지 최대 64개의 Warp(- 2,048개의 Thread)에 대한 실행 Context를 동시에 유지 가능)

> Compute Cores

· SIMD 기능 Unit(CUDA Core) : 32개의 Unit이 하나의 그룹으로 묶여 제어 신호를 공유하며 SIMD 방식으로 동일한 명령어 수행

· 각 Unit은 Clock당 1개의 곱셈-덧셈 연산 및 Load/Store 연산 수행 가능

· SFU(Special Function Unit) : sin/cos 등 연산 비용이 높은 특수 수학 함수를 전담하여 처리하는 Unit

> Memory Hierarchy

· Register : 실행 중인 Warp들의 데이터를 임시 보관하는 가장 빠른 메모리 (초대 64개의 Warp Context를 위해 총 256KB 할당)

· Shared Memory : 96KB 크기로, Programmer가 코드 상에서 직접 할당하고 관리할 수 있는 Cache 형태의 메모리

→ 같은 Block 내의 Thread들이 데이터를 빠르게 주고받을 때 사용

· L1 Cache : 48KB 크기로, CPU의 Cache처럼 HW가 자동으로 관리하는 고속 메모리 공간

<NVDIA GTX 1080 20 SMs>

· 20개의 SM 존재

** 1개의 SM는 1 Clock Cycle에 최대 4개의 Warp Selector가 4개의 Warp(128개의 Thread)만 사용하여 물리적인 연산을 수행

→ 나머지 최대 60개의 Warp는 Register에 자신의 상태와 변수를 그대로 유지한 채 대기

→ GPU는 이런 SM이 여러 개 존재 (GTX 1080 기준 20 SMs)

∴ 순간에 물리적으로 연산 중인 Thread는 20(SMs) x 4(Warps) x 32(Threads) = 2,560개 & 언제든 연산 Core에 투입될 수 있도록 Context Switching 비용 없이 꽉 찬 상태로 대기하는 Thread는 20(SMs) x 60(Warps) x 32(Threads) = 38,400개

→ GPU 칩 하나가 최대 40,960(= 2,560 + 38,400)개의 Thread를 동시에 유지(상주) 가능

▣ Thread Blocks & SM

<Thread Blocks & SM>

· SM 할당 : 하나의 Block 안에 있는 Thread들은 반드시 같은 SM에 할당되어 함께 실행

· 실행 순서 보장 불가 : 여러 개의 Block이 개념적으로는 함께 실행되지만, Block들 사이에는 어떠한 실행 순서도 보장 X

· 확장성 : 순서가 독립적 ∴ GPU가 좋으면 Code를 수정하지 않아도 여러 Block이 더 많은 SM에 나뉘어 들어가 자동으로 성능 향상

▣ Thread Scheduling

· 하나의 Warp에 속한 32개의 Thread들은 HW의 제어 Logic을 완벽하게 공유 → 32개의 Thread 모두 동시에 동일한 명령어 실행

∴ Programmer가 아무리 많은 Thread를 생성해도 HW는 결국 이를 32개씩 묶어 똑같은 명령 수행

· 만약 Programmer가 Block 크기를 32의 배수가 아닌 30의 배수로 설정

→ HW는 1개의 Warp를 할당하지만 실제로는 30개의 Thread만 일하고 나머지 2개의 CUDA Core는 IDLE 상태

· SM에 Thread가 256개인 Block 3개 할당

→ Block 1개당 Warp 수 = 256 / 32 = 8개의 Warp

→ SM 전체의 Warp 수 = 8(# of Warp) x 3(# of Block) = 24개의 Warp

※ Warp Selector : SM 내부에 상주하고 있는 여러 Warp 중 준비된 Warp를 찾아 연산 Unit에 전달하고 실행

※ Ready : 메모리에서 가져와야 할 데이터가 느린 외부 GPU 메모리(DRAM)가 아니라, 빠른 Register에 모두 도착해 당장 연산 가능한 상태

※ 동시 실행과 IDLE 상태 : GTX 1080의 SM에는 Warp Selector가 4 x 20개 있으므로, 매 Clock마다 최대 80개의 Warp가 동시에 실행 가능

→ 만약 Register에 데이터가 도착한 준비된 Warp가 단 하나도 없다면 SM은 IDLE 상태

> Zero-Overhead Warp Scheduling (Prevent GPU's IDLE State)

1) Context Switching 비용 = 0

· SM은 실행 대기 중인 모든 Warp의 정보(변수, 상태 등)를 거대한 Register 공간에 전부 올려둠

∴ CPU와 달리 A Warp에서 B Warp로 작업을 전환 시 기존 데이터를 저장하고 새 데이터를 불러오는 Context Switching 비용 = 0

2) 우선순위 기반 실행

· Warp Selector는 다음 명령어의 피연산자가 준비된 Warp들을 실행 후보로 올리고, 우선순위 Scheduling 정책에 따라 Warp를 선택하여 실행

① "TB1 W1"이 명령어 {1, 2, 3, 4, 5, 6}을 실행하다가 Stall

② GPU는 여기서 멈춰서 기다리지 않고, 즉각적으로 데이터가 준비되어 있는 "TB2 W1"을 선택해 명령어 {1, 2} 실행 후 Stall

③ 데이터가 준비되어 있는 "TB3 W1", "TB3 W2"의 명령어 {1, 2} 실행 후 "TB3 W2" Stall

④ 데이터가 준비된 "TB2 W1" 명령어 {3, 4} 실행

⑤ "TB1 W1", "TB1 W2", "TB1 W3" 실행

** Stall 원인

· 메모리에서 데이터 가져오기

· 연산 지연 : 이전 명령어의 계산 결과가 나와야 다음 명령어를 실행 가능 ∴ 해당 결과를 기다리는 동안 Warp는 Stall

· 실행 Unit 부족 : 특수 연산 또는 Load/Store 작업은 HW적으로 해당 Unit의 개수가 부족해서 이미 다른 Warp가 사용 중이라면 Stall

· 동기화 장벽 : "__syncthreads( )" 같은 동기화 함수를 만나면 하나의 TB에 속한 모든 Thread(모든 Warp)가 해당 Code Line에 도달할 때까지 먼저 도착한 Warp들은 Stall

→ GPU는 철저하게 주어진 데이터를 바탕으로 연산만 수행

∴ I/O interrupt 처리는 CPU의 몫으로 GPU Kernel이 실행되는 중간에 Stall X

** 1 Warp = 32 Threads & Warp 내에서는 동일한 명령어 처리

· 일반 연산 명령어가 들어오면 32개의 일반 연산 Unit 동시 실행

· Load/Store 명령어가 들어오면  8개의 Unit을 사용하여 4 Clock Cycle에 나누어 처리

→ 비어있는 일반 연산 Unit에 대기하고 있던 다른 Warp를 즉시 불러와서 덧셈/곱셈 명령어 수행

▣ Threads, Warps, Thread Blocks

· Thread Block들 사이에는 통신 불가 & 어떤 Block이 먼저 실행될지 순서 보장 X

· 같은 Block 안에 속한 Thread들은 Shared Memory를 통해 데이터를 통신

· "__syncthreads( )" 내장 함수 사용 시, Block 내의 모든 Thread가 특정 Code Line에 도달할 때까지 대기

** 1개의 Thread Block이 가질 수 있는 최대 Thread 수 = 1,024개 → 32 Warps

→ SM은 최대 64개의 Warps를 가질 수 있으므로 1,024개의 Thread를 가진 Block이라면 1개의 SM에 2개의 Block 배정 가능

· 1024개의 Thread로 구성된 Thread Block은 HW에 들어가면 32개의 Thread 단위인 Warp로 분리

** 하나의 Warp 안에 묶은 32개의 Thread는 무조건 모두 같은 명령어를 실행

** Warp들은 0번부터 순서대로 실행 X → 데이터가 준비된 Warp가 먼저 Warp Selector의 선택을 받음

▣ Branch Divergence

· Warp는 모든 Core가 똑같은 명령어를 수행해야 하지만, Branch는 if/else에 의해 서로 다른 Code 실행

> Prediction

· 서로 다른 구문을 실행해야 하는 경우, GPU는 두 구문을 모두 차례대로 실행하되, 해당하지 않는 Thread는 잠깐 Off(Masking)

· if/else 경로를 동시에 실행할 수 없으므로, 시간차를 두고 순차적으로 실행

→ 연산 Unit들이 100% 가동되지 못하고 놀게 되어 성능 효율 하락

> Avoiding Branch Divergence

· Thread ID로 분리 : 하나의 Warp 안에는 Thread가 순차적으로 들어있으므로 모든 Warp에서 무조건 분기 발산 발생

· Warp ID로 분리 : 같은 Warp에 속한 32개의 Thread는 모두 동일한 Warp ID → 동일한 True/False ∴ 분기 발산 발생 X

▣ Latency Hiding

· 각 Warp는 연산을 수행하는 상태와 메모리에서 데이터를 가져오느라 아무것도 못하고 기다리는 Stall 상태를 반복하다가 Done

① Warp 0이 데이터를 기다리며 Stall 상태에 빠지자마자, GPU는 멈추지 않고 즉시 Warp 1을 실행

② Warp 1이 Stall에 빠지면 Warp 2 실행

③ Warp 2가 Stall에 빠지면 Warp 3 실행

④ Warp 3이 Stall에 빠지면 데이터가 준비되어 실행 가능한 Warp 0 실행

⑤ 위 과정 반복

→ 개별 Warp가 겪는 Stall이 없는 것처럼 실행 시간으로 덮어짐

· 최신 GPU는 빠른 속도로 메모리와 데이터를 주고받을 수 있는 Pipeline 보유

→ 가능한 한 가장 많은 수의 Thread(Warp)를 SM 안에 꽉 채워 넣어 대역폭 100% 활용

(적은 Thread로는 메모리 Pipeline을 채울 수 없지만, 많은 Thread가 동시에 메모리 접근 요청 시 대역폭 한계치 사용 가능)

→ Thread(Warp)가 적은 경우 모든 Warp가 Stall에 빠졌을 때 교체해 줄 Warp가 없어 낭비 발생

▣ SM Limits

· SM 안에 항상 실행 가능한 Thread(Warp)를 꽉 채워 대역폭을 100% 활용해야 하는데 SM 한계 존재

· Max # of Threads : 2,048

· Max # of Warps : 64

· Max # of Thread Blocks : 32

· Registers : 65,536

· Shared Memory : 96KB

· L1 Cache : 48KB

Ex. 32 Threads per Block(Block당 Thread를 너무 적게 할당한 경우)

· Max # of Thread Blocks : 32

· Total Threads : 32 x 32 = 1,024

→ 활용도는 50%

Ex. 768 Threads per Block(Block당 Thread를 너무 많이 할당한 경우)

· 2 Blocks = 1,536 Threads / 3 Blocks = 2,304 Threads(2,048 초과)

· Total Threads = 1,536

→ 활용도 75%

Ex. 256 Threads per Block(Block당 Thread를 최적으로 할당한 경우)

· 256(Threads) x 8(Blocks) = 2,048 Threads

→ 활용도 100%

** 실제 프로그램에서는 Thread 개수와 Block 개수 제한뿐만 아니라, Shared Memory와 Register 사용량이라는 복잡한 변수 존재

** NVIDIA의 CUDA Occupancy Calculator에서 주황색 칸에 "Block당 Thread 수", "Thread당 Register 사용량", "Block당 Shared Memory 사용량"을 입력하면 파란색 칸에 현재 설정으로 낼 수 있는 실제 점유율을 자동으로 계산

** Code Compile 시 "--ptxas-options=-v"라는 옵션을 추가하면 해당 Kernel이 사용하는 "Thread당 Register 수"와 "Shared Memory 사용량"에 대한 수치가 출력됨

[Reference]

· 4_cuda_mgp_2026 (MGP) - Yongjun Park