CUDA 프로그래밍 - 벡터의 합 구하기(1)

 

CUDA 프로그래밍 - 벡터의 합 구하기(1)

 

목차

• 벡터의 합 구하기
• CUDA 알고리즘 성능 측정
• 벡터 합 예제 전체 코드

 

 

안녕하세요. 지난 시간까지 CUDA 프로그래밍에서 제공하는 기초 메모리 API에 대해 알아봤습니다. 해당 내용이 궁금하시면 이전 글을 참고하시는 것을 추천드립니다. 오늘은 이어서 병렬 연산에 대해 알아보고 벡터의 합을 구하는 프로그램을 작성하도록 하겠습니다.

 

 

CUDA 프로그래밍 - 기초 메모리 API(1)

 

 

벡터의 합 구하기

두 벡터의 합은 각 원소를 더하는 연산입니다. 예를 들어, 아래는 두 벡터 $a$와 $b$를 더해 $c$를 구하는 예제입니다. 컴퓨터 연산 관점에서 두 벡터의 합은 길이가 $n$인 두 배열을 더하는 것과 유사합니다.

 

 

$$a = (a_{1}, a_{2}, a_{3}, ..., a_{n})$$

$$b = (b_{b}, b_{2}, b_{3}, ..., b_{n})$$

$$c = (a_{1}+b_{1}, a_{2}+b_{2}, a_{3}+b_{3}, ... , a_{n}+b_{n})$$

 

 

$a$, $b$ 두 벡터의 합

 

 

C언어를 활용해 두 벡터의 합을 구해보겠습니다. 아래는 크기가 SIZE_VECTOR인 두 벡터의 합을 구하는 예제입니다. 벡터 $a$, $b$, 합 $c$가 저장될 배열을 미리 할당하고, $a$, $b$에 0 ~ 99 사이 임의의 값으로 초기화합니다. 이후, for문을 통해 두 배열 원소의 합을 계산합니다.

 

 

예제 코드: C언어 활용 두 벡터 합 구하기

#include <stdio.h>
#include <stdlib.h>
#include <string.h>

#define SIZE_VECTOR 1024


int main() {
  int* a;
  int* b;
  int* c;

  /* Allocate memory for each array */
  a = (int*)malloc(SIZE_VECTOR * sizeof(int));
  b = (int*)malloc(SIZE_VECTOR * sizeof(int));
  c = (int*)malloc(SIZE_VECTOR * sizeof(int));
  memset(a, 0, SIZE_VECTOR * sizeof(int));
  memset(b, 0, SIZE_VECTOR * sizeof(int));
  memset(c, 0, SIZE_VECTOR * sizeof(int));

  /* Initialize the arrays */
  for (int i = 0; i < SIZE_VECTOR; i++) {
    a[i] = rand() % 100;	// 0 to 99
    b[i] = rand() % 100;	// 0 to 99
  }

  /* Addition */
  for (int i = 0; i < SIZE_VECTOR; i++) {
    c[i] = a[i] + b[i];
  }

  delete[] a;
  delete[] b;
  delete[] c;

  return 0;
}

 

 

위의 예제를 GPU를 활용해 계산하려면 어떻게 해야 할까요? CUDA 프로그래밍을 통해 두 벡터의 합을 구해보도록 하겠습니다. GPU를 활용해 두 벡터의 덧셈 연산을 하기 위해 크게 아래의 단계를 수행합니다. 단계를 하나씩 따라가면서 설명하도록 하겠습니다.

 

 

1. 호스트, 디바이스 메모리 할당
2. 입력 데이터 복사 (호스트 메모리 to 디바이스 메모리)
3. 벡터 합 커널 호출
4. 결과 데이터 복사 (디바이스 메모리 to 호스트 메모리)
5. 호스트, 디바이스 메모리 해제

 

 

호스트, 디바이스 메모리 할당

벡터의 덧셈 연산을 수행하기 전, 데이터를 저장할 메모리를 할당해야 합니다. 기초 메모리 API 포스팅에서 설명드린 내용과 같이 GPU 디바이스에 메모리를 할당하기 위해 CUDA에서 제공하는 API를 사용합니다. 메모리를 할당하는 코드 스니펫(code snippet)은 다음과 같습니다.

 

 

/* Host data */
int32_t* p_host_vec1;
int32_t* p_host_vec2;
int32_t* p_host_vec_out;

/* Device data */
int32_t* p_device_vec1;
int32_t* p_device_vec2;
int32_t* p_device_vec_out;

