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

 

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

 

목차

• CUDA 프로그램 연산 흐름
• CUDA 기초 메모리 API (할당, 초기화, 해제)

 

 

CUDA 프로그램 연산 흐름

CUDA 프로그램은 호스트 코드와 디바이스 코드로 구성되어 있습니다. 호스트 코드는 프로그램의 기본 연산을 수행하고, GPU와 같은 장치를 사용할 때, 커널을 호출합니다. 컴퓨터 시스템의 기본 메모리 공간은 CPU가 사용하는 시스템 메모리(system memory)입니다. 시스템 메모리는 일반적으로 메인보드에 장착되어 있는 DRAM 영역으로 메인 메모리(main memory)라고도 하고, CUDA 프로그램 관점에서 호스트 메모리로 지칭합니다.

 

 

CUDA 프로그램, 호스트에서 디바이스로 데이터 복사

 

 

CPU와 GPU는 서로 독립된 장치로, 사용하는 메모리 영역이 다릅니다. GPU가 사용하는 메모리는 디바이스 메모리지만, 모든 데이터는 기본적으로 호스트 메모리에 저장되어 있습니다. 따라서, GPU를 활용해 데이터를 처리하기 위해서 호스트 메모리의 데이터를 디바이스 메모리로 복사해야 합니다.

 

일반적인 CUDA 프로그램의 첫 번째 단계는 위 그림과 같이 호스트 메모리에서 디바이스 메모리로 데이터를 복사하는 것입니다. 하드디스크, SSD 등 외부 장치나 네트워크를 통해 받은 데이터도 기본적으로 호스트 메모리에 적재되기 때문에 호스트 메모리에서 디바이스 메모리로 복사하는 과정이 필요합니다.

 

 

CUDA 프로그램, GPU 연산

 

 

연산에 필요한 데이터를 디바이스 메모리에 로드한 후, GPU 연산을 수행합니다. 커널 호출을 통해 GPU 연산을 시작하고 모든 데이터는 디바이스 메모리에서 관리됩니다. 디바이스 메모리에 저장된 결과는 다시 호스트 메모리에 전달해야 합니다. 따라서, 아래의 그림과 같이 디바이스 메모리에서 호스트 메모리로 결과를 복사합니다.

 

 

CUDA 프로그램, 디바이스에서 호스트로 데이터 복사

 

 

CUDA 기초 메모리 API (할당, 초기화, 해제)

CUDA 프로그램을 활용해 연산하기 위해 디바이스 메모리 할당, 메모리 사이의 데이터 복사 등 메모리 연산이 필요합니다. 지금부터 디바이스 메모리 할당, 해제, 초기화를 위한 CUDA 기초 메모리 API(Application Programming Interface)에 대해 알아보도록 하겠습니다.

 

 

디바이스 메모리 할당 및 해제: cudaMalloc(), cudaFree()

호스트의 데이터를 복사하기 위해 디바이스 메모리에 사용할 공간을 할당해야 합니다. CUDA 프로그램에서는 cudaMalloc() 함수를 사용해서 디바이스 메모리 공간을 할당합니다. C언어의 malloc() 함수와 상당히 유사하며, 함수의 원형은 아래와 같습니다.

 

cudaError_t cudaMalloc (void ** ptr, size_t size)

 

 

cudaMalloc() 함수의 첫 번째 인자인 ptr은 디바이스 메모리 공간의 시작 주소이고, 두 번째 인자는 할당할 공간의 크기(byte)를 나타내는 size입니다. 아래의 예제는 int형 데이터 32개를 저장할 공간을 할당하는 예제입니다. cudaMalloc() 함수를 통해 device_data_ptr가 가리키는 주소를 시작으로 메모리 공간을 할당합니다. cudaMalloc()을 통해 할당된 메모리는 디바이스 메모리 주소이기 때문에 호스트 코드에서 접근할 수 없습니다.

 

 

#include <stdio.h>

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


int main() {
  int* device_data_ptr;
  cudaError_t cuda_status;

  /* Allocate memory on the device with cudaMalloc */
  cuda_status = cudaMalloc(&device_data_ptr, sizeof(int) * 32);
  printf("cudaMalloc returned: %s\n", cudaGetErrorName(cuda_stats));

  /* Free the memory on the device with cudaFree */
  cuda_status = cudaFree(p_device_data);
  printf("cudaFree returned: %s\n", cudaGetErrorName(cuda_status));
	

  return 0;
}

 

 

할당한 메모리의 사용을 마치면 메모리를 해제해야 합니다. 메모리 자원을 제대로 반환하지 않는다면 메모리 누수(memory leak)가 발생하고, 사용 가능한 메모리가 감소합니다. CUDA 프로그램에서는 C언어의 free() 함수와 유사한 cudaFree() 함수를 사용해서 메모리를 해제합니다. cudaFree() 함수의 원형은 아래와 같습니다.

 

cudaError_t cudaFree(void* ptr)

 

 

cudaMalloc(), cudaFree() 함수 모두 수행 결과인 cudaError_t 열거형을 반환합니다. 메모리 할당, 해제에 성공하면 cudaSuccess(=0)을 반환하고, 문제가 발생하면 에러코드를 반환합니다.

 

 

디바이스 메모리 초기화: cudaMemset()

cudaMalloc()을 통해 메모리 공간을 할당받으면 메모리 공간의 값(garbage value)이 그대로 남아있습니다. CUDA 프로그램에서 메모리를 특정 값으로 초기화하는 함수는 cudaMemset()입니다. cudaMemst() 함수의 인자 ptr은 초기화할 메모리 공간의 시작 주소이고, value는 초기화할 값, size는 메모리 공간의 크기입니다.

 

