대규모 벡터 연산

왜 데이터 규모를 늘리는가

이전 예제에서 1백만 개 원소로 GPU 병렬 처리의 기본을 배웠다. 하지만 측정 결과가 실망스러웠다. 메모리 전송 오버헤드가 계산 시간보다 훨씬 커서, 전체적으로는 CPU보다 느렸다.

GPU의 진가는 대규모 데이터에서 발휘된다. 데이터가 커질수록 메모리 전송 비용의 상대적 비중이 줄어들고, 병렬 계산의 이점이 극대화된다. 실제 응용에서는 수억, 수십억 개의 데이터 포인트를 처리하는 경우가 흔하다.

머신러닝 모델 학습

• 수백만 개의 파라미터를 반복적으로 업데이트한다
• 배치당 수만 개의 샘플을 처리한다

과학 시뮬레이션

• 유체 역학, 기후 모델링에서 수억 개의 격자점을 계산한다
• 매 시간 단계마다 모든 격자점을 업데이트한다

이미지/비디오 처리

• 4K 영상은 프레임당 약 830만 픽셀이다
• 초당 60프레임이면 5억 픽셀을 처리해야 한다

이번 장에서는 데이터 크기를 256배 늘려 GPU의 스케일링 특성을 확인한다.

데이터 규모 확장

규모를 키우기로 결정했다. 구체적으로 얼마나 늘릴지, 그리고 그에 따른 메모리 요구사항을 계산한다.

256M 원소로 확장

이전 예제는 1M 원소 (1024 × 1024 = 1,048,576)를 처리했다. 이번에는 256M 원소로 확장한다.

const unsigned SIZE = 256 * 1024 * 1024;  // 256M elements = 268,435,456

const unsigned SIZE = 256 * 1024 * 1024;  // 256M elements = 268,435,456

계산 과정은 다음과 같다.

256M = 256 × 1024 × 1024
     = 256 × 1,048,576
     = 268,435,456 원소

256M = 256 × 1024 × 1024
     = 256 × 1,048,576
     = 268,435,456 원소

1M에서 256M으로 늘어났으므로 정확히 256배다.

메모리 요구사항 계산

float 타입은 4바이트를 차지한다. 세 개의 벡터 A, B, C가 필요하므로 메모리 계산은 다음과 같다.

1개 벡터 = 268,435,456 원소 × 4 bytes = 1,073,741,824 bytes = 1 GB
3개 벡터 = 3 GB

Host 메모리: 3 GB
Device 메모리: 3 GB (복사본)
총 메모리 요구량: 6 GB

1개 벡터 = 268,435,456 원소 × 4 bytes = 1,073,741,824 bytes = 1 GB
3개 벡터 = 3 GB

Host 메모리: 3 GB
Device 메모리: 3 GB (복사본)
총 메모리 요구량: 6 GB

GeForce RTX 2070은 8GB VRAM을 가지므로 충분히 처리 가능하다. 시스템 RAM도 8GB 이상이어야 한다.

실제 응용과의 비교

256M 원소는 실제로는 그리 큰 데이터가 아니다.

GPT-3 모델

• 1750억 개의 파라미터를 가진다
• 이는 256M보다 약 650배 크다

4K 비디오

• 3840 × 2160 픽셀 = 약 830만 픽셀/프레임
• 256M 원소는 약 30프레임에 해당한다
• 초당 60프레임 비디오는 2초면 이 데이터량을 초과한다

기후 시뮬레이션

• 1km 해상도 지구 표면 모델은 약 5억 격자점이 필요하다
• 3D 대기 모델은 수십억 격자점을 사용한다

256M은 GPU 성능을 테스트하기에 적당한 중간 규모다.

기본 성능 비교

데이터 크기와 메모리 요구사항을 확인했다. 이제 실제로 실행하여 CPU와 GPU 성능을 비교한다.

CPU 순차 처리

먼저 기준이 될 CPU 성능을 측정한다. 동일한 for 루프 방식으로 256M 원소를 처리한다.

ELAPSED_TIME_BEGIN(0);
for (register unsigned i = 0; i < SIZE; ++i) {
    vecC[i] = vecA[i] + vecB[i];
}
ELAPSED_TIME_END(0);