/* Allocate memory(Host) */
p_host_vec1 = (int32_t*)malloc(SIZE_VECTOR * sizeof(int32_t));
p_host_vec2 = (int32_t*)malloc(SIZE_VECTOR * sizeof(int32_t));
p_host_vec_out = (int32_t*)malloc(SIZE_VECTOR * sizeof(int32_t));
memset(p_host_vec1, 0, SIZE_VECTOR * sizeof(int32_t));
memset(p_host_vec2, 0, SIZE_VECTOR * sizeof(int32_t));
memset(p_host_vec_out, 0, SIZE_VECTOR * sizeof(int32_t));

/* Allocate memory(Device) */
cudaMalloc(&p_device_vec1, SIZE_VECTOR * sizeof(int32_t));
cudaMalloc(&p_device_vec2, SIZE_VECTOR * sizeof(int32_t));
cudaMalloc(&p_device_vec_out, SIZE_VECTOR * sizeof(int32_t));
cudaMemset(p_device_vec1, 0, SIZE_VECTOR * sizeof(int32_t));
cudaMemset(p_device_vec2, 0, SIZE_VECTOR * sizeof(int32_t));
cudaMemset(p_device_vec_out, 0, SIZE_VECTOR * sizeof(int32_t));

 

 

malloc(), memset() 함수로 호스트 메모리 공간을 할당하고, 0으로 초기화합니다. GPU 디바이스 메모리에도 동일한 작업을 하기 위해 CUDA API의 cudaMalloc(), cudaMemset() 함수를 사용합니다.

 

 

입력 데이터 복사 (호스트 to 디바이스 메모리)

벡터의 덧셈을 위해 호스트 메모리에 임의의 값을 할당하고, 디바이스 메모리로 복사합니다. 예제에서는 0부터 99 사이 임의의 값을 생성합니다. 임의의 값을 생성하기 위해 rand() 함수를 수행하고 % 나머지 연산을 수행합니다. 호스트 메모리에 값이 할당되면 cudaMemcpy() 함수를 호출하여 디바이스 메모리로 값을 복사합니다.

 

/* Initialize host data */
for (int32_t i=0; i<SIZE_VECTOR; i++) {
  p_host_vec1[i] = rand() % 100;
  p_host_vec2[i] = rand() % 100;
}

/* Copy host data to device */
cudaMemcpy(p_device_vec1, p_host_vec1, SIZE_VECTOR * sizeof(int32_t), cudaMemcpyHostToDevice);
cudaMemcpy(p_device_vec2, p_host_vec2, SIZE_VECTOR * sizeof(int32_t), cudaMemcpyHostToDevice);

 

 

cudaMemcpy()로 메모리 간 복사 연산을 할 때, 복사 방향을 입력 인자로 전달해야 합니다. 위의 코드에서 호스트 메모리의 데이터를 디바이스로 복사하기 때문에 복사 방향을 cudaMemcpyHostToDevice로 설정했습니다.

 

 

입력 데이터 복사 (호스트 to 디바이스 메모리)

 

 

벡터의 합 커널 호출

디바이스 메모리에 입력 데이터가 복사되면, CUDA 코어를 활용해 벡터의 합을 계산합니다. CUDA 코어는 $a$, $b$ 벡터에 할당된 값을 읽고(read) 더한 후 $c$ 벡터에 기록(write)합니다.

 

 

CUDA 코어를 활용한 벡터의 합

 

 

아래는 벡터의 덧셈 연산을 수행하는 커널 함수입니다. __global__ 키워드는 명시적으로 함수의 호출 관계(호스트에서 디바이스)를 나타냅니다. 커널 함수는 p_device_vec 인자를 통해 입력된 값을 더해 p_device_vec_out에 저장합니다. 자신의 스레드(thread)의 인덱스를 전달받아 인덱스에 해당하는 원소의 합을 계산합니다. 연산을 수행하는 과정에서 사전에 할당된 메모리의 범위를 벗어나지 않도록 예외 조건을 검사합니다.

 

__global__ void AddVectors(const int32_t* p_device_vec1,
                           const int32_t* p_device_vec2,
                           int32_t* p_device_vec_out,
                           const int32_t size_vector) {
  int32_t i = threadIdx.x;
  if (i < size_vector) {
    p_device_vec_out[i] = p_device_vec1[i] + p_device_vec2[i];
  }
}

 

 

이번 예제 코드의 벡터 크기 SIZE_VECTOR는 1024입니다. 1024개의 벡터를 1024개의 스레드로 연산하기 위해 아래와 같이 커널 함수를 호출합니다. 커널 함수의 입력 인자로 입력 벡터 p_device_vec1, p_device_vec2, 출력 벡터 p_device_vec_out과 벡터의 크기 SIZE_VECTOR를 전달합니다.

 

/* Kernel call */
AddVectors<<<1, SIZE_VECTOR>>>(p_device_vec1, p_device_vec2, p_device_vec_out, SIZE_VECTOR);

 

 

