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

最新推荐文章于 2023-05-04 13:44:38 发布
最新推荐文章于 2023-05-04 13:44:38 发布
阅读量488 收藏 2
点赞数
分类专栏： CUDA 文章标签： cuda 矩阵乘法

文章转自：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 plain copy print ?
  1. //CUDAtest.cu
  3. #include "cuda_runtime.h"
  4. #include "device_launch_parameters.h"
  6. #define TILE_WIDTH 64
  8. // 核函数
  9. // __global__ static void MatrixMulKernel(const float* Md,const float* Nd,float* Pd,int Width)
  10. __global__ void MatrixMulKernel(const float* Md,const float* Nd,float* Pd,int Width)
  11. {
  12. //计算Pd和Md中元素的行索引
  13. int Row = blockIdx.y * TILE_WIDTH + threadIdx.y; //行
  14. int Col = blockIdx.x * TILE_WIDTH + threadIdx.x; //列
  16. float Pvalue = 0.0;
  17. for (int k = 0; k < Width; k++)
  18. {
  19. Pvalue += Md[Row * Width + k] * Nd[k * Width + Col];
  20. }
  21. //每个线程负责计算P中的一个元素
  22. Pd[Row * Width + Col] = Pvalue;
  23. }
  25. // 矩阵乘法(CUDA中)
  26. // 在外部调用，使用extern
  27. extern "C" void MatrixMultiplication_CUDA(const float* M,const float* N,float* P,int Width)
  28. {
  29. cudaSetDevice(0); //设置目标GPU
  31. float *Md, *Nd, *Pd;
  32. int size = Width * Width * sizeof(float);//字节长度
  34. cudaMalloc((void**)&Md, size);
  35. cudaMalloc((void**)&Nd, size);
  36. cudaMalloc((void**)&Pd, size);
  38. //Copies a matrix from the memory* area pointed to by src to the memory area pointed to by dst
  39. cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice);
  40. cudaMemcpy(Nd, N, size, cudaMemcpyHostToDevice);
  42. //
  43. dim3 dimGrid(Width / TILE_WIDTH, Width / TILE_WIDTH); //网格的维度
  44. dim3 dimBlock(TILE_WIDTH, TILE_WIDTH); //块的维度
  45. MatrixMulKernel<<< dimGrid, dimBlock >>>(Md, Nd, Pd, Width);
  47. cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost);
  48. //释放设备上的矩阵
  49. cudaFree(Md);
  50. cudaFree(Nd);
  51. cudaFree(Pd);
  52. } 
//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中)
// 在外部调用，使用extern
extern "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 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);
}

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

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


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

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

