Scientific Computing을 위한 CUDA 사용법 -7

Chapter 7

Parallel Reduction

1. Problem of Parallel Computation & Solution

Reduction Problem

• 교환, 결합 법칙이 허용되는 연산을 벡터 성분에 적용
  ex. sum, max, product, etc...
• 데이터 레이싱: 병렬 연산 중 연산이 적용되지 않은 값을 읽거나 하는 문제로 연산 결과가 누적되지 않는 현상

Divide and Conquer

• N개의 스레드 -> N등분 (Divide)
• 각 스레드가 부분 계산 & 부분 결과 이용 최종 계산 (Conquer)

2. Pairwise Implementation

Iterative Pairwise Implementation: 2 types

• Neighbored: 인접한 데이터 2개씩 묶기
• Interleaved: 반갈 후 동일한 인덱스의 데이터 묶기

Serial Code

• recursive function으로 구현됨
• Interleaved와 비슷한 형태

CUDA Code

• Pairing은 같은 블록 내에서 처리
  다른 블록 접근 시 latency 발생

3. Example

Neighbored

• stride가 1부터 2배씩 커지며 Neighbor부터 순차적으로 취합됨
• __syncthreads() 사용으로 데이터 레이싱 방지

Interleaved

• stride가 blockDim/2에서 반절씩 줄어들며 같은 인덱스를 가진 데이터끼리 취합됨
• __syncthreads() 사용으로 데이터 레이싱 방지

4. Neighbored vs Interleaved

Interleaved Recommended!

Neighbored 방식의 stride

• stride가 두배씩 커지는 형태 -> 32를 넘어가게 되면 하나의 Warp 내에서의 처리가 아닌 2개의 Warp 필요
• warp 내에 비활성화 스레드 생김 -> Warp Divergence

Interleaved 방식의 stride

• stride가 반절씩 줄어드는 형태
• iteration을 돌아도 계속 한 Warp 내에서의 연산
• 99%에 상응하는 Occupancy 달성 가능