프로그래밍 마스시브하게 병렬 처리기: 실습 중심 접근

📚 콘텐츠 요약

이 과정은 CUDA C 환경을 활용한 GPU 컴퓨팅과 병렬 프로그래밍에 대한 포괄적인 소개를 제공합니다. GPU 아키텍처, 데이터 병렬성, 스레드 관리, 메모리 최적화 및 고급 성능 고려 사항을 다루며, MRI 재구성 및 분자 시각화와 같은 실제 사례 연구를 통해 설명됩니다.

실용적이고 실습 중심의 가이드를 통해 고성능 병렬 컴퓨팅의 기술을 숙달하세요.

저자: 데이비드 B. 키르크, 웬메이 웨이 후

감사의 말: 아이언 버크, 존 닉올스, NVIDIA DevTech 팀, 젠슨 황, 데이비드 루브케, 빌 비안, 사이먼 그린, 마크 해리스, 맨주 헤지, 나데엠 모하메드, 브렌트 오스터, 피터 쉴러, 에릭 영, 시릴 제러.

🎯 학습 목표

1. 멀티코어 CPU와 멀티코어 GPU의 설계 철학과 성능 전개 방식을 구분합니다.
2. 현대적인 GPU 아키텍처의 핵심 구성 요소(예: 스트리밍 멀티프로세서(SM) 및 메모리 구조)를 식별합니다.
3. 암다할의 법칙을 적용하여 이론적 속도 향상을 계산하고 순차적 블로킹의 영향을 파악합니다.
4. 고정 기능 파이프라인과 프로그래머블 통합 프로세서 배열 간의 아키텍처 차이를 비교합니다.
5. GPGPU가 중간 단계로서의 역할과 초기 샷터 프로그래밍 모델의 제약 조건을 설명합니다.
6. 원자 연산, 장벽 동기화, 이중 정밀도 지원과 같은 하드웨어 기능이 확장 가능한 일반 목적 컴퓨팅으로의 전환을 가능하게 했음을 분석합니다.
7. 행렬-행렬 곱셈 알고리즘 내에서 데이터 병렬성을 식별하고 활용합니다.
8. 할당, 호스트와 디바이스 간의 데이터 전송, 할당 해제를 포함한 디바이스 메모리 관리를 구현합니다.
9. 적절한 스레드 인덱싱과 그리드/블록 구성으로 CUDA 커널을 생성하고 실행합니다.
10. 복잡한 데이터 구조를 GPU 하드웨어에 매핑하기 위해 다차원 스레드 계층 구조(그리드 및 블록)를 설계합니다.

🔹 수업 1: 병렬 컴퓨팅과 GPU 아키텍처 소개

개요: 본 수업은 CPU와 GPU의 서로 다른 설계 철학에 의해 촉발된 순차적 컴퓨팅에서 병렬 컴퓨팅으로의 근본적인 전환을 탐색합니다. 학생들은 "멀티코어"와 "멀티코어"의 경로를 비교하고, GPU가 거대한 처리량을 달성할 수 있도록 해주는 하드웨어 아키텍처를 이해하며, 암다할의 법칙을 통해 속도 향상의 수학적 제약 조건을 배웁니다.

학습 결과:

• 멀티코어 CPU와 멀티코어 GPU의 설계 철학과 성능 전개 방식을 구분합니다.
• 스트리밍 멀티프로세서(SM)와 메모리 구조를 포함한 현대적인 GPU 아키텍처의 핵심 구성 요소를 식별합니다.
• 암다할의 법칙을 적용하여 이론적 속도 향상을 계산하고 순차적 블로킹의 영향을 파악합니다.

🔹 수업 2: GPU 컴퓨팅의 진화와 미래

개요: 본 수업은 그래픽스 처리 장치(GPU)의 아키텍처 여정을 따라가며, 삼각형 렌더링을 위한 전용 고정 기능 하드웨어에서 출발해 현재의 강력하고 통합된 일반 목적 병렬 프로세서로 발전한 과정을 탐색합니다. 학생들은 엄격한 그래픽스 파이프라인에서 프로그래머블 샷터로의 전환, GPGPU 운동의 등장, 그리고 현재 과학 및 공학 시뮬레이션을 주도하는 현대적 확장 가능한 아키텍처를 탐구합니다.

학습 결과:

• 고정 기능 파이프라인과 프로그래머블 통합 프로세서 배열 간의 아키텍처 차이를 비교합니다.
• GPGPU가 중간 단계로서의 역할과 초기 샷터 프로그래밍 모델의 제약 조건을 설명합니다.
• 원자 연산, 장벽 동기화, 이중 정밀도 지원과 같은 하드웨어 기능이 확장 가능한 일반 목적 컴퓨팅으로의 전환을 가능하게 했음을 분석합니다.

🔹 수업 3: CUDA 프로그램 구조와 메모리 관리

