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
#include
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include

#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

, matrixb

, matrixc

, gpuresult

, matrixd

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