CUDA 스트림

CUDA 스트림은 GPU에서 실행되는 작업들의 비동기적 실행 순서를 관리하는 논리적 실행 단위이다. 호스트(CPU)가 디바이스(GPU)로 보내는 명령, 즉 커널 실행이나 메모리 복사 등의 작업이 차례대로 들어가는 큐(Queue)로 이해할 수 있다. 각 스트림은 독립적인 작업 흐름을 가지며, 비동기 실행을 통해 GPU 활용률을 극대화하고 애플리케이션의 전반적인 성능을 향상시키는 핵심 역할을 한다.

CUDA 프로그래밍에서 스트림을 사용하지 않으면 모든 작업은 암시적으로 기본 스트림(Default Stream, 또는 널 스트림)에서 실행된다. 기본 스트림 내의 작업들은 호스트에 대해 비동기적으로 실행되지만, 서로 간에는 순차적으로 실행되어 중첩 실행의 이점을 얻을 수 없다. 반면, 프로그래머가 명시적으로 생성한 논-널 스트림(Non-Null Stream)을 여러 개 사용하면, 서로 다른 스트림에 속한 작업들(예: 한 스트림의 호스트-장치 데이터 전송과 다른 스트림의 커널 실행)은 GPU 하드웨어가 지원하는 범위 내에서 동시에 실행될 수 있다. 이는 병렬 컴퓨팅의 효율성을 크게 높인다.

스트림의 주요 목적은 데이터 전송과 컴퓨테이션을 중첩시켜 GPU의 유휴 시간을 최소화하는 것이다. 예를 들어, 큰 데이터를 한 번에 보내고 연산하는 대신, 데이터를 여러 부분으로 나누어 각각을 별도의 스트림에 할당하면, 첫 번째 데이터 부분의 연산이 진행되는 동안 두 번째 데이터 부분을 전송하는 등의 파이프라이닝이 가능해진다. 이러한 동시성은 고성능 컴퓨팅 애플리케이션의 처리 속도를 획기적으로 개선하는 전략이다.

비동기 실행은 CUDA 스트림의 핵심 개념으로, 호스트(CPU)가 디바이스(GPU)에 작업을 지시한 후 그 완료를 기다리지 않고 즉시 다음 작업을 수행할 수 있게 하는 방식을 의미한다. 이는 병렬 컴퓨팅의 효율성을 극대화하는 데 중요한 역할을 한다. 기본적으로 스트림은 작업의 실행 순서를 관리하는 논리적 큐이며, 하나의 스트림 내에서는 작업이 순차적으로 실행된다. 그러나 서로 다른 스트림에 배치된 작업들 사이에는 실행 순서가 보장되지 않아, GPU의 스케줄러에 따라 비동기적이고 동시에 실행될 가능성이 있다.

이러한 비동기 실행의 가장 큰 장점은 호스트와 디바이스 간의 작업 중첩을 통한 성능 향상이다. 예를 들어, 한 스트림에서 호스트에서 디바이스로의 데이터 전송(cudaMemcpyAsync)이 진행되는 동안, 다른 스트림에서 커널 실행이 동시에 이루어질 수 있다. 이렇게 하면 데이터 전송 시간이 연산 시간에 숨겨져 전체 애플리케이션의 처리 속도를 높일 수 있다. 또한, 서로 다른 커널을 다른 스트림에서 실행함으로써 GPU의 컴퓨팅 유닛 활용률을 높일 수도 있다.

비동기 실행을 구현하려면 명시적으로 스트림을 생성(cudaStreamCreate)하고, 데이터 전송 시에는 cudaMemcpyAsync 함수를, 커널 실행 시에는 스트림을 인자로 지정하여 호출해야 한다. 중요한 점은 비동기 데이터 전송을 위해서는 페이지 잠금 메모리(Pinned Memory)를 사용해야 한다는 것이다. 일반적인 페이지 가능 메모리는 운영체제에 의해 디스크로 스왑될 수 있어, DMA 방식으로 직접 접근하는 GPU의 데이터 전송 엔진과 호환되지 않기 때문이다. 작업이 완료된 시점을 확인하려면 cudaStreamSynchronize 함수를 사용하여 특정 스트림의 모든 작업이 끝날 때까지 대기하거나, CUDA 이벤트를 활용하여 스트림 간 의존성을 관리할 수 있다.

