스트림과 비동기 실행

학습 목표

• •CUDA Stream의 개념을 이해한다
• •비동기 데이터 전송과 커널 실행을 배운다
• •여러 스트림을 사용한 파이프라인을 구현한다
• •동기화 기법을 익힌다

CUDA의 동기 vs 비동기 실행

기본 동작 (동기 실행)

cuda
복사
// 기본적으로 순차적 실행
cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);  // 완료까지 대기
kernel<<<grid, block>>>(d_a);                         // 실행 시작
cudaMemcpy(h_a, d_a, size, cudaMemcpyDeviceToHost);  // 완료까지 대기

시간 흐름 (동기)

비동기 실행의 이점

비동기 실행:
- CPU는 GPU 작업을 기다리지 않고 다른 작업 수행
- 여러 GPU 작업을 동시에 실행 (overlap)
- 전체 실행 시간 단축

CUDA Stream

Stream은 순서대로 실행되는 작업 큐입니다.

Stream의 특징

Default Stream (Stream 0)

cuda
복사
// 스트림을 지정하지 않으면 기본 스트림 사용
kernel<<<grid, block>>>();  // default stream
cudaMemcpy(...);            // default stream, 동기적

// 모든 작업이 순차적으로 실행됨

Stream 생성 및 사용

Stream 생성

cuda
복사
cudaStream_t stream;
cudaStreamCreate(&stream);  // 스트림 생성

// 작업 수행...

cudaStreamDestroy(stream);  // 스트림 해제

커널 실행에 스트림 지정

cuda
복사
// <<<그리드, 블록, 공유메모리, 스트림>>>
kernel<<<grid, block, 0, stream>>>(args);
//                    ↑     ↑
//            sharedMem   stream

비동기 메모리 복사

cuda
복사
// cudaMemcpyAsync: 비동기 메모리 복사
cudaMemcpyAsync(d_a, h_a, size, cudaMemcpyHostToDevice, stream);

// 주의: Host 메모리는 반드시 pinned memory여야 함!
// 일반 malloc()은 비동기 복사 불가

Pinned Memory (Page-locked Memory)

왜 필요한가?

일반 메모리 (pageable):
- OS가 페이지 스왑 가능
- GPU가 직접 접근 불가
- 내부 버퍼 거쳐 복사 → 비동기 불가

Pinned 메모리 (page-locked):
- OS 스왑 불가, 고정 위치
- GPU가 DMA로 직접 접근
- 비동기 복사 가능!

Pinned Memory 할당

cuda
복사
float* h_pinnedData;

// Pinned memory 할당
cudaMallocHost(&h_pinnedData, size);

// 또는
cudaHostAlloc(&h_pinnedData, size, cudaHostAllocDefault);

// 사용 후 해제
cudaFreeHost(h_pinnedData);

주의사항

✓ 비동기 전송에 필수
✓ 일반 memcpy보다 빠름

✗ 과도한 사용은 시스템 메모리 부족 야기
✗ 필요한 만큼만 할당

다중 스트림 파이프라인

단일 스트림 (비효율적)

다중 스트림 (효율적)

다중 스트림 구현

cuda
복사
const int nStreams = 4;
cudaStream_t streams[nStreams];

// 스트림 생성
for (int i = 0; i < nStreams; i++) {
    cudaStreamCreate(&streams[i]);
}

// 데이터를 청크로 분할
int chunkSize = n / nStreams;
size_t chunkBytes = chunkSize * sizeof(float);

// 각 스트림에서 비동기 작업 실행
for (int i = 0; i < nStreams; i++) {
    int offset = i * chunkSize;

    // 비동기 H→D 복사
    cudaMemcpyAsync(d_a + offset, h_a + offset, chunkBytes,
                    cudaMemcpyHostToDevice, streams[i]);

    // 커널 실행
    kernel<<<gridSize, blockSize, 0, streams[i]>>>(d_a + offset, chunkSize);

    // 비동기 D→H 복사
    cudaMemcpyAsync(h_result + offset, d_result + offset, chunkBytes,
                    cudaMemcpyDeviceToHost, streams[i]);
}

// 모든 스트림 완료 대기
cudaDeviceSynchronize();

// 스트림 해제
for (int i = 0; i < nStreams; i++) {
    cudaStreamDestroy(streams[i]);
}

동기화 기법

전체 Device 동기화

cuda
복사
// 모든 스트림의 모든 작업 완료 대기
cudaDeviceSynchronize();

