junyeokk
Blog
Parallel Computing·2025. 10. 23

GPU 벡터 덧셈으로 보는 병렬 처리

왜 벡터 덧셈부터 시작하는가

병렬 프로그래밍을 배울 때 가장 먼저 다루는 예제가 벡터 덧셈이다. 복잡한 머신러닝 모델도, 과학 시뮬레이션도 결국 기본 연산의 조합이다. 벡터 덧셈은 이런 기본 연산 중 가장 단순하면서도 병렬화의 핵심 원리를 담고 있다.

병렬 프로그래밍의 첫 걸음은 문제를 독립적인 작은 조각으로 나누는 것이다. 벡터 덧셈은 이 개념을 가장 명확하게 보여준다. 각 원소의 계산이 완전히 독립적이므로, 병렬화를 배우기에 이상적이다.

벡터(Vector)란?
수학에서 벡터는 크기와 방향을 가진 양이다. 하지만 프로그래밍에서는 같은 타입의 데이터를 일렬로 나열한 1차원 배열을 의미한다. [1, 2, 3, 4, 5]처럼 연속된 메모리 공간에 저장된다.

문제 정의

벡터 덧셈은 두 배열의 같은 위치에 있는 원소끼리 더하는 연산이다. 수학적으로는 다음처럼 표기한다.

A=[a0,a1,a2,…,an−1]B=[b0,b1,b2,…,bn−1]C=A+B=[a0+b0,a1+b1,a2+b2,…,an−1+bn−1]\begin{aligned} A &= [a_0, a_1, a_2, \ldots, a_{n-1}] \\ B &= [b_0, b_1, b_2, \ldots, b_{n-1}] \\ C &= A + B = [a_0+b_0, a_1+b_1, a_2+b_2, \ldots, a_{n-1}+b_{n-1}] \end{aligned}

코드로 표현하면 C[i]=A[i]+B[i]C[i] = A[i] + B[i]다. 모든 인덱스 ii에 대해 이 연산을 수행한다. 벡터 덧셈의 핵심은 데이터 독립성이다. C[0]C[0]을 계산할 때 C[1]C[1]의 값이 필요하지 않다. C[100]C[100]을 계산할 때도 다른 원소가 필요 없다. 모든 원소가 독립적이므로 동시에 계산할 수 있다.

반대 예시로 누적 합을 보자. sum[1]=sum[0]+A[1]\text{sum}[1] = \text{sum}[0] + A[1]처럼 이전 결과에 의존하므로 단순 병렬화가 불가능하다. 각 단계가 이전 단계의 완료를 기다려야 한다. 이것이 벡터 덧셈을 첫 예제로 선택한 이유다.

문제 규모 설정

병렬화할 연산이 정해졌다. 이제 얼마나 큰 데이터로 실험할지 결정한다. 이번 예제는 1,048,5761{,}048{,}576개 (=1024×1024= 1024 \times 1024) 원소를 사용한다. 이 숫자를 선택한 이유는 병렬화 효과를 체감할 수 있으면서도 GPU 하드웨어와 잘 맞아떨어지기 때문이다. 작은 데이터는 메모리 전송 오버헤드에 묻혀버린다. 반면 1024는 GPU 블록 크기의 배수이고, 대부분 GPU에서 블록당 최대 1024개 스레드를 허용한다. 1024×1024=1,048,5761024 \times 1024 = 1{,}048{,}576개 스레드로 완벽하게 분할된다.

메모리 사용량을 계산하면 float 타입은 4바이트이므로, 1개 벡터가 약 4MB다. 세 개의 벡터 A, B, C가 필요하므로 CPU 메모리 12MB, GPU 메모리 12MB로 총 24MB를 사용한다. 현대 GPU는 수 GB VRAM을 가지므로 여유롭게 처리 가능하다.

테스트 데이터 준비

