CUDA. 4 Asynchronous Execution
Asynchronous Execution
What is Asynchronous Concurrent Execution?
Concurret Execution이 가능한 것들
- Host에서의 연산
- Device 연산
- Data transfer
- Host -> Device
- Device -> Host
- Device 내 메모리 간 전송
- Device간 전송
데이터 전송과 커널 연산을 overlap해 실행시키는 것과 같이, 여러 작업들을 overlap 시켜서 데이터 전송에 걸리는 오버헤드를 줄이거나 없애자
Asynchronous Interface
CUDA의 동시성은 asynchronous interface를 통해 구현된다.
- 즉시 return: GPU가 작업을 끝내기는 커녕 시작하기 전에도 함수가 즉시 return 됨 -> GPU가 일 끝내는걸 기다리지 않고 CPU가 다른 작업 가능
Synchronization
동기화 방법
- Blocking: 작업이 끝날때까지 CPU block 시켜
- Non-blocking/polling: 즉시 return 되면서 작업의 상태만 확인해
- Callback: 작업이 완료되면 자동으로 실행될 함수를 미리 등록해
CUDA Stream & Events, CUDA Graphs 로 구현 가능
CUDA Streams
GPU에게 시킬 일들은 순서대로 담아두는 queue (cudaMemcpy, kernel 실행들이 다 들어가)
- Stream 내부에서는 순서가 보장
- Stream 간의 순서는 Runtime이 GPU 자원 상황을 보고 알아서 scheduling 해, 자원 상태에 따라 stream이 overlap 될수도 있어
- CPU와 Stream의 관계는 Asynchronous: GPU가 일을 끝낼때까지 기다리지 않고 즉시 return 돼, 일을 다 끝낸는지 확인하고 싶으면
cudaStreamSynchronize(),cudaDeviceSynchronize()를 명시적으로 호출해야해, (stream level 동기화, device level 동기화)
- Default Stream: kernel이나 memcp 할때 stream을 지정하지 않으면 자동으로 default stream에 할당 됨, 얘는 동기화 규칙이 좀 달라
Creating and Destroying CUDA Streams
1
2
3
4
cudaStream_t stream;
cudaStreamCreate(&stream); // 생성
// ... 작업 수행 ...
cudaStreamDestroy(stream); // 소멸 (작업 완료 대기 후 삭제됨)
Stream 안에 처리 중인 작업이 남았있때, cudaStreamDestroy를 호출하면 오류가 나거나 작업이 최소되지 않아
GPU에 남아있는 작업을 모두 끝마친 뒤에 stream resource 해제.
단순한 작업이더라도 Stream 쓰는게 이득
memcp -> kernel 실행이라 할때, async하면 memcp 후에 cpu에서 kernel 실행 명령의 통신 latency가 없어져
Launching Kernels in CUDA Streams
Tripple chevron의 인자로 넣어주자 kernel<<<grid, block, shared_mem_size, stream>>>(...);
Launching Memory Transfers in CUDA Streams
// Copysize bytes from src to dst in stream stream cudaMemcpyAsync(dst, src, size, cudaMemcpyHostToDevice, stream);
cudaMemcpyAsync가 진짜로 비동기(Asynchronous)로 동작하려면, CPU 쪽 메모리(src)가 반드시 Pinned Memory (Page-locked)여야 한다.
- cudaMallocHost()로 할당한 메모리를 사용해야 합니다.
- 일반 malloc()으로 만든 메모리를 쓰면, 비동기 함수를 호출해도 동기적(Synchronous)으로 작동
Stream Synchronization
Asynchronous한 작업이 끝났는지 확인하는데 2가지 방법 존재
cudaStreamSynchronize(): CPU가 해당 스트림의 모든 작업이 끝날 때까지 멈춰서 기리기 (Blocking)cudaStreamQuery(): CPU가 멈추지 않고, GPU 작업이 다 끝났는지 상태만 확인 (Unblocking/Polling)
1
2
3
4
5
6
7
8
9
10
// Blocking
cudaStreamSynchronize(stream); // 여기서 딱 대기
// 이 줄부터는 스트림 작업 완료 보장됨
// Unblocking / Polling
if (cudaStreamQuery(stream) == cudaSuccess) {
// 완료되었을 때 처리
} else {
// 아직 안 끝났으니 다른 일 수행
}
CUDA Events
- 개념: Stream에 있는 작업들 사이에 끼워놓는 marker(tracer), 하나의 작업처럼 queue에 들어가
- 어떨때 쓰냐: Stream 단위, Device 단위의 Synchronization은 해당 stream의 모든 kernel이 완료되기를 기다려야 해
- 예시)
Stream: [Kernel 1] → [Event A] → [Kernel 2]kernel 1이 끝났는지 확인하려면 CPU는 Event A가 실행됐나만 확인 (cudaEventSynchronize)
- Profiling 할떄도 사용할수 있어: 두 이벤트가 기록된 시점의 타임스탬프 차이를 계산하여, 커널이 정확히 몇 ms 동안 실행됐는지 정밀하게 측정 (cudaEventElapsedTime)