호스트와 디바이스

호스트

호스트는 일반적으로 CPU를 의미한다. 따라서 호스트 코드는 CPU에서 실행되는 코드를 의미한다. 또한 호스트 메모리는 CPU가 사용하는 시스템 메모리이다. (DRAM)

디바이스

디바이스는 일반적으로 GPU를 의미한다. 따라서 디바이스 코드는 GPU에서 실행되는 코드를 의미한다. 또한 디바이스 코드는 GPU가 사용하는 GPU 메모리이다.

CUDA 프로그램

CUDA 프로그램은 호스트 코드와 디바이스 코드로 구성된다. 프로그램 실행 시 처음 호출되는 코드는 CPU에서 프로세스를 할당하기 때문에, 호스트 코드가 통상 같이 있다. CUDA 프로그램에서 호스트 코드는 gcc와 같은 컴파일러로, 디바이스 코드는 NVCC 컴파일러로 컴파일한다. C++에서 소스 파일은 “.cpp”, 헤더 파일은 “.h” 에 대응하면, CUDA는 소스 파일은 “.cu”, 헤더 파일은 “.cuh” 이다.

# 간단한 CUDA 프로그램 예시

#include "cuda_runtime.h"
#include "device_launch_paramters.h"
#include <stdio.h>

__global__ void helloCUDA(void) {
  printf("GPU에서 CUDA 프로그램 실행 \n");
}

int main(void) {
  printf("CPU에서 CUDA 프로그램 실행 \n");
  helloCUDA <<<1, 10>>>();
  return 0;
}
  • host
    CPU가 함수를 호출하고, CPU에서 실행한다.
  • device
    GPU가 함수를 호출하고, GPU에서 실행한다.
  • global
    CPU가 함수를 호출하고, GPU에서 실행한다.
  • 커널 함수
    디바이스 함수는 디바이스 코드만 호출 가능하다. 따라서 호스트에서 호출할 수 없다. 이 경우에 특별한 함수가 필요한데 global 키워드이다. 이 키워드로 작성된 함수는 호스트에서 호출하고, 디바이스에서 실행한다.
  • 커널 실행과 그 구성
    커널은 호스트가 디바이스에게 연산을 명령하는 수단이다. CUDA 스레드들의 동작을 정의하는 함수이며, 커널 호출 시에 몇 개의 블록과 CUDA 스레드가 연산을 수행할지 미리 명시하여야 한다. (이후 설명)
# 커널함수<<블록 수, 스레드 수>>
# 아래 함수는 블록 수 1, 스레드 수 10

helloCUDA <<<1, 10>>>();

CUDA 프로그램의 구조와 흐름

메모리의 구분

  • 시스템 메모리 = 메인 메모리 = 호스트 메모리
  • GPU 메모리 = 디바이스 메모리
  • CPU와 GPU는 독립된 장치로서, 사용하는 메모리 영역이 다르다. GPU가 데이터를 읽어들일 때는 GPU 메모리를 사용하지만, 기본적으로 CPU가 프로세스를 시스템에 할당한다. 따라서 모든 데이터는 사용 시에 호스트 메모리에 저장한다.
    1. 호스트 메모리에서 디바이스 메모리로 데이터 복사 (cudaMemcpy)
      이때 다른 저장공간 및 통신으로 받는 데이터로 호스트 메모리에 저장하였다가 디바이스 메모리로 복사한다.
    2. 호스트가 CUDA 커널을 호출하여 GPU에게 프로세스를 명령한다. (GPU 연산)
      이때 모든 데이터는 디바이스 메모리 상에서 관리한다.
    3. 연산 결과 데이터는 다시 디바이스 메모리에서 호스트 메모리로 옮긴다.

CUDA 메모리 API

  • cudaMalloc()
    데이터를 호스트메모리에서 디바이스 메모리로 복사한다.
# **ptr: 디바이스 메모리 공간의 시작 주소를 담을 포인터 변수의 주소
# size: 할당한 메모리 공간의 크기
# cudaError_t cudaMalloc (void** ptr, size_t size)

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

int main(void) {
  int *dDataPtr;
  cudaMalloc(&dDataPtr, sizeof(int)*32);
}
  • cudaFree()
    할당받은 디바이스 메모리의 사용을 마치면, 메모리를 해제하여 메모리 자원을 반환한다. 더 이상 사용하지 않는 데이터를 디바이스 메모리에서 계속 점유하고 있으면, 사용 가능한 메모리가 줄어든다.
# *ptr: 해제할 메모리 공간을 가리키는 포인터 변수
# cudaError_t cudaFree (void* ptr)

cudaFree(dDataPtr);
  • cudaMemeset()
    디바이스 메모리 공간을 특정 값으로 초기화해줄 때 사용한다.
# *ptr: 초기화할 메모리 공간의 시작 주소
# size: 초기과할 메모리 공간의 초기화할 값
# size: 초기화할 메모리 공간의 크기
# cudaError_t cudaMemst (void* ptr, int value, size_ size)

cudaMemst(dDatPtr, 0, sizeof(int)*32);
  • cudaGetErrorName()
    CUDA API 반환값은 cudaError_t을 나열한 것이 많다. 에러의 경우가 매우 많다. 따라서 오류를 cudaGetErrorName()을 통해 확인하면 좋다.
# error: cudaError_t로 나열된 에러 메세지
# __host__ __device__ const char* cudaGetErrorName (cudaError_t error)

errorCode = cudaFree(dDataPtr)
cudaGetErrorName(errorCode);
  • 디바이스 메모리 할당/초기화/해제 예제

호스트 - 디바이스 메모리 간 데이터 복사 API

  • 장치 간 데이터 복사 예제
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>

__global__ void printData(int* _dDatPtr) {
  printf("%d", _dDataPtr[threadidx.x]);
}

__global__ void setData(int* _dDataPtr) {
  _dDataPtr[theradIdx.x] = 2;
}

int main(void) {
  # data 변수에 1을 할당
  int data[10] = { 0 };
  for (int i = 0; i < 10; i++) {
    data[i] = 1;
  }

  # 데이터를 저장할 포인터 변수 선언
  # 메모리를 할당하고 0으로 초기화
  int* dDataPtr;
  cudaMalloc(&dDataPtr, sizeof(int) * 10);
  cudaMemset(dDataPtr, 0, sizeof(int) * 10);
  
  printf("Data in Device: ");
  printData <<<1, 10>>>(dDataPtr);

  # 호스트 메모리의 데이터를 디바이스 메모리로 카피
  cudaMemcpy(dDataPtr, data, sizeof(int) * 10, cudaMemcpyHostToDevice);
  printf("\nHost -> Device: ");
  printData <<<1, 10>>>(dDataPtr);

  setData<<1, 10>>(dDataPtr);

  # 디바이스 메모리의 데이터를 호스트 메모리로 카페
  cudaMemcpy(data, dDataPtr, sizeof(int)*10, cudaMemcpyDeviceToHost);
  printf("\nDevice -> Host: ");
  for (int i = 0; i < 10; i++) {
    printf("%d", data[i]);
  }

  cudaFree(dDataPtr);
}
  • 기타 데이터 복사 API
# CUDA 문서에서 데이터 복사와 관련된 API를 찾아보면 다양한 함수가 있다.

cudaMemcpy2D()
cudaMemcpy3D()
... ...