[Multi-Core & GPU Programming] CUDA - DNN

"CUDA - DNN"

[DNN(Deep Neural Network)]

▣ 기본 단위 Neuron

· 하나의 Neuron : Input을 받아 Weight를 곱하고, Bias를 더한 후, 비선형 함수를 적용하여 계산하는 작은 회로

** 비선형성이 있어야 복잡한 Pattern을 학습 가능

▣ Building Block

· FC(Fully Connected) Layer : 모든 Input Node가 모든 출력 Node와 연결

· Sparsely(Locally) Connected Layer : 일부 Input만 Output에 연결되는 구조 (CNN의 Convolution Layer)

> FC(Fully Connected) Layer

· 학습 : Model이 표현하는 함수를 올바르게 바꾸기 위해 연결선에 부여된 W(Weight)를 최적의 값으로 조정하는 과정

▣ Image

[Convolution]

▣ 2D Convolution

▼

▼

※ Filter Mask : 이미지의 특징을 찾아내는 Pattern Detector (N개의 Filter 사용 시 N개의 출력 생성)

> Single Channel

· Input : W x H (흑백 이미지)
· Filter : N

→ Output : W x H x N

> Multi Channel

· Input : W x H x 3 (컬러 이미지 → RGB)

· Filter : 96

→ Output : W x H x 3 x 96

▣ Higher-Level Feature

· Layer를 깊게 쌓을수록, 앞쪽 Layer는 단순한 것을 추출하고, 뒤쪽 Layer는 고차원적이고 추상적인 특징을 추출

▣ Sparsely Connected Layer

· Convolution Layer는 출력 하나가 입력 전체가 아니라 작은 지역 영역만 영향 → Local Connectivity

· 같은 Layer에 속한 Unit들은 위치에 상관없이 같은 Filter를 사용 → Parameter Sharing

∴ FC Layer보다 훨씬 적은 Parameter로 이미지를 처리 가능

** Random Filter로 시작하여 Backpropagation을 통해 Filter 스스로 학습하여 최적의 Weight로 Update

▣ Convolution

<LeNet-5>

· 손글씨 숫자를 인식하기 위한 초기 CNN Model (INPUT : 32x32)
· C1, C3, C5 : Convolution Layer로 Feature Map 추출

· S2, S4 : Subsampling/Pooling Layer로 데이터 크기 감소

· F6, OUTPUT : FC Layer를 거쳐 최종적으로 10개의 숫자 카테고리 중 하나를 결정

> Convolution Layer의 Forward 연산 코드

· 출력 Y[n, m, h, w]를 계산하기 위해 Input Channel C 전체와 K x K Filter를 모두 곱해서 더함

(n : 이미지 번호 / m : Output Feature Map 번호 / c : Input Channel / h, w : 출력 위치 / p, q : Filter 내부 위치)

> CNN Dataflow

· Convolution Layer의 Forward 연산은 7개의 중첩 Loop (N → M → H → W → C → K₁ → K₂)
(N : Mini-Batch 크기 / M : Output Feature Map 수 / H, W : Input 높이, 너비 / C : Input Channel / K₁, K₂ : Filter 높이, 너비)

 - C = 3 / H = 4 / W = 4 / K = 3

· Input Channel : 3 / 각 Channel마다 4x4 입력 / 각 Channel마다 3x3 Filter

<X[0, 0, 0:3, 0:3] x W[0, 0, 0:3, 0:3] = 18 / X[0, 0, 0:3, 1:4] x W[0, 0, 0:3, 0:3] = 23>

▼

<X[0, 1, 0:3, 0:3] x W[0, 1, 0:3, 0:3] = 13 / X[0, 1, 0:3, 1:4] x W[0, 1, 0:3, 0:3] = 14>

▼

<X[0, 2, 0:3, 0:3] x W[0, 2, 0:3, 0:3] = 20 / X[0, 2, 0:3, 1:4] x W[0, 2, 0:3, 0:3] = 13>

▼

<Y[0, 0, 0, 0] = 51 / Y[0, 0, 0, 1] = 50>

▣ Parallelization

for (int n = 0; n < N; n++)        

        for (int m = 0; m < M; m++)

· 각 Thread가 Mini-Batch(N) 하나 또는 Output Channel(M) 하나를 담당

→ 병렬 처리할 수 있는 단위가 너무 적어서 GPU의 많은 Core를 충분히 활용 X

for (int h = 0; h < H_out; h++)        

        for (int w = 0; w < W_out; w++)

· 각 Thread가 출력 이미지의 개별 Pixel(H, W) 하나를 담당하여 계산

→ 뒤쪽 Layer로 갈수록 Feature Map 크기 H_out, W_out가 작아지면 Output Pixel 수가 줄어들어 병렬성 부족

<32x32 → 5x5>

for (int c = 0; c < C; c++)                

for (int p = 0; p < K; p++)

                for (int q = 0; q < K; q++)

· 각 Thread가 Input Channel(C) 하나 또는 Filter Weight(K₁, K₂) 하나를 담당

→ 여러 Thread가 같은 출력 위치 값에 동시에 더하므로 Atomic Operation 또는 Tree Reduction 필요

→ 병렬성은 늘릴 수 있지만 동기화 비용 발생

> Tiling

· 각 Thread Block이 Output Feature Map의 일정 구역(Tile)을 담당하도록 설계

· Grid : [Mini-Batch 내 이미지 수(N), Output Feature Map 수(M), Tile 번호]

