QtでCUDAを使う

CUDA

NVIDIAのGPUを汎用計算に使えるCUDAという開発環境がある。これを使うとCPUと比べて数十倍高速化することがある。CPUで数十倍高速化するのは何年先か分からないので、結構すごい。ただし、並列処理が可能なプログラムしか高速化しないので、CUDAで高速化するプログラムの方が珍しいのが残念なところ。

グレースケール化

前回OpenCVを使って画像のグレースケール化をしたけど、今回はCUDAを使ってグレースケール化をしてみよう。
NTSCのグレースケール化は人は青に対しては鈍感という特性を考慮して、青の係数を小さくしてより自然に見せる手法らしい。

このような式でグレースケール値を算出できる。

QtでCUDAを使う

CUDAは.cuファイルを専用のnvccでコンパイルするので、g++は直接使用できない。どうやってCUDAのプログラムを使うかというと、.cuファイルにある関数を.cppファイルで前方宣言すれば良い。.cuファイルをnvccでコンパイルしてできたオブジェクトファイルをQtのプロジェクトフォルダに入れてから、Qt Creatorでコンパイルすると動くはず。
でも、わざわざCUDAのオブジェクトファイルを移動させるのは面倒なのでQt Creatorで自動的にコンパイルさせたい。調べてみると、forumに方法があったけど、うちの環境だと上手くいかなかったので、適当に書き換えてみた。

# CUDA .proファイルの下の方に書く
CUDA_SOURCES = $$system(echo $$SOURCES | grep -oe "\\\S*\\\.cu")
SOURCES = $$system(echo $$SOURCES | grep -oe "\\\S*\\\.cpp")
unix {
  CUDA_DIR = /usr/local/cuda
  CUDA_CC = $$CUDA_DIR/bin/nvcc
  message("nvcc resides in :" $$CUDA_CC)

  INCLUDEPATH += $$CUDA_DIR/include

  INCLUDEPATH += /home/foo/Program/NVIDIA_GPU_Computing_SDK/C/common/inc
  QMAKE_LIBDIR += $$CUDA_DIR/lib64
  QMAKE_LIBDIR += /home/foo/Program/NVIDIA_GPU_Computing_SDK/C/lib
  LIBS += -lcudart -lcutil

  cuda.output = ${QMAKE_FILE_BASE}.o
  cuda.commands = $$CUDA_CC -c -use_fast_math -arch sm_20  $$join(INCLUDEPATH,'" -I "','-I "','"') ${QMAKE_FILE_NAME} -o ${QMAKE_FILE_OUT}
}
cuda.input = CUDA_SOURCES
QMAKE_EXTRA_UNIX_COMPILERS += cuda

これでCtrl+Rだけでコンパイルできるようになった。
ただ、実行時にlibcudart.so.3: cannot open shared object file:とエラーが出たので、Qt Creatorのプロジェクトのビルド時の環境変数に変数LD_LIBRARY_PATH、値/usr/local/cuda/lib64を追加した。

プログラム

前回のプログラムをCUDA対応に変更する。
デザインはこんな感じ。QGraphicsViewを使ってる。

widget.cpp

#include "widget.h"
#include "ui_widget.h"
#include "opencv2/highgui/highgui.hpp"
#include "opencv2/imgproc/imgproc.hpp"
#include <vector_types.h>
#include <QImage>
#include <QDebug>
//#include <omp.h>

namespace gpu {
void runGrayScale(unsigned char *color_image,
                  const int rows,
                  const int cols,
                  unsigned char *gray_image);
void runThreadSyncronize();
} //namespace gpu

Widget::Widget(QWidget *parent) :
    QWidget(parent),
    ui(new Ui::Widget)
{
    ui->setupUi(this);

    // 元画像の表示
    cv::Mat color_image = cv::imread("Parrots.bmp");

    QImage image(color_image.data,
                 color_image.cols,
                 color_image.rows,
                 QImage::Format_RGB888);
    image = image.rgbSwapped();

    ui->graphicsView->setScene(&scene_);
    pixmap_item_ = scene_.addPixmap(QPixmap::fromImage(image));
}

Widget::~Widget()
{
    delete ui;
}

// グレースケール化
void Widget::on_pushButton_clicked()
{
    cv::Mat color_image = cv::imread("Parrots.bmp");
    int rows = color_image.rows;
    int cols = color_image.cols;
    cv::Mat gray_image(rows, cols, CV_8U);

    gpu::runThreadSyncronize(); // CUDAは最初の命令に時間がかかる
    //double start, end;
    //start = omp_get_wtime();
    if (ui->radioButton->isChecked()) {
        // CUDA
        gpu::runGrayScale(color_image.data, rows, cols, gray_image.data);
    } else {
        // OpenCV
        cv::cvtColor(color_image, gray_image, CV_BGR2GRAY);
    }
    //end = omp_get_wtime();
    //qDebug("Processing time: %.4lf [ms]\n", (end-start)*1000.0);

    QImage image(gray_image.data,
                 gray_image.cols,
                 gray_image.rows,
                 QImage::Format_Indexed8);
    image = image.convertToFormat(QImage::Format_RGB32);
    pixmap_item_->setPixmap(QPixmap::fromImage(image));
}

肝心のCUDAのプログラム
grayscale.cu

namespace gpu {

__global__ void grayScale(uchar3 *color_image,
                          const int image_size,
                          unsigned char *gray_image)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;

    while (idx < image_size) {
        // BGRの順で格納されている?
        float b = color_image[idx].x * 0.114f;
        float g = color_image[idx].y * 0.587f;
        float r = color_image[idx].z * 0.299f;
        gray_image[idx] = b + g + r;
        idx += gridDim.x * blockDim.x;
    }
}

void runGrayScale(unsigned char *color_image,
                  const int rows,
                  const int cols,
                  unsigned char *gray_image)
{
    const int image_size = rows * cols;
    uchar3 *d_color_image;
    cudaMalloc<uchar3>(&d_color_image, image_size*sizeof(uchar3));
    cudaMemcpy(d_color_image, color_image, image_size*sizeof(uchar3), cudaMemcpyHostToDevice);

    unsigned char *d_gray_image;
    cudaMalloc<unsigned char>(&d_gray_image, image_size*sizeof(unsigned char));

    const int threads_per_block = 2*32;
    const int blocks_per_grid = 1*12;
    grayScale<<<blocks_per_grid, threads_per_block>>>(d_color_image, image_size, d_gray_image);
    cudaThreadSynchronize();

    cudaMemcpy(gray_image, d_gray_image, image_size*sizeof(unsigned char), cudaMemcpyDeviceToHost);
    cudaFree(d_color_image);
    cudaFree(d_gray_image);
}

void runThreadSyncronize()
{
    cudaThreadSynchronize();
}

} // namespace gpu

CPUだと0.1[ms]、GPUだと0.5[ms]だった。。
でも画像が大きいとGPUの方が速くなる。