基于GPU的GEMM矩阵相乘运算优化

原创 已于 2023-06-20 11:37:10 修改 · 1.9k 阅读 · 13 ·
CC 4.0 BY-SA版权
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
文章标签：

#矩阵 #算法 #c++

于 2023-06-12 10:40:27 首次发布
文章详细介绍了GEMM（GeneralMatrixMultiply）操作的优化方法，包括矩阵的分块处理（如A和B的竖切和横切），以及如何通过线程块和CUDA架构提高计算效率。关键在于将矩阵乘法转换为向量乘法，通过块点乘提升运算访存比，优化内存访问和计算性能。文章还讨论了不同级别的BLAS操作和矩阵划分策略。

从上图中我们可以看到三种处理方法。第一种是将A和B矩阵分块（竖切和横切），第二种方法是将C和B矩阵分块（竖切和竖切），第三种方法是将C和A矩阵分块（横切和横切）:

 GEMM的子任务是GEPP或GEMP；最小粒度的任务是GEBP或GEPB或点乘。 

这里面M表示横向和纵向维度都很大的矩阵，P表示横向或纵向有一个维度很小的矩阵（或者就是一个向量），B表示横向和纵向维度都很大的矩阵（或者就是只有一个元素的矩阵或向量）。

其各自计算方法大致如下：

对于内存和线程结构：