ELAPSED_TIME_BEGIN(0);
for (register unsigned i = 0; i < SIZE; ++i) {
    vecC[i] = vecA[i] + vecB[i];
}
ELAPSED_TIME_END(0);

코드는 1M 원소일 때와 완전히 동일하다. SIZE 상수만 바뀌었다.

Intel Core i5-3570에서 459,271 usec가 소요된다. 약 0.46초다.

1M 원소일 때와 비교하면 다음과 같다.

1M 원소: 1,845 usec
256M 원소: 459,271 usec
증가 비율: 459,271 / 1,845 = 249 배

1M 원소: 1,845 usec
256M 원소: 459,271 usec
증가 비율: 459,271 / 1,845 = 249 배

데이터 크기가 256배 늘었는데 시간은 249배 늘었다. 거의 완벽한 선형 스케일링이다. 약간의 차이는 캐시 효율 차이 때문이다.

CPU는 데이터 크기와 관계없이 일정한 성능을 보인다. 1개 원소를 처리하든 10억 개 원소를 처리하든, 단위 시간당 처리량이 일정하다.

1M 원소: 1,845 usec → 원소당 1.76 ns
256M 원소: 459,271 usec → 원소당 1.71 ns

1M 원소: 1,845 usec → 원소당 1.76 ns
256M 원소: 459,271 usec → 원소당 1.71 ns

원소당 처리 시간이 거의 동일하다. CPU는 예측 가능한 성능을 제공한다.

GPU 단일 스레드

CPU 기준 성능을 확인했다. GPU로 넘어가기 전에 단일 스레드로 먼저 측정한다. GPU 단일 스레드 성능은 두 가지 정보를 제공한다.

병렬화 효과 측정

• 단일 스레드 대비 병렬 버전의 가속 비율을 알 수 있다
• GPU 코어 개수 대비 효율을 평가한다

메모리 전송 오버헤드 분리

• 단일 스레드도 동일한 메모리 전송을 수행한다
• 순수 계산 시간과 전송 시간을 구분할 수 있다

__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];
    }
}

singleKernelVecAdd <<<1, 1>>>(dev_vecC, dev_vecA, dev_vecB);

__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];
    }
}

singleKernelVecAdd <<<1, 1>>>(dev_vecC, dev_vecA, dev_vecB);

<<<1, 1>>>로 1개 블록, 1개 스레드로 실행한다. 내부는 CPU 버전과 동일한 순차 루프다.

GeForce RTX 2070에서 측정한 결과는 다음과 같다.

순수 커널 실행: 11,913,504 usec (약 11.9초)
메모리 전송 포함: 12,651,351 usec (약 12.7초)

순수 커널 실행: 11,913,504 usec (약 11.9초)
메모리 전송 포함: 12,651,351 usec (약 12.7초)

CPU 459,271 usec와 비교하면 GPU가 약 26배 느리다.

이유를 분석하면 다음과 같다.

클럭 속도 차이

• CPU: 3.4 GHz (Intel Core i5-3570)
• GPU: 1.62 GHz (GeForce RTX 2070)
• 클럭만 보면 2.1배 차이다

아키텍처 차이

• CPU는 순차 처리에 최적화되어 있다
• 분기 예측, 파이프라인, 캐시 등이 발달했다
• GPU는 병렬 처리에 최적화되어 단순한 코어를 사용한다

메모리 대역폭

• GPU 메모리는 병렬 접근에 최적화되어 있다
• 단일 스레드가 순차 접근하면 비효율적이다

이것이 GPU를 순차 처리에 사용하면 안 되는 이유다. GPU의 설계 철학은 "단순한 코어를 많이"다.

GPU 병렬 처리

단일 스레드의 참담한 성능을 확인했다. 이제 GPU의 본질인 병렬 처리를 적용한다. 256M 원소를 처리하려면 그리드와 블록을 재계산해야 한다.

dim3 dimBlock(1024, 1, 1);
dim3 dimGrid((SIZE + dimBlock.x - 1) / dimBlock.x, 1, 1);

kernelVecAdd <<<dimGrid, dimBlock>>>(dev_vecC, dev_vecA, dev_vecB, SIZE);