개요: 본 수업은 CUDA 프로그램의 기본 아키텍처를 다루며, 호스트(CPU)와 디바이스(GPU) 실행 간의 차이점을 강조합니다. 학생들은 행렬 연산에서 데이터 병렬성을 식별하고, CUDA API를 사용하여 별도의 메모리 공간을 관리하며, 단일 프로그램, 다수 데이터(SPMD) 스타일을 통해 그리드, 블록, 스레드의 계층 구조를 통해 병렬 실행을 조직합니다.

학습 결과:

• 행렬-행렬 곱셈 알고리즘 내에서 데이터 병렬성을 식별하고 활용합니다.
• 할당, 호스트와 디바이스 간의 데이터 전송, 할당 해제를 포함한 디바이스 메모리 관리를 구현합니다.
• 적절한 스레드 인덱싱과 그리드/블록 구성으로 CUDA 커널을 생성하고 실행합니다.

🔹 수업 4: 고급 CUDA 스레딩 및 스케줄링

개요: 본 수업은 CUDA 내 스레드의 계층적 조직을 탐색하며, 다차원 인덱싱이 물리적 데이터 및 하드웨어 리소스에 어떻게 매핑되는지를 집중적으로 다룹니다. 장벽 동기화와 투명한 확장성의 메커니즘을 상세히 설명하고, 고성능 컴퓨팅에서 지연 시간을 견딜 수 있도록 하는 스레드 할당 및 워프 기반 스케줄링의 아키텍처 원리를 마무리합니다.

학습 결과:

• 복잡한 데이터 구조를 GPU 하드웨어에 매핑하기 위해 다차원 스레드 계층 구조(그리드 및 블록)를 설계합니다.
• 내장된 CUDA 변수(blockIdx, threadIdx, blockDim)를 사용하여 정밀한 데이터 인덱싱을 구현합니다.
• 장벽 동기화를 적용하여 다양한 GPU 아키텍처에서 데이터 무결성을 유지하면서도 투명한 확장성을 보장합니다.

🔹 수업 5: 메모리 최적화 및 공유 메모리 타일링

개요: 본 수업은 메모리 대역폭과 자원 제약이 병렬 컴퓨팅에서 주요 블로킹 요소임을 탐색합니다. 타일링을 통해 글로벌 메모리 트래픽을 줄이는 방법을 상세히 설명하고, 동기화 장벽(__syncthreads())의 중요한 역할과 레지스터와 공유 메모리 사이의 전략적 선택이 성능 최적화에 미치는 영향을 설명합니다.

학습 결과:

• 레지스터와 공유 메모리 한계가 커널 내 병렬성(점유율) 수준을 결정함을 분석합니다.
• 타일링 기법을 통해 글로벌 메모리 대역폭 소비가 얼마나 감소하는지 정량화합니다.
• 공유 메모리 접근 시 데이터 무결성을 유지하기 위해 동기화 함수가 필수적임을 식별합니다.

🔹 수업 6: 성능 분석 및 SIMT 실행

개요: 본 수업은 CUDA 커널 최적화를 위한 아키텍처적 및 알고리즘적 고려 사항을 탐색합니다. 단순한 실행 모델, 특히 단일 명령, 다수 스레드(SIMT) 유닛과 워프 분할에서 시작하여, 메모리 코일링, 타일링된 행렬 곱셈, 그리고 스트리밍 멀티프로세서(SM) 리소스의 동적 분할과 같은 고급 성능 조정 기술로 전환합니다.

학습 결과:

• 다차원 스레드 블록이 하드웨어의 선형 워프 실행 순서에 어떻게 매핑되는지를 분석합니다.
• 병렬 축소 알고리즘에서 제어 흐름의 분열을 평가하고 최소화합니다.
• 메모리 코일링 및 타일링된 데이터 접근 패턴을 구현하여 글로벌 메모리 대역폭을 최적화합니다.

🔹 수업 7: 부동소수점 산술과 수치 정확도

개요: 본 수업은 부동소수점 숫자의 기본 아키텍처를 다루며, 특히 IEEE 754 표준의 구성 요소인 부호, 초과 인코딩된 지수, 정규화된 맨티사에 초점을 맞춥니다. 학생들은 이러한 비트 패턴이 이산 수선 위에 어떻게 매핑되는지 탐색하고, 이 표현 방식의 제약이 대규모 합산과 같은 복잡한 알고리즘의 정확도에 미치는 영향을 살펴봅니다.

학습 결과:

• 정규화된 표현과 초과 인코딩을 사용하여 비트 패턴에서 수치 값을 계산하기 위해 부동소수점 형식을 분해합니다.
• 수직선 위에서 표현 가능한 수의 분포를 시각화하고, 지수와 맨티사에 할당된 비트 수의 영향을 설명합니다.
• ULP를 사용하여 수치 부정확도를 정량화하고, 다양한 반올림 모드가 오류에 기여하는 방식을 식별합니다.

🔹 수업 8: 사례 연구: MRI 재구성의 병렬화

