CUDA开发矩阵乘法测试你的GPU效率

vs2008上创建cuda项目，新建test.cu文件，将如下代码拷贝进去，编译执行，能很清楚地看到GPU跑矩阵乘法和CPU的效率区别。在我的pc机上执行得到如下结果，可见矩阵乘法的GPU效率大概提高了一个数量级（相对应CPU而言）,开发环境VS2008+cuda5.x开发包+GT520M显卡。

程序代码（参考程序员下载程序，进行修改：http://download2.pudn.com/downloads245/sourcecode/windows/csharp/05872102CUDAMatrixMul.rar）：

///
#include "stdio.h"
#include "cuda_runtime.h"
#include "device_launch_parameters.h"



#if __DEVICE_EMULATION__
bool InitCUDA(void){return true;}
#else
bool InitCUDA(void)
{
	int count = 0;
	int i = 0;
	cudaGetDeviceCount(&count);
	if(count == 0)
	{
		fprintf(stderr, "There is no device.\n");
		return false;
	}
	for(i = 0; i < count; i++)
	{
		cudaDeviceProp prop;
		if(cudaGetDeviceProperties(&prop, i) == cudaSuccess)
		{
			if(prop.major >= 1)
			{
				break;
			}
		}
	}
	if(i == count)
	{
		fprintf(stderr, "There is no device supporting CUDA.\n");
		return false;
	}
	cudaSetDevice(i);
	printf("CUDA initialized.\n");
	cudaDeviceProp prop;
	cudaGetDeviceProperties(&prop,i);
	printf("Device : \" %s \" \n\n", prop.name);
	return true;
}
#endif

#define aW 855
#define aH 511
#define bW 1013
#define blocknum 32//32
#define threadnum 256//256

typedef struct
{
	int width;
	int height;
	int *element;
}Matrix;
Matrix InitMatrix(int w, int h)
{
	Matrix t;
	t.element = (int *)malloc(w * h * sizeof(int));
	for ( int i = 0 ; i < w * h ; i ++)
	{
		t.element[i]= rand() % 10;
	}
	t.width  = w;
	t.height = h;
	return t;
}
Matrix MM( Matrix a, Matrix b)
{
	Matrix t;
	t.element = (int *) malloc (a.height * b.width * sizeof(int));
	t.width  = b.width;
	t.height = a.height;
	int x;
	int y;
	for (int i =0 ; i < t.width * t.height ; i ++)
	{
		x = i / t.width * a.width;
		y = i - i / t.width * t.width;
		t.element[i] = 0;
		for (int k = 0 ; k < a. width ; k ++)
		{
			t.element[i] += a.element[x + k] * b.element [y +b.width * k];
		} 
	}
	return t;
}

__global__ static void MatrixMul(int *ma , int *mb , int *mc , int *mp)
{
	int aw = mp[0];
	int bw = mp[2];
	int cw = mp[4];
	int ch = mp[5];
	const int bid = blockIdx.x;
	const int tid = threadIdx.x;
	int i , x , y  ;

	for (i = bid * threadnum + tid ; i < cw * ch ; i += threadnum * blocknum)
	{
		x = i / cw * aw;
		y = i - i / cw * cw;
		mc[i] = 0;
		for (int k = 0 ; k < aw ; k ++)
		{
			mc[i] += ma[ x + k ] * mb[ y + k * bw ];
		}
	}
}


int main(int argc, char* argv[])
{
	if(!InitCUDA())
	{
		return 0;
	}
	//定义矩阵
	//int matrixa[N][N] , matrixb[N][N] , matrixc[N][N] , gpuresult[N][N] , matrixd[N][N] ;
	Matrix matrixa = InitMatrix(aW,aH);
	Matrix matrixb = InitMatrix(bW,aW);
	Matrix matrixc;
	Matrix gpuresult = InitMatrix(bW,aH);

	int matrixprop[6];

	//为CPU运算计时
	unsigned int timer1 = 0;
	//CPU矩阵相乘
	int start = clock();
	matrixc = MM(matrixa,matrixb);
	int finish = clock();
	printf("cpu time = %d\n", finish - start);
	start = clock();
	matrixprop[0] = matrixa.width;
	matrixprop[1] = matrixa.height;
	matrixprop[2] = matrixb.width;
	matrixprop[3] = matrixb.height;
	matrixprop[4] = matrixc.width;
	matrixprop[5] = matrixc.height;

	//申请显存
	int *ma, *mb, *mc, *mp;
	cudaMalloc( (void**)&ma , sizeof(int) * matrixa.width * matrixa.height );
	cudaMalloc( (void**)&mb , sizeof(int) * matrixb.width * matrixb.height );
	cudaMalloc( (void**)&mc , sizeof(int) * matrixc.width * matrixc.height );
	cudaMalloc( (void**)&mp , sizeof(int) * 6 );
	//将数据复制到显存内
	cudaMemcpy( ma ,  matrixa.element , sizeof(int) * matrixa.width * matrixa.height ,cudaMemcpyHostToDevice );
	cudaMemcpy( mb ,  matrixb.element , sizeof(int) * matrixb.width * matrixb.height ,cudaMemcpyHostToDevice );
	cudaMemcpy( mp , matrixprop , sizeof(int) * 6 , cudaMemcpyHostToDevice );
	unsigned int timer2 = 0;
	//调用CUDA函数
	MatrixMul <<< blocknum , threadnum , 0 >>>( ma , mb , mc , mp);
	cudaThreadSynchronize();
	//cutilCheckError( cutStopTimer( timer2));
	//将数据从显存中复制出来
	cudaMemcpy( gpuresult.element , mc , sizeof(int) * gpuresult.width * gpuresult.height , cudaMemcpyDeviceToHost );
	finish = clock();
	printf("gpu time = %d\n", finish-start);

	for (int i =0 ; i < gpuresult.width * gpuresult.height ; i ++)
	{
		//printf("%d -- %d\n",matrixc.element[ i ],gpuresult.element[ i ]);
		if (matrixc.element[i] != gpuresult.element[i])
		{
			printf("ERROR");
		}
	}

	cudaFree(ma);
	cudaFree(mb);
	cudaFree(mc);
	cudaFree(mp);

	return 0;
}

http://vivianlaogong.blog.163.com/blog/static/213726121201210611590366/

http://www.cnblogs.com/viviman/archive/2012/11/06/2775098.html