CUDA 스트림(stream)을 활용한 동시성 향상

  CUDA 스트림(stream)은 호스트 코드에서 기술된 순서대로 디바이스에서 동작하는 일련의 연속된 연산을 의미한다. 스트림 안의 연산들은 순차적인 실행이 보장되지만, 다른 스트림의 연산과는 함께 동작할 수 있다. 특히 GPU의 실행 리소스가 허락된다면, 서로 다른 스트림의 커널도 동시에(concurrently) 실행될 수 있다.  일반적으로 소개되고 있는 예제에서는 데이터 전송(data transfer)을 1)호스트에서 동작 중인 연산, 2)디바이스에서 동작 중인 연산, 3)또 다른 데이터 전송과 중첩되게 하기 위해 CUDA 스트림을 사용하는 경우를 보여준다.

 

  CUDA의 모든 디바이스 연산은 스트림 안에서 실행되어야 한다. 따라서 스트림이 지정되지 않는다면, 널 스트림(null stream)이라고도 하는 기본 스트림(default stream) 안에서 실행된다.

그런데 기본 스트림은 디바이스의 연산과 동기화 되는 스트림이다. 이전에 시작된 연산이 디바이스에서 끝나지 않는다면, 기본 스트림의 연산은 시작될 수 없다. 기본 스트림이 아닌 스트림을 사용하려면 cudaStreamCreate()을 이용하여 스트림을 생성하고, cudaStreamDestroy()를 사용해 스트림을 종료한다.

          cudaStream_t s1;          

          cudaStreamCreate(s1);

          cudaStreamDestroy(s1);

이렇게 생성된 스트림 식별자(stream identifier) s1을 이용해 비동기적 데이터 전송을 하려면 cudaMemcpyAsync()를 사용한다. ( 2차원/3차원 배열을 위해서는 각각 cudaMemcpy2DAsync()와 cudaMemcpy3DAsync()가 사용된다)

          cudaMemcpyAsync( d_a, a, numBytes, cudaMemcpyHostToDevice, s1);

다른 스트림의 커널 실행을 비동기적으로 실행하려면, 스트림 식별자를 네 번째 매개변수로 전달하면 된다. 세 번째 매개변수는 공유 메모리를 얼마나 할당할 지를 나타낸다.

          kernel<<<1, N, 0, s1>>>(d_a);

기본 스트림과 달리 다른 스트림은 non-blocking 스트림이므로, 위의 함수들이 실행되자마자 실행 제어권한은 다시 호스트 코드 쪽으로 넘어간다. 스트림 안의 어떠한 연산도 호스트 코드의 실행을 가로막지 못한다.

따라서 만약 스트림 안의 연산과 호스트 코드를 동기화할 필요가 있다면, cudaDeviceSynchronize()와 같은 동기화함수를 명시적으로 사용해야 한다. cudaDeviceSynchronize()함수는 디바이스에서 이전에 시작된 모든 연산이 종료될 때까지 호스트 코드를 블록한다. 따라서 이 함수의 사용은 다른 디바이스 스레드나 호스트 스레드를 멈추게(stall)하여서 성능 저하를 가져오게 된다.

cudaStreamSynchronize(stream) 함수는 특정 스트림 안에서 시작된 모든 연산이 끝날 때까지 호스트 스레드를 블록하는 방식이다. cudaStreamQuery(stream)는 특정 스트림 안의 모든 연산의 종료 여부를 테스트 할 수 있다.

cudaEventSynchronize(event)와 cudaEventQuery(event)는 특정 이벤트의 기록 여부를 이용하여 위의 두 함수와 유사한 기능을 수행할 수 있다. cudaStreamWaitEvent(event)를 이용하면 스트림 안의 연산들을 특정 이벤트와 동기화시킬 수 있다.

첫번째 참고 링크에서는 전체 데이터를 스트림 개수만큼 나누어 데이터 전송과 커널 실행이 다른 스트림과 중첩될 수 있는 방식을 보여주고 있다. 

1) 스트림마다의 데이터 전송과 커널 실행을 하나의 배치로 묶어서 반복하는 방식과

2) 모든 스트림에 대한 데이터 전송을 반복하고, 전송된 데이터에 대한 스트림의 커널 실행을 반족하는 방식

2가지가 제시되었다. 두 방식 모두 똑같이 옳은 결과를 가져오지만, 블로그가 작성되는 시점의 테슬라 GPU에서 하드웨어 데이터 전송 엔진의 개수 및 방식에 따라 각기 성능 차이를 보였다. 데이터 전송 엔진이 하나만 있는 경우는 두 번째 방식이 높은 성능을 보였지만, 데이터 전송 엔진이 디바이스 방향과 호스트 방향 엔진으로 두 개 있는 경우 첫 번째 방식이 높은 성능을 보였다.

그러나 어느 방식도 기본 스트림만 사용하는 경우보다는 월등한 성능을 가져오기 때문에, 개발 단계에서는 기본 스트림을 사용하더라도 최종단계에서는 스트림을 명시적으로 특정할 필요가 있다. 특히 라이브러리를 개발하는 경우 기본 스트림을 사용한다면, 라이브러리 사용자가 데이터 전송을 라이브러리 커널 실행과 겹칠 수 없게하는 한계가 존재한다고 밝히고 있다.

---

그런데 스트림을 사용하였다고 하더라도 중간에 기본 스트림을 사용하는 커널이 들어가게 되면, 동기화를 위해 블록하는 기본 스트림의 특성때문에 암묵적인 동기화가 적용되게 된다.

 

위의 코드에서는 아무런 작업도 하지 않는 더미 커널을 기본 스트림에서 시작시켰다. 실제 작업을 하는 커널은 모두 다른 스트림에 할당되었기때문에 원래라면 실행이 중첩될 수 있어야하지만, 기본 스트림의 더미 커널로 인해 모두 순차적으로 실행되게 된다.

이런 의도하지 않은 코드 작성자의 실수를 방지하기 위해, CUDA 7부터는 모든 호스트 스레드마다 독립적인 기본 스트림을 사용할 수 있게 하는 옵션을 제공한다. 이런 per-thread default stream 기능을 활용하려면, 컴파일시에 --default-stream per-thread 옵션을 지정하면 된다. 또는 원본 코드에서 cuda.h나 cuda_runtime.h 등의 CUDA 헤더파일을 포함시키기 전에 CUDA_API_PER_THREAD_DEFAULT_STREAM을 전처리 매크로로 선언하면 된다.

           1) nvcc --default-stream per-thread ./test.cu -o test

           2) #define CUDA_API_PER_THREAD_DEFAULT_STEAM 포함

위의 옵션대로 실행할 경우, 호스트 스레드의 기본 스트림으로 Stream14가 할당되었고, 다른 스트림의 연산들은 이 스트림과 동시에 실행될 수 있기 때문에 동시성이 향상 되었음을 알 수 있다. 

이렇게 할당된 기본 스트림은 cudaStreamPerThread 핸들을 이용해서 접근할 수 있다. 또한 만약 기존의 기본 스트림에 접근하려면 cudaStreamLegacy 핸들을 사용해야된다. 또 이 경우 cudaStreamLegacy의 사용은 스레드 당 기본 스트림에서 사용되는 경우 암시적인 동기화를 발생시키게 된다.

 

참고) How to Overlap Data Transfers in CUDA C/C++

        GPU Pro Tip: CUDA 7 Streams Simplify Concurrency