CUDA入門 > OpenCVとCUDA > 拡大縮小:Bilinear

拡大縮小:Bilinear


拡大縮小:Bilinear
拡大縮小の手法の一つ、Bilinear(バイリニア補間)のCUDAでの実装例です。

サンプルプログラム
#include <stdio.h> #include "cv.h" #include "cxcore.h" #include "highgui.h" #pragma comment(lib,"cv.lib") #pragma comment(lib,"cxcore.lib") #pragma comment(lib,"highgui.lib") #include <cutil.h> // 倍率を設定 #define RATE 2 __global__ void Bilinear8Kernel(unsigned char *in, unsigned char *out, int width_in, int height_in, int channels, int width_out, int height_out){ int x = blockDim.x * blockIdx.x + threadIdx.x; int y = blockDim.y * blockIdx.y + threadIdx.y; // 拡大・縮小率に合わせて入力画像の座標を調整 int x_in = (int)(x * ((float)width_in / width_out)); int y_in = (int)(y * ((float)height_in / height_out)); // バッファオーバーフローしないように調整 if(x < width_out && y < height_out){ if((int)x_in >= width_in){ x_in = width_in - 1; } if((int)y_in >= height_in){ y_in >= height_in - 1; } out[y * width_out + x] = (unsigned char)(((int)x_in + 1 - x_in) * ((int)y_in + 1 - y_in) * (double)in[(int)y_in * width_in + (int)x_in] + (x_in - (int)x_in) * ((int)y_in + 1 - y_in) * (double)in[(int)y_in * width_in + (int)x_in + 1] + ((int)x_in + 1 - x_in) * (y_in - (int)y_in) * (double)in[((int)y_in + 1) * width_in + (int)x_in] + (x_in - (int)x_in) * (y_in - (int)y_in) * (double)in[((int)y_in + 1) * width_in + (int)x_in + 1]); } } __global__ void Bilinear24Kernel(unsigned char *in, unsigned char *out, int width_in, int height_in, int channels, int width_out, int height_out){ int x = blockDim.x * blockIdx.x + threadIdx.x; int y = blockDim.y * blockIdx.y + threadIdx.y; // 拡大・縮小率に合わせて入力画像の座標を調整 int x_in = (int)(x * ((float)width_in / width_out)); int y_in = (int)(y * ((float)height_in / height_out)); // バッファオーバーフローしないように調整 if(x < width_out && y < height_out){ if((int)x_in >= width_in){ x_in = width_in - 1; } if((int)y_in >= height_in){ y_in >= height_in - 1; } out[y * width_out * 3 + x * 3] = (unsigned char)(((int)x_in + 1 - x_in) * ((int)y_in + 1 - y_in) * (double)in[(int)y_in * width_in * 3 + (int)x_in * 3] + (x_in - (int)x_in) * ((int)y_in + 1 - y_in) * (double)in[(int)y_in * width_in * 3 + (int)(x_in + 1) * 3] + ((int)x_in + 1 - x_in) * (y_in - (int)y_in) * (double)in[((int)y_in + 1) * width_in * 3 + (int)x_in * 3] + (x_in - (int)x_in) * (y_in - (int)y_in) * (double)in[((int)y_in + 1) * width_in * 3 + (int)(x_in + 1) * 3]); out[y * width_out * 3 + x * 3 + 1] = (unsigned char)(((int)x_in + 1 - x_in) * ((int)y_in + 1 - y_in) * (double)in[(int)y_in * width_in * 3 + (int)x_in * 3 + 1] + (x_in - (int)x_in) * ((int)y_in + 1 - y_in) * (double)in[(int)y_in * width_in * 3 + (int)(x_in + 1) * 3 + 1] + ((int)x_in + 1 - x_in) * (y_in - (int)y_in) * (double)in[((int)y_in + 1) * width_in * 3 + (int)x_in * 3 + 1] + (x_in - (int)x_in) * (y_in - (int)y_in) * (double)in[((int)y_in + 1) * width_in * 3 + (int)(x_in + 1) * 3 + 1]); out[y * width_out * 3 + x * 3 + 2] = (unsigned char)(((int)x_in + 1 - x_in) * ((int)y_in + 1 - y_in) * (double)in[(int)y_in * width_in * 3 + (int)x_in * 3 + 2] + (x_in - (int)x_in) * ((int)y_in + 1 - y_in) * (double)in[(int)y_in * width_in * 3 + (int)(x_in + 1) * 3 + 2] + ((int)x_in + 1 - x_in) * (y_in - (int)y_in) * (double)in[((int)y_in + 1) * width_in * 3 + (int)x_in * 3 + 2] + (x_in - (int)x_in) * (y_in - (int)y_in) * (double)in[((int)y_in + 1) * width_in * 3 + (int)(x_in + 1) * 3 + 2]); } } __global__ void Bilinear32Kernel(unsigned char *in, unsigned char *out, int width_in, int height_in, int channels, int width_out, int height_out){ int x = blockDim.x * blockIdx.x + threadIdx.x; int y = blockDim.y * blockIdx.y + threadIdx.y; // 拡大・縮小率に合わせて入力画像の座標を調整 int x_in = (int)(x * ((float)width_in / width_out)); int y_in = (int)(y * ((float)height_in / height_out)); // バッファオーバーフローしないように調整 if(x < width_out && y < height_out){ if((int)x_in >= width_in){ x_in = width_in - 1; } if((int)y_in >= height_in){ y_in >= height_in - 1; } out[y * width_out * 4 + x * 4] = (unsigned char)(((int)x_in + 1 - x_in) * ((int)y_in + 1 - y_in) * (double)in[(int)y_in * width_in * 4 + (int)x_in * 4] + (x_in - (int)x_in) * ((int)y_in + 1 - y_in) * (double)in[(int)y_in * width_in * 4 + (int)(x_in + 1) * 4] + ((int)x_in + 1 - x_in) * (y_in - (int)y_in) * (double)in[((int)y_in + 1) * width_in * 4 + (int)x_in * 4] + (x_in - (int)x_in) * (y_in - (int)y_in) * (double)in[((int)y_in + 1) * width_in * 4 + (int)(x_in + 1) * 4]); out[y * width_out * 4 + x * 4 + 1] = (unsigned char)(((int)x_in + 1 - x_in) * ((int)y_in + 1 - y_in) * (double)in[(int)y_in * width_in * 4 + (int)x_in * 4 + 1] + (x_in - (int)x_in) * ((int)y_in + 1 - y_in) * (double)in[(int)y_in * width_in * 4 + (int)(x_in + 1) * 4 + 1] + ((int)x_in + 1 - x_in) * (y_in - (int)y_in) * (double)in[((int)y_in + 1) * width_in * 4 + (int)x_in * 4 + 1] + (x_in - (int)x_in) * (y_in - (int)y_in) * (double)in[((int)y_in + 1) * width_in * 4 + (int)(x_in + 1) * 4 + 1]); out[y * width_out * 4 + x * 4 + 2] = (unsigned char)(((int)x_in + 1 - x_in) * ((int)y_in + 1 - y_in) * (double)in[(int)y_in * width_in * 4 + (int)x_in * 4 + 2] + (x_in - (int)x_in) * ((int)y_in + 1 - y_in) * (double)in[(int)y_in * width_in * 4 + (int)(x_in + 1) * 4 + 2] + ((int)x_in + 1 - x_in) * (y_in - (int)y_in) * (double)in[((int)y_in + 1) * width_in * 4 + (int)x_in * 4 + 2] + (x_in - (int)x_in) * (y_in - (int)y_in) * (double)in[((int)y_in + 1) * width_in * 4 + (int)(x_in + 1) * 4 + 2]); out[y * width_out * 4 + x * 4 + 3] = (unsigned char)(((int)x_in + 1 - x_in) * ((int)y_in + 1 - y_in) * (double)in[(int)y_in * width_in * 4 + (int)x_in * 4 + 3] + (x_in - (int)x_in) * ((int)y_in + 1 - y_in) * (double)in[(int)y_in * width_in * 4 + (int)(x_in + 1) * 4 + 3] + ((int)x_in + 1 - x_in) * (y_in - (int)y_in) * (double)in[((int)y_in + 1) * width_in * 4 + (int)x_in * 4 + 3] + (x_in - (int)x_in) * (y_in - (int)y_in) * (double)in[((int)y_in + 1) * width_in * 4 + (int)(x_in + 1) * 4 + 3]); } } int main(int argc, char **argv){ CUT_DEVICE_INIT( argc, argv ); IplImage *imgin, *imgout; unsigned char *d_in, *d_out; int memsize_in, memsize_out; float width_out, height_out; imgin = cvLoadImage("ファイル名", 1); // 出力画像の幅を高さを求める width_out = (int)(imgin->width * RATE); height_out = (int)(imgin->height * RATE); imgout = cvCreateImage(cvSize(width_out, height_out), imgin->depth, imgin->nChannels); printf("width = %d, height = %d, nChannels = %d\n", imgout->width, imgout->height, imgin->nChannels); memsize_in = imgin->width * imgin->height * imgin->nChannels; memsize_out = imgout->width * imgout->height * imgout->nChannels; CUDA_SAFE_CALL(cudaMalloc((void**)&d_in, memsize_in)); CUDA_SAFE_CALL(cudaMalloc((void**)&d_out, memsize_out)); CUDA_SAFE_CALL(cudaMemcpy(d_in, imgin->imageData, memsize_in, cudaMemcpyHostToDevice)); dim3 threads(32,8); // 全画素処理するため、小数点は切り上げる dim3 blocks((int)((float)imgout->width / threads.x + 0.5), (int)((float)imgout->height / threads.y + 0.5)); // 前の情報を一旦初期化(処理されなかったメモリは、前のデータが表示されるため) // デバッグ用。通常は必要なし cudaMemset(d_out, 0, memsize_out); CUDA_SAFE_CALL( cudaThreadSynchronize() ); switch(imgin->nChannels){ case 1: Bilinear8Kernel<<< blocks, threads >>> (d_in, d_out, imgin->width, imgin->height, imgin->nChannels, imgout->width, imgout->height); break; case 3: Bilinear24Kernel<<< blocks, threads >>> (d_in, d_out, imgin->width, imgin->height, imgin->nChannels, imgout->width, imgout->height); break; case 4: Bilinear32Kernel<<< blocks, threads >>> (d_in, d_out, imgin->width, imgin->height, imgin->nChannels, imgout->width, imgout->height); break; } CUDA_SAFE_CALL(cudaMemcpy(imgout->imageData, d_out, memsize_out, cudaMemcpyDeviceToHost)); cvNamedWindow("in", 1); cvNamedWindow("out", 1); cvShowImage("in", imgin); cvShowImage("out", imgout); // キーを何か押したら終了 cvWaitKey(-1); cvDestroyWindow("in"); cvDestroyWindow("out"); cvReleaseImage(&imgin); cvReleaseImage(&imgout); cudaFree(d_in); cudaFree(d_out); return 0; }