업무자동화코딩연구소

Home » 인사이트 » KISTI N-Ways to GPU Programming Bootcamp 2022 (With 한국과학기술정보연구원xNVIDIA)

KISTI N-Ways to GPU Programming Bootcamp 2022 (With 한국과학기술정보연구원xNVIDIA)

by 오토소장(김효복)
부트 캠프 로고

2022년 5월 16일부터 18일까지 약 3일간 한국과학기술정보연구원(KISTI)엔비디아(NVidia)에서 주최한 부트 캠프에 참여하였습니다. 이번 부트 캠프에서는 C 언어 기반의 다양한 서드파티(third party)1 모듈을 활용해서 개발자가 작성한 연산 수행 알고리즘을 CPU가 아닌 GPU에게 수행하게 하는 방법이라는 주제로 진행되었습니다. 이번 부트 캠프를 통해 GPU를 가속화하는 여러가지 방법에 대해서 알게 된 뜻 깊은 시간이었습니다.

GPU 가속화 기술이 보급되기 이전에는 아래 그림과 같이 CPU를 중심으로 한 병렬 처리 기술들이 많이 연구 되었습니다. 처음에는 CPU 하나를 이용하였지만, 데이터가 점점 많아 지고 연산에 대한 비중이 점점 커지면서 다중 CPU를 이용한 병렬 처리 기술들이 연구되기 시작했습니다. 또한 자원의 한계를 벗어나기 위해 여러 대의 컴퓨터를 활용하는 메세지 전달 인터페이스(Message Passing Interface, MPI) 기술이 제안 되었는데, 이를 병렬 및 분산 처리 시스템(Distributed MPI)이라고 말합니다. 쉽게 말해 다수의 컴퓨터를 하나의 네트워크로 묶어 하나의 시스템처럼 운용하는 기술입니다.

CPU를 이용한 연산 처리 기술의 변천사

하지만 아래 그림처럼 CPU를 이용한 대량의 연산 처리는 일정 수준 이상이 되면 성능의 증가가 정체 되는 부분이 발생하게 되는데, 이는 CPU 설계 구조에 대한 한계점으로 Core 개수가 많지 않기 때문에 나타나는 현상으로 생각이 됩니다. 이런 성능 증가가 둔화되는 법칙을 무어의 법칙(Moore’s Law)2에 보통 비유하고 있습니다. 따라서 이런 CPU의 한계점으로 인해 GPU를 이용한 방법들이 제안 되기 시작하였고, 현재는 미국의 앤비디아(NVidia)가 선두 주자로 나서고 있습니다.

CPU(파란색)와 GPU(초록색)의 성능 차이

전통적인 C 프로그래밍은 CPU를 대부분 활용하지만, GPU 병렬 처리 프로그래밍을 통해 아래 그림과 같이 많은 연산을 처리하는 부분을 GPU에게 위임하여 빠르게 처리할 수 있습니다.

GPU 가속화 방법

시스템 프로그래밍에 강한 언어인 C 언어를 중심으로 한 GPU 병렬 처리 방법은 크게 4가지가 있습니다.

  • 표준 템플릿 라이브러리 (Standard Template Library, STL)
  • OpenACC
  • OpenMP (Open Multi-Processing)
  • CUDA(Compute Unified Device Architecture)

표준 템플릿 라이브러리 (Standard Template Library, STL)

표준 템플릿 라이브러리 (Standard Templete Library, 이하 STL)는 C++을 위한 표준 라이브러리입니다. STL은 컨테이너(Container), 반복자(Iterator), 알고리즘(Algorithm)과 같이 자주 사용되는 요소들을 템플릿으로 구조화하여 제공하고 있습니다. C++17부터는 GPU 병렬 연산 처리를 위한 문법(std::par)을 지원하고 있습니다. C++98과 비교하여 C++17에서는 아래 코드와 같이 std::par를 인자로 넣어주면 sort() 알고리즘에서 처리해야 하는 연산을 CPU가 아닌 GPU에게 위임할 수 있습니다.

C++98: std::sort(c.begin(), c.end());
C++17: std::sort(std::execution::par, c.begin(), c.end());

OpenACC

OpenACC는 크레이, CAPS, 엔비디아, PGI가 개발한 병렬 컴퓨팅을 위한 프로그래밍 표준 입니다. 이 표준은 이기종 CPU/GPU 시스템의 병렬 프로그래밍을 단순하게 만들기 위해 설계된 것입니다. OpenACC는 #Pragma 전처리 문법을 사용하여 GPU에게 연산을 위임할 수 있습니다.

#pragma acc data copy(A) create(Anew)
while ( error > tol  &&  iter  <  iter_max )  {
  error = 0.0;
#pragma acc kernels
  {
#pragma acc loop independent collapse(2) reduction(max:error)
  for (  int  j = 1; j < n-1;  j++ )  {
    for (  int  i = 1; i < m-1; i++ )  {
       Anew [j] [i] = 0.25 * ( A [j] [i+1] + A [j] [i-1] +
                                      A [j-1] [i] + A [j+1] [i]);
       error = max ( error, fabs (Anew [j] [i] - A [j] [i]));
      }
    }
    ...
  } 
}