데이터 크기가 정해졌으니 실제 배열에 채울 값을 준비한다. 벡터 연산의 정확성을 검증하려면 의미 있는 테스트 데이터가 필요하다. 모든 원소를 0으로 초기화하면 버그를 발견하기 어렵다. C 표준 라이브러리의 rand() 함수로 [0.000, 1.000] 범위의 난수를 생성한다.

c
#include <stdlib.h>

float num = (rand() % 1000) / 1000.0F;

rand()는 0부터 RAND_MAX까지의 정수를 반환한다. rand() % 1000으로 0999 사이의 정수를 구하고, / 1000.0F로 나누면 0.0000.999 사이의 실수가 된다. F 접미사는 float 리터럴을 의미한다. 이것이 없으면 double 타입으로 계산되어 불필요한 형변환이 발생한다.

의사 난수(Pseudo-Random Number)란?
진짜 난수가 아니라 수학적 알고리즘으로 생성한 난수다. 같은 시드 값으로 시작하면 항상 같은 수열을 생성한다. 이 특성 덕분에 디버깅과 결과 재현이 가능하다.

srand() 함수로 난수 생성기의 시드를 설정할 수 있다. 디버깅 시에는 고정된 시드를 사용하여 결과를 재현한다. 프로덕션 환경에서는 time(NULL)을 시드로 사용하여 매번 다른 값을 생성한다.

배열 전체를 난수로 채우는 헬퍼 함수를 작성한다.

c
void setRandomData(float* dst, int size) {
    while (size--) {
        *dst++ = (rand() % 1000) / 1000.0F;
    }
}

포인터를 증가시키며 배열을 순회한다. size--는 후위 감소 연산자로, 0이 되면 루프가 종료된다. *dst++는 현재 위치에 값을 쓴 후 포인터를 다음 위치로 이동한다.

CPU 버전 구현

테스트 데이터 준비가 끝났다. 먼저 CPU에서 순차적으로 벡터를 더하는 기준 구현을 작성한다. 이 결과가 GPU 버전의 정확성을 검증하는 기준이 된다. CPU에서는 for 루프를 사용해 순차적으로 벡터를 더한다. 단일 코어는 한 번에 하나의 덧셈만 수행할 수 있다.

c
const unsigned SIZE = 1024 * 1024;

int main(void) {
    float* vecA = new float[SIZE];
    float* vecB = new float[SIZE];
    float* vecC = new float[SIZE];

    srand(0);
    setRandomData(vecA, SIZE);
    setRandomData(vecB, SIZE);

    chrono::system_clock::time_point time_begin = chrono::system_clock::now();
    for (register unsigned i = 0; i < SIZE; ++i) {
        vecC[i] = vecA[i] + vecB[i];
    }
    chrono::system_clock::time_point time_end = chrono::system_clock::now();

    float sumA = getSum(vecA, SIZE);
    float sumB = getSum(vecB, SIZE);
    float sumC = getSum(vecC, SIZE);
    float diff = fabsf(sumC - (sumA + sumB));

    delete[] vecA;
    delete[] vecB;
    delete[] vecC;

    return 0;
}

메모리 할당 후 난수로 배열을 채운다. chrono 라이브러리로 시간 측정을 시작한다. for 루프가 순차적으로 모든 원소를 처리한다. 시간 측정을 종료하고 결과를 검증한다.

결과 검증은 각 벡터의 모든 원소를 합한 값을 비교한다. sumC가 sumA + sumB와 일치하면 계산이 올바르게 수행된 것이다. 부동소수점 연산은 정확하지 않으므로, diff가 0이 아니라 매우 작은 값(예: 0.0001 이하)인지 확인한다.

Intel Core i5-3570에서 실행하면 약 1,845 usec가 소요된다. 이것이 GPU 가속의 기준점이 된다.

GPU 단일 코어 실험

CPU 기준 구현이 완성되었다. 이제 GPU로 넘어가는데, 본격적인 병렬 구현 전에 단일 코어로 먼저 실험한다. 병렬 버전을 작성하기 전에 GPU 단일 코어 성능을 측정한다. 이 실험으로 메모리 전송 오버헤드가 얼마나 큰지 파악하고, GPU 단일 코어 성능을 확인하며, 병렬화 이전에 기본 로직을 검증한다.

