GPU矩阵点积代码实现及分析

最新推荐文章于 2024-07-06 02:15:17 发布

seeCrazy

最新推荐文章于 2024-07-06 02:15:17 发布

阅读量1.8k

点赞数

分类专栏： gpu 文章标签： gpu cuda 高性能缓存优化

本文链接：https://blog.csdn.net/seeCrazy/article/details/41693653

版权

本文通过设计CUDA代码实现矩阵乘法，并分析了不同算法的效率。对比了不使用共享内存、使用16*16和8*8块大小的共享内存实现，指出共享内存的使用能有效减少全局存储器访问次数，提高计算效率。实验结果表明，随着分块大小增加，性能提升与全局存储器访问次数减少成正比。然而，在特定运行次数下，如10000次，任务a与任务b的效率接近，这可能得益于GPU的硬件优化技术，如缓存或合并访问技术。

摘要由CSDN通过智能技术生成

实验内容

设计实现矩阵乘法A*B=C的CUDA代码，并分析使用不同算法所产生的不同效率的原因。

不使用共享内存的实现方式。
使用共享内存的实现方式，分块实现块大小为16*16
使用共享内存的实现方式，分块实现块大小为8*8

注：因共享内存有限，故对矩阵的乘法采用分块实现。

实验分析

任务a：每一个线程负责产生结果矩阵C中的一个元素，C中每一个元素的产生都需要A中一行元素与B中一列元素进行点积运算，最终的问题就是在代码中将线程标识与其所要计算的元素联系起来。因为C中的元素个数与线程总数相等，所以可以先求出每个线程在整个grid空间中的二维坐标，自然而然地与问题空间的二维坐标对应了起来。Kernel函数部分的代码如下：

__global__ void MatrixMulKernel_v1(const float* d_Md, const float* d_Nd, float* d_Pd, const int Width)
{
1	int col = blockIdx.x * blockDim.x + threadIdx.x;
2	int row = blockIdx.y * blockDim.y + threadIdx.y;
3	int value = 0;
4	for(int i = 0; i < Width; i++)
5	{
6		value += d_Md[row * Width + i] * d_Nd[i * Width + col];
7	}
8	d_Pd[row * Width + col] = value;
}

在上面的代码实现中，C中每一个元素的计算都需要访问A中的一行与B中的一列，而A与B中的元素是存储在全局存储器中的，若不考虑其他硬件技术的支持，假设每取一个元素访问一次存储单元，对于一个宽度与高度都为N的矩阵来说，每一个线程要访问全局存储器2*N+1次，则完成整个矩阵计算任务需要访问全局存储器N*N*(2*N+1)次，即O(N^3)次。这样，访问存储器的时间延迟将成为计算效率的瓶颈。而且可以注意到，每个线程所访问的元素之间并不是完全独立的，比如row值相同的线程会访问A中的同一行，col值相同的线程会访问B中的同一列。所以可以对算法进行改进，通过将数据先缓存到共享存储器，然后再进行计算，可以有效地减少访问全局存储器的次数。