OpenMP (Open Multi-Processing)

OpenMP API는 C/C++ 및 Fortran에서 다중 플랫폼 공유 메모리 병렬 프로그래밍을 지원합니다. OpenMP API는 데스크톱에서 슈퍼컴퓨터에 이르는 플랫폼에서 병렬 응용 프로그램을 개발하기 위한 간단하고 유연한 인터페이스로 이식 가능하고 확장 가능한 모델을 정의합니다. OpenMP도 OpenACC와 마찬가지로 #pragma 전처리 문법을 통해서 사용할 수 있습니다. OpenACC와 OpenMP를 사용하시는 분이라면 문법적인 요소가 비슷해서 혼동되는 경우가 자주 발생합니다.

#include <omp.h>
#include <stdio.h>
#define NT 4
#define thrd_no omp_get_thread_num

int main()
{
	#pragma omp parallel for num_threads(NT) // PRAG 1
	for(int i=0; i<NT; i++) printf("thrd no %d\n",thrd_no());

	#pragma omp parallel for \
		num_threads(NT) // PRAG 2
	for(int i=0; i<NT; i++) printf("thrd no %d\n",thrd_no());

	#pragma omp parallel num_threads(NT) // PRAG 3-4
	#pragma omp for
	for(int i=0; i<NT; i++) printf("thrd no %d\n",thrd_no());

	#pragma omp parallel num_threads(NT) // PRAG 5
	{
		int no = thrd_no();
		if (no%2) { printf("thrd no %d is Odd \n",no);}
		else { printf("thrd no %d is Even\n",no);}

		#pragma omp for
		for(int i=0; i<NT; i++) printf("thrd no %d\n",thrd_no());
	}
}

CUDA

CUDA는 그래픽 처리 장치(GPU)에서 수행하는 (병렬 처리) 알고리즘을 C 프로그래밍 언어를 비롯한 산업 표준 언어를 사용하여 작성할 수 있도록 하는 GPGPU 기술입니다. CUDA는 엔비디아가 개발해오고 있으며 이 아키텍처를 사용하려면 엔비디아 GPU와 엔비디아 SDK가 필요합니다. C언어에서 사용하는 CUDA 문법은 C문법과 매우 유사합니다.

cudaArray* cu_array;
texture<float, 2> tex;

// 행렬 할당
cudaMallocArray(&cu_array, cudaCreateChannelDesc<float>(), width, height);

// 이미지 데이터를 행렬로 복사
cudaMemcpy(cu_array, image, width*height, cudaMemcpyHostToDevice);

// 행렬을 텍스처에 연결한다.
cudaBindTexture(tex, cu_array);

// 커널을 실행한다
dim3 blockDim(16, 16, 1);
dim3 gridDim(width / blockDim.x, height / blockDim.y, 1);
kernel<<< gridDim, blockDim, 0 >>>(d_odata, width, height);
cudaUnbindTexture(tex);

__global__ void kernel(float* odata, int height, int width)
{
   unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
   unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
   float c = texfetch(tex, x, y);
   odata[y*width+x] = c;
}

CUDA의 경우 STL, OpenACC, OpenMP와는 다르게 매우 세부적으로 GPU를 제어할 수 있습니다. 이로 인해 CUDA를 학습하는것은 C언어를 처음 배우는 것과 같이 매우 힘이 들수도 있습니다. 하지만 이러한 단점을 극복하면 아래 그림과 같은 성능을 얻을 수 있습니다.

CUDA 성능

이번 부트 캠프를 통해 C언어를 이용한 GPU 병렬 처리 기법을 4가지를 알게 된 뜻 깊은 시간이었습니다. CUDA 책은 사놓기만하고 있었는데 이 캠프를 통해 한번 경험해보고 책을 읽어보니 매우 수월하게 이해가 되었습니다. GPU 기술에 대한 큰 인사이트를 얻게 된 시간이었습니다.

수료증

본 포스팅에 사용한 이미지는 부트 캠프간 진행된 엔비디아(NVidia) 강의 자료를 일부 발췌 하였습니다.

  1. 일반적으로 하드웨어 또는 플랫폼((platform) 생산자와 소프트웨어 개발자의 관계를 나타내는 용어[]
  2. 무어의 법칙은 반도체 집적회로의 성능이 24개월마다 2배로 증가한다는 법칙이다. 경험적인 관찰에 바탕을 두고 있다. 인텔의 공동 설립자인 고든 무어가 1965년에 내 놓은 것이다.[]

You may also like

이 글에 대한 댓글 구독
알림
guest
0 Comments
Inline Feedbacks
모든 댓글 보기
0
Would love your thoughts, please comment.x
[]