[Multi-Core & GPU Programming] Optimizing Parallel Reduction in CUDA

"Optimizing Parallel Reduction in CUDA"

[Parallel Reduction]

· 각 Thread Block은 Tree 기반의 접근 방식을 사용하여 배열의 일부를 Reduction한 후 Partial Result 생성

· 큰 배열을 처리하려면 여러 Thread Block 필요

→ Thread Block들 사이의 Partial Result들을 어떻게 통신?

▣ Global Synchronization

· 모든 Thread Block 간에 동기화가 가능하다면 거대한 배열도 쉽게 Reduction 가능

· But, CUDA는 HW 비용과 Core 효율성 문제(Deadlock)로 인해 전역 동기화 지원 X

  (한 Block 안에서는 "__syncthreads( )"로 동기화 가능 / 여러 Block 전체를 한 번에 동기화 불가)

→ 여러 개의 Kernel로 분할(Decomposition)하여 Kernel Launch 자체를 전역 동기화 지점으로 활용

** Kernel : GPU용 작업 단위로, GPU에서 실행되는 함수를 의미

▣ Kernel Decomposition

· 한 번의 Kernel에서 모든 Reduction을 끝내려 하지 않고, 여러 Kernel을 반복 호출

① 1번째 Kernel(Level 0)에서 여러 Block이 각각 Partial Sum 생성

② 2번째 Kernel(Level 1)에서 해당 Partial Sum들을 다시 Reduction

<반복적인 Kernel 실행 구조>

① Global Memory의 큰 배열을 여러 Thread Block이 나누어 Read

② 각 Block 안에서 Shared Memory를 사용해 Reduction

③ 각 Block의 결과가 다시 Global Memory에 저장

④ 다음 Kernel Launch에서 해당 결과들을 다시 Read하여 Reduction 진행

⑤ 결과가 하나 남을 때까지 반복

[Optimize]

· Reduction 연산은 Memory에서 값을 하나 Read하여 1번의 연산만 수행 → 연산량보다 Memory 접근량이 중요

∴ 연산 속도(GFLOP/s)보다는 Memory Bandwidth 최적화 중요

▣ Kernel 1 - Interleaved Addressing(교차 주소 지정) Code

① 각 Thread가 Global Memory에서 하나의 값을 읽어 Shared Memory(sdata)에 저장

② "s"를 증가시키면서 Thread별로 교차하여 값을 더하는 연산 수행

③ Thread 0번이 Block의 Reduction 결과를 Global Memory에 기록

> Kernel 1 동작 방식

① Stride 1 : 바로 옆(1칸 떨어진) 값과 덧셈 연산 진행 (Thread 0, 2, 4, 6, 8, 10, 12, 14만 동작)

② Stride 2 : 2칸 떨어진 값과 덧셈 연산 진행 (Thread 0, 4, 8, 12만 동작)

③ Stride 4 : 4칸 떨어진 값과 덧셈 연산 진행 (Thread 0, 8만 동작)

④ Stride 8 : 8칸 떨어진 값과 덧셈 연산 진행 (Thread 0만 동작)

→ Thread들이 띄엄띄엄 동작하기 때문에 Warp 안에서 일부 Thread만 동작하고 일부는 쉼

> 문제점

· 같은 Warp 안의 Thread들이 모두 같은 명령을 실행해야 효율↑

· But, "if (tid % (2 * s) == 0)" 때문에 일부 Thread만 실행(Warp Divergence↑) & "%" 연산은 GPU에서 상대적으로 느림

▣ Kernel 2 - 분기 및 % 연산 제거

▼

· Thread Id 자체를 "%"로 검사하지 않고, 각 Thread가 담당할 Shared Memory Index를 계산

> Kernel 2 동작 방식

· Thread 0, 1, 2, 3, …이 연속적으로 동작하여 연산에 참여

> 문제점

· Shared Memory 위치가 (Stride 2 = 0, 2, 4, 6, …), (Stride 3 = 0, 4, 8, …)처럼 간격 증가