CUDA 스트림을 생성하려면 cudaStream_t 타입의 변수를 선언한 후 cudaStreamCreate() 함수를 호출한다. 이 함수는 생성된 스트림 객체에 대한 포인터를 인자로 받아 새로운 비동기 스트림을 초기화한다. 생성된 스트림은 이후 커널 실행이나 cudaMemcpyAsync()와 같은 비동기 메모리 복사 작업을 지시할 때 사용된다.

스트림 사용이 끝나면 cudaStreamDestroy() 함수를 호출하여 자원을 해제해야 한다. 이 함수는 호출 즉시 반환되며, 만약 해당 스트림에 아직 실행 중인 작업이 있다면 그 작업이 완료된 후 내부 자원이 정리된다. 이를 통해 메모리 누수을 방지할 수 있다. 스트림을 명시적으로 생성하고 관리하는 것은 기본 스트림만 사용할 때 발생할 수 있는 암묵적인 동기화를 피하고, 여러 작업을 중첩시켜 GPU 활용률을 높이는 데 필수적이다.

CUDA 스트림에서 작업의 실행 순서를 보장하거나, 다른 작업과의 의존성을 관리하기 위해 동기화 메커니즘이 필요하다. 스트림 동기화는 특정 스트림 내의 모든 비동기 작업이 완료될 때까지 호스트 스레드의 실행을 차단하거나, 서로 다른 스트림 간의 작업 순서를 조정하는 역할을 한다.

주요 동기화 함수로는 cudaStreamSynchronize()가 있다. 이 함수는 인자로 지정된 스트림에 큐에 들어간 모든 작업(예: cudaMemcpyAsync 또는 커널 실행)이 완료될 때까지 호출한 호스트 스레드의 실행을 차단한다. 이는 특정 스트림의 작업이 끝난 후에만 그 결과를 사용해야 하는 경우에 필수적이다. 또한, cudaDeviceSynchronize() 함수는 특정 GPU 장치에 대한 모든 스트림의 작업이 완료될 때까지 대기하는 데 사용된다.

보다 세밀한 제어를 위해 CUDA 이벤트(cudaEvent_t)를 활용한 동기화도 가능하다. 이벤트를 스트림 내 특정 지점에 기록(cudaEventRecord())하고, 다른 스트림이 해당 이벤트가 완료될 때까지 대기(cudaStreamWaitEvent())하도록 설정할 수 있다. 이를 통해 서로 다른 스트림 간의 작업 순서와 의존성을 명시적으로 정의할 수 있으며, 데이터 생산자와 소비자 관계를 안전하게 구성하는 데 유용하다.

단일 스트림을 활용하는 방식은 CUDA 프로그래밍의 기본이 된다. 이는 명시적으로 스트림을 생성하지 않고, 암시적으로 존재하는 기본 스트림을 사용하는 경우를 의미한다. 모든 커널 실행과 메모리 복사 작업이 이 단일 스트림에 순차적으로 큐잉되고, GPU는 이 명령들을 호스트에 의해 기술된 순서대로 하나씩 실행한다.

이 방식은 구현이 간단하고 디버깅이 용이하다는 장점이 있다. 기본 스트림은 동기화 스트림으로 작동하여, 한 작업이 완전히 끝난 후에 다음 작업이 시작되도록 보장한다. 이는 호스트와 디바이스 간의 실행 순서를 명확하게 하여 프로그램의 정확성을 쉽게 유지할 수 있게 해준다. 따라서 초보자가 CUDA를 학습하거나 간단한 프로토타입을 개발할 때 유용하게 사용된다.