cudaError_t cudaMemset (void * ptr, int value, size_t size)

 

 

에러 코드 확인: cudaGetErrorName()

cudaMemset() 예제를 살펴보기 전에, 에러 코드(cudaError_t 열거형)를 확인하는 cudaGetErrorName()에 대해 알아보겠습니다. CUDA API의 반환값 대부분은 에러 코드입니다. 개발자가 모든 에러 코드를 알고 개발하거나 버전 호환성을 유지하는 것은 어렵습니다. 따라서, API의 에러 코드를 처리하는 함수를 사용하는 것이 개발에 용이하고 안전합니다. CUDA API에서는 에러 코드를 확인할 수 있는 cudaGetErrorName() 함수를 제공합니다.

 

__host__ __device__ const char* cudaGetErrorName (cudaError_t error)

 

 

메모리를 초기화 함수 cudaMemset()과 에러 코드 확인 함수 cudaGetErrorName()를 활용해서 예제 코드를 완성해 보겠습니다. 아래의 예제는 디바이스 메모리를 할당, 초기화, 해제하는 예제입니다. 아래의 예제는 int형 정수 (1024 * 1024) 개를 담을 수 있는 공간(총 4 MiB)을 할당합니다. 할당된 메모리 공간을 0으로 초기화하고, 할당된 메모리 공간을 해제합니다.

 

CheckDeviceMemory() 함수에서 cudaMemGetInfo() 함수를 활용해 디바이스 메모리의 가용 가능한 크기, 총크기를 확인합니다. 만약 메모리 정보를 얻는 과정에서 문제가 발생하면 오류를 콘솔에 출력하고 리턴합니다.

 

 

#include <stdio.h>

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


void CheckDeviceMemory() {
  size_t free_byte;
  size_t total_byte;

  cudaError_t cuda_status = cudaMemGetInfo(&free_byte, &total_byte);
  if (cuda_status != cudaSuccess) {
    printf("Error: cudaMemGetInfo failed\n");
    return;
  }

  printf("GPU memory free: %lu bytes, total: %lu bytes\n\n", free_byte, total_byte);
}

/*
  We can allocate memory on the device using the cudaMalloc function.
    - In host code, we can't directly access device memory.
*/
int main() {
  int *p_device_data;
  cudaError_t cuda_status;

  printf("Before any memory allocation\n");
  CheckDeviceMemory();

  /* 1. Allocate memory on the device with cudaMalloc */
  cuda_status = cudaMalloc(&p_device_data, sizeof(int) * 1024 * 1024);
  // cuda_status = cudaMalloc(&p_device_data, sizeof(int) * 1024 * 1024 * 1024 * 12);  // Error case
  printf("cudaMalloc returned: %s\n", cudaGetErrorName(cuda_status));
  CheckDeviceMemory();

  /* 2. Memsset the device memory with cudaMemset */
  cuda_status = cudaMemset(p_device_data, 0, sizeof(int) * 1024 * 1024);
  printf("cudaMemset returned: %s\n", cudaGetErrorName(cuda_status));
  CheckDeviceMemory();

  /* Free the memory on the device with cudaFree */
  cuda_status = cudaFree(p_device_data);
  printf("cudaFree returned: %s\n", cudaGetErrorName(cuda_status));
  CheckDeviceMemory();

  return 0;
}

 

 

아래 그림은 예제 코드의 수행 결과입니다. 코드 수행 초기, 총 사용 가능한 메모리의 크기는 약 4 GiB 정도이고, 사용 가능한 메모리의 크기는 약 2.914 GiB입니다. cudaMalloc() 함수를 호출한 후, 사용 가능한 메모리의 크기는 약 2.910 GiB로 할당한 4 MiB만큼 감소한 것을 확인할 수 있습니다. cudaFree()를 통해 메모리를 해제하면 사용 가능한 메모리 크기가 2.914 GiB로 돌아오는 것을 확인할 수 있습니다.

 

 

예제 코드 수행 결과 (정상 케이스)

 

 

아래의 결과는 예제 코드의 Error case 주석을 해제하고 수행한 결과입니다. 메모리의 총크기보다 큰 용량의 메모리를 할당하는 케이스로 CUDA API 함수에서 에러 코드를 반환하는 것을 확인할 수 있습니다. cudaMalloc()를 통해 (1024 * 1024* 1024 * 12)의 용량, 48 GiB의 메모리 할당을 시도하자 cudaErrorMemoryAllocation 에러 코드를 리턴합니다. 이어서 cudaMemset() 또한, 할당되지 않은 메모리에 대해 초기화를 시도하자 cudaErrorInvalideValue()를 반환하는 것을 확인할 수 있습니다.

 

 

예제 코드 수행 결과(에러 케이스)

 

 

마치며

오늘 포스팅에서는 CUDA 프로그램이 GPU를 사용할 때, 프로그램이 어떤 흐름으로 제어되는지 간단하게 확인했습니다. GPU 디바이스 메모리를 활용하기 위해 CUDA API를 활용해 메모리 할당, 초기화, 해제를 수행했습니다. 오늘 소개해 드리지 못한 호스트, 디바이스 간 메모리 복사는 다음 포스팅에 이어서 설명하도록 하겠습니다.

 

오늘 포스팅이 도움이 되셨기를 바라면서 포스팅 마치겠습니다. 좋은 하루 보내시길 바랍니다. 고맙습니다!

 

 

Reference

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