CUDA矩阵乘法——VS2010中使用CUDA示例

来源：互联网发布：染色体分析软件编辑：程序博客网时间：2024/05/18 12:01

文章转自：http://blog.csdn.net/augusdi/article/details/12617177

各工具或库的版本：

IDE：VS2008

VS2010 (使用MSVC编译器)

CUDA：5.5

下面以宽为1024的举证乘法为例，VS2010工程。

1.新建VS2010下VC++ Console工程

2.设置项目属性

设置Project属性自定义配置，支持CUDA。

然后添加链接库：项目——>属性——>链接器——>常规，在附加库目录中添加ToolKit和SDK目录里的lib，在输入的附加库目录下添加需要用到的lib文件。这一步和在单独使用CUDA时的做法是一样的，详见http://www.cnblogs.com/Romi/archive/2012/04/20/2459669.html

3.编写CUDA文件(.cu)

在项目中新建一个.cu的文件，加上如下代码，完成在GPU设备上进行矩阵乘法：

[cpp] view plaincopyprint?

//CUDAtest.cu
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#define TILE_WIDTH 64
// 核函数
// __global__ static void MatrixMulKernel(const float* Md,const float* Nd,float* Pd,int Width)
__global__ void MatrixMulKernel(constfloat* Md,const float* Nd,float* Pd,int Width)
{
//计算Pd和Md中元素的行索引
int Row = blockIdx.y * TILE_WIDTH + threadIdx.y;//行
int Col = blockIdx.x * TILE_WIDTH + threadIdx.x;//列
float Pvalue = 0.0;
for (int k = 0; k < Width; k++)
{
Pvalue += Md[Row * Width + k] * Nd[k * Width + Col];
}
//每个线程负责计算P中的一个元素
Pd[Row * Width + Col] = Pvalue;
}
// 矩阵乘法(CUDA中)
// 在外部调用，使用extern
extern "C"void MatrixMultiplication_CUDA(constfloat* M,const float* N,float* P,int Width)
{
cudaSetDevice(0); //设置目标GPU
float *Md, *Nd, *Pd;
int size = Width * Width * sizeof(float);//字节长度
cudaMalloc((void**)&Md, size);
cudaMalloc((void**)&Nd, size);
cudaMalloc((void**)&Pd, size);
//Copies a matrix from the memory* area pointed to by src to the memory area pointed to by dst
cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice);
cudaMemcpy(Nd, N, size, cudaMemcpyHostToDevice);
//
dim3 dimGrid(Width / TILE_WIDTH, Width / TILE_WIDTH); //网格的维度
dim3 dimBlock(TILE_WIDTH, TILE_WIDTH); //块的维度
MatrixMulKernel<<< dimGrid, dimBlock >>>(Md, Nd, Pd, Width);
cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost);
//释放设备上的矩阵
cudaFree(Md);
cudaFree(Nd);
cudaFree(Pd);
}

//CUDAtest.cu#include "cuda_runtime.h"  #include "device_launch_parameters.h"#define TILE_WIDTH 64// 核函数// __global__ static void MatrixMulKernel(const float* Md,const float* Nd,float* Pd,int Width)__global__ void MatrixMulKernel(const float* Md,const float* Nd,float* Pd,int Width){//计算Pd和Md中元素的行索引int Row = blockIdx.y * TILE_WIDTH + threadIdx.y; //行int Col = blockIdx.x * TILE_WIDTH + threadIdx.x; //列float Pvalue = 0.0;for (int k = 0; k < Width; k++){Pvalue += Md[Row * Width + k] * Nd[k * Width + Col];}//每个线程负责计算P中的一个元素Pd[Row * Width + Col] = Pvalue;}// 矩阵乘法(CUDA中)// 在外部调用，使用externextern "C" void MatrixMultiplication_CUDA(const float* M,const float* N,float* P,int Width){cudaSetDevice(0);  //设置目标GPU float *Md, *Nd, *Pd;int size = Width * Width * sizeof(float);//字节长度cudaMalloc((void**)&Md, size);cudaMalloc((void**)&Nd, size);cudaMalloc((void**)&Pd, size);//Copies a matrix from the memory* area pointed to by src to the memory area pointed to by dstcudaMemcpy(Md, M, size, cudaMemcpyHostToDevice);cudaMemcpy(Nd, N, size, cudaMemcpyHostToDevice);//dim3 dimGrid(Width / TILE_WIDTH, Width / TILE_WIDTH);//网格的维度dim3 dimBlock(TILE_WIDTH, TILE_WIDTH);//块的维度MatrixMulKernel<<< dimGrid, dimBlock >>>(Md, Nd, Pd, Width);cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost);//释放设备上的矩阵cudaFree(Md);cudaFree(Nd);cudaFree(Pd);}

