CPU와 GPU가 함께 작업을 실행하는 방법

지난 글에서는 커널의 스레드들이 블록과 워프로 묶이고, GPU 안에서 실행되는 과정을 살펴봤습니다. 이번에는 시야를 프로그램 전체로 넓혀 보겠습니다. GPU에서 실행할 커널이 준비되었다면, 누가 입력을 마련하고 실행을 요청할까요? 계산 결과가 필요한 다음 작업은 언제 시작할 수 있을까요?

이번 글은 CPU가 작업을 요청하는 시점과 GPU가 그 작업을 완료하는 시점을 구별하는 것에서 출발합니다. 먼저 CPU의 호출 코드와 GPU의 커널 본문을 연결하고, 입력과 결과를 어디에 보관하는지 살펴보겠습니다. 이어서 작업을 순서대로 등록하는 방법과 CPU가 결과를 기다려야 하는 때를 알아보겠습니다. 마지막으로 독립적인 작업을 여러 스트림에 나누고, 결과를 주고받는 작업 사이에는 필요한 실행 순서를 연결하겠습니다.

입력 배열에 1을 더하는 계산 A와, 그 결과를 두 배로 만드는 계산 B를 글 전체에서 사용하겠습니다. 마지막에는 A·B와 별도 데이터를 사용하는 계산 C를 추가합니다.

그림과 코드에는 NVIDIA GPU에 작업을 요청할 때 사용하는 CUDA가 등장합니다. CUDA 문법을 몰라도 이 글을 읽는 데 지장은 없습니다. 본문과 그림으로 실행 흐름을 설명하고, 코드에는 각 단계가 하는 일을 주석으로 붙이겠습니다. 코드는 주석을 따라 준비 → 복사 → 계산 요청 → 결과 사용의 순서만 살펴봐도 충분합니다. ...은 세부 내용을 생략했다는 표시이며, 예시는 그대로 실행하는 완성된 프로그램이 아닙니다.

CPU 코드에서 GPU 커널 실행하기

입력 x가 [1, 2, 3, 4]라고 해보겠습니다. 계산 A는 각 원소에 1을 더해 중간값 u를 [2, 3, 4, 5]로 만들고, 계산 B는 u의 각 원소를 두 배로 만들어 결과 y를 [4, 6, 8, 10]으로 만듭니다. B는 A가 만든 u를 읽으므로, A의 결과가 준비된 뒤 B가 실행되어야 합니다.

이 계산을 GPU에 맡기는 프로그램에는 CPU에서 실행되는 코드와 GPU에서 실행되는 코드가 함께 있습니다. CPU 쪽 코드는 입력을 준비하고, GPU에 메모리 공간을 요청하고, 복사와 커널 실행을 호출합니다. GPU에서는 커널의 스레드들이 각자 맡은 원소를 계산합니다.

그림 1의 왼쪽에 있는 A 호출은 GPU에 계산 A의 실행을 요청하는 CPU 코드입니다. 오른쪽의 커널 본문은 그 요청을 받아 GPU가 수행하는 계산입니다. 지난 글에서 본 것처럼 GPU 스레드마다 자신이 담당하는 원소를 읽고, 1을 더한 값을 중간값 u에 씁니다.

핵심은 실행을 요청하는 쪽은 CPU이고, 요청받은 계산을 수행하는 쪽은 GPU라는 점입니다. 이번 글에서 메모리 공간 확보와 복사를 요청하는 코드도 커널 밖의 CPU 코드에 둡니다. 함수 이름이나 괄호 안의 세부 내용보다 이 역할 구분에 집중하면 됩니다. CUDA의 커널 정의와 실행 호출

GPU에 입력과 결과 공간 준비하기

GPU가 x를 읽고 u와 y를 쓰려면 각 데이터를 놓을 공간이 필요합니다. 데이터를 담는 메모리 공간을 버퍼라고 부릅니다. 이번 예시에서는 CPU의 입력·결과 버퍼와 GPU의 x·u·y 버퍼를 별도로 둡니다.

다음은 그림 1의 CPU 쪽 흐름을 간단히 옮긴 코드입니다. CPU에서 입력 x를 준비한 뒤 시작한다고 가정합니다. A와 B는 앞서 설명한 계산이며, A<<<...>>>(...)는 GPU에 A의 실행을 요청하는 부분입니다. 각 호출의 세부 인자와 오류 처리, 자원 해제는 생략했습니다.