dim3 dimBlock(1024, 1, 1);
dim3 dimGrid((SIZE + dimBlock.x - 1) / dimBlock.x, 1, 1);

kernelVecAdd <<<dimGrid, dimBlock>>>(dev_vecC, dev_vecA, dev_vecB, SIZE);

계산 과정은 다음과 같다.

블록 크기: 1024 스레드
블록 개수: ceil(268,435,456 / 1024) = 262,144 블록
총 스레드: 262,144 × 1024 = 268,435,456 개

블록 크기: 1024 스레드
블록 개수: ceil(268,435,456 / 1024) = 262,144 블록
총 스레드: 262,144 × 1024 = 268,435,456 개

262,144개 블록이 생성된다. 이는 16진수로 0x40000이다.

왜 (SIZE + dimBlock.x - 1) / dimBlock.x인가?
단순히 SIZE / dimBlock.x를 사용하면 정수 나눗셈으로 내림 처리된다. 크기가 블록 크기의 정확한 배수가 아니면 일부 원소가 누락된다.

잘못된 방법:
SIZE = 1000, dimBlock.x = 1024
dimGrid.x = 1000 / 1024 = 0 블록
→ 아무것도 처리되지 않음!

올바른 방법:
dimGrid.x = (1000 + 1024 - 1) / 1024 = 2023 / 1024 = 1 블록
→ 1024개 스레드 생성, 1000개만 사용

잘못된 방법:
SIZE = 1000, dimBlock.x = 1024
dimGrid.x = 1000 / 1024 = 0 블록
→ 아무것도 처리되지 않음!

올바른 방법:
dimGrid.x = (1000 + 1024 - 1) / 1024 = 2023 / 1024 = 1 블록
→ 1024개 스레드 생성, 1000개만 사용

(a + b - 1) / b는 올림 나눗셈의 정수 버전이다. 수학적으로 ceil(a / b)와 동일하다.

커널 함수는 1M 원소 버전과 동일하다.

__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];
    }
}

__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];
    }
}

각 스레드가 자신의 인덱스를 계산하고 해당 원소만 처리한다. if (i < n) 조건은 마지막 블록에서 범위를 벗어나는 스레드를 걸러낸다.

코드의 아름다움은 데이터 크기와 무관하게 동일한 코드가 작동한다는 점이다. 1백만 개든 10억 개든 커널 함수는 바뀌지 않는다.

GeForce RTX 2070에서 측정한 결과는 다음과 같다.

순수 커널 실행: 8,136 usec (약 8.1 ms)
메모리 전송 포함: 746,448 usec (약 0.75초)

순수 커널 실행: 8,136 usec (약 8.1 ms)
메모리 전송 포함: 746,448 usec (약 0.75초)

CPU 459,271 usec와 비교하면 다음과 같다.

커널만: 459,271 / 8,136 = 56.4 배 빠름
전체: 459,271 / 746,448 = 0.62 배 (CPU가 더 빠름)

커널만: 459,271 / 8,136 = 56.4 배 빠름
전체: 459,271 / 746,448 = 0.62 배 (CPU가 더 빠름)

순수 계산 성능은 CPU 대비 56배 빠르다. 하지만 메모리 전송을 포함하면 오히려 CPU가 빠르다.

메모리 전송 시간을 계산하면 다음과 같다.

전체 시간: 746,448 usec
커널 시간: 8,136 usec
전송 시간: 746,448 - 8,136 = 738,312 usec

전송 비율: 738,312 / 746,448 = 98.9%
계산 비율: 8,136 / 746,448 = 1.1%

전체 시간: 746,448 usec
커널 시간: 8,136 usec
전송 시간: 746,448 - 8,136 = 738,312 usec

전송 비율: 738,312 / 746,448 = 98.9%
계산 비율: 8,136 / 746,448 = 1.1%

전체 시간의 99%가 메모리 전송이다. 계산은 1%에 불과하다.

1M 원소일 때와 256M 원소일 때를 비교하면 흥미로운 패턴이 보인다.

1M 원소    256M 원소    비율
커널:      118 usec   8,136 usec   69배
전송+커널: 5,944 usec 746,448 usec 126배

1M 원소    256M 원소    비율
커널:      118 usec   8,136 usec   69배
전송+커널: 5,944 usec 746,448 usec 126배