c
__global__ void singleKernelVecAdd(float* c, const float* a, const float* b) {
    for (register unsigned i = 0; i < SIZE; ++i) {
        c[i] = a[i] + b[i];
    }
}

__global__ 키워드는 이 함수가 GPU에서 실행되는 커널임을 나타낸다. 내부 로직은 CPU 버전과 동일한 for 루프다. 포인터 매개변수는 GPU 메모리 주소를 가리킨다.

커널 호출은 다음처럼 한다.

c
ELAPSED_TIME_BEGIN(0);
singleKernelVecAdd <<<1, 1>>>(dev_vecC, dev_vecA, dev_vecB);
cudaDeviceSynchronize();
ELAPSED_TIME_END(0);

<<<1, 1>>> 구문은 실행 설정을 지정한다. 첫 번째 인자는 그리드의 블록 개수, 두 번째는 블록의 스레드 개수다. <<<1, 1>>>은 1개 블록, 1개 스레드를 의미한다. cudaDeviceSynchronize()는 커널 실행이 완료될 때까지 CPU를 대기시킨다. CUDA 커널은 비동기로 실행되므로, 명시적으로 동기화해야 정확한 시간을 측정할 수 있다.

GeForce RTX 2070에서 60,436 usec가 소요된다. CPU보다 32배 느리다. 이유는 GPU 단일 코어의 낮은 클럭(1.6 GHz vs CPU 3.4 GHz)과 메모리 전송 오버헤드(총 16MB 전송) 때문이다. GPU는 병렬 처리에서만 의미가 있다.

GPU 병렬 버전 설계

단일 코어 실험으로 기본 로직과 메모리 전송을 확인했다. 이제 GPU의 진짜 장점인 병렬 처리를 적용한다. 벡터 덧셈의 병렬화 전략은 명확하다. 하나의 원소를 하나의 스레드가 처리한다. 모든 스레드가 동시에 실행되면, 백만 번의 덧셈이 한 번에 완료된다.

1백만 개 원소를 처리하려면 1백만 개 스레드가 필요하다. kernelVecAdd <<<1, SIZE>>>처럼 한 블록에 모두 담고 싶지만, 이 코드는 "invalid configuration argument" 에러를 발생시킨다.

SM(Streaming Multiprocessor)란?
GPU의 물리적 연산 유닛이다. 각 SM은 제한된 수의 스레드만 동시에 실행할 수 있다. 대부분의 GPU에서 블록당 최대 1024개 스레드만 허용한다. 이는 하드웨어 제약이므로 소프트웨어로 변경할 수 없다.

단일 블록에 1백만 개 스레드를 배치할 수 없다. 여러 블록으로 나누어야 한다. 1백만 개 스레드를 1024개 블록으로 나누고, 각 블록에 1024개 스레드를 배치한다.

c
kernelVecAdd <<<SIZE/1024, 1024>>>(dev_vecC, dev_vecA, dev_vecB, SIZE);

그리드는 1024개 블록을 가지고, 각 블록은 1024개 스레드를 가진다. 총 1,048,576개 스레드가 생성된다. 각 블록이 독립적으로 SM에 할당되어 실행된다. GPU가 여러 SM을 가지므로, 여러 블록이 동시에 처리된다.

스레드 인덱스 계산

블록과 스레드 구조가 정해졌다. 각 스레드가 어떤 배열 원소를 처리할지 결정하는 인덱스 계산이 필요하다. 각 스레드는 자신이 처리할 배열 인덱스를 계산해야 한다.