그러나 단일 스트림 활용의 주요 한계는 성능이다. 호스트-디바이스 데이터 전송과 커널 실행이 서로 중첩되어 동시에 실행될 수 없기 때문에, PCI 익스프레스 버스를 통한 데이터 이동 시간 동안 GPU가 유휴 상태로 대기하거나, 그 반대의 상황이 발생할 수 있다. 이는 고성능 컴퓨팅 애플리케이션에서 전체 실행 시간을 증가시키는 주요 요인이 된다.

따라서 단일 스트림 활용은 성능 최적화보다는 코드의 간결성과 안정성에 중점을 둔 경우에 적합하다. 보다 높은 GPU 활용률과 성능 향상을 위해서는 다중 스트림을 생성하여 데이터 전송과 연산을 중첩시키는 비동기 실행 기법으로 발전시켜야 한다.

다중 스트림 활용은 GPU의 병렬 컴퓨팅 능력을 극대화하기 위한 핵심 기법이다. 단일 스트림을 사용할 경우 호스트와 디바이스 간의 데이터 전송과 커널 실행이 순차적으로 이루어져 GPU나 호스트 중 하나가 대기 상태에 머무르는 시간이 발생한다. 반면, 여러 개의 독립적인 스트림을 생성하여 작업을 분배하면, 서로 다른 스트림에 속한 작업들은 비동기적으로 실행될 수 있으며, 특정 조건 하에서 동시에 수행되어 전체 실행 시간을 단축할 수 있다.

다중 스트림을 효과적으로 활용하는 전형적인 패턴은 데이터를 여러 부분으로 나누어 각각 다른 스트림에 할당하는 것이다. 예를 들어, 대용량 데이터를 처리할 때 첫 번째 스트림에서 데이터의 첫 번째 부분을 GPU로 전송하는 동시에, 두 번째 스트림에서는 이미 전송된 두 번째 부분에 대한 커널 실행을 시작할 수 있다. 이렇게 하면 데이터 전송 엔진과 커널 실행 엔진이 동시에 가동되어 하드웨어 활용률을 높일 수 있다. 이러한 기법은 특히 PCI 익스프레스 버스를 통한 데이터 전송 시간이 큰 오버헤드로 작용하는 경우에 성능 향상 효과가 두드러진다.

다만, 모든 작업이 동시에 실행될 수 있는 것은 아니다. 동시 실행 가능성은 GPU 아키텍처와 성능에 따라 달라진다. 일반적으로 서로 다른 스트림 간의 호스트에서 디바이스로의 메모리 복사, 디바이스에서 호스트로의 메모리 복사, 그리고 서로 다른 커널 실행이 중첩되어 수행될 가능성이 있다. 이를 위해서는 비동기 데이터 전송 함수인 cudaMemcpyAsync를 사용해야 하며, 해당 데이터가 페이지 잠금 메모리에 상주해야 한다.

성공적인 다중 스트림 활용을 위해서는 작업 간의 의존성을 신중하게 관리해야 한다. 한 스트림의 결과가 다른 스트림의 입력으로 필요한 경우, cudaStreamWaitEvent와 같은 동기화 함수를 사용하여 명시적으로 실행 순서를 제어해야 한다. 또한, 너무 많은 수의 스트림을 생성하면 스케줄링 오버헤드가 증가하거나 GPU 리소스 경쟁으로 인해 오히려 성능이 저하될 수 있으므로, 실제 하드웨어에서의 동시 실행 가능 스트림 수를 고려하여 적절한 개수를 선정하는 것이 중요하다.

호스트-장치 데이터 전송과 커널 실행의 중첩은 CUDA 스트림을 활용한 핵심적인 성능 최적화 기법이다. 기본적으로 호스트와 GPU 사이의 데이터 전송(PCI 익스프레스 버스를 통해 이루어짐)과 커널 실행은 순차적으로 이루어지며, 이 과정에서 한쪽이 작업할 때 다른 쪽은 대기 상태가 되어 자원 활용률이 낮아진다. 스트림을 이용하면 이러한 작업들을 비동기적으로 스케줄링하고 중첩시켜 전체 실행 시간을 단축할 수 있다.