개요: 본 수업은 고급 자기공명영상(MRI) 재구성의 GPU 병렬화를 탐색합니다. 비카르테시안 궤도에 대한 반복적 재구성 과정에 초점을 맞추며, F^H d 커널의 계산 비용이 매우 큰 부분을 차지하는 점을 최적화하기 위해 루프 변환, 상수 메모리 관리, 데이터 레이아웃 재구성, 하드웨어 가속 트리거 함수의 사용을 다룹니다.

학습 결과:

• 카르테시안 FFT 기반 재구성에서 비카르테시안 k-space 데이터를 위한 반복적 선형 해법 기반 알고리즘으로의 전환을 이해합니다.
• 루프 분할과 루프 교환을 적용하여 순차적 C 코드를 대규모 CUDA 스레드 매핑에 적합한 구조로 변환합니다.
• 상수 메모리 청크화와 구조의 배열(AoS) 데이터 레이아웃을 사용하여 메모리 대역폭을 최적화합니다.

🔹 수업 9: 사례 연구: 분자 시각화와 멀티-GPU 실행

개요: 본 수업은 GPU 컴퓨팅의 실용적 응용을 탐색하며, 전기적 잠재력 지도를 계산하기 위해 직접 쿨롱 합산(DCS) 방법을 사용하는 분자 시각화를 중심으로 합니다. 학생들은 기본 커널 구현에서 시작하여 명령 미리 풀기, 메모리 코일링, 패딩을 활용한 고도로 최적화된 버전으로 진전합니다.

학습 결과:

• CUDA 상수 메모리와 글로벌 메모리 지연 숨기기 기법을 사용하여 직접 쿨롱 합산(DCS) 커널을 구현합니다.
• 명령 미리 풀기 및 공통 좌표 계산의 재사용을 통해 커널 성능을 최적화합니다.
• 메모리 코일링 및 패딩 전략을 적용하여 GPU 글로벌 메모리 접근을 최대 대역폭으로 정렬합니다.

🔹 수업 10: 계산적 사고와 병렬 알고리즘 선택

개요: 본 수업은 순차적 사고에서 병렬 문제 해결으로의 전환을 탐색하며, 병렬 프로그래밍의 목표와 알고리즘의 전략적 선택에 초점을 맞춥니다. 학생들은 문제를 병렬화 가능한 단위로 분해하고, 도메인 과학과 하드웨어 아키텍처 사이의 격차를 다루기 위한 계산적 사고를 적용하며, 알고리즘 성능을 평가합니다.

학습 결과:

• 병렬 프로그래밍의 주요 목표를 식별하고, 암다할의 법칙을 사용하여 이론적 속도 향상을 계산합니다.
• 작업 수준과 데이터 수준의 분해를 구분하고, 원자 중심(산란)과 그리드 중심(수집) 전략을 적용합니다.
• 메모리 대역폭, 계산 복잡도, 아키텍처 제약 조건 등의 기준에 따라 병렬 알고리즘을 평가하고 선택합니다.

🔹 수업 11: OpenCL 프로그래밍 모델 소개

개요: 본 수업은 이종 병렬 컴퓨팅을 위한 프레임워크인 OpenCL을 소개하며, 데이터 병렬성 모델과 계층적 하드웨어 추상화에 초점을 맞춥니다. 학생들은 OpenCL의 NDRange와 메모리 구조를 CUDA에 해당하는 것과 매핑하고, 동적 컴파일 모델을 통해 호스트 측 장치 관리를 익힙니다.

학습 결과:

• OpenCL의 병렬성과 메모리 계층 구조를 CUDA 특화 아키텍처(예: 워크 그룹 → 블록, 로컬 메모리 → 공유 메모리)로 매핑합니다.
• OpenCL 커널 함수를 구현하고, 컨텍스트와 명령 큐를 사용하여 호스트 측 실행 환경을 관리합니다.
• 런타임에 소스 코드로부터 커널을 빌드하는 동적 컴파일 워크플로우를 실행합니다.

🔹 수업 12: 현대 GPU 기능과 미래 전망

개요: 본 수업은 GPU의 아키텍처적 및 기능적 진화를 탐색하며, 복잡한 메모리 관리, 향상된 커널 실행 능력, 증가된 코어 성능으로의 전환에 초점을 맞춥니다. 학생들은 통합 장치 메모리 공간과 커널 수준의 함수 호출과 같은 기능이 GPU를 일반 목적 프로세서로 전환하는 방식을 검토합니다.

학습 결과:

• 메모리 아키텍처 진화의 의미와 64비트 통합 장치 메모리 공간으로의 이동을 설명합니다.
• 강화된 원자 연산과 커널 수준의 함수 호출이 복잡한 데이터 구조와 알고리즘 구현을 가능하게 하는 방식을 분석합니다.
• 동시 커널 실행, 이중 정밀도 속도 향상, 제어 흐름 효율성의 성능 영향을 현대적인 GPU 환경에서 평가합니다.