· Thread Block : TILE_WIDTH x TILE_WIDTH 크기로 설정

> Basic Convolution CUDA Kernel Code

· 각 Thread는 하나의 출력값 Y[n, m, h, w]를 계산, 내부에서는 모든 Input Channel C와 K x K Filter를 돌면서 누적합 계산

· Thread 0 : Y[0, 0, 0, 0] 계산

· Thread 1 : Y[0, 0, 0, 1] 계산

· Thread 2 : Y[0, 0, 1, 0] 계산

· Thread 3 : Y[0, 0, 1, 1] 계산

** h = (blockIdx.z / W_grid) * TILE_WIDTH + threadIdx.y

** w = (blockIdx.z % W_grid) * TILE_WIDTH + threadIdx.x

> Shared Memory

· 인접한 Thread들은 Input 영역이 많이 겹치므로 같은 데이터를 중복해서 Read

∴ Block 단위로 필요한 Input Tile과 Filter를 Shared Memory에 Load → 여러 Thread가 재사용하여 Global Memory 접근 횟수↓

① Filter(W)를 Shared Memory에 Load → W_sh

② "__syncthreads( )"로 동기화

③ Input Tile(X)을 Shared Memory에 Load → X_sh

④ "__syncthreads( )"로 동기화

⑤ Shared Memory에서 Convolution 계산

⑥ "__syncthreads( )"로 동기화

⑦ Channel C에 대해 반복

** ② 동기화를 진행하지 않고, Input Tile(X)까지 Load한 이후에 ④ 동기화에서 한 번에 진행 가능

** Input Tile > Output Tile

[im2col]

· 복잡한 Convolution 연산을 단순한 행렬 곱 연산으로 변환    ∵ GPU : 행렬 곱셈(GEMM) 연산에 극도로 최적화

· Convolution = im2col + Matmul의 Pipeline으로 구성

▣ im2col Diagram

· Input Feature Map에서 Convolution에 사용될 작은 Patch들을 하나씩 펼쳐서 Column으로 변경

· Filter도 2D 형태가 아니라 1D Vector로 펼침

→ W' x X_unrolled = Y (Convolution 연산이 행렬 곱 연산으로 변경)

※ W' : 펼친 Filter 행렬 / X_unrolled : 입력 Patch들을 펼친 행렬 / Y : Output Feature Map

▣ im2col Kernel Code

· Input(X)을 Unrolled 행렬(X_col)로 변환하기 위해 출력 이미지의 해상도인 H_out과 W_out 계산

    → H_out & W_out : Output Feature Map 크기이자, X_col의 Column 개수

    → Convolution을 행렬 곱으로 변경하려면 Filter(W)만 펼치는 것 X

    → Input(X)의 각 K x K Window도 하나의 Column으로 펼쳐야 함

· filter_length : Filter 하나가 가지는 원소의 총 개수 (Ex. 3x3 Filter = 9)

    → Input Channel이 C개라면, 한 Output Pixel을 계산하는 데 필요한 입력 Patch의 전체 길이는 C x K x K

· C_base : 현재 연산 중인 Input Channel(c)에 따라 펼쳐진 행렬 세로축에서 몇 번째 Block부터 데이터를 써야 하는지 시작점 계산

· y : Filter 내부의 Local 2차원 좌표(p, q)를 펼쳐진 행렬의 최종 세로축 Index(y)로 변환

    → 현재 Thread가 담당할 Input Channel과 Filter 내 위치(p, q) 기반 Row Index 계산

· x : 원래 이미지에서 추출할 Window의 중심 위치를 펼쳐진 행렬의 가로축 Index(x)로 일렬로 나열

    → Output Feature Map에서의 위치 기반 Column Index 계산

∴ X_col[y, x] 위치에 원래 이미지의 Pixel 값을 배치하여 행렬 곱셈이 가능한 형태로 정렬

** 각 Column은 하나의 Convolution Window를 의미 & 각 Row는 해당 Window 안의 특정 Channel/Filter 위치를 의미

△ GPU에 매우 최적화된 GEMM(Matrix Multiplication) 연산 가능

▼ X_unrolled를 만들면서 Input 데이터가 중복 저장되기 때문에 메모리 사용량 증가

▣ Pooling

· CNN에서 데이터의 공간 크기를 축소 (위 코드는 K x K 영역의 평균을 내는 Average Pooling을 구현)

→ Input Feature Map(Y)에서 K x K 영역을 하나로 줄여서 출력 S를 생성

▣ Convolution 연산

> Direct Convolution

· Padding과 Stride를 고려한 다중 중첩 Loop를 사용하여 Input Window와 Filter를 직접 곱하고 더하는 방식

→ 메모리 낭비는 없지만 GPU 내부 연산 장치를 100% 활용 X

> GEMM vi im2col

· im2col을 이용해 Convolution을 행렬 곱 형태로 변경 후, Matrix Multiplication으로 계산

→ 데이터가 중복 저장되는 메모리 손해가 존재하지만 GPU 내부 연산 장치 활용↑

> FFT-Based Convolution

· 고속 Fourier 변환을 이용해 Convolution을 Frequency Domain에서 Point-Wise Multiplication으로 빠르게 계산 후 역변환

> Winograd Transform

· 수치해석적으로 연산식을 재조합(중간 계산을 재배치)하여, HW 장치에서 가장 연산 비용이 비싼 곱셈의 횟수를 감소

→ 덧셈 횟수↑

[Reference]

· 7_cuda_dnn_mgp_2026 (MGP) - Yongjun Park