CUDA 프로그래밍 Study Ch 3. CUDA 프로그램의 기본 흐름

김재만·2023년 10월 2일

CUDA Programming Study

목록 보기
3/12
  • 들어가기에 앞서
    : 해당 게시물은 책 'CUDA 기반 GPU 병렬 처리 프로그래밍' (김덕수 지음, 비제이퍼블릭) 을 통해 CUDA 프로그래밍을 공부하면서 정리한 것이다.

1. CUDA 프로그램의 구조 및 흐름

  • CPU가 운영체제와 같은 컴퓨터 시스템의 기본 연산 장치이며, GPU등 다른 연산 장치를 사용하기 위해서는 호스트 코드에서 커널을 호출해야함.

  • CPU와 GPU는 서로 독립된 장치
    : 사용하는 메모리 영역도 다르다.

  • 모든 데이터는 기본적으로 호스트 메모리에 저장
    : GPU를 이용해서 데이터를 처리하기 이해서는 호스트 메모리에 있는 데이터를 디바이스 메모리로 복사해주어야 함.

CUDA 프로그램의 흐름

  1. 호스트 메모리에서 디바이스 메모리로의 입력 데이터 복사.
  2. GPU 연산.
    : 커널 호출 통해 시작, 모든 데이터는 디바이스 메모리에서 관리.
    : GPU가 수행한 연산 결과 역시 디바이스 메모리에 저장.
  3. 디바이스 메모리에서 호스트 메모리로의 결과 데이터 복사.

2. CUDA 기초 메모리 API

2.1 디바이스 메모리 공간 할당 및 초기화 API

cudaMalloc()

  • 디바이스 메모리 공간 할당.

    cudaError_t cudaMalloc (void ** ptr, size_t size)

    : 첫 번째 인자인 ptr은 할당된 디바이스 메모리 공간의 시작 주소를 담을 포인터 변수의 주소
    : 두 번째 인자인 size는 할당할 공간의 크기로 단위는 byte.
    : size_t 형에 대해 다시 짚고 넘어가자면...

    size_t can store the maximum size of a theoretically possible object of any type (C99 원문)

    : 반환값은 cudaError_t enumeration 으로, API 호출에 성공하면 cudaSuccess (=0)가 반환.
    : 실패시 대표적인 에러 코드는 cudaErrorMemoryAllocation (=2)

  • 디바이스 메모리 공간에 int형 데이터 32개를 담을 공간이 필요하다면, 

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

 int main(void)
 {
 	int *dDataPtr;
 	coudaMalloc(&dDataPtr, sizeof(int)*32);
 }

    : 할당된 메모리 공간의 시작 주소가 dDataPtr 포인터 변수에 저장. -> dDataPtr 포인터 변수를 통해 할당된 메모리 영역에 접근
    : 디바이스 메모리상의 주소는 호스트 코드에서 직접 접근할 수 없다.

cudaFree()

  • 디바이스 메모리 해제

    	cudaError_t cudaFree (void* ptr)

    : 인자 ptr은 해제할 메모리 공간을 가리키는 포인터 변수
    : 반환값은 cudaMalloc()과 같이 cudaError_t enumeration

  • cudaMalloc() 예제에서 할당한 디바이스 멤리 공간을 해제하려면

    cudaFree(dDataPtr);

cudaMemset()

  • 디바이스 메모리 초기화
    : cudaMalloc()을 통해 메모리 공간을 할당 -> 해당 메모리 공간에 남아 있던 garbage value가 그대로 남아있기 때문에, 디바이스 메모리 공간을 특정 값으로 초기화하는 cudaMemset() 함수가 필요하다. 

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

    : 첫 번째 인자 ptr은 값을 초기화 할 메모리 공간의 시작 주소
    : 마지막 인자 size는 초기화 할 메모리 공간의 크기
    : 두 번째 인자 value는 해당 공간의 각 바이트를 초기화 할 값을 의미.

  • cudaMalloc() 예제에서 할당된 메모리 공간의 값을 모두 0으로 초기화하고 싶다면,

    cudaMemset(dDataPtr, 0, sizeof(int)*32)

cudaGetErrorName()

  • 에러 코드 확인

    	__host__ __device__ const char* cudaGetErrorName() (cudaError_t error)

    : 원형 앞에 __host__ 와 __device__ 가 모두 붙어있는데, 이는 호스트와 디바이스 코드 모두에서 사용 가능함을 의미

메모리 할당/초기화/해제 예제

  • 예제 코드
 %%cu

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

void checkDeviceMemory(void)
{
    size_t free, total;

    cudaMemGetInfo(&free, &total);
    printf("Device memory (free / total) = %lld / %lld bytes\n", free, total);
}

