OpenGL de プログラミング

以下の画像を読み込んで、２値化せよ。

pgm画像

その他条件：

[1]
　閾値 127より大きい　→　255
　それ以外　→ 0
とする。

[2]
　テクスチャ座標へのアクセスは、正規化しない。→tex.normalized = false;
　テクスチャフィルタ　→　cudaFilterModePoint


→ CUDA::2値化::複数枚

#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);
}



unsigned int size = imgA->width * imgA->height * sizeof(char); //char型でサイズ確保

/** GPU用にデータ確保 **/
char* d_data = NULL;
CUDA_SAFE_CALL( cudaMalloc( (void**) &d_data, size));


/** 配列のメモリを確保しデータをコピー **/
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(8, 0, 0, 0, cudaChannelFormatKindSigned);//符号付きchar
cudaArray* cu_array;

CUDA_SAFE_CALL( cudaMallocArray( &cu_array, &channelDesc, imgA->width, imgA->height ));//メモリ確保
CUDA_SAFE_CALL( cudaMemcpyToArray( cu_array, 0, 0, imgA->imageData, size, cudaMemcpyHostToDevice)); //コピー


/** テクスチャのパラメータ設定 **/
tex.addressMode[0] = cudaAddressModeWrap;
tex.addressMode[1] = cudaAddressModeWrap;
tex.filterMode = cudaFilterModePoint; //フィルタ
tex.normalized = false; // 正規化された座標でアクセス



/** 配列をテクスチャにバインド **/
CUDA_SAFE_CALL( cudaBindTextureToArray( tex, cu_array, channelDesc));

dim3 dimBlock(8, 8, 1); //Db
dim3 dimGrid(imgA->width / dimBlock.x, imgA->height / dimBlock.y, 1); //Dg

main_kernel<<<dimGrid, dimBlock, 0>>>(d_data,imgA->width,imgA->height);

CUT_CHECK_ERROR("Kernel execution failed");
CUDA_SAFE_CALL( cudaThreadSynchronize() ); //同期をとる




CUDA_SAFE_CALL( cudaMemcpy( imgA->imageData, d_data, size, cudaMemcpyDeviceToHost) );


CUDA_SAFE_CALL(cudaFree(d_data));
CUDA_SAFE_CALL(cudaFreeArray(cu_array));

cout << "OK ! ..."<<endl;


cvNamedWindow("window",CV_WINDOW_AUTOSIZE);
cvShowImage("window",imgA);

cvWaitKey(0); // ０秒待つ => ずっと入力待ち
cvDestroyWindow("window");

cvSaveImage("out.jpg",imgA);

cvReleaseImage( &imgA );imgA=NULL;



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<signed char, 2, cudaReadModeElementType> tex;


__global__ void main_kernel(char *o_data,int width,int height)
{
unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;


if( unsigned(tex2D(tex, x, y)) > unsigned char(127))
{
o_data[y*width + x] = char(255);
}else
{
o_data[y*width + x] = char(0);
}


}

#endif