这里使用extern以声明函数可以在外部被调用。如果是在调用该函数的原文件中使用include “XXX.cu”，我这会出现编译错误，暂没有解决，所以使用extern

设置.cu文件属性，支持CUDA编译。

4.在工程源文件中添加CUDA的引用

添加CPP文件，调用CUDA文件内容，调试CPU与GPU运行效率。

[cpp] view plaincopyprint?

#include <time.h>
#include <stdlib.h>
#include <stdio.h>
//这里不要忘了加引用声明
extern "C" void MatrixMultiplication_CUDA(const float* M, const float* N,float* P, int Width);
//构造函数...
//析构函数...
// 产生矩阵,矩阵中元素0~1
void matgen(float* a,int Width)
{
int i, j;
for (i = 0; i < Width; i++)
{
for (j = 0; j < Width; j++)
{
a[i * Width + j] = (float)rand() / RAND_MAX + (float)rand() / (RAND_MAX*RAND_MAX);
}
}
}
//矩阵乘法(CPU验证)
void MatrixMultiplication(constfloat* M,const float* N,float* P,int Width)
{
int i,j,k;
for (i = 0; i < Width; i++)
{
for (j = 0; j < Width; j++)
{
float sum = 0;
for (k = 0; k < Width; k++)
{
sum += M[i * Width + k] * N[k * Width + j];
}
P[i * Width + j] = sum;
}
}
}
double MatrixMul_GPU()
{
float *M, *N, *Pg;
int Width = 1024; //1024×1024矩阵乘法
M = (float*)malloc(sizeof(float) * Width * Width);
N = (float*)malloc(sizeof(float) * Width * Width);
Pg= (float*)malloc(sizeof(float) * Width * Width);//保存GPU计算结果
srand(0);
matgen(M, Width); //产生矩阵M
matgen(N, Width); //产生矩阵N
double timeStart, timeEnd; //定义时间，求时间差用
timeStart = clock();
MatrixMultiplication_CUDA(M, N, Pg, Width); //GPU上计算
timeEnd = clock();
free(M);
free(N);
free(Pg);
return timeEnd - timeStart;
}
double MatrixMul_CPU()
{
float *M, *N, *Pc;
int Width = 1024; //1024×1024矩阵乘法
M = (float*)malloc(sizeof(float) * Width * Width);
N = (float*)malloc(sizeof(float) * Width * Width);
Pc= (float*)malloc(sizeof(float) * Width * Width);//保存CPU计算结果
srand(0);
matgen(M, Width); //产生矩阵M
matgen(N, Width); //产生矩阵N
double timeStart,timeEnd; //定义时间，求时间差用
timeStart = clock();
MatrixMultiplication(M, N, Pc, Width); //CPU上计算
timeEnd = clock();
free(M);
free(N);
free(Pc);
return timeEnd - timeStart;
}
//////////////////////////////////////////////////////////////////////////
void main()
{
printf("CPU use time %g\n", MatrixMul_CPU());
printf("GPU use time %g\n", MatrixMul_GPU());
}