[cpp] view plain copy print ?
  1. #include <time.h>
  2. #include <stdlib.h>
  3. #include <stdio.h>
  5. //这里不要忘了加引用声明
  6. extern "C" void MatrixMultiplication_CUDA(const float* M, const float* N, float* P, int Width);
  8. //构造函数...
  9. //析构函数...
  11. // 产生矩阵,矩阵中元素0~1
  12. void matgen(float* a, int Width)
  13. {
  14. int i, j;
  15. for (i = 0; i < Width; i++)
  16. {
  17. for (j = 0; j < Width; j++)
  18. {
  19. a[i * Width + j] = (float)rand() / RAND_MAX + (float)rand() / (RAND_MAX*RAND_MAX);
  20. }
  21. }
  22. }
  24. //矩阵乘法(CPU验证)
  25. void MatrixMultiplication(const float* M,const float* N,float* P,int Width)
  26. {
  27. int i,j,k;
  28. for (i = 0; i < Width; i++)
  29. {
  30. for (j = 0; j < Width; j++)
  31. {
  32. float sum = 0;
  33. for (k = 0; k < Width; k++)
  34. {
  35. sum += M[i * Width + k] * N[k * Width + j];
  36. }
  37. P[i * Width + j] = sum;
  38. }
  39. }
  40. }
  42. double MatrixMul_GPU()
  43. {
  44. float *M, *N, *Pg;
  45. int Width = 1024; //1024×1024矩阵乘法
  46. M = (float*)malloc(sizeof(float) * Width * Width);
  47. N = (float*)malloc(sizeof(float) * Width * Width);
  48. Pg= (float*)malloc(sizeof(float) * Width * Width); //保存GPU计算结果
  50. srand(0);
  52. matgen(M, Width); //产生矩阵M
  53. matgen(N, Width); //产生矩阵N
  55. double timeStart, timeEnd; //定义时间，求时间差用
  56. timeStart = clock();
  57. MatrixMultiplication_CUDA(M, N, Pg, Width); //GPU上计算
  58. timeEnd = clock();
  60. free(M);
  61. free(N);
  62. free(Pg);
  63. return timeEnd - timeStart;
  64. }
  66. double MatrixMul_CPU()
  67. {
  68. float *M, *N, *Pc;
  69. int Width = 1024; //1024×1024矩阵乘法
  70. M = (float*)malloc(sizeof(float) * Width * Width);
  71. N = (float*)malloc(sizeof(float) * Width * Width);
  72. Pc= (float*)malloc(sizeof(float) * Width * Width); //保存CPU计算结果
  74. srand(0);
  76. matgen(M, Width); //产生矩阵M
  77. matgen(N, Width); //产生矩阵N
  79. double timeStart,timeEnd; //定义时间，求时间差用
  80. timeStart = clock();
  81. MatrixMultiplication(M, N, Pc, Width); //CPU上计算
  82. timeEnd = clock();
  84. free(M);
  85. free(N);
  86. free(Pc);
  87. return timeEnd - timeStart;
  88. }
  90. //
  91. void main()
  92. {
  93. printf("CPU use time %g\n", MatrixMul_CPU());
  94. printf("GPU use time %g\n", MatrixMul_GPU());
  95. } 
#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(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