→ 여러 Thread가 같은 Memory Bank에 동시에 접근하면서 Shared Memory 접근 속도↓ = Bank Conflict

** Kernel 1 : 띄엄띄엄 있는 Thread가 띄엄띄엄 있는 주소에 접근 → Bank Conflict 영향↓

** Kernel 2 : 연속된 Thread가 띄엄띄엄 있는 주소에 접근 = 같은 Warp 안에서 Bank 겹치기 쉬움 → Bank Conflict 영향↑

▣ Kernel 3 - Sequential Addressing

▼

· "s = blockDim.x / 2"로 Stride가 증가(1 → 2 → 4 → 8)하지 않고 감소(8 → 4 → 2 → 1)

> Kernel 3 동작 방식

· 앞쪽 절반의 Thread들이 뒤쪽 절반의 값과 덧셈 연산

→ 역방향 구조를 취하여 Thread가 연속적인 메모리 공간에 접근하도록 하여 Bank Conflict 방지

> 문제점

· 첫 반복이 "s = blockDim.x / 2"이므로 Thread의 절반만 동작, 절반은 IDLE 상태 → GPU 자원 낭비

▣ Kernel 4 - First Add During Load

▼

· Global Memory에서 Shared Memory로 Load할 때 1번째 덧셈을 같이 수행

  (기존 : 각 Thread가 하나만 Load / 개선 : 각 Thread가 2개를 Load하며 동시에 바로 덧셈)

→ Block 수를 절반으로 줄이고, Shared Memory Reduction의 첫 단계를 Global Load 과정에서 처리 가능

> 문제점

· Kernel 4까지 최적화했지만 이론적 Bandwidth에는 많이 못 미침

· Reduction은 연산량이 적고 Memory 접근이 중요한 작업

→ 이 정도 성능이면 Memory 자체가 병목이라기 보다는 명령어 Overhead가 병목일 가능성 존재

** 명령어 Overhead : 실제 연산에 쓰이는 명령어가 아닌, 주소 계산, 반복문 조건 검사, Iteration 증가 같은 부가 명령어

· Reduction이 진행될수록 Active Thread 수 감소

· s ≤ 32가 되면 Active Thread들이 하나의 Warp 안에 들어감

→ 하나의 Warp 내부 Thread들은 기본적으로 같은 명령을 동기화되어 실행되므로, 이 구간에서는 "__syncthreads( )" 필요 X

∴ 마지막 단계를 반복문 없이 직접 코드로 펼칠 수 있음

▣ Kernel 5 - Unroll the Last Warp

· "s > 32" 조건을 활용하여 하나의 Warp만 남기 전까지 Loop 진행

· "s ≤ 32"를 만족하는 경우 Unrolling
→ 반복문 Overhead(Iteration 증가, 조건 검사)와 불필요한 동기화↓

** 마지막 Warp뿐만 아니라 Compile 시점에 전체 반복 횟수를 알 수 있다면 코드를 완전히 Unrolling 가능

· GPU의 Block 크기는 512 Thread로 한정되어 있고, 보통 2의 거듭 제곱 형태를 사용

→ 고정된 크기에 대해서 코드를 Unrolling 가능

→ 다만, 사용자가 지정하는 Block 크기에 유연하게 대처하기 위해 C++의 Template 기능을 활용

<C++ Template>

· Block 크기를 Template Parameter로 지정하여 선언하는 구조

→ Compiler가 "blockSize >= 512" 같은 조건을 미리 판단하고 필요 없는 코드 제거

▣ Kernel 6 - Completely Unrolled

· "s = 256, 128, 64" 단계를 반복문 없이 Unrolling (다만, "s > 32" 구간에서는 여러 Warp가 참여 → 동기화 필요)

· Block Size ≥ 512 : 앞의 256개 Thread가 뒤의 256개의 값과 덧셈 연산

· Block Size ≥ 256 : 앞의 128개 Thread가 뒤의 128개의 값과 덧셈 연산

· Block Size ≥ 128 : 앞의 64개 Thread가 뒤의 64개의 값과 덧셈 연산

