Intro to Parallel Programming, Lesson 1

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)はできない、なぜならブロック当たりスレッドの数の最大は1024)
CPUの場合、4コアで8スレッドだからスレッド当たり多数のピクセルが処理されてしまい、
結構処理速度が速いとしても、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で画像を処理するkernel(cudaプログラム)を実行し、結果をHost側に送信する(cudaMemcpy, cudaMemcpyDeviceToHost)。

GridとBlockとThread

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倍程度早くなった。