데이터가 256배 늘었는데 커널 시간은 69배만 늘었다. 전송 포함 시간은 126배 늘었다.

데이터가 커질수록 GPU 병렬 처리의 효율이 높아진다. 더 큰 데이터에서는 CPU를 추월할 것이다.

커널 내부 시간 측정

호스트에서 측정한 전체 커널 실행 시간을 확인했다. 이제 커널 내부의 세부적인 시간을 측정하는 방법을 알아본다.

clock() 함수를 이용한 측정

지금까지는 호스트에서 ELAPSED_TIME_BEGIN()과 ELAPSED_TIME_END()로 시간을 측정했다. 이는 전체 커널 실행 시간을 알려준다. 하지만 커널 내부의 특정 구간 시간은 알 수 없다.

복잡한 커널에서는 어느 부분이 느린지 파악해야 최적화할 수 있다. 이를 위해 GPU 내부에서 시간을 측정하는 방법이 필요하다.

clock() 함수란?
GPU 내부의 사이클 카운터를 읽는 CUDA 내장 함수다. clock_t clock(void)는 32비트, long long int clock64(void)는 64비트 카운터를 반환한다. CPU의 clock() 함수와는 다른 별개의 함수다. __device__와 global 함수에서만 호출 가능하다.

__global__ void kernelVecAdd(float* c, const float* a, const float* b,
                              unsigned n, long long* times) {
    clock_t start = clock();

    unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }

    clock_t end = clock();

    if (i == 0) {  // 스레드 0번만 결과 저장
        times[0] = (long long)(end - start);
    }
}

__global__ void kernelVecAdd(float* c, const float* a, const float* b,
                              unsigned n, long long* times) {
    clock_t start = clock();

    unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }

    clock_t end = clock();

    if (i == 0) {  // 스레드 0번만 결과 저장
        times[0] = (long long)(end - start);
    }
}

커널 시작 시 clock()으로 시작 클럭을 기록한다. 계산을 수행한 뒤 다시 clock()으로 종료 클럭을 기록한다. 차이를 계산하여 경과 클럭 수를 구한다.

스레드 0번만 결과를 저장하는 이유는 무엇일까? 모든 스레드가 저장하면 26만 개 스레드가 동시에 같은 메모리에 쓰기를 시도한다. 불필요한 경쟁이 발생하므로, 대표로 하나만 저장한다.

클럭-시간 변환

clock()이 반환하는 값은 클럭 틱 수다. 실제 시간으로 변환하려면 GPU 클럭 주파수가 필요하다.

int clk_freq = 1;
cudaDeviceGetAttribute(&clk_freq, cudaDevAttrClockRate, 0);
float elapsed_usec = clk_ticks * 1000.0F / clk_freq;

int clk_freq = 1;
cudaDeviceGetAttribute(&clk_freq, cudaDevAttrClockRate, 0);
float elapsed_usec = clk_ticks * 1000.0F / clk_freq;

cudaDeviceGetAttribute()는 GPU 속성 정보를 가져온다. cudaDevAttrClockRate는 클럭 주파수를 kHz 단위로 반환한다.

경과 시간 계산식은 다음과 같다.

경과 시간(usec) = (클럭 틱 수) × 1000 / (주파수 kHz)

경과 시간(usec) = (클럭 틱 수) × 1000 / (주파수 kHz)

왜 1000을 곱하는가?
주파수가 kHz 단위이므로 Hz로 변환하려면 1000을 곱해야 한다. 1 클럭 틱의 시간은 1 / (주파수 Hz) 초다. 이를 마이크로초로 바꾸면 1,000,000 / (주파수 Hz)다. kHz 단위를 사용하면 1,000,000 / (주파수 kHz × 1000) = 1000 / (주파수 kHz)가 된다.

측정 결과 분석

GeForce RTX 2070의 클럭은 약 1.62 GHz = 1,620,000 kHz다.

측정 결과 스레드 0번이 1,743 클럭 틱을 소비했다. 시간으로 변환하면 다음과 같다.

경과 시간 = 1,743 × 1000 / 1,620,000 = 1.075 usec

경과 시간 = 1,743 × 1000 / 1,620,000 = 1.075 usec

