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

 

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

 

목차

• CUDA 기초 메모리 API (호스트, 디바이스 데이터 복사)
• 호스트, 디바이스 간 데이터 복사 예제

 

 

안녕하세요. 지난 포스팅에서 CUDA 프로그램의 기본 연산 흐름과 디바이스(GPU) 메모리를 사용하기 위한 기초 API에 대해 소개했습니다. 이번 포스팅에서는 지난 글에 이어 호스트, 디바이스 간 데이터를 복사하는 방법에 대해 소개합니다. 지난 내용을 확인하고 싶으시면 아래의 글을 읽어보시는 것을 추천드립니다!

 

 

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

 

 

CUDA 기초 메모리 API (호스트, 디바이스 데이터 복사)

호스트 메모리와 디바이스 메모리는 서로 독립된 공간입니다. 다른 장치에 있는 데이터가 필요하면 명시적으로 데이터를 복사하는 과정이 필요합니다. 지금부터 데이터 복사를 위한 CUDA 기초 메모리 API에 대해 알아보도록 하겠습니다.

 

 

장치 간 데이터 복사: cudaMemcpy()

CUDA 프로그램에서 장치 간 데이터 복사를 위해 사용하는 API는 cudaMemcpy() 함수입니다. 이름에서 볼 수 있듯이 C언어의 memcpy()의 사용법과 동작 방식이 유사합니다. cudaMemcpy()는 데이터 복사를 수행하고, 에러 코드 cudaError_t를 반환합니다.

 

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

 

 

cudaMemcpy() 첫 번째 인자 dst는 데이터를 복사할 메모리 공간(destination)의 시작 주소 포인터입니다. 두 번째 인자 src는 복사할 원본 데이터(source) 메모리 공간의 시작 포인터입니다. 세 번째 인자 size는 복사할 데이터의 크기(바이트 단위)이고, 마지막 인자 kind는 데이터의 복사 방향을 설정하는 cudaMemcpyKind입니다.

 

cudaMemcpyKind는 데이터 복사 방향을 나타내는 열거형 변수입니다. 접두어로 cudaMemcpy, 복사 방향(OOOToOOO)으로 구성되어 있어, 어떤 방향으로 복사할지 나타냅니다. 아래의 표는 cudaMemcpyKind 열거형의 항목을 나타냅니다. 

 

 

항목	복사 방향
cudaMemcpyHostToHost	호스트 메모리 to 호스트 메모리
cudaMemcpyHostToDevice	호스트 메모리 to 디바이스 메모리
cudaMemcpyDeviceToHost	디바이스 메모리 to 호스트 메모리
cudaMemcpyDeviceToDevice	디바이스 메모리 to 디바이스 메모리
cudaMemcpyDefault	dst와 src의 포인터 값에 의해 결정 (unified virtual addessing을 지원하는 시스템에서만 사용 가능)

 

 

 

cudaMemcpy()의 첫 번째, 두 번째 인자인 dst, src가 가리키는 메모리 영역의 주소는 kind로 지정한 복사 방향과 일치해야 합니다. 예를 들어 복사 방향(kind)이 cudaMemcpyHostToDevice인 경우, src는 호스트 메모리의 주소, dst는 디바이스 메모리의 주소를 전달해야 합니다. 복사 방향과 포인터 값이 가리키는 메모리 주소가 일치하지 않으면 실행 시 오류가 발생하거나 잘못된 데이터 복사가 발생할 수 있습니다.

 

복사 방향으로 cudaMemcpyDefault를 설정하는 경우, dst와 src로 전달한 주소 값을 통해 실행 시간에 방향성을 판단합니다. 단, cudaMemcpyDefault는 호스트 메모리, 디바이스 메모리를 하나의 메모리 공간처럼 가상화하는 unified virtual addressing을 지원하는 시스템에서만 사용 가능합니다.

 

cudaMemcpyDefault를 사용하면 실행 시간에 복사 방향을 유연하게 처리할 수 있는 장점이 있지만, 명시적으로 방향을 설정하는 것에 비해 코드의 안정성, 가독성 측면에서 불리할 수 있습니다. 따라서, CUDA 프로그램의 유연성, 유지보수성 등을 고려해서 선택적으로 설계해야 합니다.

 

 

호스트, 디바이스 간 데이터 복사 예제

아래는 호스트, 디바이스 사이에서 데이터를 복사하는 예제입니다. 코드를 차근차근 따라가면서 설명하겠습니다. 예제 코드는 다음의 순서로 동작합니다.

 

1. 호스트, 디바이스 메모리 할당 및 초기화 (0으로 초기화)
2. 디바이스 데이터 업데이트(각 원소에 인덱스 정보 저장)
3. 디바이스 데이터를 호스트에 복사
4. 호스트 데이터 0으로 초기화
5. 호스트 데이터를 디바이스에 복사

 

예제 코드

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

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


/* Print the device data */
__global__ void PrintData(int* p_device_data, int size) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < size) {
        printf("Device Data[%d] = %d\n", idx, p_device_data[idx]);
    }
}

/* Set the device data */
__global__ void SetData(int* p_device_data, int size) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < size) {
        p_device_data[idx] = idx;
    }
}