最低0.47元/天 解锁文章
从 GEMM 到 DeepGEMM：GPU 矩阵计算原理与优化
关于我的学习记录
07-29 1995
1.1 GEMM简介GEMM通用矩阵的矩阵乘法）是BLAS（Basic Linear Algebra Subprograms）库中的一个函数，用于实现矩阵与矩阵之间的乘法运算。BLAS库是一组用于执行基础线性代数运算的子程序库，包括向量加法、数乘、点积、矩阵相乘等。GEMM 的数学原理GEMM 的基本公式：如上图所示，简单来说，GEMM就是将两个矩阵相乘，得到一个输出矩阵的过程。m < M;k++) {
CUDA算子优化：矩阵乘GEMM优化(二)
qq_56047026的博客
06-13 2480
A_TILE_ROW_START代表在这个维度为bm*bk的数据块中，当前线程需要搬运的数据的竖向坐标，而A_TILE_COL代表需要搬运的数据的横向坐标。那么我们需要开启（2048/128）*（2048/128）=256个block，每个block里面有（128/8）*（128/8）=256个线程，每个线程需要负责计算C矩阵中8*8=64个元素的结果，每个block负责256*64=16384个元素的结果。在一轮迭代中需要使用bm*bk的数据，为了加快后续的访存，所以需要进行一次转置。
GEMM详解
最新发布
m0_49133355的博客
09-22 1121
GEMM是的缩写，即通用矩阵乘法。A, B, C: 这三个都是矩阵。A 的维度是m x k，B 的维度是k x n，C 的维度是m x n。α (alpha)和β (beta): 这两个是标量（也就是普通的数字），用于对矩阵运算的结果进行缩放和累加。A * B: 这是标准的矩阵乘法，结果是一个m x n的矩阵。: 将矩阵乘法的结果 A*B 中的每个元素都乘以标量 α。β * C: 将原始矩阵 C 中的每个元素都乘以标量 β。: 最后，将前两项的结果相加，更新矩阵 C。标准矩阵乘法。
矩阵乘法(GEMM)优化算法在GPU上的实现
m0_72410588的博客
05-26 1928
矩阵乘法的基本思路是将两个矩阵相应位置的元素相乘，并累加得到结果矩阵的对应位置的元素。设AAA为形状为m×km\times km×k的矩阵，BBB为形状为k×nk \times nk×n的矩阵，则得到的矩阵CCC形状为m×nm \times nm×n，其中第iii行第jjj列的元素为Cij∑r1kAirBrjCij​r1∑k​Air​Brj​。
一个用GPU写的矩阵乘法
m0_38036750的博客
08-18 2846
这个过程相当繁琐,个人认为有优化的可能: 先说一下思路,矩阵相乘A矩阵乘B矩阵相当于A矩阵和B矩阵的转置做内积.所以我就先把B矩阵做了转置,再做内积.其中有两个核函数是在主函数中执行的,先执行转置,再执行乘法.再乘法函数中又嵌套了一个内积函数.这样充分的利用了并行化. 如图所示: 以3*3矩阵为例.我先开3*3个线程做内积运算,然后在每个线程中又开了1*3个线程做内积运算.其中求和部分我...
GPU矩阵计算
leon1171的博客
05-22 4923
学习GPU加速 参考例子和网上博文写的GPU矩阵计算 //矩阵转置 __global__ void transposeGPU(double *outdata, double *indata, int width, int height) { __shared__ double block[BLOCK_DIM][BLOCK_DIM + 1]; unsigned int xIndex...
通用矩阵乘（GEMM）优化与卷积计算
qq_35985044的博客
12-28 4483
​ 本文简要介绍通用矩阵乘（GEMM，General Matrix Multiplication）优化的基本概念和方法、神经网络量化中矩阵乘的优化方法。
实现CUDA核函数——二维矩阵相乘
02-19 1132
初学CUDA，自己练手实现二维矩阵乘法，且并发线程数与矩阵计算数据量自适应。计算结果，与调用numpy、torch库的计算方法进行比较验证通过。 环境： win10，PyCUDA: (2021, 1)，Python: 3.9.7 说明：调用时传入的grid、block参数，不能超过GPU的限制，所以需要先调用pycuda.driver库进行相关操作获取，网上有很多资料。 代码如下： # -*- coding: utf-8 -*- """ Created on Wed Feb 19 21:33:17
细粒度任务并行GPU通用矩阵乘.pdf
09-25
GEMM是两个矩阵相乘的操作，对于大规模计算可以提供高效的性能。 现有的GPU库如cuBLAS和MAGMA已经实现了GEMM操作，并通过内核函数在大型GEMM计算中达到高性能。然而，这些实现对大批量的小型GEMM计算效率并不高，且...
FlashAttention3：“GEMM”就是比较快！
鲁班模锤
07-15 2524
FlashAttention-3比使用FP16的FlashAttention-2 快1.5-2.0倍，即H100理论最大FLOPS利用率为 75%。使用FP8 时，FlashAttention-3 达到接近 1.2 PFLOPS。
用gpu实现矩阵相乘，加速比达500
08-19
用GPU实现两矩阵相乘，与串行算法相比，加速比达到了500多。
opencl 多gpu矩阵数组运算
12-06
1：opencl 多gpu运算； 2：opencl 矩阵运算； 3：opencl 数组运算；
GPU上 多种方法的矩阵乘法
11-16
GPU上使用多种方法的矩阵乘法，让你在用到 矩阵乘的时候有更为简洁的方法
GPU上矩阵乘法快速实现
07-14
GPU上矩阵乘法的设计与实现，对于在cuda上做编程优化的同学有一定的参考价值。
矩阵乘法——基于GPU的并行编程模型CUDA程序设计
千色千寻的博客
08-21 4470
CUDA并行编程——矩阵乘法 目录CUDA并行编程——矩阵乘法1 题目描述2 设计思路实验环境3 源码3.1 串行程序3.2 并行程序3.3 性能对比与分析 1 题目描述 题目1：编写一个矩阵乘法的GPU并行程序，并且与对应规模的串行程序进行运行时间的比对（n=500，1000，1500，2000，3000，5000），画出规模和时间对比图。 矩阵A（n，n） 矩阵B（n，n） C = A x B 要求：   1、完成程序的开发并验证其正确性，完成一个实验报告（程序源代码、变量和语句的详细说明；   2、
GPU矩阵并行相乘
qq_45769678的博客
05-24 2701
内部类 理解 把类A定义在类B内部，类A就被称为内部类 内部类语法格式 public class OuterClass{ } 内部类访问特点 内部类可以直接访问外部类的成员，包括私有 外部类要想访问内部类的成员，必须创建对象 内部类分类 成员内部类 分类 静态内部类（static，属于外部类本身） 非静态内部类（private） num 先查找局部变量 this.num 再查找内部类的成员 Outer.this.name; 最后查找外部类对象 局部内部类 局部内部类访问局部变量，必须加f
gpu cuda 矩阵乘法
闪电彬彬的博客
10-24 712
gpu矩阵乘法
如何在CPU上优化GEMM（下）
吴建明wujianming_110117
12-09 738
如何在CPU上优化GEMM（下） Array Packing 另一个重要的技巧是数组打包。这个技巧是对数组的存储维度进行重新排序，将某个维度上的连续访问模式在平滑后转换为顺序模式。 如上图所示，在阻塞计算之后，可以观察到B的数组访问模式（扁平化后），它是规则的但不连续的。期望经过一些转换，可以得到连续访问模式。可以将[16][16]数组重新排序为[16/4][16][4]数组，这样当从压缩数组中获取相应的值时，B的访问模式将是顺序的。 We have to re-write the algorithm s
cublas矩阵相乘
02-27
### 使用 cuBLAS 实现矩阵相乘 为了使用 cuBLAS 进行高效的矩阵相乘操作，需要初始化 CUDA 和 cuBLAS 库，并调用相应的 API 函数完成计算。下面是一个完整的例子说明如何通过 cuBLAS 完成两个浮点数矩阵之间的乘法。 #### 初始化环境与资源分配 在执行任何 cuBLAS 调用之前，先要创建一个上下文句柄用于后续的操作： ```cpp #include <cuda_runtime.h> #include <cublas_v2.h> // 错误处理宏定义简化错误检测过程 #define cudaErrorCheck(call) \ do { \ cudaError_t err = call; \ if (err != cudaSuccess) { \ fprintf(stderr, "CUDA Error: %s\n", cudaGetErrorString(err)); \ exit(EXIT_FAILURE); \ } \ } while(0) int main() { cublasHandle_t handle; cublasStatus_t status; // 创建 cuBLAS 句柄 status = cublasCreate(&handle); if (status != CUBLAS_STATUS_SUCCESS){ printf("Failed to create cuBLAS handle.\n"); return -1; } } ``` 此部分代码负责设置必要的运行时环境以及异常情况下的报错机制[^1]。 #### 数据准备 接着准备好待运算的数据集并将它们传输到 GPU 设备上： ```cpp const int m = 2; // 行数 const int n = 3; // 列数 const int k = 4; // 中间维度大小 float h_A[m*k] = {...}; // 主机端输入矩阵A数据填充... float h_B[k*n] = {...}; // 主机端输入矩阵B数据填充... float *d_A, *d_B, *d_C; size_t size_A = m * k * sizeof(float); size_t size_B = k * n * sizeof(float); size_t size_C = m * n * sizeof(float); // 分配设备内存空间给三个矩阵变量 cudaErrorCheck(cudaMalloc((void**)&d_A, size_A)); cudaErrorCheck(cudaMalloc((void**)&d_B, size_B)); cudaErrorCheck(cudaMalloc((void**)&d_C, size_C)); // 将主机上的初始值复制到对应的设备指针位置处 cudaErrorCheck(cudaMemcpy(d_A, h_A, size_A, cudaMemcpyHostToDevice)); cudaErrorCheck(cudaMemcpy(d_B, h_B, size_B, cudaMemcpyHostToDevice)); ``` 这里假设 `h_A` 和 `h_B` 已经被适当赋值为所需的测试数值。 #### 执行矩阵乘法 现在可以调用 cuBLAS 提供的 gemm 接口来进行实际的矩阵乘法运算了: ```cpp // 设置 alpha=1.0f,beta=0.0f 参数控制线性组合系数 const float alf = 1.f; const float bet = 0.f; // 调用 cuBLAS 的 GEMM 方法实现 C := α*A*B + β*C status = cublasSgemm(handle, CUBLAS_OP_N /* op(A)=A */, CUBLAS_OP_N /* op(B)=B */, n /* Number of rows in matrix op(B) and C*/, m /* Number of columns in matrix op(A) and C*/, k /* Number of columns/rows in matrices A/B */, &alf, d_B, n /* Leading dimension of B */, d_A, k /* Leading dimension of A */, &bet, d_C, n /* Leading dimension of C */ ); if(status != CUBLAS_STATUS_SUCCESS){ printf("GEMM failed!\n"); return -1; } // 把结果从设备拷贝回主机侧以便验证正确性 float h_C[m*n]; cudaErrorCheck(cudaMemcpy(h_C, d_C, size_C, cudaMemcpyDeviceToHost)); ``` 这段程序片段展示了怎样配置参数并通过一次函数调用来高效地完成大规模矩阵间的乘积运算。 #### 清理工作 最后不要忘记释放所占用的所有硬件资源： ```cpp // 解除对 cuBLAS 上下文对象的引用并销毁它 cublasDestroy(handle); // 释放先前申请过的显存区域 cudaFree(d_A); cudaFree(d_B); cudaFree(d_C); ``` 上述流程构成了基于 NVIDIA cuBLAS 库的一个典型应用案例，能够显著提升涉及大量密集型代数变换任务的工作效率。
评论 3
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

抵扣说明：

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

余额充值