GPGPU 프로그래밍

서광열  kwangyul.seo@gmail.com·현재 (주)노매드커넥션의 CTO로 일하고 있다. 현재 동영상 미디어 플랫폼 까멜레오(http://www.chameleo.org)를 개발하고 있다. 그리고 블로그 서광열의 소프트웨어 이야기(http://skyul.tistory.com)을 통해서 소프트웨어 개발, 프로그래밍 언어에 대한 생각을 정리하고 있다.

GPGPU 프로그래밍은 GPU를 보조 CPU처럼 일반적인 연산 수행에도 사용하려는 취지에서 출발했다. GPU는 복잡한 3D 게임을 지원하기 위해 빠른 속도로 발전했지만 최신 게임이 아닌 이상 GPU의 파워를 100% 활용하는 경우는 드물다. GPGPU는 남은 GPU 연산 능력을 CPU를 보조하는 데 쓰자는 취지다. GPU는 원래 그래픽 애플리케이션의 연산을 돕는 목표로 설계됐지만, GPGPU로 인해 지금은 복잡한 과학 연산이나 암호 알고리즘 계산 등 그래픽 이외의 영역에서도 많이 사용되고 있다.

하지만 GPU를 비그래픽 애플리케이션에 활용하려는 시도는 쉽지 않았다. GPU는 CPU와 달리 아키텍처나 인스트럭션에 대한 정보가 드물고, 일반 개발자 입장에서는 OpenGL이나 DirectX와 같은 API를 통하지 않고 직접 하드웨어를 제어하기가 무척 어렵다. 이 때문에 GPGPU 방법론은 OpenGL, DirectX 등 3D 그래픽 API를 변용하거나 GLSL, Cg, HLSL 등 셰이딩 언어를 사용하는 방법이 주종을 이루고 있다. 따라서 그래픽 애플리케이션 작성에 익숙지 않은 개발자들은 GPU 활용이 상당히 까다로울 수밖에 없다.

최근 애플은 맥 OS X의 차세대 버전인 스노우 레오파드(Snow Leopard)의 프리뷰를 선보이면서 OpenCL(Open Computing Language) 프로젝트를 언급했다. OpenCL은 그래픽 애플리케이션이 아닌 일반 애플리케이션 개발에도 GPU의 컴퓨팅 파워를 쉽게 사용할 수 있도록 만드는 것이 목표다. OpenCL은 OpenGL 같은 3D 그래픽 API를 사용하지 않고서도, CPU의 보조 장치로 GPU를 일반 계산에 사용할 수 있도록 해준다. 하지만 아쉽게도 아직 OpenCL에 대한 자세한 정보는 공개되지 않았다.

다행히 엔비디아와 ATI가 이와 비슷한 노력을 하고 있다. 특히, 엔비디아의 쿠다(CUDA: Compute Unified Device Architecture) 프로젝트는 OpenCL에 직접적인 영향을 미친 선구적인 프로젝트이다. 앞으로 나올 OpenCL의 모습은 쿠다를 통해 어느 정도 예측해 볼 수 있는 셈이다. 쿠다는 엔비디아의 GPU를 GPGPU로 활용할 수 있도록 C 언어를 확장한 언어와 런타임을 제공한다. 쿠다 개발자는 C 언어로 애플리케이션을 개발하면서 일부 코드를 GPU에서 실행하도록 지정할 수 있다. 엔비디아 컴파일러(nvcc)는 이 코드를 GPU에서 실행할 수 있도록 컴파일 해준다. 그래픽 API를 전혀 몰라도 GPU를 사용할 수 있다는 면에서 쿠다는 GPGPU의 새로운 획을 그었다.

쿠다의 특징

쿠다는 쉽게 말해 C 언어로 GPU 프로그래밍을 할 수 있게 해주는 도구다. 여러분이 C 개발자라면 곧바로 GPU 프로그래밍을 할 수 있다. 물론 몇 가지 사전 지식은 필요하지만 그래픽 프로그래밍의 고수일 필요가 없다는 뜻이다.  쿠다의 특징은 다음과 같다.

- C 언어로 작성된 코드를 직접 GPU에서 수행
- 그래픽 API나 GPU 프로그래밍에 대한 사전 지식이 필요 없음
- GPU 고유 명령어와 메모리에 대한 접근
- 윈도우와 리눅스에서 무료로 사용 가능
- 엔비디아의 안정적인 지원과 풍부한 문서 및 예제 코드

쿠다를 배우는 이유

쿠다는 엔비디아에서 개발되었기 때문에 엔비디아 GPU 위에서만 동작한다. 또한, 비교적 고사양의 그래픽 카드인 GeFor ce 8xxx, Telsa, Quadro 시리즈에서만 동작한다. 일부 독자는 특정 하드웨어에서만 동작하는 프로그래밍 모델이나 API를 습득할 필요가 있느냐는 질문을 할 수도 있을 것이다. 그에 대한 대답은 쿠다가 엔비디아에 종속된 프로젝트임에도 불구하고 쿠다는 충분히 배워둘 가치가 있다는 것이다.

일단 쿠다는 스트림 컴퓨팅(stream computing)과 관련해서 가장 최신의 API다. 즉, 스트림 컴퓨팅의 마켓 리더이며, 앞으로 등장할 제품에도 막대한 영향을 미칠 것이 분명하다. 맥 OS X의 OpenCL 팀 역시 쿠다의 존재를 알고 있으며, 쿠다와 유사한 프로그래밍 모델을 내놓을 가능성이 크다. 또한 엔비디아의 시장이 적지 않기 때문에 쿠다가 장착된 GPU가 2007년 12월 기준으로 이미 4,000만 대 이상 팔렸다. 원래 고성능 컴퓨팅은 특정 하드웨어를 목표로 잡는 경우가 많았다는 점도 무시할 수 없다. 마지막으로, 엔비디아가 앞으로 쿠다 지원을 멈출 가능성이 높지 않다는 점을 들 수 있다.

쿠다 프로그래밍 모델

구체적인 API를 살펴보기에 앞서 쿠다의 프로그래밍 모델을 정확히 이해하는 것이 가장 중요하다. 쿠다 프로그래밍 모델은 쿠다 프로그래밍 시에 프로그래머가 GPU를 바라보는 관점을 명확히 설명해 주기 때문이다.

쿠다 프로그래밍 모델에서 GPU는 CPU의 계산을 도와주는 보조프로세서(coprocessor) 역할을 한다. GPU의 G가 그래픽스를 뜻하지만 쿠다는 GPU를 CPU와 조금 다른 특징을 가진 또 하나의 단순 연산 장치로 생각한다고 보면 된다. 특히, 데이터-병렬의 계산양이 많은 함수는 CPU가 아닌 GPU에서 실행하도록 지정할 수 있다. 같은 코드를 데이터만 조금 다르게 해서  여러 번 수행하는 부분이 주 대상이 되는데, 일반적인 프로그래밍에서는 for 문이 대표적이다.

쿠다는 일반 C 함수를 디바이스(쿠다에서는 GPU를 보통 디바이스라고 부르고, CPU를 호스트로 부른다)에서 실행하도록 표시해줄 수 있는데, 이런 함수를 커널(kernel)이라고 부른다. 커널은 디바이스에서 실행되며, 수많은 스레드가 동시에 실행할 수 있다. 여기서 스레드는 GPU 상의 스레드를 의미한다. GPU는 사실상 거대한 ALU의 집합체이기 때문에 구조상 데이터-병렬 프로그래밍에 최적화되어 있다. 쿠다는 이런 GPU의 속성을 스레드로 표현한 것이다.

쿠다는 <그림 1>처럼 스레드 블록과 스레드로 구성된다. 커널을 실행할 때는 몇 개의 블록과 스레드를 할당할 것인지 지정할 수 있다. 또한, 행렬 연산이나 복잡한 데이터 연산을 돕기 위해 그리드의 레이아웃은 1, 2, 3차원이 모두 가능하다. 각각의 블록은 유일한 block ID를 가지며, 각 스레드는 유일한 thread ID를 가진다.

쿠다는 또한 호스트(CPU) 메모리와 디바이스(GPU) 메모리를 명확히 구분한다. 메모리가 어느 영역에 속해있느냐에 따라 성능 차이가 크게 날 수 있기 때문에, 쿠다는 메모리 할당 시에 어떤 메모리를 할당할 것인지 명확히 구분한다. 대신, 호스트와 디바이스 간의 메모리를 복사할 수 있다.

쿠다 프로그램 예제

이해를 돕기 위해 간단한 N×N 행렬 두 개를 받아서 더하는 간단한 C 함수를 쿠다 프로그램으로 변경해 보자. 일단 C 프로그램은 <리스트 1>과 같다. 이 프로그램은 N×N 행렬을 float * 타입으로 나타냈으며, 이중으로 for 루프를 돌면서 a와 b의 각 원소를 더해서 c에 저장하는 간단한 프로그램이다.

<리스트 1>을 쿠다 프로그램으로 변경하면 <리스트 2>와 같아진다. 참고로 코드를 단순화하기 위해 main()에서 a, b, c, blocksize, N의 선언은 생략했다. 완전한 코드는 쿠다 메모리 모델을 설명한 후에 제시하도록 하겠다.

add_matrix 함수 속성을 __global__로 선언해 주었는데, __global__은 해당 함수를 CPU가 아닌 GPU로 수행하라고 컴파일러에게 알려주는 역할을 한다. 커널은 호출할 때 몇 개의 블록과 스레드로 실행할 것인지 알려줘야 하는데 함수 이름과 파라미터 사이에 << >>를 사용한다. dim3는 쿠다의 확장 데이터 타입으로 세 개의 정수를 받는 벡터 타입이다. dimBlock과 dimGrid는 각각 두 개의 정수만 지정했으므로 2차원이 된다.

예를 들어, N을 20으로 잡고 blocksize가 10이라고 가정하면, main 함수는 add_matrix를 10×10 블록을 2×2 그리드로 배치해서 호출하라는 뜻이 된다. 스레드는 동시에 실행되기 때문에 add_matrix는 최대 40개의 스레드에서 병렬로 수행된다. add_matrix 함수는 자신이 현재 어떤 블록의 어떤 스레드에서 수행되고 있는지 알아야 하는데, blockIdx와 threadIdx는 이 때 사용되는 내장 변수(built-in variable)이다. block Idx는 현재 커널을 수행 중인 block의 인덱스를, threadIdx는 블록 내에서 thread의 인덱스를 뜻한다. blockIdx와 thread Idx는 구조체로 차원에 따라 x, y, z, w를 멤버로 가진다. blockIdx.x * blockDim.x + threadIdx.x는 blockIdx.x와 threadIdx.x의 값에 따라 정확히 20개의 행렬의 행(row) 중에 하나를 가리키게 된다. 쉽게 말해 i, j가 각각 [0,20)으로 index가 [0, 400)인 400개의 스레드가 동시에 수행된다는 뜻이다.

스레드는 병렬로 수행되기 때문에 변경된 add_matrix는 CPU가 아닌 GPU에서 수행될 뿐만 아니라 병렬로 수행된다. 이 때문에 쿠다 프로그래밍은 병렬 프로그래밍과도 밀접한 관계를 맺고 있다. 스레드와 스레드 블록, 그리드의 관계를 다시 살펴보면 <그림 2>와 같다. 스레드는 스레드 블록에 속하고, 스레드 블록은 그리드에 속한다. 병렬화를 여러 단계를 거쳐서 하는 셈인데, 각 단계는 메모리 모델에 차이가 있다. 하나의 스레드 블록은 최대 512개의 스레드를 가질 수 있고, 스레드들은 공유 메모리를 통해 통신한다. 반대로 스레드 블록 그리드는 공유 메모리를 통해 통신한다.

쿠다 메모리 모델

쿠다는 GPU의 메모리 모델을 상당 부분 프로그래머에게 노출시키고 있다. 메모리 모델을 정확히 이해해야만 최적의 성능을 이끌어낼 수 있기 때문이다. GPU 메모리는 cuda Malloc(), cudaFree() 같은 쿠다 함수를 이용해 명시적으로 할당하고 해제한다. GPU 메모리에 대해 포인터를 사용할 수 있다는 것이 특징이다. 더불어, CPU와 GPU 사이의 메모리 복사는 매우 비용이 큰 연산이므로 CPU-GPU 간 메모리 복사를 줄이기 위해 많은 노력을 했다.

쿠다에는 다양한 메모리 종류가 존재하는데, <그림 3>처럼 레지스터, 로컬(local) 메모리, 공유(shared) 메모리, 전역(global) 메모리, 상수(constant) 메모리, 텍스처(texture) 메모리 등이 있다.
레지스터와 로컬 메모리는 스레드마다 하나씩 존재하며, 읽기-쓰기가 가능하다. 공유 메모리 역시 읽기-쓰기가 가능하고 블록 별로 존재하여 여러 스레드가 공유한다. 전역 메모리는 그리드 별로 존재하며 읽기-쓰기가 가능하지만 캐시가 되지 않는다. 상수 메모리와 텍스처 메모리 역시 그리드 별로 존재하며 상수 메모리는 캐 쉬되고 텍스처 메모리는 일부만 캐시 된다. 입문자들은 CPU 보다 더 복잡한 메모리 모델에 놀랄 수도 있겠다. 하지만 일단 전역 메모리를 사용하고 필요에 따라 최적화하는 방법이 일반적이다.

앞서 살펴본 예제 프로그램을 완전히 동작하도록 수정해 보자. add_matrix 커널은 전과 동일하기 때문에 생략했다.

행렬 a, b, c는 호스트 메모리에 잡혀 있기 때문에 커널에서 사용하기 위해서는 GPU 메모리로 복사가 되어야 한다. <리스트 3>을 보면 ad, bd, cd를 cudaMalloc()으로 할당했는데, cudaMalloc()은 GPU 메모리를 할당한다. 행렬 a를 ad로 복사하기 위해서는 호스트->디바이스 메모리 복사가 이루어져야 하므로, 일반적인 memcpy가 아닌 cudaMemcpy의 cudaMemcpyHostToDevice 옵션을 사용한다. 반대로 연산이 끝나고 나면 다시 GPU에서 CPU로 메모리를 복사해 와야 하는데, 이때는 cudaMemcpy의 cudaMemcpyDeviceTo Host 옵션을 사용한다.

쿠다 API와 C 확장

쿠다는 기존 C 프로그래머가 쉽게 사용할 수 있을 것을 목표로 하기 때문에, C 언어에 최소한의 확장만을 추가했다. 쿠다는 크게 C 확장과 런타임 라이브러리로 나뉘는데, 런타임 라이브러리는 크게 세 가지 콤포넌트로 나뉜다.

1) 디바이스(GPU) 컴포넌트

