AI 시스템 성능 엔지니어링 (Korean Edition)

제10장. 커널 내부 파이프라이닝, 워프 전문화 및 협동 스레드 블록 클러스터

이 작품은 AI를 사용하여 번역되었습니다. 여러분의 피드백과 의견을 환영합니다: [e-mail]

이전 장에서는 메모리 액세스 조정, 병렬성 극대화, 계산과 데이터 전송 중첩, 점유율 향상, 워프 정지 최소화 등 기본적인 최적화 기법을 다루었습니다. 이러한 기법들은 지연 시간을 숨기고 병목 현상을 제거하는 데 도움이 되었습니다. 그러나 현대 GPU는 기본적인 최적화 기법을 한 단계 더 발전시킬 수 있는 고급 하드웨어 기능과 실행 모델을 제공합니다.

이 장에서는 워프 특화 파이프라인, 그리드 수준 및 클러스터 수준 동기화를 지원하는 협력 그룹, 동적 작업 큐를 반복하는 지속적 커널, 분산 공유 메모리(DSMEM 또는 DSM) 및 텐서 메모리 가속기(TMA) 멀티캐스트를 사용하는 스레드 블록 클러스터( 협동 스레드 배열 클러스터 [CTA]라고도 함)와 같은 보다 고급 CUDA 기법을 소개합니다. 스레드 블록 클러스터는 높은 수준에서 동시 실행이 보장되는 스레드 블록 그룹입니다. 이들은 DSMEM을 사용하여 서로의 공유 메모리를 읽고, 쓰고, 원자 연산을 수행할 수 있습니다.

이러한 방법을 통해 호스트 개입 없이 메모리 액세스와 연산 작업을 중첩할 수 있습니다. 또한 스레드 블록 간에 온칩 데이터 공유가 가능하며, 모든 SM을 완전히 활용할 수 있습니다.

이러한 현대적인 GPU 실행 모델을 이해하면, CUDA 스트림을 활용한 커널 간 파이프라인 탐구를 통해 이러한 최적화를 더욱 확장하는 다음 장으로 나아갈 준비가 될 것입니다. 다음 장에서는 본 장에서 논의된 커널 내 최적화를 기반으로 커널 간 파이프라인을 구축합니다.

커널 내 파이프라인 기법

커널 내 파이프라이닝 ( )은 단일 커널 실행 내에서 메모리 작업과 연산을 중첩시키는 일련의 기법을 의미합니다. (다음 장에서는 서로 다른 스트림에서 실행되는 여러 커널 간 작업을 중첩시키는 커널 간 파이프라이닝을 살펴볼 것입니다.)

핵심 아이디어는 커널을 동시 실행 단계로 구조화하여, 한 데이터 조각이 로드되거나 저장되는 동안 이전에 로드된 데이터가 처리되도록 하는 것입니다. 이러한 단계들은 서로 다른 타일 또는 데이터 청크에 대해 병렬로 작동합니다. 이는 처리량을 향상시키고 지연 시간을 효율적으로 숨깁니다.

기존 GPU는 지연 시간을 숨기기 위해 워프 수준 멀티스레딩에 의존했습니다. 한 워프가 메모리 로드에서 정지하는 동안 다른 워프들은 계산을 진행합니다. 이는 실행 모델에서 단일 명령 다중 스레드(SIMT) 지연 시간 숨김의 기초입니다.

커널 내 파이프라이닝은 동일한 워프 또는 커널 내에서 메모리 작업과 연산을 중첩시켜 이를 한 단계 발전시킵니다. 세밀한 조정을 통해 메모리 로드와 연산을 시간차를 두고 수행하며, 때로는 단일 워프 내에서 이를 구현합니다.

CUDA 파이프라인 API를 활용한 커널 내 파이프라이닝은 __syncthreads() 없이 비동기 메모리 전송과 계산을 중첩합니다. ...