CUDA Stream & Concurrent Execution

CUDA Stream
→ Host에서 Device로 명령을 보내는 통로(Host에서 호출하는 명령들이 차례대로 들어감)

Stream 개념

Types of CUDA Stream

• NULL stream → 암묵적으로 선언된 stream
  • 사용할 stream을 명시하지 않은 경우(default stream)
• Non-NULL stream → 명시적으로 선언된 stream
  • 명시적으로 생성 및 사용

하나의 stream에 들어온 명령은 순서대로 Device에 의해 처리된다.
→ Synchronous execution

C() → B() → A()

Synchronous execution

하지만 서로 다른 stream에 있는 명령들은 순서가 정해져 있지 않다.(비동기적)
즉, stream들이 명시되어 있으면 비동기적인 stream이 될 수 있다.

Asynchronous execution

 

Concurrent Execution
→ 서로 다른 stream에 있는 작업 중 일부는 동시에 수행될 수 있다.

동시에 수행(overlap)이 가능한 작업들

• Host computation and device computation
• Host computation and host-device data transfer
• Device computation and host-device data transfer
• Concurrent device compuatations

* 동시 수행 능력은 GPU마다 다를 수 있다.

동시에 작업이 가능하기 때문에 이를 활용하면 높은 성능을 낼 수 있다.
아래 그림을 살펴보자.

Concurrent Execution

stream을 3개로 나눴을 때 다음과 같이 동시에 작업이 처리될 수 있다.

 

그렇다면 어떻게 non-NULL Stream을 만드는지 알아보자.

• Create/Destory
  • Data type : cudaStream_t
  • cudaError_t cudaStreamCreate(cudaStream_t*)
  • cudaError_t cudaStreamDestory(cudaStream_t)

cudaStream_t stream;
cudaStreamCreate(&stream);
myKernel<<<dimGrid, dimBlock, 0, stream>>>(...)
cudaStreamDestory(stream);

cudaMemcpy(d_a, a, memSize, cudaMemcpyHostToDevice);
cudaMemcpy(b, d_b, memSize, cudaMemcpyDeviceToHost);