cpu gpu做矩阵乘法效率比对,虽然如此,但是对需要自己做的算法是否能如此高效的提高还是未知
2014-09-25 12:16
477 查看
vs2010上创建cuda项目,kernel.cu,将如下代码拷贝进去,编译执行,能很清楚地看到GPU跑矩阵乘法和CPU的效率区别。在我的pc机上执行得到如下结果,开发环境VS2010+cuda4.1开发包+Tesla c2075显卡
//代码借用别人的,自己修改后测试通过
///////////////////////
#include <stdio.h>
#include <cstdlib>
#include <iostream>
#include <cutil.h>
//#include "cuda_class.h"
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
using namespace std;
#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);
system("pause");
return 0;
}
测试的结果如下:仅供参考
//代码借用别人的,自己修改后测试通过
///////////////////////
#include <stdio.h>
#include <cstdlib>
#include <iostream>
#include <cutil.h>
//#include "cuda_class.h"
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
using namespace std;
#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);
system("pause");
return 0;
}
测试的结果如下:仅供参考
相关文章推荐
- 取模、乘法和除法运算在CPU和GPU上的效率
- GPU上大规模稀疏矩阵特征值计算高效算法之二——稀疏矩阵
- 算法提高 矩阵乘法 区间DP
- GPU与CPU版本的矩阵乘法对比
- 蓝桥 ADV-232 算法提高 矩阵乘法 【区间DP】
- CUDA开发矩阵乘法测试你的GPU效率
- 算法提高 矩阵乘法
- 蓝桥杯 算法提高 矩阵乘法 (区间dp)
- 算法提高 矩阵乘法 (区间dp)
- 蓝桥杯 算法提高 矩阵乘法 区间dp
- GPU上大规模稀疏矩阵特征值计算高效算法之一——GPU介绍
- 【蓝桥】算法提高 矩阵乘法
- 本来从动态壁纸预览页面设置一个动态壁纸回到桌面便可以看到桌面动态壁纸,可以观察得到自己的动态壁纸是否设置成功了(必须知道设置是否成功的结构,因为还有一些操作需要完成)! 但是现在是要在自己的应用中进入
- 算法提高 矩阵乘法
- cuda开发矩阵乘法测试你的GPU效率
- 折腾二维数组(哎,真的把自己折腾死了,虽然在做的时候错了好多地方,需要认真改错,但是看到最后的结果,真的很开心)
- 对象数组的使用及this指针(自己改的,虽然只是一小点,但是还是有成就感)
- 华为U8800 虽然系统升级到2.3以上但是app还是无法移动到SD卡上,需要刷ROM到U8800+才能解决
- 算法提高 矩阵乘法
- 算法提高 矩阵乘法