// GPU에 입력 x를 담을 공간을 확보합니다.
cudaMalloc(...);
// GPU에 중간값 u를 담을 공간을 확보합니다.
cudaMalloc(...);
// GPU에 최종 결과 y를 담을 공간을 확보합니다.
cudaMalloc(...);

// CPU에서 준비한 입력 x를 GPU로 복사합니다.
cudaMemcpy(...);

// GPU에 계산 A를 요청합니다: x에 1을 더해 u를 만듭니다.
A<<<...>>>(...);
// GPU에 계산 B를 요청합니다: u를 두 배로 만들어 y를 만듭니다.
B<<<...>>>(...);

// GPU의 결과 y를 CPU로 복사하고, 완료될 때까지 기다립니다.
cudaMemcpy(...);
// 이제 CPU에서 결과 y를 사용할 수 있습니다.

공간을 확보하는 것과 그곳에 데이터를 채우는 것은 별개의 작업입니다. GPU에 입력 x의 공간을 마련한 뒤에는 CPU가 준비한 입력을 그곳으로 복사해야 합니다. 중간값 u와 결과 y의 공간에는 각각 A와 B가 계산한 값을 써넣습니다. 원소 네 개는 이 흐름을 따라가기 위한 작은 예시입니다.

그림 1의 아래쪽에서 중간값 u는 GPU 메모리에 계속 남아 있습니다. A와 B를 모두 CPU가 호출하더라도, u를 CPU로 가져왔다가 다시 GPU로 보낼 필요는 없습니다. 다음 GPU 계산에 필요한 데이터는 GPU에 유지하고, CPU에서 사용할 최종 결과만 가져올 수 있습니다.

스트림에 작업을 등록하고 순서대로 실행하기

CPU 코드에 A 호출 다음 B 호출이 있다고 해서, CPU가 A의 계산 완료까지 기다린 뒤 B를 호출하는 것은 아닙니다. 커널 실행 호출은 GPU 작업을 요청한 뒤, 그 계산이 끝나기 전에 CPU 코드가 계속 진행할 수 있는 비동기 호출입니다. 함수 호출이 끝나 다음 줄로 돌아오는 것을 ‘반환’이라고 하는데, 커널 호출의 반환은 GPU 계산의 완료를 뜻하지 않습니다. CUDA API의 동기화 동작

그렇다면 A가 끝나기 전에 B를 호출해도 순서가 지켜지는 이유는 무엇일까요? CUDA에서는 복사나 커널 실행 같은 작업을 스트림(stream)에 등록합니다. 스트림은 등록된 GPU 작업을 순서대로 실행하도록 지정하는 흐름입니다. 앞의 코드처럼 스트림을 별도로 지정하지 않은 호출은 기본 스트림을 사용합니다.

그림 2는 입력 복사가 완료된 뒤 A와 B를 실행하는 부분입니다. 위쪽은 CPU가 작업을 등록하는 시간, 아래쪽은 GPU가 실제로 계산하는 시간입니다. 세로선이 가리키는 시점에는 CPU가 B의 등록까지 마쳤지만, GPU는 아직 A를 실행하고 있습니다. B는 같은 스트림에서 자신의 차례를 기다립니다.

이 예시에서 같은 스트림의 B는 A가 끝난 뒤 실행되므로, A가 만든 u를 읽을 수 있습니다. 입력 복사도 앞에 등록했다면 복사 → A → B라는 순서로 연결됩니다. CPU가 각 작업의 완료를 하나씩 확인해야 이 순서가 유지되는 것은 아닙니다. CUDA 스트림의 실행 순서

여기서는 CPU와 GPU의 진행이 겹치는 것과 GPU의 A와 B가 겹쳐 실행되는 것을 구별해야 합니다. CPU는 다음 작업을 등록하거나 다른 일을 할 수 있지만, 같은 스트림의 A와 B는 차례로 실행됩니다.

CPU에서 결과가 필요할 때 기다리기

GPU 안에서 다음 계산을 이어갈 때는 스트림의 순서로 결과가 준비되는 시점을 연결할 수 있습니다. 이번에는 CPU가 최종 결과 y의 실제 값을 읽어 사용해야 한다고 해보겠습니다. 그러려면 B의 계산이 끝나고, 그 결과가 CPU 메모리로 복사되어야 합니다.