이 기법을 구현하려면 먼저 여러 개의 논-널 스트림을 생성해야 한다. 그런 다음, 처리할 전체 데이터를 여러 개의 청크로 나누어 각 스트림에 할당한다. 각 스트림 내에서는 cudaMemcpyAsync 함수를 사용해 호스트의 페이지 잠금 메모리에서 장치로 데이터를 비동기적으로 복사하고, 이어서 커널을 실행한 후, 다시 결과를 장치에서 호스트로 비동기적으로 복사하는 작업 파이프라인을 구성한다. 서로 다른 스트림에 속한 작업들은 GPU의 DMA 엔진과 컴퓨트 엔진이 지원하는 범위 내에서 가능한 한 동시에 실행되어, 데이터 전송 시간의 일부를 커널 실행 시간 아래로 "숨기는" 효과를 낸다.

이러한 중첩의 효율성을 극대화하기 위해서는 몇 가지 조건이 필요하다. 첫째, 비동기 메모리 복사에 사용되는 호스트 메모리는 반드시 cudaMallocHost나 cudaHostAlloc으로 할당된 페이지 잠금 메모리여야 한다. 둘째, 작업의 부하를 균등하게 나누고 충분한 양의 작업을 각 스트림에 제공하여 파이프라이닝 효과를 얻어야 한다. 마지막으로, cudaStreamSynchronize나 CUDA 이벤트를 사용해 모든 스트림의 작업이 완료될 때까지 호스트가 적절히 대기하도록 동기화를 관리해야 한다.

페이지 잠금 메모리는 호스트의 가상 메모리 시스템에서 운영체제에 의해 디스크로 스왑 아웃되지 않도록 고정된 메모리 영역이다. CUDA에서 비동기 메모리 복사를 수행하려면 반드시 페이지 잠금 메모리를 사용해야 한다. 이는 DMA 방식으로 동작하는 PCI 익스프레스 버스를 통한 데이터 전송이 물리적 메모리 주소를 필요로 하기 때문이다. 일반적인 페이지 가능 메모리를 사용하면 GPU 드라이버가 먼저 임시의 페이지 잠금 메모리를 할당하고 데이터를 중간에 복사하는 추가 오버헤드가 발생하지만, 페이지 잠금 메모리를 직접 사용하면 이 과정이 생략되어 전송 성능이 향상된다.

페이지 잠금 메모리는 cudaMallocHost() 또는 cudaHostAlloc() 함수를 사용하여 할당하고, cudaFreeHost()로 해제한다. cudaMemcpyAsync()와 같은 비동기 복사 함수는 소스 또는 목적지가 페이지 잠금 메모리일 때만 정상적으로 동작하며, 이를 통해 데이터 전송과 커널 실행을 중첩시키는 성능 최적화가 가능해진다. 그러나 페이지 잠금 메모리는 물리적 램에 상주해야 하므로 과도하게 사용하면 시스템 전체의 메모리 부족을 초래할 수 있으며, 가상 메모리의 장점을 활용할 수 없다는 단점도 있다.

다중 스트림을 이용한 성능 향상은 GPU의 병렬 컴퓨팅 능력을 극대화하기 위한 핵심 기법이다. 기본적으로 단일 스트림을 사용하면 호스트와 디바이스 간의 데이터 전송과 커널 실행이 순차적으로 이루어지기 때문에, 데이터를 보내는 동안 GPU는 대기하고, 연산을 하는 동안 데이터 전송 경로는 유휴 상태가 된다. 다중 스트림을 활용하면 이러한 작업들을 여러 개의 독립적인 실행 큐에 분배하여, 서로 다른 스트림의 작업들이 가능한 한 동시에 실행되도록 할 수 있다. 특히 호스트-장치 데이터 전송과 커널 실행을 서로 다른 스트림에 배치하면, 한 스트림에서 데이터를 전송하는 동안 다른 스트림에서 커널을 실행하는 식으로 작업을 중첩시켜 전체 실행 시간을 단축할 수 있다.

