CUDA矩阵乘法——VS2010中使用CUDA示例
来源:互联网 发布:js将对象转换成字符串 编辑:程序博客网 时间:2024/05/18 14:45
各工具或库的版本:
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设备上进行矩阵乘法:
//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运行效率。
#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- CUDA矩阵乘法——VS2010中使用CUDA示例
- CUDA矩阵乘法——VS2010中使用CUDA示例
- CUDA矩阵乘法——VS2010中使用CUDA示例
- 矩阵乘法——CUDA 优化记录
- CUDA之矩阵乘法——globalmemory
- CUDA之矩阵乘法——复数
- CUDA矩阵乘法
- CUDA 矩阵乘法优化
- CUDA:矩阵乘法原理
- CUDA: 矩阵乘法优化
- cuda 矩阵乘法
- cuda编程------矩阵乘法
- CUDA: 矩阵乘法优化
- cuda中的矩阵乘法
- CUDA矩阵乘法
- CUDA矩阵乘法
- CUDA矩阵乘法——利用共享存储器
- CUDA之矩阵乘法——TILE&sharedmemory
- 多维矩阵转一维矩阵
- ffmpeg 在ubuntu 编译
- 第八周项目4-给r,h,求各项数据
- 例解三层交换原理
- 使用 JCaptcha 开发图形和声音验证码
- CUDA矩阵乘法——VS2010中使用CUDA示例
- android jni开发
- 开发日志:注意struts标签的注释
- 将博客搬至CSDN
- CAS项目登录流程介绍(二)
- 《Oracle Applications DBA 基础》- 7-8 OAM及系统管理
- makefile概述
- spring的定时任务
- 使用MediaRecorder录制音频