앞의 코드에서 마지막 단계가 이 GPU→CPU 결과 복사입니다. 여기서 사용한 복사 호출(cudaMemcpy)은 이 방향의 복사가 완료될 때까지 CPU를 기다리게 합니다. 따라서 CPU 코드가 그다음 줄로 진행하면 CPU 메모리에 도착한 y를 사용할 수 있습니다. 복사 방법에 따라 CPU가 기다리는 방식은 달라질 수 있으며, 여기서는 결과 복사가 끝난 뒤 다음 줄로 진행하는 경우를 살펴봅니다. 복사 방향별 반환 조건

그림 3에서 CPU가 결과 복사를 호출했을 때 GPU는 아직 A를 계산하고 있습니다. 복사 작업은 B의 결과가 필요하므로 A와 B가 끝난 뒤 진행됩니다. 따라서 CPU가 복사 호출 안에 머무는 시간에는 실제 데이터 전달 시간뿐 아니라 앞선 GPU 계산이 끝나기를 기다리는 시간도 포함될 수 있습니다.

이때 기다리는 것은 해당 호출을 실행한 CPU 스레드입니다. GPU는 그동안 A와 B를 계산하고 데이터를 전달합니다. 따라서 CPU 대기가 길게 보인다는 사실만으로 GPU가 쉬고 있었다고 판단할 수는 없습니다. CPU의 대기 시간과 GPU의 실행 시간도 서로 겹치므로 둘을 더해 전체 시간을 구할 수 없습니다.

완료를 기다리는 동작과 데이터를 복사하는 동작도 구별해야 합니다. CPU가 특정 스트림에 앞서 등록한 작업이 모두 끝날 때까지 기다리게 할 수도 있습니다. 이렇게 완료를 기다리기만 해서는 GPU 데이터가 CPU 메모리로 복사되지 않습니다. CPU에서 데이터를 읽으려면 필요한 복사도 요청해야 합니다. CUDA에서 스트림 완료 대기에 사용하는 함수가 cudaStreamSynchronize입니다. 스트림 완료 대기 API

여러 스트림에 독립적인 작업 나누기

A와 B 사이에는 지켜야 할 순서가 있습니다. 하지만 별도 입력과 출력 버퍼를 사용하는 계산 C는 A의 결과를 읽지 않고, A·B가 사용하는 데이터도 변경하지 않는다고 해보겠습니다. 그러면 C를 진행하기 위해 A가 끝날 때까지 기다릴 필요는 없습니다.

이런 작업은 서로 다른 스트림에 배치할 수 있습니다. 스트림을 만들고 각 작업을 어느 스트림에 등록할지 정하는 주체는 개발자 또는 프레임워크·라이브러리입니다. 이번 예시에서는 CPU 코드가 같은 GPU에 사용할 스트림 S₀와 S₁을 만들고, S₀에 A를, S₁에 C와 B를 배치하겠습니다. GPU가 계산 내용을 분석해서 자동으로 두 스트림으로 나누는 과정은 아닙니다.

스트림 생성은 GPU의 SM을 둘로 나누거나 각 스트림의 전용 자원을 예약하는 동작이 아닙니다. 프로그램은 작업의 흐름을 지정하고, CUDA 실행 시스템은 실행 가능한 작업과 GPU 자원 상태에 따라 실행을 진행합니다. 커널의 블록이 실제 SM에 배정되는 과정은 지난 글에서 살펴본 실행 구조로 이어집니다.

서로 다른 스트림의 작업은 필요한 의존성이 충족되고 하드웨어와 실행 자원이 허용하면 겹쳐 실행될 수 있습니다. 먼저 실행된 커널이 SM의 상주 가능 공간이나 레지스터·공유 메모리 등을 많이 사용한다면 다른 커널이 기다릴 수도 있습니다. 스트림을 여러 개 만들었다고 동시 실행이 보장되지는 않습니다. 여러 커널의 동시 실행과 조건

겹쳐 실행되는 경우에도 성능이 반드시 좋아지는 것은 아닙니다. 두 작업이 같은 메모리 대역폭을 경쟁하면 각 작업의 시간이 길어질 수 있습니다. 독립적인 작업을 나누어 실행할 여지를 만들고, 실제로 전체 완료 시간이 줄었는지 확인해야 합니다.

이제 한 가지 조건이 남았습니다. C는 A와 독립적이지만, S₁의 B는 여전히 A의 결과 u를 필요로 합니다. 다른 스트림에 놓인 A와 B 사이의 순서는 따로 연결해야 합니다.

이벤트로 스트림 사이의 순서 연결하기