huikong2013
关注
专栏目录
cublas,tensorcore矩阵乘法基本介绍
forrestguang的博客
11-14 434
CαopAopBβCCαopAopB))βC，这里的运算符op决定是否转置。这里首先说明，cublasSgemm专门处理float类型的二维矩阵乘法，针对其他类型的矩阵，cublas有类似的函数接口cublasDgemm：处理double类型矩阵cublasCgemm：处理cuComplex类型矩阵cublasZgemm：处理cuDoubleComplex类型矩阵cublasHgemm：处理__half类型的矩阵。
一篇文章理解CUDA架构、编程与进阶使用
qq_37764141的博客
01-26 5224
一篇文章理解CUDA架构、编程与进阶使用一、CUDA架构二、CUDA编程基础1.矩阵加法2.矩阵乘法三、CUDA进阶——利用共享内存加速访存1.CUDA内存读写速度比较2.申请共享内存 一、CUDA架构 二、CUDA编程基础 从矩阵加法和矩阵乘法来认识CUDA编程 1.矩阵加法 __global__ void vecAddKernel(float* A_d, float* B_d, float* C_d, int n) { int i = threadIdx.x + blockDim.x * bl
CUDA学习--矩阵乘法的并行运算
热门推荐
Yi_M的博客
08-24 1万+
1. CUDA学习步骤 CPU实现 a*b = c 的矩阵乘法（矩阵尺寸是n*m的，n和m大于1000） 下载 https://developer.nvidia.com/cuda-downloads，安装好cuda 将cpu代码移植到cuda。将CPU值传入GPU，使用cuda计算，与cpu结果对比。 优化思路1：将矩阵分块进行计算 优化思路2：使用share memory进行优化 优化思路3：将数
基于CUDA的矩阵相乘
身是菩提树、心是明镜台、时时勤拂拭、莫使惹尘埃
03-09 9435
这几天研究了一下CUDA，发现其并行的思想和普通的CPU多线程思想不太一致，但还是挺不错。主要是将任务划分成一个个block，然后每个block里面再划分成细的线程。然后每个线程做自己做的 事情。这种并行思想很适用于像矩阵运算这些元素与元素之间的运算并不耦合得很厉害，但整体数据很大的情况，这只是我对CUDA的初步感觉。 矩阵相乘的CPU程序如下： //C = A*B void
2.3CUDA矩阵乘法
Z_Paper的博客
10-19 469
CPU 矩阵乘法 能相乘的两个矩阵，必须满足一个矩阵的行数和第二个矩阵的列数相同. A(N*P) * B(P*M) = C(N*M).　其中P是行数，N是列数，　从宽高的角度来说，即 A的宽度和B的高度是相同的．C矩阵 = ha * wb. 其中C(i,j) = A矩阵中的i行和B矩阵中的j列进行点乘得到该点的值． //C = A*B void MatrixMulCPU(float* _C,const float *_A,const float *_B,int _wa,int _ha,in..
cuda 矩阵乘法
todayq的专栏
12-04 1163
原始矩阵乘法 #include #include #define BLOCK_SIZE 1 typedef struct{ int width; int height; float* elements; }Matrix; __global__ void MatMulKernel(const Matrix,const Matrix,Matrix); void MatMul(const Mat
使用cuda C完成矩阵相乘算法详解
sinat_38368658的博客
03-26 5194
矩阵相乘大家应该都不陌生。 设有两个矩阵M和N，假设M和N都是方阵，维度均为width × width 如果M和N均为1000 × 1000的矩阵，总共要进行1000000次点乘。其中，每次点乘有1000次乘法和1000次加法。 Matirx Multiply：CPU实现 先来看看使用普通的c代码在CPU端如何实现 void MatrixMulOnHost(float* M,float* N,...
CUDA程序性能分析-矩阵乘法
張行之的专栏
12-02 3882
矩阵乘法非常适合在GPU上并行运行，但是使用GPU并行后能获得多大的性能提升？本文将通过一些实验分析CUDA程序的性能。本文主要记录了本人测试CUDA程序性能的结果，并对结果进行了分析，从测试结果和分析可以为并行程序和优化性能带来一些启示。
CUDA矩阵乘法及优化【参加CUDA线上训练营】
Mr_8j的博客
02-09 705
CUDA矩阵乘法及优化。本文是参加2023春季CUDA线上训练营学习笔记。
CUDA矩阵乘法计算
08-26
CUDA实现的矩阵乘法，运用了共享内存，纹理内存。
一个老外写的可扩展C++矩阵类（包括实例，VS2010可用）
07-23
一个老外写的可扩展C++矩阵类（包括实例，VS2010可用）
VS环境下的矩阵运算库，包含源代码
01-08
VS环境下的矩阵运算库，包含源代码和实例程序，很好用，推荐给大家
CUDA矩阵乘法
taifyang的博客
11-23 366
【代码】CUDA矩阵乘法。
CUDA编程之矩阵乘法
最新发布
qq_43398345的博客
05-04 2867
对于平铺矩阵算法，对于每一个Block，首先将其存储在共享内存中, 花费 * Block_Size ^ 2，对BLock_SIze中每一个元素进行读取需要花费: 2 * k / Block_Size 个单位时间，一共有Block_Size ^ 2个元素， 一共有 (m * n) / (block_size * block_size)个Block，所以总时间就是这些值进行相乘。n])的矩阵乘法运算，每个线程都要读取A的一整行和B的一整列。• 共享内存属于线程块，可以被一个线程块内的所有线程访问。
cuda矩阵乘法(简单理解)
Super5311的博客
07-18 1135
cuda矩阵乘法
cuda编程------矩阵乘法
ICT_100的专栏
03-21 1485
cuda + VS 2010 安装：http://www.cnblogs.com/xing901022/archive/2013/08/09/3248469.html 本文主要介绍如何使用CUDA并行计算矩阵乘法： //头文件 #include "cuda_runtime.h" #include "device_launch_parameters.h" #include __glob
Visual Studio 2019 CUDA配置教程：详解步骤与矩阵乘法示例
- 提供了矩阵乘法的CUDA C++代码片段，展示了如何编写并行计算的函数`MatrixMulKernel`。使用CUDA并行计算模型，利用`__global__`关键字定义kernel（线程块），以及数据块大小、线程索引等概念。 - 主函数`main()`...
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值