GPU는 어떻게 작동할까 + CUDA programming

문상준·2024년 11월 25일

CPU vs GPU

CPU : 고성능 프로세서가 몇 개 정도 들어있다.
GPU : 성능이 높지 않은 프로세서가 수천개 정도 들어있다.

GPU는 왜 많은 프로세서를 가지는가?

한장의 프레임(사진)에 있는 픽셀 수 $= 1920 * 1080 = 2,073,600$

동영상에서 1초에 30프레임을 사용한다면, 컴퓨터가 1초에 계산해야되는 픽셀 수 $= 2,073,600 * 30 = 62,208,000$

==> 만약 이렇게 많은 픽셀을 CPU가 처리하는 것은 교수에게 쉬운 수학 문제를 많이 풀게 하는 것과 같다.

==> 비효율

==> GPU 사용

GPU의 기본 연산

GPU는 원래 컴퓨터 그래픽 작업을 가속화하기 위해 설계되었다. 그래픽 작업에서는 화면에 이미지를 렌더링하기 위해 많은 벡터와 행렬 연산이 필요하다.

예를들어, 3D 모델 변환의 경우, 객체의 회전, 이동, 스케일링과 같은 변환을 수행해야 하는데, 이는 모두 행렬 곱셈으로 표현된다.

그렇기에 GPU의 기본 연산은 행렬 곱셈인 것이다.

GPU의 병렬 처리

GPU는 위에서 설명했듯이, 각 픽셀 값을 계산하기 위해 사용되었다. 옆의 픽셀을 계산하는데, 위쪽의 픽셀 값이 사용되지는 않는다. 즉, 각각의 계산이 병렬적이라는 것이다.

이런 GPU의 특징을 활용해 대량의 데이터를 병렬 처리하고자 한다.

CUDA, OpenCL

위의 요구에 의해 대량의 데이터를 그래픽이 처리되는 형태로 바꾸어 GPU연산을 하고 싶다.
But 그래픽이 처리되는 형태는 행렬이다.
즉, 대량의 데이터의 처리를 행렬의 곱셈 방식으로 쉽게 처리하고자, CUDA, OpenCL탄생.

이렇게 그래픽 처리에만 GPU를 사용하는 것이 아니기에, GPGPU(General Purpose GPU)라고도 한다.

CUDA 코어와 프로그램이 데이터를 처리하는 방식

GPU가 처리해야 할 100개의 작업이 있다고 하자.
GPU 코어가 처리해야 할 하나의 작업을 쓰레드라고 하면,
서로 독립적인 100개의 쓰레드를 처리하는 가장 이상적인 방법은 100개의 코어에 100개의 쓰레드를 각각 하나씩 할당하는 SIMD 방법이다.

But 현실에서 쓰레드 >>> 코어

==> SIMT(Single Instruction Multiple Threading)

서로 독립적인 100만개의 쓰레드를 1000개의 코어에서 계산한다고 하자.

이때, 100만개의 쓰레드를 1000개의 코어중 어딘가에서 처리해야한다면, 그 분배의 과정이 매우 복잡해진다.

만약 100만개의 쓰레드를 1000개의 묶음으로 묵고, 1000개의 코어를 10개의 묶음으로 묶는다면, 1000개의 쓰레드 묶음을 10개의 코어 묶음에 할당하면 되니, 할당의 문제가 더 쉬워진다.

==> 하나의 코어 그룹에서 하나의 쓰레드 그룹을 처리하게 된다.

==> 그 쓰레드 그룹에는 1000개의 쓰레드가 있으므로, 하나의 코어 그룹에는 100개의 코어가 있으므로, 한번에 처리 X

==> 쓰레드 그룹에서 쓰레드를 100개씩 묶어주면, 다시 10개의 작은 그룹이 생성되는데 이를 warp라고 한다.

==> 하나의 warp에는 100개의 쓰레드가 있기에, 하나의 코어 그룹(100개의 코어)에서 처리 可能

==> 하나의 warp안의 쓰레드들은 하나의 작업처럼 처리된다.

RTX 3090 GPU에서의 예시를 기반으로 GPU의 병렬처리 이해하기

RTX 3090 GPU에는 만개가 넘는 코어가 있다.
코어는 128개씩 그룹으로 묶인다.
warp안에는 128개씩의 쓰레드가 있다.

==> 하나의 쓰레드가 하나의 코어에서 연산 可能

코어가 연산을 수행하다가 메모리로부터 데이터를 읽어와야 할 때가 있다. 이때, 지연 시간이 생긴다.

CPU의 경우, 지연시간을 줄이기위해,
==> 캐시 메모리 사용.

GPU의 경우, 지연시간을 줄이기위해, 지연시간이 발생하면,
그저 다른 warp를 계산하면 된다.
==> 쓰레드 ↑↑↑이면, 효율 ↑↑↑

그저 다른 warp를 계산하는 것이 효율적인 이유

cpu에서 multi-threading을 하기 위해서는 software적인 overhead가 ↑

But

GPU에서는 context switching(thread간 switching (= warp간 switching))이 hardware적으로 구현되어, overhead ↓

CUDA의 indexing을 예시를 기반으로 이해하기

for(int i=0; i<100; i++)
	C[i] = A[i] + B[i];

라는 단순한 코드를 GPU로 parallel programming하게 된다고 하자.
이때, thread는 100개, GPU core는 10개라고 하자.
thread를 10개의 block으로 묶는다면, 각 block에는 2개의 warp를 갖게된다.

그렇다면 각각의 데이터를 어떤 thread로 처리할 것인가?
아래의 코드로 확인해 볼 수 있다.

__global__ VectorAdd(A,B,C)
{
	C[blockIdx*10 + threadIdx] = A[blockIdx*10 + threadIdx]
    							+B[blockIdx*10 + threadIdx];
}

==> block은 10개 있고, 각 block에 thread가 10개 있으므로,
==> block index, thread index, 두개의 index로 모든 쓰레드를 넘버링하여, 데이터를 처리할 수 있는 것이다.

그렇다면 A,B,C arrays는 어디에 있나?

위 VectorAdd함수를 실행하기 이전에는 당연히 Host(CPU)의 Main memory에 있다.

따라서 GPU에서 연산을 수행하기 위해서는 다음의 과정을 거쳐야한다.

Device의 global memory에 공간을 할당하고 (cudaMalloc)
Host의 memory 내용을 Device의 global memory로 옮기고 (cudaMemcpy)
연산을 수행하고 (Kernel launch) // gpu에서 수행될 함수를 kernel이라고 부른다.
Device의 연산 결과를 Host의 main memory로 옮기는 (cudaMemcpy)
Device의 global memory에 할당된 공간 삭제 (cudaFree)

과정이 필요할 것이다.

Warp divergence란?

GPU는 쓰레드를 대량으로 warp라는 단위로 묶어 실행하여 병렬 처리한다고 했다.

이때 하나의 warp에 속한 모든 쓰레드는 동일한 pc값을 갖는다.
즉, 동일한 명령어를 동시에 실행한다.

워프는 한 번에 하나의 공통 명령어를 실행하므로, 워프의 32개 쓰레드가 동일한 실행 경로에 동의할 때 최대 효율을 달성한다.
만약 워프의 쓰레드가 데이터 의존적인 조건부 분기를 통해 경로를 나누는 경우, 워프는 선택된 각 분기 경로를 실행하며, 해당 경로에 있지 않은 스레드는 비활성화된다.

즉, 분기 다이버전스(branch divergence) 는 워프 내에서만 발생하며, 다른 워프는 공통 또는 별개의 코드 경로를 실행하든 관계없이 독립적으로 실행된다.