CUDA의 이벤트(event)는 스트림 안의 특정 지점까지 작업이 완료되었는지 확인하는 데 사용할 수 있습니다. 이번에는 S₀의 A 바로 뒤에 이벤트 E를 기록하도록 요청합니다. 이 이벤트가 완료되면 A까지 끝났다는 것을 알 수 있습니다. CPU가 이벤트 기록을 요청한 순간과 GPU 작업이 진행되어 이벤트가 완료되는 순간은 다릅니다. CUDA 이벤트 기록

S₁에는 C 다음에 이벤트 E 완료 대기를 넣고, 그 뒤에 B를 등록합니다. 그러면 C는 먼저 진행할 수 있고, B는 E가 완료된 뒤에만 실행될 수 있습니다.

스트림 S₀: 계산 A → 이벤트 E 기록
스트림 S₁: 계산 C → 이벤트 E 완료 대기 → 계산 B

그림 4의 위쪽은 CPU가 등록하는 작업의 순서이고, 아래쪽은 GPU에서 실제로 실행되는 시간을 나타냅니다. 아래에서 C가 먼저 끝났지만 A는 아직 실행 중입니다. 따라서 S₁은 이벤트 E 완료를 기다립니다. A가 끝나고 E가 완료되면 대기 조건이 충족되어 B가 실행될 수 있습니다.

다음 코드는 이 순서를 CPU에서 어떻게 요청하는지 보여줍니다. 앞의 단일 스트림 예시와 별개로, 두 스트림과 이벤트를 만들고 GPU 버퍼와 입력까지 준비한 상태에서 시작합니다. C는 A·B와 별도 데이터를 사용합니다. 스트림 선택을 포함한 커널 호출의 세부 인자는 생략했으므로, 주석에 적힌 스트림과 요청 순서를 따라 읽으면 됩니다.

// S₀에 계산 A를 요청합니다.
A<<<...>>>(...);

// S₀에서 A까지 끝났음을 표시할 이벤트 E를 기록하도록 요청합니다.
cudaEventRecord(E, S0);

// S₁에 독립적인 계산 C를 요청합니다.
C<<<...>>>(...);

// S₁의 뒤쪽 작업이 이벤트 E 완료를 기다리도록 지정합니다.
// CPU가 E 완료까지 기다리는 호출은 아닙니다.
cudaStreamWaitEvent(S1, E, ...);

// S₁에 계산 B를 요청합니다. B는 E가 완료된 뒤 실행됩니다.
B<<<...>>>(...);

// CPU는 E 완료를 기다리지 않고 다음 작업을 계속합니다.

코드에서 중요한 것은 함수 이름보다 대상을 구분하는 것입니다. 이벤트 기록은 S₀의 A가 끝났음을 표시하고, 이벤트 대기는 S₁의 B가 시작할 조건을 정합니다. CPU는 이 조건을 등록하므로 E가 완료되기 전에도 B를 등록하고 다음 CPU 작업을 계속할 수 있습니다. 스트림의 이벤트 대기

이벤트가 u를 복사하는 것도 아닙니다. u는 같은 GPU의 메모리에 남아 있고, 이벤트는 B가 u를 읽기 전에 A가 쓰기를 마치도록 실행 순서를 연결합니다. 그림의 기다리는 구간이 짧아지거나 없어져도 이 관계는 유지됩니다. 예를 들어 C가 끝나기 전에 E가 이미 완료되었다면, B는 C 뒤에서 E를 추가로 기다릴 필요가 없습니다.

그림 3과 비교하면 대기의 주체가 분명해집니다. 그림 3에서는 결과를 사용할 CPU 스레드가 복사 완료까지 기다렸습니다. 그림 4에서는 CPU가 실행 조건을 등록하고, 그 조건의 영향을 받는 GPU 후속 작업 B가 기다립니다. 나중에 CPU에서 y를 읽거나 버퍼를 재사용·해제할 때는 그 용도에 필요한 작업 완료와 데이터 전달을 별도로 보장해야 합니다.

이제 CPU의 호출, GPU의 실행, 결과 사용을 위한 완료 확인을 하나의 흐름으로 연결할 수 있습니다. 이 구분은 성능을 볼 때도 필요합니다. CPU 호출에 걸린 시간만으로 GPU 계산 시간을 판단할 수 없고, 여러 스트림을 사용했다는 사실만으로 병렬 실행의 효과를 알 수도 없습니다. 다음 글에서는 이 실행 흐름 위에서 계산량과 데이터 이동, 중간 결과의 재사용을 살펴보며 GPU 작업을 더 효율적으로 만드는 방법으로 이어가겠습니다.