- 일반 sin 함수에 비해 덜 정확하지만 수행 속도가 빠른 수학 함수 예) __sinf(x)
- __syncthreads() : 블록 내의 모든 스레드가 이 지점에 도달하기를기다리는 함수
- 타입 변환 함수(라운딩 모드 지원)
- 타입 캐스팅 함수
- 텍스처 함수
- 아토믹 함수(Compute Capability 1.1 이상)

2) GPU을 제어할 호스트(CPU) 컴포넌트- 디바이스 관리 : 디바이스 속성 정보, 다중 GPU 제어

3) 공통 컴포넌트- 내장 벡터 타입: float1, float2, int3, ushort4 등.  int2 i = make_int2(i, j)

- 수학 함수: CPU에서는 표준 math.h, GPU에서는 HW
- 벤치마크를 위한 시간 함수
- 텍스처 레퍼런스

C 언어에도 몇 가지 확장을 가했는데, 앞서 예제에서 살펴본 __global__도 그 중 하나이다. 이 외에도 GPU 내에서만 호출이 가능하다고 표시하는 __device__, CPU에서 호출한다는 뜻의 __host__ 등의 함수 타입 한정자(type qualifier)가 있다. 변수를 선언할 때도 메모리 종류를 지정할 수 있는데, __device__, __constant__, __shared__ 변수 타입 한정자가 있다. 커널을 호출할 때는 foo<<GridDim, BlockDim>>(...) 형태를 사용한다. blockIdx, threadIdx, blockDim 등 그리드/블록, 블록/스레드 크기를 나타내는 내장 변수도 추가되었다.