약 1 마이크로초다. 단일 스레드가 256M 원소 중 1개만 처리하므로 매우 짧다.

전체 커널 실행 시간 8,136 usec와 비교하면 흥미롭다. 26만 개 스레드가 동시에 실행되지 않고, 여러 warp로 나뉘어 시분할 실행되기 때문에 개별 스레드 시간과 전체 시간이 크게 차이난다.

런타임 크기 설정 (Argument 처리)

다양한 크기로 실험하려면 매번 소스를 수정하고 재컴파일해야 한다. 실행 시 크기를 지정할 수 있도록 개선한다.

명령행 인자 파싱

지금까지는 벡터 크기를 소스 코드에 하드코딩했다.

const unsigned SIZE = 256 * 1024 * 1024;

const unsigned SIZE = 256 * 1024 * 1024;

다른 크기를 테스트하려면 소스를 수정하고 다시 컴파일해야 한다. 불편하고 시간이 걸린다.

실행 시 명령행 인자로 크기를 지정하면 편리하다.

unsigned vecSize = 256 * 1024 * 1024;  // 기본값

int main(const int argc, const char* argv[]) {
    char* pEnd = nullptr;

    switch (argc) {
        case 1:
            // 인자 없음, 기본값 사용
            break;
        case 2:
            // 인자 하나, 크기 파싱
            vecSize = strtol(argv[1], &pEnd, 10);
            break;
        default:
            // 인자가 너무 많음
            printf("usage: %s [size]\\n", argv[0]);
            exit(EXIT_FAILURE);
            break;
    }

    printf("vecSize = %u\\n", vecSize);
    // ...
}

unsigned vecSize = 256 * 1024 * 1024;  // 기본값

int main(const int argc, const char* argv[]) {
    char* pEnd = nullptr;

    switch (argc) {
        case 1:
            // 인자 없음, 기본값 사용
            break;
        case 2:
            // 인자 하나, 크기 파싱
            vecSize = strtol(argv[1], &pEnd, 10);
            break;
        default:
            // 인자가 너무 많음
            printf("usage: %s [size]\\n", argv[0]);
            exit(EXIT_FAILURE);
            break;
    }

    printf("vecSize = %u\\n", vecSize);
    // ...
}

argc는 명령행 인자 개수다. 프로그램 이름을 포함하므로 최소값은 1이다. argv는 인자 문자열 배열이다.

strtol()은 문자열을 long 정수로 변환한다. 세 번째 인자 10은 10진수를 의미한다. pEnd는 파싱이 끝난 위치를 가리킨다.

실행 예시는 다음과 같다.

$ ./giga-add-arg.exe 512000000
vecSize = 512000000
dimBlock = 1024
dimGrid = 500000
total # thread = 512000000
elapsed wall-clock time[0] = 15422 usec

$ ./giga-add-arg.exe 512000000
vecSize = 512000000
dimBlock = 1024
dimGrid = 500000
total # thread = 512000000
elapsed wall-clock time[0] = 15422 usec

512M 원소 (512,000,000)를 처리하는 데 15,422 usec가 소요되었다.

256M일 때 8,136 usec였으므로 비율을 계산하면 다음과 같다.

데이터 증가: 512M / 256M = 2배
시간 증가: 15,422 / 8,136 = 1.9배

데이터 증가: 512M / 256M = 2배
시간 증가: 15,422 / 8,136 = 1.9배

거의 완벽한 선형 스케일링이다. GPU 병렬 처리는 데이터 크기에 비례하여 시간이 증가한다.

800M 원소를 시도하면 어떻게 될까?

$ ./giga-add-arg.exe 800000000
cudaMalloc error: out of memory

$ ./giga-add-arg.exe 800000000
cudaMalloc error: out of memory

에러가 발생한다. 이유는 메모리 부족이다.

800M 원소 × 4 bytes × 3 벡터 = 9.6 GB
RTX 2070 VRAM = 8 GB
부족한 메모리 = 1.6 GB

800M 원소 × 4 bytes × 3 벡터 = 9.6 GB
RTX 2070 VRAM = 8 GB
부족한 메모리 = 1.6 GB

GPU 메모리 용량을 초과했다. 이것이 GPU 프로그래밍의 현실적 제약이다.

