OpenGL de プログラミング

512×512の画像を読み込んで、線形メモリでテクスチャを作成、反転処理を行え。

【main.cu】

#include <iostream>
#include <cutil_inline.h>

#pragma comment (lib,"cudart.lib")

#ifdef _DEBUG
#pragma comment (lib,"cutil32D.lib")
#else
#pragma comment (lib,"cutil32.lib")
#endif

texture<float,1,cudaReadModeElementType> texRef;

//------------- CUDAでの処理 ---------------//
__global__ void cuInverse(float *gpuAry,int widthStep)
{
	//現在地の取得
	unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
	unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
	unsigned int Idx = y*widthStep+x;

	gpuAry[Idx] = 1.f - tex1Dfetch(texRef,Idx);//floatの場合1〜0の間で扱う
}



void printDeviceInfomation(const cudaDeviceProp &devInfo)
{
	std::cout<<"GPUの総メモリ数 : "<<devInfo.totalGlobalMem/1024/1024<<" MB\n";
	std::cout<<"各ブロックに割り当てられる最大 shared memory : "<<devInfo.sharedMemPerBlock/1024<<" KB\n";
	std::cout<<"1ブロックあたりの最大スレッド数:"<<devInfo.maxThreadsPerBlock<<"\n";
	std::cout<<"X方向：スレッドサイズのとれる値:"<<devInfo.maxThreadsDim[0]<<"\n";
	std::cout<<"Y方向：スレッドサイズのとれる値:"<<devInfo.maxThreadsDim[1]<<"\n";
	std::cout<<"X方向：グリッドサイズのとれる値:"<<devInfo.maxGridSize[0]<<"\n";
	std::cout<<"Y方向：グリッドサイズのとれる値:"<<devInfo.maxGridSize[1]<<"\n";
	std::cout<<"Z方向：グリッドサイズのとれる値:"<<devInfo.maxGridSize[2]<<"\n";
}

//------------- メイン関数 --------------------//
int main(int argc, char **argv)
{

	CUT_DEVICE_INIT(argc,argv);

	float *imgData = NULL;//NULLポインタで必ず初期化。
	unsigned int w,h;

	if( CUTTrue != cutLoadPGMf("lena.pgm",&imgData,&w,&h) )
	{
		std::cerr << "Can't Load Image\n";
		CUT_EXIT(argc, argv);//終了
	}



	//GPUの情報を取得
	cudaDeviceProp devInfo;
	CUDA_SAFE_CALL(cudaGetDeviceProperties(&devInfo, 0)); //0番目のGPUの情報を取得
	printDeviceInfomation(devInfo);



	//GPUのメモリ確保
	int dataSize = sizeof(float)* w * h;
	float *deviceAry;
	CUDA_SAFE_CALL(cudaMalloc( reinterpret_cast<void**>(&deviceAry),dataSize ));//GPUのメモリ確保


	//GPUへデータをコピー(Host -> GPU)
	CUDA_SAFE_CALL( cudaMemcpy( deviceAry, imgData, dataSize , cudaMemcpyHostToDevice));

	//バインド
	cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
	CUDA_SAFE_CALL(	cudaBindTexture(NULL,texRef,deviceAry,channelDesc)) ;


	//512×512画像 = 32*32 × 16*16
	//ブロックとスレッドを用意
	dim3 gridSize(32,32,1);//maxGridSize[0〜1]の上限に注意
	dim3 blockSize(16,16,1);//maxThreadsPerBlock,maxThreadsDim[0〜2]の上限に注意

	cuInverse<<<gridSize,blockSize>>>(deviceAry,w);

	CUDA_SAFE_CALL( cudaThreadSynchronize() ); //同期

	//GPUからデータを取り出す(GPU -> HOST)
	CUDA_SAFE_CALL( cudaMemcpy( imgData,deviceAry, dataSize , cudaMemcpyDeviceToHost));
	CUT_SAFE_CALL( cutSavePGMf( "InveseImage.pgm", imgData, w, h));


	//各種メモリ解放
	cutFree(imgData);
	cudaFree(deviceAry);
	CUT_EXIT(argc, argv);//終了

	return 0;
}