CUDA Stream

스트림 : CUDA의 비동기 구현을 위한 객체로서, 데이터 입력을 받으면서  GPU 프로세서가 함께 계산하도록 지원해준다.

일반적인 CUDA 프로그램의 구조는 3단계로 구성되어 있다.

1. Host 에서 Device로 입력 데이터 전송
2. 커널 함수에서 데이터 처리
3. 결과 데이터 Device 에서 Host로 전송

위의 3단계 과정을 처리할 때 1, 3 단계에서는 GPU 프로세서는 대기하고 2 단계에서만 GPU가 계산을 처리한다.
만약 데이터를 좀 더 작게 나누고 입력 데이터 전송이 완료된 것부터 GPU에서 계산을 함과 동시에 다음 데이터를 전송하면 전체적으로 효율이 좋아진다.

[-] Collapse

    cudaStream_t stream1; // Stream 개체 선언
    cudaStream_t stream2;

    cudaStreamCreate(&stream1); // Stream 개체 생성
    cudaStreamCreate(&stream2);

    cudaMemcpyAsync(deviceInput1, hostInput1, dataSize1, cudaMemcpyHostToDevice, stream1); // 비동기로 Data 전송
    cudaMemcpyAsync(deviceInput2, hostInput2, dataSize2, cudaMemcpyHostToDevice, stream2);

    kernel<<<nBlocks, nThreads, 0, stream1>>>(deviceInput1, deviceOutput1); // Kernel 함수 실행
    kernel<<<nBlocks, nThreads, 0, stream2>>>(deviceInput2, deviceOutput2);

    cudaStreamDestroy(stream1); // Stream 종료
    cudaStreamDestroy(stream2);

위의 Code를 실행하면 stream1 의 cudaMemcpyAsync() 의 실행이 완료되면 stream1의 kernel()이 실행됨과 동시에 stream2의 cudaMemcpyAsync()가 실행된다. stream1의 kernel()의 실행 완료와 stream2의 cudaMemcpyAsync()가 모두 완료되면 stream2의 kernel()이 실행된다.

일반적인 CUDA 프로그램에 스트림을 적용 할 수 있는데, 그 코드를 정형화 하면 다음과 같다.
[-] Collapse

    const int N; // 데이터 개수
    const int nStreams; // 분할 개수
    const int nBlocks; // 블록 개수
    const int nThreads; // 스레드 갯수
    int offset = 0;

    cudaStream_t streams = (cudaStream_t *)malloc(nStreams * sizeof(cudaStream_t));

    for (int i = 0; i < nStreams; i++)
        cutilSafeCall(cudaStreamCreate(&(streams[i])));

    for (int i = 0; i < nStreams; i++)
    {
        offset = i * N / nStreams;
        cudaMemcpyAsync(dev_In + offset, InputData + offset, Size, cudaMemcpyHostToDevice, streams[i]);
    }

    for (int i = 0; i < nStreams; i++)
    {
        offset = i * N / nStreams;
        kernel<<<nBlocks/nStreams, nThreads, 0, streams[i]>>>(dev_In + offset, dev_Out + offset);
    }

    for (int i = 0; i < nStreams; i++)
    {
        offset = i * N / nStreams;
        cudaMemcpyAsync(host_Out + offset, dev_Out + offset, Size, cudaMemcpyDeviceToHost, streams[i]);
    }

    for (int i = 0; i < nStreams; i++)
        cudaStreamDestory(&(streams[i]));