· Block Size ≤ 64 : Active Thread 수 = 32개 → 하나의 Warp

** 빨간색 조건들은 Compile 시점에 평가 → Runtime에는 Overhead X

· Template을 사용하여 가능한 Block Size 별 Kernel 버전을 미리 만들어두고, 실행 시 Switch 구문으로 선택

> Complexity

· Reduction은 매 단계마다 원소 수가 절반으로 감소 → Step Complexity = O(log N) (Big-O 표기법에서는 밑 생략)

· 전체 덧셈 횟수 = (N - 1)번 → Work Complexity = O(N)

    → 순차 Reduction과 같은 수의 덧셈 수행 → Resource 낭비가 없는 Work-Efficient

· 실제 병렬 처리에서는 P개의 Processor 존재 → Time Complexity = O(N/P + log N)

    → 한 Thread Block 내부에서 Thread 수가 충분히 많다면 Reduction 단계는 O(logN)으로 볼 수 있음

> Cost

· Cost = (# of Processor) x (Time Complexity)

· 데이터 개수 N개와 동일한 Thread N개를 사용 & Reduction 시간이 O(log N)

    → Cost = O(N log N) (순차 Algorithm = O(N))

    → Cost-Inefficient

· Brent's Theorem : Thread 수를 O(N / log N)개로 줄이고, 각 Thread가 O(log N)만큼의 순차적인 덧셈 작업을 먼저 수행

    → Cost = O((N / log N) x log N) = O(N)

    → Cost-Efficient

> Algorithm Cascading

① 각 Thread가 Global Memory에서 여러 원소를 Load하고 더함

② 더한 결과를 Shared Memory에 저장

③ Shared Memory 안에서 Tree 기반 Reduction 수행

→ 순차 Reduction(담당 원소들끼리 덧셈 연산)과 병렬 Reduction(Partial Sum끼리 더해서 Reduction 연산) 방식을 결합

** 완전 병렬만 사용 시

    → Thread 수↑

    → Synchronization↑

    → Kernel Launch↑

    → Instruction Overhead↑

** Thread 하나의 작업량 증가 시(순차 처리↑)

    → Thread 수↓

    → Synchronization↓

    → Reduction 단계↓

    → Memory Latency Hiding↑

▣ Kernel 7 - Multiple Adds / Threads

▼

· 각 Thread가 2개만 Load & Sum하는 것에서 각 Thread가 While Loop를 통해 여러 원소를 Load & Sum하는 것으로 변경

→ 하나의 Thread가 배열의 여러 위치를 맡아서 부분합을 만든 후 Shared Memory Reduction에 참여

** i += gridSize : Memory Coalescing을 유지하기 위함

    → Stride를 설정해야 같은 Warp 내의 Thread들이 Global Memory에 접근할 때 연속된 주소 공간을 Read

▣ Performance for 4M Element Reduction

▣ Final Optimized Kernel

· Template으로 Block Size 고정

· 각 Thread가 여러 원소를 더함

· Shared Memory에 Partial Sum 저장

· Block 내부 Reduction 전체 Unrolling

· 마지막 Warp Unrolling

· 마지막에 Thread 0이 Block 결과를 Global Memory에 저장

▣ Types of Optimization

> Algorithmic Optimizations

· 주소 접근 방식 변경

· Algorithm Cascading

> Code Optimizations

· Loop Unrolling

[Conclusion]

▣ Optimizing CUDA Performance

· Memory Coalescing

· Divergent Branching

· Shared Memory Bank Conflict

· Latency Hiding

· 병목이 Memory인지, 연산인지, Instruction Overhead인지 구분하기

· Parallel Algorithm Complexity 이해하기

· Algorithm을 먼저 최적화하고, 그 다음 Loop Unrolling 적용하기

· Template Parameter로 Compiler가 최적 Code를 만들도록 하기

→ GPU에서는 단순히 병렬화하는 것만으로 충분하지 않고, Memory Access Pattern과 Thread 동작 방식을 잘 설계해야 함

[Reference]

· 8_cuda_reduction_mgp_2026 (MGP) - Yongjun Park