CUDAの勉強がてら書いてみました。2枚の同サイズの画像を読み込んで、平均画像を作成して保存するだけのプログラムです。環境はUbuntu 16.04です。
#include <chrono>
#include <cstdio>
#include <iostream>
#include <opencv2/core/core.hpp>
#include <opencv2/imgproc/imgproc.hpp>
#include <opencv2/highgui/highgui.hpp>
const int THREADS_PER_BLOCK = 16;
static void HandleError( cudaError_t err, const char *file, int line ) {
if (err != cudaSuccess) {
printf( "%s in %s at line %d\n", cudaGetErrorString( err ), file, line );
exit( EXIT_FAILURE );
}
}
#define HANDLE_ERROR( err ) (HandleError( err, __FILE__, __LINE__ ))
__global__ void use_gpu_sub(unsigned char *a, unsigned char *b, unsigned char *c, int N) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < N) {
c[tid] = (a[tid] + b[tid]) / 2;
}
}
void use_gpu(unsigned char* a, unsigned char* b, unsigned char* c, int height, int width, int channel) {
unsigned char *dev_a, *dev_b, *dev_c;
const int N = height * width * 3;
HANDLE_ERROR( cudaMalloc( (void**)&dev_a, N * sizeof(unsigned char) ) );
HANDLE_ERROR( cudaMalloc( (void**)&dev_b, N * sizeof(unsigned char) ) );
HANDLE_ERROR( cudaMalloc( (void**)&dev_c, N * sizeof(unsigned char) ) );
HANDLE_ERROR( cudaMemcpy( dev_a, a, N * sizeof(unsigned char),
cudaMemcpyHostToDevice ) );
HANDLE_ERROR( cudaMemcpy( dev_b, b, N * sizeof(unsigned char),
cudaMemcpyHostToDevice ) );
use_gpu_sub<<<(N + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>( dev_a, dev_b, dev_c, N );
HANDLE_ERROR( cudaMemcpy( c, dev_c, N * sizeof(unsigned char),
cudaMemcpyDeviceToHost ) );
HANDLE_ERROR( cudaFree( dev_a ) );
HANDLE_ERROR( cudaFree( dev_b ) );
HANDLE_ERROR( cudaFree( dev_c ) );
}
void use_cpu(unsigned char* a, unsigned char* b, unsigned char* c, int height, int width, int channel) {
int N = height * width * channel;
for (int i = 0; i < N; ++i) {
c[i] = (a[i] + b[i]) / 2;
}
}
void average_two_images() {
cv::Mat img1 = cv::imread("adventure-asia-bali-1090551.jpg");
cv::Mat img2 = cv::imread("beautiful-close-up-color-943903.jpg");
assert(img1.size() == img2.size());
assert(img1.type() == img2.type());
std::cout << "image size: " << img1.size() << std::endl;
cv::Mat img3_gpu = cv::Mat(img1.size(), img1.type());
cv::Mat img3_cpu = cv::Mat(img1.size(), img1.type());
auto s1 = std::chrono::system_clock::now();
use_gpu(img1.data, img2.data, img3_gpu.data, img1.rows, img1.cols, img1.channels());
auto e1 = std::chrono::system_clock::now();
std::cout << "gpu time: " << std::chrono::duration_cast<std::chrono::milliseconds>(e1 - s1).count() << " ms" << std::endl;
auto s2 = std::chrono::system_clock::now();
use_cpu(img1.data, img2.data, img3_cpu.data, img1.rows, img1.cols, img1.channels());
auto e2 = std::chrono::system_clock::now();
std::cout << "cpu time: " << std::chrono::duration_cast<std::chrono::milliseconds>(e2 - s2).count() << " ms" << std::endl;
imwrite("averaged.jpg", img3_gpu);
assert(cv::countNonZero(img3_gpu != img3_cpu) == 0);
}
int main()
{
average_two_images();
return 0;
}
これを以下のコマンドでビルドします。
nvcc -O2 -std=c++11 average_two_images.cu $(pkg-config --cflags --libs opencv)
以下は実行結果です。GPUは転送速度がネックになっていてCPUより遅いです。
image size: [6000 x 4000]
gpu time: 180 ms
cpu time: 51 ms
コード中で定数にしている THREADS_PER_BLOCKの値を変えると処理時間が変わりました。どうやって最適な値を決めればよいのでしょう?
176ms: 512スレッド
172ms: 256スレッド
158ms: 128スレッド
158ms: 64スレッド
180ms: 32スレッド
180ms: 16スレッド