#include <time.h>#include <stdlib.h>#include <stdio.h>//这里不要忘了加引用声明extern "C" void MatrixMultiplication_CUDA(const float* M, const float* N, float* P, int Width);//构造函数...//析构函数...// 产生矩阵,矩阵中元素0~1void matgen(float* a, int Width){    int i, j;    for (i = 0; i < Width; i++)    {        for (j = 0; j < Width; j++)        {            a[i * Width + j] = (float)rand() / RAND_MAX + (float)rand() / (RAND_MAX*RAND_MAX);        }    }}//矩阵乘法(CPU验证)void MatrixMultiplication(const float* M,const float* N,float* P,int Width){    int i,j,k;    for (i = 0; i < Width; i++)    {        for (j = 0; j < Width; j++)        {            float sum = 0;            for (k = 0; k < Width; k++)            {                sum += M[i * Width + k] * N[k * Width + j];            }            P[i * Width + j] = sum;        }    }}double MatrixMul_GPU(){     float *M, *N, *Pg;    int Width = 1024;//1024×1024矩阵乘法    M = (float*)malloc(sizeof(float) * Width * Width);    N = (float*)malloc(sizeof(float) * Width * Width);    Pg= (float*)malloc(sizeof(float) * Width * Width); //保存GPU计算结果    srand(0);    matgen(M, Width);//产生矩阵M    matgen(N, Width);//产生矩阵N    double timeStart, timeEnd;//定义时间，求时间差用    timeStart = clock();    MatrixMultiplication_CUDA(M, N, Pg, Width);//GPU上计算    timeEnd = clock();    free(M);    free(N);    free(Pg);return timeEnd - timeStart;}double MatrixMul_CPU(){    float *M, *N, *Pc;    int Width = 1024;//1024×1024矩阵乘法    M = (float*)malloc(sizeof(float) * Width * Width);    N = (float*)malloc(sizeof(float) * Width * Width);    Pc= (float*)malloc(sizeof(float) * Width * Width);//保存CPU计算结果    srand(0);    matgen(M, Width);//产生矩阵M    matgen(N, Width);//产生矩阵N    double timeStart,timeEnd; //定义时间，求时间差用    timeStart = clock();    MatrixMultiplication(M, N, Pc, Width);//CPU上计算    timeEnd = clock();    free(M);    free(N);    free(Pc);return timeEnd - timeStart;}//////////////////////////////////////////////////////////////////////////void main(){printf("CPU use time %g\n", MatrixMul_CPU());printf("GPU use time %g\n", MatrixMul_GPU());}

5.测试结果

测试时开了其他的应用程序，另外本机配置很戳，看看吧，使用CUDA进行加速甩了使用传统方法几条街呢

参考一下文章改写：

后注：代码中有点问题，测试结果也不对，后来发现了，改过的结果见该文http://www.cnblogs.com/Romi/archive/2012/05/17/2506787.html

http://www.cnblogs.com/Romi/archive/2012/05/09/2492363.html

-----------------------------------------------------------------------------------------------------------

上篇中http://www.cnblogs.com/Romi/archive/2012/05/09/2492363.html，出了点问题，也是后来才发现的，意识到每个块中最多只能有512个线程，而该文的块大小为64*64，显然超过了512，因此此篇将其改为16，即TILE_WIDTH=16。其他代码还是和上篇一样。

矩阵计算模型的数组元素索引如下图所示

测试结果如下：

GPU计算时间变长了，上篇那样数组中的元素并没有全计算到。可以看到GPU计算时间虽然也有点多，但还是比CPU串行计算要快。

此文中数据保存在全局存储器，进行计算时，从全局存储区取数据进行计算，而从全局存储器取数据的速度是很慢的，而且取矩阵元素有很多重复，即一个矩阵元素取了好多次，这些都会对计算性能产生影响，因此还可以进一步优化。

http://www.cnblogs.com/Romi/archive/2012/05/17/2506787.html

0 0