결과 데이터 복사 (디바이스 메모리 to 호스트 메모리)

커널 연산이 완료되면 p_device_vec_out에 결과 값이 저장됩니다. p_device_vec_out의 값은 디바이스 메모리에 있는 값이므로 결과를 호스트에 복사해야 합니다. 결과 값 복사도 입력 값 복사와 같이 cudaMemcpy()를 호출해 수행합니다.

 

cudaMemcpy(p_host_vec_out, p_device_vec_out, SIZE_VECTOR * sizeof(int32_t), cudaMemcpyDeviceToHost);

 

 

단, 복사의 방향이 디바이스에서 호스트이므로 복사 방향을 cudaMemcpyDeviceToHost로 설정합니다.

 

결과 데이터 복사 (디바이스 to 호스트 메모리)

 

 

호스트, 디바이스 메모리 해제

프로그램의 모든 연산을 완료하면 할당한 모든 메모리를 해제해야 합니다. 할당된 호스트 메모리, 디바이스 메모리는 delete와 cudaFree() 함수를 통해 해제됩니다. 아래의 코드 스니펫은 할당한 호스트, 디바이스 메모리를 해제하는 예제입니다.

 

/* Free memory */
delete[] p_host_vec1;
delete[] p_host_vec2;
delete[] p_host_vec_out;
cudaFree(p_device_vec1);
cudaFree(p_device_vec2);
cudaFree(p_device_vec_out);

 

 

CUDA 알고리즘 성능 측정

프로그램 전체 코드를 보여드리기 전에 연산 성능 측정 코드를 소개하겠습니다. 연산 성능 측정 기능은 프로그램이 연산하는 데 소요되는 시간을 측정합니다. 연산시간을 측정하기 위해 C++의 <chrono> 라이브러리를 사용합니다. <chrono> 라이브러리는 C++ 표준 라이브러리의 일부로, 시간 관련 기능을 다루기 위한 헤더 파일입니다.

 

아래의 코드는 <chrono> 라이브러리를 활용해 연산 시간을 측정하는 코드입니다. 연산 시작 전, 후 시점의 time point를 저장하고, time point를 활용해 연산에 소요된 시간(microseconds)을 계산합니다.

 

#include <chrono>

std::chrono::steady_clock::time_point start, end;

start = std::chrono::steady_clock::now();  	// 측정 시작 시점의 time point
/*
	측정 대상 코드
*/
end = std::chrono::steady_clock::now(); 	// 측정 종료 시점의 time point

// 측정 시간을 microseconds로 환산
auto elapsed_us = std::chrono::duration_cast<std::chrono::microseconds>(end - start).count();

 

 

벡터 합 예제 전체 코드

아래의 코드는 벡터의 합을 구하는 전체 코드입니다. 예제 코드를 실행해서 벡터의 합 성능을 비교해 보시는 것을 추천드립니다.  벡터의 크기(1024)가 작기 때문에 GPU의 연산 성능이 두드러져 보이지 않을 수 있습니다. 오히려, CPU의 연산 성능이 더 좋게 나올 수도 있습니다. 하지만, 현재 예제에서 크기를 크게 설정하면 벡터의 합을 제대로 구하지 못합니다. 이 부분에 대해서는 추후 포스팅에서 CUDA 스레드 계층을 소개해 드리면서 설명드리도록 하겠습니다.

 

#include <cassert>
#include <chrono>
#include <cstdint>
#include <cstring>
#include <iostream>
#include <random>
#include <string>
#include <thread>

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#define SIZE_VECTOR 1024


__global__ void AddVectors(const int32_t* p_device_vec1,
                           const int32_t* p_device_vec2,
                           int32_t* p_device_result,
                           const int size_vector) {
  int index = threadIdx.x;
  if (index < size_vector) {
    p_device_result[index] = p_device_vec1[index] + p_device_vec2[index];
  }
}

void PrintPerformance(const std::string message,
                      const std::chrono::steady_clock::time_point start,
                      const std::chrono::steady_clock::time_point end) {
  auto elapsed_us = std::chrono::duration_cast<std::chrono::microseconds>(end - start).count();
  std::cout << message << " took: " << elapsed_us << " us" << std::endl << std::endl;
}

/*
  * This program adds two vectors using CUDA.
  * It measures the time it takes to add the vectors.
  * Then, compares the result with the CPU version.
  * SIZE_VECTOR is small, so the CPU version should be faster.
*/