int main() {
    int size = 10;

    /* ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
        1. Set the host, device data
      ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ */
    int* p_host_data = (int*)malloc(size * sizeof(int));
    memset(p_host_data, 0, size * sizeof(int));

    int* p_device_data;
    cudaMalloc((void**)&p_device_data, size * sizeof(int));
    cudaMemset(p_device_data, 0, size * sizeof(int));

    /* ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
        2. Set the host, device data
       ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ */
    printf("Set the host, device data\n");
    SetData <<<(size + 255) / 256, 256 >>> (p_device_data, size);
    for (int i = 0; i < size; i++) {
        printf("Host Data[%d] = %d\n", i, p_host_data[i]);
    }
    PrintData <<<(size + 255) / 256, 256 >>> (p_device_data, size);

    /* Wait for the kernel to finish */
    cudaDeviceSynchronize();
    

    /* ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
        3. Copy the device data to host
       ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ */
    printf("\nCopy the device data to host\n");
    cudaMemcpy(p_host_data, p_device_data, size * sizeof(int), cudaMemcpyDeviceToHost);
    for (int i = 0; i < size; i++) {
        printf("Host Data[%d] = %d\n", i, p_host_data[i]);
    }
    PrintData <<<(size + 255) / 256, 256 >>> (p_device_data, size);

    /* Wait for the kernel to finish */
    cudaDeviceSynchronize();

    /* ~~~~~~~~~~~~~~~~~~~~~~~~~~
        4. Set the host data as 0
       ~~~~~~~~~~~~~~~~~~~~~~~~~~ */
    printf("\nSet the host data as 0\n");
    memset(p_host_data, 0, size * sizeof(int));

    /* ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
        5. Copy the host data to device
       ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ */
    printf("\nCopy the host data to device\n");
    cudaMemcpy(p_device_data, p_host_data, size * sizeof(int), cudaMemcpyHostToDevice);
    for (int i = 0; i < size; i++) {
        printf("Host Data[%d] = %d\n", i, p_host_data[i]);
    }
    PrintData <<<(size + 255) / 256, 256 >>> (p_device_data, size);

    /* Wait for the kernel to finish */
    cudaDeviceSynchronize();

    /* Free the device data */
    cudaFree(p_device_data);
    free(p_host_data);

    return 0;
}

 

 

호스트, 디바이스 메모리 할당 및 초기화

호스트 메모리 할당은 malloc(), 디바이스 메모리 할당은 cudaMalloc() 함수를 통해 수행합니다. 사전에 정의한 size 변수의 값을 크기로 int형 배열을 할당합니다. 메모리가 할당되면 memset()과 cudaMemset()을 호출해 메모리를 0으로 초기화합니다.

 

디바이스 데이터 업데이트

디바이스 데이터 메모리를 새로운 값으로 업데이트하기 위해 SetData() 커널 함수를 사전에 정의했습니다. 해당 커널 함수는 디바이스 메모리 원소에 접근하여 인덱스 값으로 업데이트합니다. 커널 함수 내부에서 현재 스레드, 블록 정보를 통해 인덱스를 계산하고 원소에 업데이트합니다. 스레드, 블록, 그리드 등 레이아웃에 대한 개념은 추후 포스팅에서 자세히 다루도록 하겠습니다.

 

디바이스 데이터를 호스트에 복사

cudaMemcpy 함수를 호출해 디바이스 메모리의 정보를 호스트 메모리에 복사합니다. 사전에 정의한 size 정보를 인자로 전달하고, 복사 방향을 cudaMemcpyDeviceToHost로 설정합니다. 이후, cudaDeviceSynchronize()를 통해 GPU 연산이 모두 종료될 때까지 기다립니다.

 

호스트 데이터 0으로 초기화

호스트에서 디바이스로 데이터 복사하는 기능을 테스트하기 위해 호스트 메모리를 다시 0으로 초기화합니다. memset() 함수를 호출해 호스트 메모리의 데이터를 0으로 초기화합니다.

 

호스트 데이터를 디바이스에 복사

cudaMemcpy함수를 호출해 호스트 메모리 정보를 디바이스 메모리에 복사합니다. 이번에는 복사를 반대로 수행할 수 있도록 복사 방향을 cudaMemcpyHostToDevice로 설정합니다.

 

예제 코드를 수행한 결과는 아래와 같습니다. 결과를 확인하면 디바이스 메모리에 설정한 인덱스 정보가 호스트에 복사되고, 호스트에 초기화된 0 값이 디바이스에 복사되는 것을 확인할 수 있습니다.

 

메모리 복사 예제 수행 결과

 

 

마치며

오늘 포스팅에서는 CUDA API를 활용해서 호스트, 디바이스 간 데이터 복사 기능을 소개했습니다. GPU 디바이스를 효율적으로 사용하기 위해 호스트, 디바이스 메모리 간 데이터 전송 과정이 필수입니다. CUDA 프로그래밍 과정에서 단순히 커널 성능을 고려하는 것 이상으로, 데이터 복사 시 오버헤드를 줄이는 것 또한 중요합니다. 앞으로 CUDA 프로그램 성능을 높이기 위한 성능 최적화 방법에 대해서도 다룰 예정이니 많은 관심 부탁드립니다.

 

오늘 포스팅이 도움이 되셨기를 바라면서 글 마치겠습니다. 고맙습니다.

 

 

Reference

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