K/M 단위 지원

명령행 인자로 크기를 지정할 수 있게 되었다. 하지만 큰 숫자를 직접 입력하기 불편하다.

$ ./giga-add-arg.exe 512000000  # 0이 7개, 세기 어려움

$ ./giga-add-arg.exe 512000000  # 0이 7개, 세기 어려움

사람이 읽기 쉬운 표기법이 필요하다.

$ ./giga-add-arg.exe 512M  # 훨씬 명확함

$ ./giga-add-arg.exe 512M  # 훨씬 명확함

템플릿 함수로 구현한다.

template <typename TYPE>
TYPE procArg(const char* progname, const char* str,
             TYPE lbound = -1, TYPE ubound = -1) {
    char* pEnd = nullptr;
    TYPE value = 0;

    // float/double vs 정수 처리
    if (typeid(TYPE) == typeid(float) || typeid(TYPE) == typeid(double)) {
        value = strtof(str, &pEnd);
    } else {
        value = strtol(str, &pEnd, 10);
    }

    // 단위 접미사 처리
    switch (*pEnd) {
        case 'M':
        case 'm':
            value *= (1024 * 1024);
            break;
        case 'K':
        case 'k':
            value *= 1024;
            break;
    }

    // 범위 검사
    if (lbound != (TYPE)(-1) && value < lbound) {
        fprintf(stderr, "Error: %s must be >= %lld\\n", str, (long long)lbound);
        exit(EXIT_FAILURE);
    }
    if (ubound != (TYPE)(-1) && value > ubound) {
        fprintf(stderr, "Error: %s must be <= %lld\\n", str, (long long)ubound);
        exit(EXIT_FAILURE);
    }

    return value;
}

template <typename TYPE>
TYPE procArg(const char* progname, const char* str,
             TYPE lbound = -1, TYPE ubound = -1) {
    char* pEnd = nullptr;
    TYPE value = 0;

    // float/double vs 정수 처리
    if (typeid(TYPE) == typeid(float) || typeid(TYPE) == typeid(double)) {
        value = strtof(str, &pEnd);
    } else {
        value = strtol(str, &pEnd, 10);
    }

    // 단위 접미사 처리
    switch (*pEnd) {
        case 'M':
        case 'm':
            value *= (1024 * 1024);
            break;
        case 'K':
        case 'k':
            value *= 1024;
            break;
    }

    // 범위 검사
    if (lbound != (TYPE)(-1) && value < lbound) {
        fprintf(stderr, "Error: %s must be >= %lld\\n", str, (long long)lbound);
        exit(EXIT_FAILURE);
    }
    if (ubound != (TYPE)(-1) && value > ubound) {
        fprintf(stderr, "Error: %s must be <= %lld\\n", str, (long long)ubound);
        exit(EXIT_FAILURE);
    }

    return value;
}

함수는 템플릿으로 구현하여 int와 float 모두 처리한다. typeid로 런타임에 타입을 확인한다.

strtol()로 숫자 부분을 파싱한 뒤 pEnd가 가리키는 문자를 확인한다.

• 'M' 또는 'm': 1024 × 1024를 곱한다 (mebibyte)
• 'K' 또는 'k': 1024를 곱한다 (kibibyte)
• 그 외: 숫자 그대로 사용한다

선택적으로 범위 검사도 수행한다. lbound와 ubound를 지정하면 값이 범위 내에 있는지 확인한다.

사용 예시는 다음과 같다.

case 2:
    vecSize = procArg<unsigned>(argv[0], argv[1], 1, 1024*1024*1024);
    break;

case 2:
    vecSize = procArg<unsigned>(argv[0], argv[1], 1, 1024*1024*1024);
    break;

최소값 1, 최대값 1G로 제한한다.

$ ./giga-add-km.exe 512M
vecSize = 536870912
elapsed time = 15422 usec

$ ./giga-add-km.exe 16K
vecSize = 16384
elapsed time = 124 usec

$ ./giga-add-km.exe 512M
vecSize = 536870912
elapsed time = 15422 usec

$ ./giga-add-km.exe 16K
vecSize = 16384
elapsed time = 124 usec

내부적으로 512 × 1024 × 1024 = 536,870,912로 변환된다.