int main() {
  std::chrono::steady_clock::time_point start, end;

   /* Host data */
  int32_t* p_host_vec1;
  int32_t* p_host_vec2;
  int32_t* p_host_vec_out;

  /* Device data */
  int32_t* p_device_vec1;
  int32_t* p_device_vec2;
  int32_t* p_device_vec_out;

  /* Allocate memory(Host) */
  p_host_vec1 = (int32_t*)malloc(SIZE_VECTOR * sizeof(int32_t));
  p_host_vec2 = (int32_t*)malloc(SIZE_VECTOR * sizeof(int32_t));
  p_host_vec_out = (int32_t*)malloc(SIZE_VECTOR * sizeof(int32_t));
  memset(p_host_vec1, 0, SIZE_VECTOR * sizeof(int32_t));
  memset(p_host_vec2, 0, SIZE_VECTOR * sizeof(int32_t));
  memset(p_host_vec_out, 0, SIZE_VECTOR * sizeof(int32_t));

  /* Allocate memory(Device) */
  cudaMalloc(&p_device_vec1, SIZE_VECTOR * sizeof(int32_t));
  cudaMalloc(&p_device_vec2, SIZE_VECTOR * sizeof(int32_t));
  cudaMalloc(&p_device_vec_out, SIZE_VECTOR * sizeof(int32_t));
  cudaMemset(p_device_vec1, 0, SIZE_VECTOR * sizeof(int32_t));
  cudaMemset(p_device_vec2, 0, SIZE_VECTOR * sizeof(int32_t));
  cudaMemset(p_device_vec_out, 0, SIZE_VECTOR * sizeof(int32_t));

  /* Initialize host data */
  for (int32_t i=0; i<SIZE_VECTOR; i++) {
    p_host_vec1[i] = rand() % 100;
    p_host_vec2[i] = rand() % 100;
  }

  /* Measure time to add vectors in CPU */
  start = std::chrono::steady_clock::now();
  for (int32_t i=0; i<SIZE_VECTOR; i++) {
    p_host_vec_out[i] = p_host_vec1[i] + p_host_vec2[i];
  }
  end = std::chrono::steady_clock::now();
  PrintPerformance("Add vectors in CPU", start, end);

  /* Measure time to copy host data to device */
  start = std::chrono::steady_clock::now();
  cudaMemcpy(p_device_vec1, p_host_vec1, SIZE_VECTOR * sizeof(int32_t), cudaMemcpyHostToDevice);
  cudaMemcpy(p_device_vec2, p_host_vec2, SIZE_VECTOR * sizeof(int32_t), cudaMemcpyHostToDevice);
  end = std::chrono::steady_clock::now();
  PrintPerformance("Copy host data to device", start, end);

  /* Kernel call */
  start = std::chrono::steady_clock::now();
  AddVectors<<<1, SIZE_VECTOR>>>(p_device_vec1, p_device_vec2, p_device_vec_out, SIZE_VECTOR);

  cudaDeviceSynchronize();
  end = std::chrono::steady_clock::now();
  PrintPerformance("Kernel call", start, end);

  /* Measure time to copy device data to host */
  start = std::chrono::steady_clock::now();
  cudaMemcpy(p_host_vec_out, p_device_vec_out, SIZE_VECTOR * sizeof(int32_t), cudaMemcpyDeviceToHost);
  end = std::chrono::steady_clock::now();
  PrintPerformance("Copy device data to host", start, end);

  /* Print result */
  for (int32_t i=0; i<SIZE_VECTOR; i++) {
    /* Check result */
    assert(p_host_vec_out[i] == p_host_vec1[i] + p_host_vec2[i]);
  }

  /* Free memory */
  delete[] p_host_vec1;
  delete[] p_host_vec2;
  delete[] p_host_vec_out;
  cudaFree(p_device_vec1);
  cudaFree(p_device_vec2);
  cudaFree(p_device_vec_out);

  return 0;
}

 

 

마치며

오늘 포스팅에서는 벡터의 합을 구하는 CUDA 프로그램을 작성했습니다. CUDA API 함수를 사용해서 호스트, 디바이스 간 데이터 복사를 수행하고, 벡터의 합을 구하는 커널 함수를 작성했습니다. CPU와 GPU 간 연산 성능을 비교하기 위해 C++ STL의 <chrono> 라이브러리를 활용해 성능을 측정하는 코드를 작성했습니다.

 

성능을 비교했을 때, GPU의 성능이 의외로 뛰어나지 않은 것을 보실 수 있는데요. 이번 포스팅에서 보여드린 예제의 크기 (1024)는 GPU가 가진 대규모 병렬 처리 능력을 발휘하기에 턱없이 부족한 크기입니다. 이 부분에 대해 자세히 설명드리기 위해 다음 포스팅에서 CUDA의 스레드 계층에 대해 다루도록 하겠습니다.

 

블로그에 방문해 주셔서 감사드리고, 좋은 글로 찾아올 수 있도록 하겠습니다!

 

 

Reference

1. CUDA 기반 GPU 병렬 처리 프로그래밍 - 기초부터 성능 최적화 전략까지