int main(void)
{
    int* dDataPtr;
    cudaError_t errorCode;

    checkDeviceMemory();
    errorCode = cudaMalloc(&dDataPtr, sizeof(int)* 1024 * 1024);
    printf("cudaMalloc - %s\n", cudaGetErrorName(errorCode));
    checkDeviceMemory();

    errorCode = cudaMemset(dDataPtr, 0, sizeof(int) * 1024 * 1024);
    printf("cudaMemset - %s\n", cudaGetErrorName(errorCode));

    errorCode = cudaFree(dDataPtr);
    printf("cudaFree -%s\n", cudaGetErrorName(errorCode));
    checkDeviceMemory();
}
  • 실행 결과는 다음과 같다.
Device memory (free / total) = 15727394816 / 15835398144 bytes
cudaMalloc - cudaSuccess
Device memory (free / total) = 15723200512 / 15835398144 bytes
cudaMemset - cudaSuccess
cudaFree -cudaSuccess
Device memory (free / total) = 15727394816 / 15835398144 bytes

: checkDeviceMemory() 함수는 현재 디바이스 메모리 상태를 출력하는 함수.
: 디바이스 메모리의 총 크기와 현재 사용 가능한 공간의 크기를 반환해주는 cudaMemGetInfo() 함수를 사용하였다. 가용 메모리 크기는 free에, 사용하는 GPU가 가진 총 디바이스 메모리 크기는 total에 바이트 단위로 반환.
: 코드에서 메모리 할당 크기를 디바이스 메모리 크기보다 크게 설정하면 에러가 발생한다.

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

cudaMemcpy()

  • 장치 간 데이터 복사

    cudaError_t cudaMemcpy (void* dst, const void* src, size_t size, enum cudaMemcpyKind kind)

    : 첫 번째 인자 dst는 데이터가 복사될 메모리 공간의 시작 주소 (destination)
    : 두 번째 인자 src는 복사할 원본 데이터가 들어있는 메모리 공간의 시작 주소 (sorce)
    : 세 번째 인자 size는 복사할 데이터의 크기 (byte단위)
    : 마지막 인자 kind는 cudaMemcpyKind열거형 변수, 데이터 복사의 방향을 설정. 다음은 cudaMemcpyKind 열거형의 항목들이다.

    cudaMemcpyHostToHost Host -> Host
    cudaMemcpyHostToDevice Host -> Device
    cudaMemcpyDeviceToHost Device -> Host
    cudaMemcpyDeviceToDevice Device -> Device
    cudaMemcpyDefault dst와 src의 포인터 값에 의해 결정

    : cudaMemcpyDefault 는 호스트 메모리와 디바이스 메모리를 하나의 메모리 공간처럼 가상화 해주는 unified virtual addressing을 지원하는 시스템에서만 사용가능.
    : 복사 방향을 명시적으로 지정하는 것을 추천.

장치 간 데이터 복사 예제

  • 에제 코드
%%cu

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

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

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

int main(void)
{
    int data[10] = { 0 };
    for (int i = 0; i < 10; i++)
      data[i] = 1;
  
    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);
}
  • 코드 실행 결과는 다음과 같다.
Data in device: 0000000000
Host -> Device: 1111111111
Device -> Host: 2222222222

기타 데이터 복사 API

  • 2차원 또는 3차원 데이터의 복사를 도와주는 cudaMemcpy2D()와 cudaMemcpy3D()
  • 비 동기적 복사를 수행하는 API인 cudaMemcpyAsync() (동기화의 개념 및 비동기적 데이터 복사 등에 대한 내용은 뒤에서 다룬다.)

3. CUDA로 작성하는 벡터의 합 프로그램

3.1 Device 메모리 할당

int *da, *db, *dc

cudaMalloc(&da, memSize); cudaMemset(da, 0, memSize);
cudaMalloc(&db, memSize); cudaMemset(db, 0, memSize);
cudaMalloc(&dc, memSize); cudaMemset(dc, 0, memSize);

3.2 입력 벡터 복사

cudaMemcpy(da, a, memSize, cudaMemcpyHostToDevice);
cudaMemcpy(db, b, memSize, cudaMemcpyHostToDevice);

3.3 벡터 합 커널 호출

__global__ void vecAdd(int* _a, int* _b, int* _c)
{
	int tID = threadIdx.x;
    _c[tID] = _a[tID] + _b[tID];
}

// 호스트 코드에서 커널 호출

vecAdd <<<1, NUM_DATA>>> (da, db, dc);

3.4 결과 벡터 복사 및 디바이스 메모리 해제

cudaMemcpy(c, dc, memSize, cudaMemcpyDeviceToHost);

cudaFree(da);
cudaFree(db);
cudaFree(dc);

3.5 전체 코드

%%cu

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

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

// The size of the vector
#define NUM_DATA 1024

// Simple vector sum kernel
__global__ void vecAdd(int* _a, int* _b, int* _c)
{
	int tID = threadIdx.x;
    _c[tID] = _a[tID] + _b[tID];
}