성능 분석과 한계

다양한 크기로 실험할 수 있는 도구가 완성되었다. 지금까지 측정한 결과를 종합하여 GPU의 성능 특성과 한계를 분석한다.

가속 비율 분석

256M 원소 처리 시간을 정리하면 다음과 같다.

CPU (Intel Core i5-3570):        459,271 usec (0.46초)
GPU 단일 스레드 (RTX 2070):    11,913,504 usec (11.9초)
GPU 병렬 (커널만):                  8,136 usec (8.1 ms)
GPU 병렬 (메모리 전송 포함):      746,448 usec (0.75초)

CPU (Intel Core i5-3570):        459,271 usec (0.46초)
GPU 단일 스레드 (RTX 2070):    11,913,504 usec (11.9초)
GPU 병렬 (커널만):                  8,136 usec (8.1 ms)
GPU 병렬 (메모리 전송 포함):      746,448 usec (0.75초)

순수 커널 성능은 다음과 같다.

CPU 대비: 459,271 / 8,136 = 56.4 배 빠름
GPU 단일 스레드 대비: 11,913,504 / 8,136 = 1,464 배 빠름

CPU 대비: 459,271 / 8,136 = 56.4 배 빠름
GPU 단일 스레드 대비: 11,913,504 / 8,136 = 1,464 배 빠름

GPU 단일 스레드 대비 1,464배 빠르다. GeForce RTX 2070은 2,304개 CUDA 코어를 가진다. 이론적 최대 가속은 2,304배지만, 실제로는 63% 효율을 보인다.

데이터 크기별 트렌드를 정리하면 다음과 같다.

데이터가 커질수록 커널 가속 비율이 증가한다. 하지만 메모리 전송 오버헤드 때문에 전체 성능은 여전히 CPU에 미치지 못한다.

메모리 대역폭 병목

가속 비율을 분석했다. 이제 병목의 원인인 메모리 전송을 자세히 살펴본다. 전송 비용 비율을 계산하면 다음과 같다.

1M 원소:
  전송 비율 = (5,944 - 118) / 5,944 = 98.0%

256M 원소:
  전송 비율 = (746,448 - 8,136) / 746,448 = 98.9%

1M 원소:
  전송 비율 = (5,944 - 118) / 5,944 = 98.0%

256M 원소:
  전송 비율 = (746,448 - 8,136) / 746,448 = 98.9%

데이터가 256배 늘어도 전송 비율은 거의 동일하다. 메모리 전송이 병목임을 명확히 보여준다.

언제쯤 GPU가 CPU를 추월할까? 계산해보자.

CPU 처리 시간: T_cpu = k × N (k는 상수, N은 원소 개수)
GPU 처리 시간: T_gpu = T_transfer + T_kernel
  - T_transfer ≈ c × N (메모리 대역폭에 비례)
  - T_kernel ≈ N / 2304 (코어 개수로 나눔)

CPU 처리 시간: T_cpu = k × N (k는 상수, N은 원소 개수)
GPU 처리 시간: T_gpu = T_transfer + T_kernel
  - T_transfer ≈ c × N (메모리 대역폭에 비례)
  - T_kernel ≈ N / 2304 (코어 개수로 나눔)

데이터를 계속 늘리면 언젠가 T_gpu < T_cpu가 된다. 하지만 현재 예제에서는 전송 비용이 너무 커서 수 GB 이상 데이터가 필요하다.

다음 장에서 다룰 AXPY와 FMA는 계산 강도가 높아 더 유리하다.

스케일링 한계와 해결책

메모리 대역폭 병목을 확인했다. 데이터를 더 키우면 어떻게 될까? 실제로 한계에 부딪힌다. 800M 원소는 메모리 부족으로 실패한다.

필요 메모리: 800M × 4 bytes × 3 = 9.6 GB
RTX 2070 VRAM: 8 GB
부족: 1.6 GB

필요 메모리: 800M × 4 bytes × 3 = 9.6 GB
RTX 2070 VRAM: 8 GB
부족: 1.6 GB

더 큰 데이터를 처리하려면 다른 전략이 필요하다.

해결 방법 1: 데이터 분할

데이터를 여러 조각으로 나누어 처리한다.