이를 효과적으로 구현하기 위해서는 작업을 적절히 분할해야 한다. 예를 들어, 처리해야 할 대용량 데이터를 여러 개의 청크로 나누고, 각 청크에 대한 데이터 전송과 커널 실행을 별도의 스트림에 할당한다. 첫 번째 스트림이 첫 번째 청크의 데이터를 GPU로 전송하고, 그 전송이 완료되면 해당 스트림에서 커널을 실행한다. 그동안 두 번째 스트림은 두 번째 청크의 데이터 전송을 시작할 수 있다. 이렇게 파이프라인 방식으로 작업을 구성하면 데이터 전송과 연산이 겹쳐서 수행되어 GPU의 활용률을 높이고 대기 시간을 숨길 수 있다.

성능 향상의 정도는 하드웨어의 능력에 크게 의존한다. 많은 GPU는 호스트에서 디바이스로의 전송, 디바이스에서 호스트로의 전송, 그리고 커널 실행을 담당하는 별도의 엔진을 가지고 있어, 이러한 작업들이 서로 다른 스트림에 속해 있을 때 진정한 동시성을 달성할 수 있다. 또한, 비동기 데이터 전송을 위해서는 페이지 잠금 메모리를 사용해야 한다. 일반적인 페이지 가능 메모리를 사용하면 CUDA 드라이버가 임시의 고정 메모리를 할당하는 추가 작업이 필요하므로 성능 저하가 발생할 수 있다.

그러나 다중 스트림 사용이 항상 성능을 선형적으로 향상시키는 것은 아니다. 스트림 간의 과도한 경쟁으로 인해 리소스 부족이 발생하거나, 작업 부하가 고르지 않게 분배되면 오히려 성능이 저하될 수 있다. 또한, 모든 작업이 완료되었는지 확인하기 위해 cudaStreamSynchronize와 같은 적절한 동기화 메커니즘을 사용해야 하며, 스트림 간의 의존성이 있다면 cudaEventRecord와 cudaStreamWaitEvent를 이용해 명시적으로 관리해야 한다. 따라서 애플리케이션의 특성과 하드웨어 사양을 고려하여 최적의 스트림 개수와 작업 분할 전략을 찾는 것이 중요하다.

CUDA 스트림을 활용한 성능 최적화에는 몇 가지 주의사항과 본질적인 한계가 존재한다. 우선, 다중 스트림을 사용하더라도 GPU 하드웨어의 물리적 제약으로 인해 동시 실행 가능한 작업의 수에는 한계가 있다. 예를 들어, 일반적으로 하나의 커널 실행 엔진과 두 개의 DMA 엔진(호스트-장치, 장치-호스트 방향 각각)이 존재하므로, 이론상 최대 세 개의 작업(커널 하나와 메모리 복사 두 개)만이 진정한 의미에서 동시에 실행될 수 있다. 또한, 하나의 스트림 내에서 기본 스트림은 암묵적인 동기화 지점으로 작용할 수 있어, 다른 스트림의 작업을 블로킹하고 병렬 실행의 이점을 상쇄시킬 위험이 있다.

성능 향상을 위해서는 페이지 잠금 메모리를 사용해야 한다는 점도 중요한 제약 조건이다. cudaMemcpyAsync와 같은 비동기 데이터 전송 함수는 반드시 페이지 잠금 메모리에서만 동작한다. 페이지 잠금 메모리는 운영체제가 디스크로 스왑 아웃할 수 없도록 고정시키므로, 시스템의 가용 물리 메모리를 과도하게 점유하여 다른 응용 프로그램의 성능에 영향을 미치거나 메모리 부족 현상을 초래할 수 있다. 따라서 이 메모리는 필요한 만큼만 할당하고 사용 후 즉시 해제하는 것이 바람직하다.

마지막으로, 다중 스트림을 사용하더라도 모든 작업이 완벽하게 중첩되어 실행되리라는 보장은 없다. 작업들의 실행 시간, 데이터 의존성, GPU의 작업 스케줄러 동작 방식에 따라 성능 향상 효과는 달라질 수 있다. 특히 서로 다른 스트림에 배치된 작업들 사이에 순서 의존성이 존재할 경우, cudaStreamWaitEvent와 같은 함수를 사용하여 명시적인 동기화를 수행해야 올바른 결과를 보장할 수 있다.