int main(void)
{
    int* a, * b, * c, * hc;
    int *da, *db, *dc;

    int memSize = sizeof(int) * NUM_DATA;
    printf("%d elements, memSize = %d bytes\n", NUM_DATA, memSize);

    // memory allocation on the host_side
    a = new int[NUM_DATA]; memset(a, 0, memSize);
    b = new int[NUM_DATA]; memset(b, 0, memSize);
    c = new int[NUM_DATA]; memset(c, 0, memSize);
    hc = new int[NUM_DATA]; memset(hc, 0, memSize);

    // Data generation
    for (int i = 0 ; i < NUM_DATA; i++)
    {
        a[i] = rand() % 10;
        b[i] = rand() % 10;
    }

    //Vector sum on the host
    for (int i = 0 ; i < NUM_DATA ; i++)
      hc[i] = a[i] + b[i];
    
    // Memory allocation on te device-side
    cudaMalloc(&da, memSize); cudaMemset(da, 0, memSize);
    cudaMalloc(&db, memSize); cudaMemset(db, 0, memSize);
    cudaMalloc(&dc, memSize); cudaMemset(dc, 0, memSize);

    // Data copy : Host -> Device
    cudaMemcpy(da, a, memSize, cudaMemcpyHostToDevice);
    cudaMemcpy(db, b, memSize, cudaMemcpyHostToDevice);

    // Kernel call
    vecAdd <<<1, NUM_DATA>>> (da, db, dc);

    // Copy results : Device -> Host
    cudaMemcpy(c, dc, memSize, cudaMemcpyDeviceToHost);

    // Release device Memory
    cudaFree(da);
    cudaFree(db);
    cudaFree(dc);

    // Check results
    bool result = true;
    for (int i = 0 ; i < NUM_DATA ; i++)
    {
        if (hc[i] != c[i])
        {
            printf("[%d] The result is not matched! (%d, %d)\n", i, hc[i], c[i]);
            result = false;
        }
    }

    if (result)
      printf("GPU works well!\n");
    
    // Release host memory
    delete[] a; 
    delete[] b;
    delete[] c;

    return 0;
}
  • 실행 결과는 다음과 같다.
1024 elements, memSize = 4096 bytes
GPU works well!

  • 벡터의 크기를 늘려보면, 검증 과정에서 결과가 일치하지 않는다는 메시지가 출력된다. 커널이 연산을 제대로 수행하지 못하는 것이다.
    : 4장의 스레드 계층 구조 및 실행모델, 5장의 스레드 인덱싱을 학습하면서 이를 해결해 보자.

  • 병렬 처리 알고리즘의 설계 및 구현은 논리적 버그가 발생할 가능성이 높으며, 디버그 난이도가 훨씬 높기 때문에, 호스트 코드와 같이 검증된 코드의 연산 결과와 비교해보기를 추천한다.

4. CUDA 알고리즘의 성능 측정

  • 소요 시간으로 그 성능을 측정한다.

4.1 커널 수행시간

  • GPU 연산은 커널 호출을 통해 진행.
    : 성능 측정의 주요 지점 중 하나는 커널 부분이다.

  • 커널 호출 시 디바이스에게 명령을 전달한 후 프로그램 흐름의 제어권을 바로 호스트에게 반환한다.
    : 즉, 호스트는 디바이스에게 커널 수행을 요청하고 바로 다음 작업으로 진행한다.
    : 이는 호스트와 디바이스의 Asynchronous한 수행이 가능하게 하는 특징으로, 두 장치를 동시에 사용 가능하다는 의미.

  • 정확한 시간 측정을 위해서는, 커널 연산이 종료될 때 까지 대기한 후 시간 측정을 종료해야한다.
    : 디바이스가 수행 죽인 작업이 끝날 때까지 대기하는 CUDA Synchronous 함수는 cudaDeviceSynchronize()

  • CUDA API는 기본적으로 순차적으로 진행된다.
    : 호스트 코드를 디바이스 코드 제어에만 사용하는 경우에는 cudaDeviceSynchronize()와 같은 동기화 함수를 사용하지 않아도 된다.

    4.2 데이터 전송 시간

  • CUDA 알고리즘의 성능을 판단할 때는 반드시 데이터 전송시간도 함께 고려해야 한다. (실제로 컴퓨터 구조에서 메모리 접근 시간은 상당한 성능 저하를 일으킨다. 이를 보완하기 위해 Cache를 사용하며, Cache 실패시 메모리 접근에 의한 성능 저하를 줄이기 위한 기법들도 있다.)
    : cudaMemcpy() 함수는 호스트 코드와 동기적으로 수행된다.

이 외에도 CUDA를 통해 GPU를 사용하기 위해 데이터 변환 등과 같은 추가 작업이 있다면 해당 작업에 소요되는 시간도 포함해 성능을 측정해야 한다.

profile
김재만
Hardware Engineer가 되자
이전 포스트

CUDA 프로그래밍 Study Ch 2. CUDA 개요

다음 포스트

CUDA 프로그래밍 Study Ch 4. CUDA 스레드 계층

0개의 댓글