const unsigned CHUNK_SIZE = 256 * 1024 * 1024;
for (unsigned offset = 0; offset < SIZE; offset += CHUNK_SIZE) {
    unsigned chunk = min(CHUNK_SIZE, SIZE - offset);

    cudaMemcpy(dev_vecA, vecA + offset, chunk * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_vecB, vecB + offset, chunk * sizeof(float), cudaMemcpyHostToDevice);

    kernelVecAdd <<<dimGrid, dimBlock>>>(dev_vecC, dev_vecA, dev_vecB, chunk);

    cudaMemcpy(vecC + offset, dev_vecC, chunk * sizeof(float), cudaMemcpyDeviceToHost);
}

const unsigned CHUNK_SIZE = 256 * 1024 * 1024;
for (unsigned offset = 0; offset < SIZE; offset += CHUNK_SIZE) {
    unsigned chunk = min(CHUNK_SIZE, SIZE - offset);

    cudaMemcpy(dev_vecA, vecA + offset, chunk * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_vecB, vecB + offset, chunk * sizeof(float), cudaMemcpyHostToDevice);

    kernelVecAdd <<<dimGrid, dimBlock>>>(dev_vecC, dev_vecA, dev_vecB, chunk);

    cudaMemcpy(vecC + offset, dev_vecC, chunk * sizeof(float), cudaMemcpyDeviceToHost);
}

800M을 256M씩 3번에 나누어 처리한다. 각 반복마다 메모리 전송과 커널 실행을 반복한다.

단점은 메모리 전송 횟수가 증가한다는 것이다. 전송 오버헤드가 3배로 늘어난다.

해결 방법 2: 통합 메모리

CUDA 6.0부터 Unified Memory를 지원한다.

float* vecA;
cudaMallocManaged(&vecA, SIZE * sizeof(float));  // 자동 관리

float* vecA;
cudaMallocManaged(&vecA, SIZE * sizeof(float));  // 자동 관리

시스템 메모리와 GPU 메모리를 하나처럼 사용한다. 필요한 데이터를 자동으로 전송하므로 프로그래머가 cudaMemcpy()를 호출하지 않아도 된다.

장점은 프로그래밍이 간단해진다는 것이다. 단점은 성능 예측이 어렵고, 명시적 관리보다 느릴 수 있다는 것이다.

해결 방법 3: 스트리밍

CUDA 스트림을 사용하여 메모리 전송과 커널 실행을 중첩시킨다.

cudaStream_t stream[N_STREAMS];
for (int i = 0; i < N_STREAMS; i++) {
    cudaStreamCreate(&stream[i]);
}

for (int i = 0; i < n_chunks; i++) {
    int sid = i % N_STREAMS;
    cudaMemcpyAsync(..., stream[sid]);
    kernelVecAdd <<<grid, block, 0, stream[sid]>>>(...);
    cudaMemcpyAsync(..., stream[sid]);
}

cudaStream_t stream[N_STREAMS];
for (int i = 0; i < N_STREAMS; i++) {
    cudaStreamCreate(&stream[i]);
}

for (int i = 0; i < n_chunks; i++) {
    int sid = i % N_STREAMS;
    cudaMemcpyAsync(..., stream[sid]);
    kernelVecAdd <<<grid, block, 0, stream[sid]>>>(...);
    cudaMemcpyAsync(..., stream[sid]);
}

스트림 0이 메모리를 전송하는 동안 스트림 1이 커널을 실행할 수 있다. 이를 통해 전송과 계산을 중첩시켜 전체 시간을 단축한다.

이 기법은 고급 최적화 주제로, 나중 장에서 다룬다.

관련 문서

• CUDA C Programming Guide - Execution Configuration - Grid와 Block 크기 설정에 대한 공식 가이드
• CUDA C Programming Guide - Performance Guidelines - GPU 성능 최적화 전략과 메모리 대역폭 최적화
• CUDA Runtime API - cudaMalloc - GPU 메모리 할당 API 레퍼런스
• CUDA C Best Practices Guide - Memory Optimizations - 메모리 전송 오버헤드 최소화 기법
• NVIDIA Blog - How to Optimize Data Transfers - 데이터 전송 최적화 실전 가이드