最終更新: mikk_ni3_92 2008年04月20日(日) 19:10:31履歴
現在地 >> メニュー >> CUDA >> CUDA::画像処理::2値化 >> CUDA::画像処理::CUDA+CV
#include <stdlib.h>
#include <stdio.h>
#include <iostream>
#include <cstdio>
#include <cutil.h>
#include <cv.h>
#include <highgui.h>
#include "main_kernel.h"
using namespace std;
char *filename="test.pgm";
int main( int argc, char **argv)
{
CUT_DEVICE_INIT();
/** 画像読み込み **/
IplImage *imgA = cvLoadImage( filename, CV_LOAD_IMAGE_GRAYSCALE);//
if(imgA ==NULL)
{
printf("File not found\n");
CUT_EXIT(argc, argv);
exit(0);
}
IplImage *tmp = cvCreateImage (cvGetSize (imgA), IPL_DEPTH_32F, 1);
cvScale (imgA, tmp, 1.0/255.0, 0.0); //32ビット(float)にスケール変換
cvReleaseImage( &imgA );imgA=NULL;
unsigned int size = tmp->width * tmp->height * sizeof(float);
printf("filename [%s]: %d x %d (%d)\n", filename, tmp->width, tmp->height,size);
/** GPU用にデータ確保 **/
float* d_data = NULL;
CUDA_SAFE_CALL( cudaMalloc( (void**) &d_data, size));
/** 配列のメモリを確保しデータをコピー **/
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
cudaArray* cu_array;
CUDA_SAFE_CALL( cudaMallocArray( &cu_array, &channelDesc, tmp->width, tmp->height ));
CUDA_SAFE_CALL( cudaMemcpyToArray( cu_array, 0, 0, tmp->imageData, size, cudaMemcpyHostToDevice));
/** テクスチャのパラメータ設定 **/
tex.addressMode[0] = cudaAddressModeWrap;
tex.addressMode[1] = cudaAddressModeWrap;
tex.filterMode = cudaFilterModeLinear; //フィルタ
tex.normalized = true; // 正規化された座標でアクセス
/** 配列をテクスチャにバインド **/
CUDA_SAFE_CALL( cudaBindTextureToArray( tex, cu_array, channelDesc));
dim3 dimBlock(8, 8, 1); //Db
dim3 dimGrid(tmp->width / dimBlock.x, tmp->height / dimBlock.y, 1); //Dg
main_kernel<<<dimGrid, dimBlock, 0>>>(d_data,tmp->width,tmp->height);
CUT_CHECK_ERROR("Kernel execution failed");
CUDA_SAFE_CALL( cudaThreadSynchronize() ); //同期をとる
CUDA_SAFE_CALL( cudaMemcpy( tmp->imageData, d_data, size, cudaMemcpyDeviceToHost) );
CUDA_SAFE_CALL(cudaFree(d_data));
CUDA_SAFE_CALL(cudaFreeArray(cu_array));
cout << "OK ! ..."<<endl;
IplImage *output = cvCreateImage (cvGetSize (tmp), IPL_DEPTH_8U, 1);
cvConvertScale(tmp, output,255);
cvSaveImage("out.jpg",output);
cvReleaseImage( &tmp );tmp=NULL;
cvReleaseImage( &output );output=NULL;
cvDestroyWindow("window");
CUT_EXIT(argc, argv);
return 0;
}
#ifndef _MAIN_KERNEL_H_
#define _MAIN_KERNEL_H_
#include <cstdio>
#include <iostream>
#include <cutil.h>
#include <ctime>
#include <cstdlib>
using namespace std;
texture<float, 2, cudaReadModeElementType> tex;
__global__ void main_kernel(float *o_data,int width,int height)
{
unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
float u = x / (float) width;
float v = y / (float) height;
if( tex2D(tex, u, v) > 0.5)
{
o_data[y*width + x] = 1;
}else
{
o_data[y*width + x] = 0;
}
}
#endif
#include <stdlib.h>
#include <stdio.h>
#include <iostream>
#include <cstdio>
#include <cutil.h>
#include <cv.h>
#include <highgui.h>
#include "main_kernel.h"
using namespace std;
char *filename="test.pgm";
int main( int argc, char **argv)
{
CUT_DEVICE_INIT();
/** 画像読み込み **/
IplImage *imgA = cvLoadImage( filename, CV_LOAD_IMAGE_GRAYSCALE);//
if(imgA ==NULL)
{
printf("File not found\n");
CUT_EXIT(argc, argv);
exit(0);
}
IplImage *tmp = cvCreateImage (cvGetSize (imgA), IPL_DEPTH_32F, 1);
cvScale (imgA, tmp, 1.0/255.0, 0.0); //32ビット(float)にスケール変換
cvReleaseImage( &imgA );imgA=NULL;
unsigned int size = tmp->width * tmp->height * sizeof(float);
printf("filename [%s]: %d x %d (%d)\n", filename, tmp->width, tmp->height,size);
/** GPU用にデータ確保 **/
float* d_data = NULL;
CUDA_SAFE_CALL( cudaMalloc( (void**) &d_data, size));
/** 配列のメモリを確保しデータをコピー **/
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
cudaArray* cu_array;
CUDA_SAFE_CALL( cudaMallocArray( &cu_array, &channelDesc, tmp->width, tmp->height ));
CUDA_SAFE_CALL( cudaMemcpyToArray( cu_array, 0, 0, tmp->imageData, size, cudaMemcpyHostToDevice));
/** テクスチャのパラメータ設定 **/
tex.addressMode[0] = cudaAddressModeWrap;
tex.addressMode[1] = cudaAddressModeWrap;
tex.filterMode = cudaFilterModeLinear; //フィルタ
tex.normalized = true; // 正規化された座標でアクセス
/** 配列をテクスチャにバインド **/
CUDA_SAFE_CALL( cudaBindTextureToArray( tex, cu_array, channelDesc));
dim3 dimBlock(8, 8, 1); //Db
dim3 dimGrid(tmp->width / dimBlock.x, tmp->height / dimBlock.y, 1); //Dg
main_kernel<<<dimGrid, dimBlock, 0>>>(d_data,tmp->width,tmp->height);
CUT_CHECK_ERROR("Kernel execution failed");
CUDA_SAFE_CALL( cudaThreadSynchronize() ); //同期をとる
CUDA_SAFE_CALL( cudaMemcpy( tmp->imageData, d_data, size, cudaMemcpyDeviceToHost) );
CUDA_SAFE_CALL(cudaFree(d_data));
CUDA_SAFE_CALL(cudaFreeArray(cu_array));
cout << "OK ! ..."<<endl;
IplImage *output = cvCreateImage (cvGetSize (tmp), IPL_DEPTH_8U, 1);
cvConvertScale(tmp, output,255);
cvSaveImage("out.jpg",output);
cvReleaseImage( &tmp );tmp=NULL;
cvReleaseImage( &output );output=NULL;
cvDestroyWindow("window");
CUT_EXIT(argc, argv);
return 0;
}
#ifndef _MAIN_KERNEL_H_
#define _MAIN_KERNEL_H_
#include <cstdio>
#include <iostream>
#include <cutil.h>
#include <ctime>
#include <cstdlib>
using namespace std;
texture<float, 2, cudaReadModeElementType> tex;
__global__ void main_kernel(float *o_data,int width,int height)
{
unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
float u = x / (float) width;
float v = y / (float) height;
if( tex2D(tex, u, v) > 0.5)
{
o_data[y*width + x] = 1;
}else
{
o_data[y*width + x] = 0;
}
}
#endif