소스 파일은 반드시 쿠다 컴파일러인 nvcc로 컴파일 되어야 한다.

쿠다 구성도

쿠다 프로그램은 <그림 4>와 같이 구성된다. 가장 상위에는 CPU/GPU가 혼합된 소스 코드가 있고, BLAS나 FFT 등 미리 작성된 최적화 라이브러리가 있다. 이 소스 코드는 쿠다 C 컴파일러인 nvcc로 컴파일 된다. 호스트 코드는 일반 C 컴파일러를 거쳐서 CPU에서 수행될 수 있는 오브젝트 코드가 되고, 디바이스 코드는 nvcc가 머신 의존적이지 않은 어셈블리(PTX)로 컴파일 한다. PTX는 쿠다 드라이버를 통해 GPU에서 수행되는 코드가 된다. 보통 커널은 .cu 확장자로 저장하는 것이 일반적이다.

쿠다 최적화

일단 올바르게 동작하는 프로그램을 작성했으면 다음으로 할 일은 최적화다. 물론 섣부른 최적화는 모든 악의 근원이므로, 최적화는 성능 요구사항이 있을 때만 수행하는 것이 원칙이다. 쿠다는 clock 함수를 제공하는데, GPU 내 멀티프로세서마다 있는 카운터의 값을 리턴 한다. 카운터는 매 클럭 사이클마다 갱신되므로, 연산 시작 전에 clock을 얻어오고, 끝난 후에 얻어서 차이를 비교하면 연산 비용이 얼마나 큰지 추정할 수 있다.

