[CUDA] CUDA Streams

CUDA 프로그래밍에서 Stream은 순서대로 실행되는 일련의 명령체계이다. CUDA 어플리케이션에서 Kernel 실행과 일부 메모리 전송은 CUDA Stream 내에서 발생한다. 명시적으로 CUDA Stream에 대해 작성하지 않았을 때에는 실제로 CUDA Code는 default stream이라고 불리는 stream 내에서 kernel이 실행된다.

 

CUDA 프로그래밍에서는 default stream 외에 non-default CUDA stream을 생성하고 활용할 수 있는데, 이를 활용하면 여러 kernel을 동시에 여러 개의 다른 stream에서 실행하는 등의 여러 작업을 수행할 수 있다. 다중 stream을 사용하면 가속화된 어플리케이션에 병렬화 계층을 추가할 수 있고 어플리케이션 최적화를 위한 더 많은 기회를 얻을 수 있다.

 

아래의 규칙은 CUDA stream 동작에 관한 몇 가지 rule로써, 효과적으로 활용하기 위해 알아야 하는 내용이며, 슬라이드를 통해 자세히 확인할 수 있다.

 

• 지정된 stream 내의 작업은 순차적으로 동작한다.
• default가 아닌 stream에서의 작업은 서로에 대해 특정 순서로 작동한다고 보장할 수 없다.
• default stream은 다른 모든 stream이 완료될 때까지 대기하며 다른 stream의 실행을 차단한다.

 

아래의 코드는 non-default CUDA stream을 만들고, 활용하며, 해제하기까지의 예제이다. non-default stream에서 CUDA kernel을 실행하려면 execute configuration의 stream이 네 번째 arguement로 전달되어야 한다.

 

cudaStream_t stream;       // CUDA streams are of type `cudaStream_t`.
cudaStreamCreate(&stream); // Note that a pointer must be passed to `cudaCreateStream`.

someKernel<<<number_of_blocks, threads_per_block, 0, stream>>>(); // `stream` is passed as 4th EC argument.

cudaStreamDestroy(stream); // Note that a value, not a pointer, is passed to `cudaDestroyStream`.

위 코드의 세 번째 argument는 kernel 실행을 위해 block 당 동적으로 할당될 shared memory의 바이트 수에 대한 것이며, 블록 당 공유메모리에 할당되는 기본 바이트 수는 0이다. (추후 자세히 알아보자)