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

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

 

목차

• 벡터의 합을 위한 스레드 레이아웃
• 벡터의 합 예제 코드

 

 

'벡터의 합 구하기' 포스팅에서 CUDA 프로그램을 작성해서 벡터의 합을 구했습니다. 당시 구현한 예제 프로그램은 최대 1,024개의 스레드를 활용해 병렬적으로 벡터의 합을 계산했습니다. 사용하는 스레드가 제한적이기 때문에 병렬 컴퓨팅의 능력을 제대로 발휘하지 못했고, 1024개의 스레드를 초과했을 때, 커널이 제대로 연산을 수행하지 못했습니다. 오늘 포스팅에서는 지난 포스팅에서 다룬 스레드 계층을 활용해서 더 큰 규모의 벡터의 합을 구해보도록 하겠습니다.

 

 

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

CUDA 프로그래밍 - CUDA 스레드 계층

 

 

벡터의 합을 위한 스레드 레이아웃

'벡터의 합 구하기(1)'에서 설명드린 예제에서 벡터의 합을 구하기 위해 아래와 같이 커널 함수를 호출했습니다. AddVectors 커널 함수 호출에 사용한 스레드 레이아웃은 x-차원 길이가 SIZE_VECTOR인 1차원 블록 하나를 갖는 그리드입니다.

 

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

 

 

CUDA 공식 가이드의 compute capability에 따라 블록의 사이즈는 최대 1,024입니다. 한 블록이 가질 수 있는 최대 스레드 수는 1,024이고, 따라서, 커널 함수의 스레드 레이아웃에서 SIZE_VECTOR는 1,024를 넘을 수 없습니다. 1024보다 큰 벡터를 처리하기 위해 블록의 크기뿐만 아니라 블록의 개수도 고려해야 합니다.

 

그렇다면 몇 개의 블록을 사용해야 할까요? 위의 코드처럼 1,024개의 스레드를 사용한다고 가정하면 총데이터를 1,024로 나눈 것만큼의 블록이 필요합니다. 하지만, 스레드 계층에서 설명드린 것처럼, 그리드의 각 차원의 길이는 정수로만 정의가 가능합니다. 따라서, SIZE_VECTOR를 스레드의 개수(1,024)로 나누고 올림 처리하여 모든 데이터가 연산이 되도록 설정해야 합니다.

 

AddVectors<<<ceil(SIZE_VECTOR/1024), 1024>>>(p_device_vec1, p_device_vec2, p_device_out, SIZE_VECTOR);

 

 

예를 들어, 크기가 5,000인 두 벡터의 합을 구한다고 가정하겠습니다. 5,000을 1,024로 나누면 몫 4, 나머지 952이기 때문에 1,024개를 계산할 4개의 블록과 952개를 계산할 1개의 블록이 필요합니다. 아래 그림은 크기가 5,000인 두 벡터의 합을 구하기 위해 블록을 할당하는 것을 나타냅니다.

 

 

벡터의 합을 위한 스레드 레이아웃

 

 

위 그림이 나타내는 벡터의 합을 구하기 위한 스레드 레이아웃은 아래와 같습니다. AddVectors 커널 함수에서 블록 5개와 스레드 1,024개를 할당해서 연산을 수행합니다.

 

AddVectors<<<5, 1024>>>(p_device_vec1, p_device_vec2, p_device_out, SIZE_VECTOR);

 

 

그렇다면 AddVectors의 커널 함수는 어떻게 구성해야 할까요? 블록 내부에서는 같은 스레드 번호를 가지지 않지만 서로 다른 블록에서는 같은 스레드 번호를 가질 수 있습니다. 따라서, 블록 번호를 사용해 구분해야 합니다. 여러 블록을 사용할 때, 어떻게 독립적인 스레드를 할당해야 할까요?

 

스레드 계층에서 소개했던 것처럼 커널 함수 내부에서 블록 차원의 크기 blockDim과 인덱스 blockIdx 정보를 얻을 수 있습니다. 즉, 블록의 크기와 인덱스 정보, 스레드 정보를 알면 연산을 수행할 스레드를 특정할 수 있습니다. 아래의 그림은 연산에 사용할 스레드를 지정하는 것을 나타냅니다.

 

 

커널 함수 내부, 특정 블록의 스레드 지정하기

 

 

벡터의 합 예제 코드

아래는 16 MiB 크기를 가지는 두 벡터의 합을 구하는 예제입니다. SIZE_VECTOR는 16 * 1,024 * 1,024, BLOCK_SIZE는 1,024로 설정했습니다. 1,024개를 초과하는 벡터의 합을 구하기 위해 blockDim, blockIdx 정보를 활용해서 스레드를 지정합니다. 또한, 스레드가 SIZE_VECTOR를 초과하는 범위에 접근하지 못하도록 접근 범위를 검사합니다.

 

 

#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 16 * 1024 * 1024
#define BLOCK_SIZE 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 + blockIdx.x * blockDim.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 <<<ceil(static_cast<float>(SIZE_VECTOR)/BLOCK_SIZE), BLOCK_SIZE >>> (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;
}

 

 

아래는 예제 코드를 실행한 결과를 나타냅니다. 16 MiB 크기의 벡터를 더하기 위해 CPU는 107ms가 소요되었고, 커널 함수에서는 5ms가 소요되었습니다. 연산 속도만 고려했을 때, GPU가 약 21배 정도 빠른 것을 확인할 수 있습니다. 작은 크기의 벡터를 처리할 때와 달리 크기가 큰 벡터에 대해서는 GPU가 명백히 더 빠른 연산 성능을 보여줍니다.

 

벡터의 합 CPU, GPU 연산 성능 비교

 

 

위의 결과에서 볼 수 있듯 호스트, 디바이스 간 데이터 전송에도 시간이 소요되는 것을 확인할 수 있습니다. 호스트에서 디바이스로 데이터를 복사할 때 41ms, 디바이스에서 호스트로 복사할 때 21ms 소요되었습니다. 따라서, 데이터 복사까지 고려하면 67ms 소요되었습니다. CUDA 프로그램의 성능을 비교할 때, 커널의 연산 시간뿐만 아니라 데이터 전송 시간까지 함께 고려해야 합니다.

 

 

마치며

오늘 포스팅에서는 크기 1,024를 초과하는 벡터의 합을 구하기 위해 스레드 레이아웃을 설정하는 방법과 커널 함수에서 스레드를 인덱싱하는 방법에 대해 소개했습니다. 예제 코드를 통해 16 MiB 크기의 두 벡터의 합을 구하고 CPU, GPU의 성능을 비교했습니다. 1,024 크기 벡터의 합을 구했던 예제와 달리 크기가 커지면서 GPU가 명백히 더 빠른 연산 성능을 보여주는 것을 확인할 수 있었습니다. 다음 시간에는 연산을 더 효율적으로 수행하기 위해 스레드 인덱싱하는 방법에 대해 알아보도록 하겠습니다.

 

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

 

 

Reference

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