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

제11장. 커널 간 파이프라이닝, 동기화 및 CUDA 스트림 순서 메모리 할당

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

지금까지 우리는 단일 커널에 대해 SM을 바쁘게 유지하기 위한 커널 내부 도구들—cuda::pipeline, 더블 버퍼링, 워프 전문화(로더/컴퓨트/스토어 워프), 지속적 커널, DSMEM/TMA를 사용한 스레드 블록 클러스터—에 집중해 왔습니다. 이 장에서는 이러한 커널들을 유지한 채 CUDA 스트림, 이벤트, 스트림 순서 메모리 할당기를 사용하여 커널과 배치 간에 파이프라인을 구성하는 방법을 보여줍니다. 간단히 말해, 제10장은 커널 내 지연 시간 숨기기에 초점을 맞췄습니다. 본 장에서는 커널 간 및 GPU와 호스트 간 지연 시간을 숨기는 방법을 보여줍니다.

이러한 커널 간 동시성은 실제 워크로드에서 GPU의 모든 엔진을 바쁘게 유지하는 데 필수적입니다. 현대 GPU에서 최대 GPU 활용도를 달성하려면 GPU의 컴퓨팅 엔진과 직접 메모리 액세스(DMA) 엔진을 병렬로 바쁘게 가동해야 합니다.

CUDA 스트림은 이러한 커널 간 동시성의 기반을 제공합니다. 비동기 메모리 작업, 세밀한 동기화, CUDA 그래프(본 장에서 간략히 소개하고 다음 장에서 자세히 다룸)를 결합함으로써 호스트 측 정체를 피하는 고효율 파이프라인을 구축할 수 있습니다.

CUDA 스트림을 통한 커널 실행 중첩

CUDA 스트림은 커널 실행( ), 메모리 복사, 메모리 할당 등의 작업 순서로 실행되는 일련의 커널 실행 작업입니다. 그림 11-1과 같이 CPU에서 GPU로 2개의 스트림을 사용하여 5개의 커널을 실행한다고 가정해 보겠습니다.

그림 11-1. CPU에서 GPU 상의 두 스트림으로 다섯 개의 커널을런칭하는 과정

ker_B 여기서 커널 1( ker_1, ker_2, ker_3 )은 스트림 2에서 실행되는 커널 2( ker_A )와 병렬 실행됩니다. 하드웨어 자원이 허용하는 한 모든 커널은 서로 및 CUDA 스트림 간에 중첩될 수 있습니다.

스트림이 커널 작업을 비동기적으로 수행하는 동안 CPU는 작업을 계속 수행할 수 있습니다(cpu_code_1 와 cpu_code_2) ). 두 CUDA 스트림에서 이 다섯 개의 커널을 실행하는 코드는 다음과 같습니다:

#include <cstdio>
#include <cuda_runtime.h>

__global__ void ker_A()  { /* ... do some work ... */ }
__global__ void ker_B()  { /* ... do some work ... */ }

__global__ void ker_1()  { /* ... do some work ... */ }
__global__ void ker_2()