c
__global__ void kernelVecAdd(float* c, const float* a, const float* b, unsigned n) {
    unsigned i = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

인덱스 계산 공식을 분해하면 blockIdx.x는 현재 블록의 인덱스(01023), blockDim.x는 블록당 스레드 개수(1024), threadIdx.x는 블록 내 스레드 인덱스(01023)다. 예를 들어 블록 2의 스레드 5라면 다음과 같다.

i=2×1024+5=2053i = 2 \times 1024 + 5 = 2053

이 스레드는 C[2053]=A[2053]+B[2053]C[2053] = A[2053] + B[2053]을 계산한다.

if (i < n) 조건문은 배열 범위를 벗어나는 접근을 방지한다. 스레드 개수가 배열 크기와 정확히 일치하지 않을 수 있기 때문이다. 예를 들어 배열 크기가 1,000,0001{,}000{,}000이고 블록 크기가 1024라면 977개 블록이 필요하고, 실제 스레드는 1,000,4481{,}000{,}448개가 생성된다. 마지막 블록의 448개 스레드는 유효한 데이터가 없다. 경계 검사 없이 메모리에 접근하면 segmentation fault가 발생한다.

커널 함수의 매개변수를 살펴보자. float* c는 결과를 저장할 배열(쓰기), const float* a와 const float* b는 입력 배열(읽기 전용)이다. const 키워드는 이 포인터가 가리키는 데이터를 수정하지 않음을 나타낸다. 컴파일러가 최적화할 수 있고, 의도하지 않은 수정을 방지한다. unsigned n은 배열의 원소 개수로 경계 검사에 사용한다.

CUDA 실행 흐름

커널 함수 구현이 끝났다. 이 커널을 실행하기 위한 전체 흐름을 정리한다. CUDA 프로그램은 정해진 순서로 실행된다. CPU 메모리에 배열을 할당하고(Step 1), GPU 메모리에 배열을 할당한 뒤(Step 2), CPU에서 GPU로 데이터를 전송한다(Step 3). 그 다음 커널을 실행하고(Step 4), GPU에서 CPU로 결과를 전송한 후(Step 5), 메모리를 해제한다(Step 6).

c
// Step 1: Host 메모리 할당과 초기화
float* vecA = new float[SIZE];
float* vecB = new float[SIZE];
float* vecC = new float[SIZE];
srand(0);
setNormalizedRandomData(vecA, SIZE);
setNormalizedRandomData(vecB, SIZE);

// Step 2: Device 메모리 할당
float* dev_vecA = nullptr;
float* dev_vecB = nullptr;
float* dev_vecC = nullptr;
cudaMalloc((void**)&dev_vecA, SIZE * sizeof(float));
cudaMalloc((void**)&dev_vecB, SIZE * sizeof(float));
cudaMalloc((void**)&dev_vecC, SIZE * sizeof(float));

// Step 3: Host → Device 데이터 전송
cudaMemcpy(dev_vecA, vecA, SIZE * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(dev_vecB, vecB, SIZE * sizeof(float), cudaMemcpyHostToDevice);

// Step 4: 커널 실행
kernelVecAdd <<<SIZE/1024, 1024>>>(dev_vecC, dev_vecA, dev_vecB, SIZE);
cudaDeviceSynchronize();

// Step 5: Device → Host 결과 전송
cudaMemcpy(vecC, dev_vecC, SIZE * sizeof(float), cudaMemcpyDeviceToHost);

// Step 6: 메모리 해제
cudaFree(dev_vecA);
cudaFree(dev_vecB);
cudaFree(dev_vecC);
delete[] vecA;
delete[] vecB;
delete[] vecC;

cudaMalloc()은 GPU의 전역 메모리에 공간을 예약한다. 포인터 변수 자체는 CPU 메모리에 있지만, 가리키는 주소는 GPU 메모리다. dev_ 접두사는 이 포인터가 device 메모리를 가리킴을 나타내는 관례다.

cudaMemcpy()로 host 데이터를 device로 복사한다. 마지막 인자가 전송 방향을 지정한다. 이 전송은 PCI Express 버스를 통해 이루어진다. 버스 대역폭이 제한되어 있으므로 시간이 소요된다. GPU 프로그래밍에서 메모리 전송이 주요 병목이다.

성능 분석

전체 구현이 완료되었다. CPU 버전, GPU 단일 코어, GPU 병렬 버전의 성능을 비교한다. GeForce RTX 2070 기준으로 측정한 결과를 보자.

plain
CPU (Intel Core i5-3570):        1,891 usec
CUDA 단일 코어:                 60,436 usec
CUDA 1K 블록 (커널만):             118 usec
CUDA 1K 블록 (메모리 전송 포함):   5,944 usec

순수 커널 실행 시간만 비교하면 CPU 대비 16배 빠르다.

1,891118=16.0\frac{1{,}891}{118} = 16.0

하지만 메모리 전송까지 포함하면 오히려 CPU보다 느려진다.

1,8915,944=0.32\frac{1{,}891}{5{,}944} = 0.32

메모리 전송 비용이 성능 향상을 모두 상쇄했다.

전체 시간 5,944 usec를 분해하면 커널 실행이 118 usec(2%), 메모리 전송이 5,826 usec(98%)다. 전체 시간의 98%가 메모리 전송에 소비된다. 계산 자체는 매우 빠르지만, 데이터를 옮기는 데 시간이 걸린다.

작은 데이터는 GPU에 부적합하다. 메모리 전송 오버헤드가 너무 크다. 데이터가 클수록 GPU의 장점이 커진다. 데이터를 한 번 전송하고 여러 연산을 수행하여 전송 비용을 여러 연산에 분산시켜야 한다. 메모리 접근 대비 계산량이 많을수록 유리하다. 벡터 덧셈은 계산 강도가 낮은 편이다.

단일 GPU 코어는 CPU보다 훨씬 느리다(32배). GPU의 장점은 단일 코어 성능이 아니라 수천 개 코어를 동시에 사용할 수 있다는 점이다. CPU는 4개 또는 8개 코어를 가진다. GPU는 수천 개 코어를 가진다. GeForce RTX 2070은 2304개 CUDA 코어를 탑재한다.

C++ 템플릿으로 일반화

기본 구현이 완성되었다. 실무에서는 다양한 데이터 타입을 처리해야 하므로 코드를 일반화한다. float 대신 double이나 int를 사용하려면 어떻게 해야 할까. 코드를 복사해서 타입만 바꾸는 것은 유지보수가 어렵다. C++ 템플릿을 사용하면 타입에 독립적인 코드를 작성할 수 있다.

c
template<typename TYPE>
__global__ void kernelVecAdd(TYPE* c, const TYPE* a, const TYPE* b, unsigned n) {
    unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

// 호출 시 타입을 명시
kernelVecAdd<float> <<<SIZE/1024, 1024>>>(dev_vecC, dev_vecA, dev_vecB, SIZE);
kernelVecAdd<double> <<<SIZE/1024, 1024>>>(dev_vecC, dev_vecA, dev_vecB, SIZE);

컴파일러가 각 타입에 대해 별도의 커널을 생성한다. 코드는 하나지만 실행 파일에는 여러 버전이 포함된다.

템플릿으로 타입을 일반화했다. 이제 코드 구조를 개선한다. 각 단계를 클래스로 캡슐화하면 코드 가독성이 높아진다.

c
class VecAdd {
protected:
    const unsigned SIZE = 1024 * 1024;
    float* vecA;
    float* vecB;
    float* vecC;
    float* dev_vecA;
    float* dev_vecB;
    float* dev_vecC;

public:
    void prepare_host(void) {
        vecA = new float[SIZE];
        vecB = new float[SIZE];
        vecC = new float[SIZE];
        setNormalizedRandomData(vecA, SIZE);
        setNormalizedRandomData(vecB, SIZE);
    }

    void prepare_device(void) {
        cudaMalloc((void**)&dev_vecA, SIZE * sizeof(float));
        cudaMalloc((void**)&dev_vecB, SIZE * sizeof(float));
        cudaMalloc((void**)&dev_vecC, SIZE * sizeof(float));
    }

    void copy_to_device(void) {
        cudaMemcpy(dev_vecA, vecA, SIZE * sizeof(float), cudaMemcpyHostToDevice);
        cudaMemcpy(dev_vecB, vecB, SIZE * sizeof(float), cudaMemcpyHostToDevice);
    }

    void execute_kernel(void) {
        kernelVecAdd<float> <<<SIZE/1024, 1024>>>(dev_vecC, dev_vecA, dev_vecB, SIZE);
        cudaDeviceSynchronize();
    }

    void copy_to_host(void) {
        cudaMemcpy(vecC, dev_vecC, SIZE * sizeof(float), cudaMemcpyDeviceToHost);
    }

    void clear(void) {
        cudaFree(dev_vecA);
        cudaFree(dev_vecB);
        cudaFree(dev_vecC);
        delete[] vecA;
        delete[] vecB;
        delete[] vecC;
    }
};

int main(void) {
    VecAdd vecAdd;
    vecAdd.prepare_host();
    vecAdd.prepare_device();
    vecAdd.copy_to_device();
    vecAdd.execute_kernel();
    vecAdd.copy_to_host();
    vecAdd.clear();
    return 0;
}

각 단계가 명확하게 분리되어 이해하기 쉽다. 템플릿 버전의 실행 시간은 129 usec로 C 버전(118 usec)과 거의 동일하다. 템플릿 오버헤드는 컴파일 타임에 해소되므로 런타임 성능에 영향이 없다.

병렬화의 핵심 원칙

벡터 덧셈 구현과 성능 분석을 마쳤다. 여기서 배운 핵심 원칙들을 정리한다. 벡터 덧셈이 병렬화에 적합한 이유는 각 원소의 계산이 완전히 독립적이기 때문이다. 병렬화 가능 여부를 판단하는 질문은 간단하다. "이 계산을 수행할 때 다른 계산의 결과가 필요한가?" 대답이 "아니오"라면 병렬화할 수 있다.

SM당 최대 스레드 수 제한은 모든 CUDA 프로그램에서 고려해야 할 요소다. 블록당 1024개 스레드는 하드웨어 한계다. 이 제약을 넘으면 에러가 발생한다. 1D, 2D, 3D 레이아웃 모두 이 제약을 만족해야 한다.

2D: blockDim.x×blockDim.y≤1024\text{2D: } blockDim.x \times blockDim.y \leq 1024 3D: blockDim.x×blockDim.y×blockDim.z≤1024\text{3D: } blockDim.x \times blockDim.y \times blockDim.z \leq 1024

메모리 전송 비용은 GPU 프로그래밍 최적화의 핵심이다. 작은 데이터를 여러 번 전송하면 오버헤드가 커진다. 여러 작은 전송보다 하나의 큰 전송이 효율적이다. cudaMemcpy 호출 횟수를 줄여야 한다. 데이터를 전송한 후 GPU에서 모든 처리를 완료하고, 중간 결과를 CPU로 가져오지 않는다. 메모리 접근 대비 계산량을 늘리고, 같은 데이터로 더 많은 연산을 수행한다.

다음 단계로

여기까지가 CUDA 병렬 프로그래밍의 기본이다. 벡터 덧셈으로 CUDA 병렬 프로그래밍의 기본 구조를 배웠다. 메모리 할당, 데이터 전송, 커널 실행, 결과 회수의 전체 흐름을 이해했다. 하지만 1백만 개 원소에서는 메모리 전송 오버헤드가 너무 크다. 다음 장에서는 데이터 크기를 10억 개로 늘려 GPU의 진가를 확인한다. 데이터가 클수록 GPU의 장점이 극대화된다. 또한 AXPY와 FMA 같은 더 복잡한 벡터 연산을 다룬다. 계산 강도가 높아지면 메모리 전송 비용이 상대적으로 줄어든다.

관련 문서