OpenGL de プログラミング

以下の画像をよみこんで、2値化せよ。
pgm画像


→CUDA::画像処理::CUDA+CV

#include <stdlib.h>
#include <stdio.h>
#include <iostream>
#include <cstdio>
#include <cutil.h>

#include "main_kernel.h"

using namespace std;


char *filename="test.pgm";


int main( int argc, char **argv)
{

CUT_DEVICE_INIT();

/** 画像読み込み **/
float* h_data = NULL;
unsigned int width, height;
if( cutLoadPGMf(filename, &h_data, &width, &height) == NULL )
{
CUT_EXIT(argc, argv);
}

unsigned int size = width * height * sizeof(float);
printf("filename [%s]: %d x %d (%d)\n", filename, width, 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, width, height ));
CUDA_SAFE_CALL( cudaMemcpyToArray( cu_array, 0, 0, h_data, 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(width / dimBlock.x, height / dimBlock.y, 1); //Dg

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

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


float* h_odata = (float*) malloc( size);
CUDA_SAFE_CALL( cudaMemcpy( h_odata, d_data, size, cudaMemcpyDeviceToHost) );

CUT_SAFE_CALL( cutSavePGMf( "output.pgm", h_odata, width, height));

CUDA_SAFE_CALL(cudaFree(d_data));
CUDA_SAFE_CALL(cudaFreeArray(cu_array));
free(h_data);
free(h_odata);
cout << "OK ! ..."<<endl;
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

floatだと0〜1の範囲で画素値をあつかうっぽい。