쿠다 최적화를 위해서는 GPU의 구조를 어느 정도 알아야 한다. GPU는 보통 여러 멀티프로세서(MP)로 구성되며, 각각의 멀티프로세서는 여러 개의 스칼라 프로세서(SP)를 가진다. 각 MP는 일련의 블록을 처리하는데, 블록 하나는 MP 하나에서만 처리된다. 각 블록은 와프(warps)라는 그룹으로 나뉘는데, 와프는 물리적으로 병렬 수행되며 스케줄러는 와프를 바꿔가며 수행한다. 와프는 보통 32개의 스레드로 구성된다.

따라서 블록의 수가 최소한 MP의 수만큼은 있어야 모든 MP가 놀지 않고 일을 하게 된다. 블록/MP의 비율이 2보다 크면 하나의 MP에서 여러 블록이 병행 수행된다. 그리고 앞으로 나올 GPU에서도 규모가변성(scalability)을 보장하려면 블록 수가 최소 100개 이상은 되어야함을 알 수 있다.

각 연산이 얼마나 많은 시간을 소비하는지 알아두는 것도 성능 향상에 큰 도움이 된다.

마지막으로 GPU에 존재하는 다양한 메모리 중에서 적당한 것을 골라 쓰는 것이 중요하다. 각각의 메모리는 상황에 따라 액세스 속도와 캐시 등에서 큰 차이가 있기 때문이다. 일례로 상수 메모리는 와프 내의 모든 스레드가 같은 위치에 접근하는 경우 레지스터만큼 빠르다. 반대로 전역 메모리는 메모리 패치를 내보는데 4 사이클이 걸리고, 지연 시간(latency)은 400-600 사이클에 달한다. 성능을 내기 위해서는 전역 메모리 접근은 다음 <그림 5>처럼 반드시 coalesce되어야 한다.

이 글에서는 GPGPU 프로그래밍의 일환으로 GPU를 보조 CPU처럼 사용할 수 있는 쿠다 프로그래밍에 대해 살펴보았다. 쿠다는 비록 엔비디아 GPU에 종속된 프로젝트지만 GPGPU 프로그래밍과 병렬 프로그래밍의 미래에 대해 많은 시사점을 주고 있다.

쿠다는 값비싼 그래픽 카드가 하는 일 없이 놀고 있는 상황이 싫은 개발자라면 한 번쯤 도전해 볼만한 가치가 있는 주제다. 지면 관계상 쿠다 프로그래밍 모델을 설명하는 데 집중했고, 쿠다 최적화와 실제 API에 대해서는 이 글에서 자세히 소개하지 못했다. 관심 있는 개발자는 쿠다 홈페이지를 방문해서 쿠다가 사용된 다양한 애플리케이션을 접해보길 바란다.