CPU와 GPU

CPU는 고수준의 프로세서가 소수 모여 생긴 것.
GPU는 저수준의 프로세서가 다수 모여 생긴 것.
GPU는 그 이름(Graphics)으로부터 알 수 있듯이 화상 처리에 특화된 것으로 화상과 같이 병렬 처리하기 쉬운 신호를 자랑으로 한다.
예를 들어, 제 경우 C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.1\extras\demo_suite\deviceQuery.exe를 실행하면 다음과 같이 출력됩니다.

output

Detected 1 CUDA Capable device(s)

Device 0: "GeForce GTX 750 Ti"
  CUDA Driver Version / Runtime Version          9.1 / 9.1
  ...
  Maximum number of threads per block:           1024
  Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
  Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)
  ...
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

우선 Maximum number of threads per block에 주목해 보면 1024가 되고 있다.
이것은 block 당 얼마나의 thread를 사용할 수 있는지를 나타낸다.
예를 들어 1024*1024 이미지의 픽셀에 1을 더하고 싶다면,
<<< grid(1024,1,1), block(1024, 1, 1)>>>로 처리하면,
픽셀 한 개에 1을 더하는 시간으로 계산이 모두 끝나 버린다.
(참고: Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
block(1024, 1024, 1) 은 할 수 없다, 왜냐하면 블록당 thread의 수의 최대는 1024)
CPU의 경우, 4 코어로 8 thread이기 때문에 thread당 다수의 픽셀이 처리되어 버려,
상당히 처리 속도가 빠르더라도, 1만 더하면 끝이라고 하는 연산에서는 GPU가 유리하다.

GPU 디자인 원칙

① 다수의 간단한 계산 유닛으로 구성되어 있다.
② 병렬 프로그래밍 모델
③지연(latency)에서 처리량(throughput)에 최적화

CUDA Program Diagram

CPU(Host)와 GPU(Device) 사이에 데이터 통신이 있다.
예를 들어, 이미지를 GPU로 처리하는 것을 생각한다.
· CPU가 Hard disk에서 이미지 파일을 읽고 메모리 (Host)에 저장합니다.
・CPU가 GPU에 화상을 처리하는 메모리(Device)를 준비하도록 명령(cudaMalloc)한다.
・GPU가 메모리에 화상 데이터가 들어가는 공간을 준비해, CPU로부터 화상 ​​데이터가 송신되는 것을 기다린다.
· CPU가 GPU에 이미지 데이터를 전송합니다 (cudaMemcpy, cudaMemcpyHostToDevice).
· GPU에서 이미지를 처리하는 커널 (cuda 프로그램)을 실행하고 결과를 Host 측으로 전송한다 (cudaMemcpy, cudaMemcpyDeviceToHost).

그리드와 블록과 스레드

Grid⊃Block⊃Thread
위 그림의 경우 kernel <<< dim3(2, 3, 1), dim3(3, 4, 1) >>>

Map 패턴

GPU가 자랑하는 패턴으로 배열의 요소를 처리하고 그대로 같은 장소에 저장한다.
다른 thread의 읽고 쓰기에 의해, 자신의 처리가 방해받지 않는 패턴으로
알고리즘 등 궁리하지 않고, 그대로 병렬 처리가 가능.

Gray화 코드 예

gray.cpp

#include <cuda.h>
#include <cuda_runtime.h>
#include <opencv2/core.hpp>
#include <opencv2/highgui.hpp>

using namespace cv;

void bgr_to_grayscale(uchar3 *d_inImg, unsigned char *d_outImg,
    size_t numRows, size_t numCols);

int main(int argc, char const *argv[])
{
    // init host memory
    Mat inImg = imread("/home/username/Pictures/lena.png");
    Mat outImg(inImg.rows, inImg.cols, CV_8UC1);
    uchar3 *h_inImg, *d_inImg;
    unsigned char *h_outImg, *d_outImg;
    h_inImg = (uchar3 *)inImg.ptr<unsigned char>(0);
    h_outImg = (unsigned char *)outImg.ptr<unsigned char>(0);

    // init device memory
    size_t numPixels = inImg.rows * inImg.cols;
    cudaMalloc(&d_inImg, sizeof(uchar3) * numPixels);
    cudaMemcpy(d_inImg, h_inImg, sizeof(uchar3) * numPixels, cudaMemcpyHostToDevice);
    cudaMalloc(&d_outImg, numPixels);

    // cuda process
    bgr_to_grayscale(d_inImg, d_outImg, inImg.rows, inImg.cols);

    // device to host
    cudaMemcpy(h_outImg, d_outImg, numPixels, cudaMemcpyDeviceToHost);

    // show result
    imshow("Gray Image", outImg);
    imshow("Original Image", inImg);
    waitKey(0);
    destroyAllWindows();

    return 0;
}

gray.cu

#include <cuda.h>
#include <cuda_runtime.h>

__global__
void cuda_bgr_to_grayscale(uchar3 *d_inImg,
    unsigned char *d_outImg,
    int numRows, int numCols) {
    int idx = threadIdx.x + blockIdx.x * numCols;
    d_outImg[idx] = float(d_inImg[idx].x) * 0.114f +
                    float(d_inImg[idx].y) * 0.587f +
                    float(d_inImg[idx].z) * 0.299f;
}

void bgr_to_grayscale(uchar3 *d_inImg, unsigned char *d_outImg,
    size_t numRows, size_t numCols) {
    cuda_bgr_to_grayscale <<<numRows, numCols>>> (d_inImg, d_outImg, numRows, numCols);
};

compile(ubuntu)

nvcc gray.cpp gray.cu `pkg-config --cflags --libs opencv`

위의 예를 사용하여 gray화하면 lena.png의 처리가 opencv의 cv::cvtColor의 gray화보다 30배 정도 빨라졌다.