특정 스트림 동기화

cuda
복사
// 해당 스트림의 작업 완료 대기
cudaStreamSynchronize(stream);

스트림 완료 확인 (비차단)

cuda
복사
// 완료 여부만 확인 (대기하지 않음)
cudaError_t status = cudaStreamQuery(stream);
if (status == cudaSuccess) {
    // 스트림 작업 완료
} else if (status == cudaErrorNotReady) {
    // 아직 실행 중
}

이벤트 기반 동기화

cuda
복사
cudaEvent_t event;
cudaEventCreate(&event);

// 스트림에 이벤트 기록
cudaEventRecord(event, stream1);

// 다른 스트림이 이벤트 완료 대기
cudaStreamWaitEvent(stream2, event, 0);
// stream2는 stream1의 이벤트가 완료될 때까지 대기

cudaEventDestroy(event);

CUDA 이벤트로 시간 측정

이벤트 기반 타이밍

cuda
복사
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

// 시작 시간 기록
cudaEventRecord(start);

// 측정할 작업
kernel<<<grid, block>>>(args);

// 종료 시간 기록
cudaEventRecord(stop);

// 완료 대기
cudaEventSynchronize(stop);

// 경과 시간 계산
float milliseconds = 0;
cudaEventElapsedTime(&milliseconds, start, stop);
printf("Kernel time: %.3f ms\n", milliseconds);

cudaEventDestroy(start);
cudaEventDestroy(stop);

CPU 타이밍 vs GPU 이벤트 타이밍

CPU 타이밍 (부정확):
- GPU 작업이 비동기이므로 실제 실행 시간과 다름
- CPU는 커널 시작 즉시 다음 코드로 진행

GPU 이벤트 타이밍 (정확):
- GPU 타임스탬프 사용
- 실제 GPU 실행 시간 측정
- 비동기 작업에도 정확

실제 파이프라인 예제

완전한 파이프라인 코드

cuda
복사
#include <cuda_runtime.h>
#include <stdio.h>

__global__ void processKernel(float* data, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        data[idx] = data[idx] * 2.0f;  // 간단한 처리
    }
}

int main() {
    const int n = 1 << 20;  // 1M elements
    const int nStreams = 4;
    const int chunkSize = n / nStreams;
    const size_t totalBytes = n * sizeof(float);
    const size_t chunkBytes = chunkSize * sizeof(float);

    // Pinned host memory
    float *h_data, *h_result;
    cudaMallocHost(&h_data, totalBytes);
    cudaMallocHost(&h_result, totalBytes);

    // Device memory
    float *d_data;
    cudaMalloc(&d_data, totalBytes);

    // Initialize data
    for (int i = 0; i < n; i++) h_data[i] = (float)i;

    // Create streams
    cudaStream_t streams[nStreams];
    for (int i = 0; i < nStreams; i++) {
        cudaStreamCreate(&streams[i]);
    }

    // Timing
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    cudaEventRecord(start);

    // Pipeline execution
    for (int i = 0; i < nStreams; i++) {
        int offset = i * chunkSize;
        cudaMemcpyAsync(d_data + offset, h_data + offset,
                        chunkBytes, cudaMemcpyHostToDevice, streams[i]);

        processKernel<<<chunkSize/256, 256, 0, streams[i]>>>
                      (d_data + offset, chunkSize);

        cudaMemcpyAsync(h_result + offset, d_data + offset,
                        chunkBytes, cudaMemcpyDeviceToHost, streams[i]);
    }

    cudaEventRecord(stop);
    cudaEventSynchronize(stop);

    float ms;
    cudaEventElapsedTime(&ms, start, stop);
    printf("Pipeline time: %.3f ms\n", ms);

    // Cleanup
    for (int i = 0; i < nStreams; i++) {
        cudaStreamDestroy(streams[i]);
    }
    cudaFreeHost(h_data);
    cudaFreeHost(h_result);
    cudaFree(d_data);
    cudaEventDestroy(start);
    cudaEventDestroy(stop);

    return 0;
}

정리

성능 향상 전략

1. 데이터를 청크로 분할
2. 여러 스트림 사용
3. 복사와 연산 오버랩
4. Pinned memory 활용
5. 적절한 동기화 지점 설정

핵심 포인트: 스트림을 활용한 파이프라인은 GPU 활용률을 극대화하고 전체 처리 시간을 단축합니다.