j ) a[ (i)*lda + (j) ] #define B( i, j ) b[ (i)*ldb + (j) ] #define C( i, j ) c[ (i)*ldc + (j) ] // gemm 因此,为了解决上一问题,gemm论文提出了矩阵分块的做法,直击核心,这篇论文针对矩阵乘法主要提出了下面6种不同的分块计算方法,如下图所示: ? 总结 这篇文章讲到的优化方法都是有理论支撑的,也就是第5节展示的gemm论文中的那个Figure4。 gemm论文我打算放到我后面的文章中进行解读,另外会再分享一些优化程度更大的算法,感兴趣的请关注一下我们的公众号,谢谢。 7. https://github.com/tpoisonooo/how-to-optimize-gemm
如何处理TensorFlow中的InternalError: Blas GEMM launch failed 摘要 大家好,我是默语,擅长全栈开发、运维和人工智能技术。 今天我们来探讨一个在使用TensorFlow时常见的问题:InternalError: Blas GEMM launch failed。 然而,在使用TensorFlow时,我们可能会遇到各种各样的错误,其中之一就是InternalError: Blas GEMM launch failed。 什么是InternalError: Blas GEMM launch failed? Q: 如何避免InternalError: Blas GEMM launch failed?
, CUBLAS_GEMM_ALGO3, CUBLAS_GEMM_ALGO4, CUBLAS_GEMM_ALGO5, CUBLAS_GEMM_ALGO6, CUBLAS_GEMM_ALGO7 , CUBLAS_GEMM_ALGO8, CUBLAS_GEMM_ALGO9, CUBLAS_GEMM_ALGO10, CUBLAS_GEMM_ALGO11, CUBLAS_GEMM_ALGO12 , CUBLAS_GEMM_ALGO13, CUBLAS_GEMM_ALGO14, CUBLAS_GEMM_ALGO15, CUBLAS_GEMM_ALGO16, CUBLAS_GEMM_ALGO17, CUBLAS_GEMM_DFALT_TENSOR_OP, CUBLAS_GEMM_ALGO0_TENSOR_OP, CUBLAS_GEMM_ALGO1 , CUBLAS_GEMM_ALGO18, CUBLAS_GEMM_ALGO19, CUBLAS_GEMM_ALGO20, CUBLAS_GEMM_ALGO21,
本文为Open AI Lab 工作人员投稿,对深度学习的核心操作GEMM进行了详细的使用介绍,欢迎对模型部署AI推断感兴趣的朋友关注Tengine。 福利来了,Tengine带来了一份超详细的gemm汇编教程。 ? GEMM简介 什么是GEMM? Why gemm is at the heart of deep learning[1]介绍了为什么GEMM在深度学习计算中如此重要,以及卷积计算中是如何使用GEMM。 ? 教程大纲 教程分为三部分: Step1: 纯C实现的gemm Step2: 调用OpenBLAS的gemm Step3: Tengine中的gemm 运行这个教程的代码,你需要: 可以执行armv8汇编的环境 Reference: [1] Why gemm is at the heart of deep learning (https://petewarden.com/2015/04/20/why-gemm-is-at-the-heart-of-deep-learning
,目前主流的卷积算法都是基于GEMM来实现的。 例如可以将A和B分解为分块矩阵,使得GEMM可以递归实现。有关GEMM的详细信息可以参见[1][2][3]。如何对GEMM进行优化,是BLAS相关工作的研究热点。 通过将卷积操作用矩阵乘法来代替,进而使用GEMM算法来间接进行卷积操作,这使得卷积操作可以在任何包含GEMM的平台上进行,并且受益于矩阵乘法的高效性,任何针对GEMM的改进和研究都能有助于卷积运算效率的提升 具体的,基于GEMM的卷积方法需要借助于 im2col或im2row buffer来内存转换,使得数据格式满足GEMM算法的输入要求,从而将卷积操作转化为GEMM操作,然而这个转换过程是一个计算开销和内存开销都比较大的过程 其中A是输入张量,B是一个常量滤波器,C是输出矩阵,在传统的im2col+GEMM算法中,通常α=1而β=0,原始GEMM操作示意图如下: ?
Deepseek开源周第三弹:DeepGEMM,一个 FP8 GEMM 库,支持密集 GEMM 和 MoE GEMM,为 V3/R1 训练和推理提供支持。 DeepGEMM 是一个库,专为干净高效的 FP8 通用矩阵乘法 (GEMM) 而设计,具有细粒度缩放,如 DeepSeek-V3 中所建议的那样。它支持普通和混合专家 (MoE) 分组 GEMM。 布局对比 接口说明 DeepGEMM 提供以下主要接口: 普通 GEMM:deep_gemm.gemm_fp8_fp8_bf16_nt,支持非分组的 FP8 GEMM。 分组 GEMM(连续布局):m_grouped_gemm_fp8_fp8_bf16_nt_contiguous,适用于 MoE 模型的训练前向传播或推理预填充。 分组 GEMM(掩码布局):m_grouped_gemm_fp8_fp8_bf16_nt_masked,适用于推理解码阶段。
今天我们从基础到进阶,系统拆解GEMM运算的原理、优化策略与工程落地方法,结合代码示例与性能监控方案,对GEMM刨根问底、一探究竟。二、GEMM运算介绍1. ,GEMM运算占总算力消耗的90%-95%,非GEMM运算(如激活函数、LayerNorm)仅占5%-10%,因此算力测算可近似围绕GEMM展开。 同时验证了单头注意力GEMM运算量理论值的准确性,帮助开发者理解高维场景下GEMM与大模型架构的深度绑定关系。 示例4:简化分块GEMM实现import torchdef tiled_gemm(A, B, tile_size=32): """ 分块GEMM运算实现 Args: A: ;GEMM耗时占比:通过PyTorch Profiler监控GEMM运算占总推理耗时的比例,理想值≥80%,占比过低说明非GEMM开销过大;显存带宽利用率:GEMM运算的显存读写带宽占GPU总带宽的比例
DeepGEMM Github地址:https://github.com/deepseek-ai/DeepGEMM 矩阵乘法GEMM 通用矩阵乘法 (GEMM) 是线性代数、机器学习、统计学和许多其他领域的常见算法 为什么GEMM 是深度学习的核心 GEMM(General Matrix Multiply,通用矩阵乘法)是深度学习神经网络优化中的一个关键函数。 为了更直观地理解GEMM的影响,我们可以参考贾扬清在其论文中的一张图表。该图表展示了在不同硬件平台上执行矩阵乘法时,GEMM优化带来的性能提升。 所有以 fc(全连接)或 conv(卷积)开头的层都是使用 GEMM 实现的,几乎所有时间(GPU 版本的 95%,CPU 版本的 89%)都花在这些层上。 那么什么是 GEMM? Y_e = XW_e 是分组 GEMM 的典型场景。
由于笔者只对GEMM的优化熟悉,这里就以优化X86的GEMM为例子来探索。 入门GEMM优化 https://github.com/flame/how-to-optimize-gemm/wiki 中介绍了如何采用各种优化方法来优化GEMM。 最终将GEMM的性能提到到原始版本的8倍以上。 假如你和我一样,对如何凑出高效的GEMM并不敏感,并且你有需要将一个GEMM算子优化到性能比较好的需求时你可以怎么做呢? 所以我想的是是否可以基于Ansor的搜索结果来指导我来编写高效的GEMM程序。 /gemm.py 。
将 fp16 GEMM 替换为 int8 GEMM 不仅可以缩短 GEMM 时间,还可以减小前后算子的输入输出位宽,从而减小读写数据的时间。 首先所有 int8 GEMM 的输入和输出都需要进行量化。由于 int8 GEMM 的 shape 限制,部分 GEMM(例如注意力分数的计算)仍然采用 float GEMM。 此外第二层 FFN 的 GEMM 采用的是 int32 的输出,因为它的 GEMM 输入是 ReLU 激活函数的输出结果,只包含正数,非对称,因此如果采用 int8 输出的 GEMM,将无法反量化为正确的浮点数结果 图中黄色矩形表示 int8 GEMM,绿色矩形表示 float GEMM。这里采用 float GEMM 是由于 shape 的限制,不适合使用 int8 GEMM 加速。 GEMM 计算。
2025 年 2 月 26 日,DeepSeek 在开源周的第三天,正式发布了高效的 FP8 通用矩阵乘法(GEMM)库 —— DeepGEMM。 在大规模模型训练和推理中,矩阵乘法(GEMM,General Matrix Multiplications)是最核心的计算操作之一,尤其是在深度学习模型的训练和推理过程中,GEMM 占据了绝大部分的计算资源 随着模型规模的不断扩大,尤其是混合专家模型(MoE)的兴起,传统的 GEMM 实现已经无法满足高效计算的需求。 然而,现有的 GEMM 库对 FP8 的支持有限,尤其是在 MoE 场景下,缺乏针对性的优化。 密集和 MoE GEMM 支持:该库不仅能够高效处理传统的密集矩阵乘法,还专门针对 MoE 模型中的 GEMM 运算进行了优化,能够满足不同类型模型的计算需求。
的文章,它主要优化了Im2Col+GEMM计算策略中的内存消耗,并且也能提升一点速度,是一个不错的卷积加速算法。 目前,对卷积层的计算一般有以下几种方式: Im2Col+GEMM。 Figure1 所以,MEC改进了Im2Col+GEMM的策略,目的是减少它的内存消耗同时提升一点速度。 3. 但是,在实际操作中,子矩阵的数量对性能的影响是很大的,在Solution1中执行了 次gemm,而Solution2中执行了 次gemm,如果使用Blas矩阵计算库,那么这两种方法在特定硬件平台如GPU 效果 测试了一下复现的MEC和原始的Im2Col+GEMM的速度和内存消耗,结果如下: ?
这里假设以一个gemm说明为例: // cutlass::gemm::kernel::GemmUniversal: ClusterTileM and ClusterTileN loops // are ::gemm::kernel::GemmUniversal Collective cutlass::gemm::collective::CollectiveMma cutlass::epilogue:: 这是一个基本的gemm示例,你可以看到它可能是所需参数的三分之一多。 using Gemm = typename gemm::kernel::DefaultGemmUniversal< half, layout::RowMajor, ComplexTransform B可以是int8,而在gemm中,我们将向上转换到BF16的操作。我们还对我们的lower alignment gemm进行了性能改进。
cudaFree(A); cudaFree(B); cudaFree(C); } template <typename T, typename S> int cublas_gemm_ex ; int end_algo = CUBLAS_GEMM_ALGO23; int start_algo_t_op = CUBLAS_GEMM_DEFAULT_TENSOR_OP; ,然后执行下面命令进行编译: nvcc test_gemm.cpp -o test_gemm -L/usr/local/cuda/lib64 -lcudart -lcuda -lcublas 最后执行 /test_gemm运行就行了。 这里计算的是 ,其中 的维度是 , 的维度是 , 的维度是 。 此外我还对比了不同的GEMM算法的效果。
将 fp16 GEMM 替换为 int8 GEMM 不仅可以缩短 GEMM 时间,还可以减小前后算子的输入输出位宽,从而减小读写数据的时间。 首先所有 int8 GEMM 的输入和输出都需要进行量化。由于 int8 GEMM 的 shape 限制,部分 GEMM(例如注意力分数的计算)仍然采用 float GEMM。 此外第二层 FFN 的 GEMM 采用的是 int32 的输出,因为它的 GEMM 输入是 ReLU 激活函数的输出结果,只包含正数,非对称,因此如果采用 int8 输出的 GEMM,将无法反量化为正确的浮点数结果 图中黄色矩形表示 int8 GEMM,绿色矩形表示 float GEMM。这里采用 float GEMM 是由于 shape 的限制,不适合使用 int8 GEMM 加速。 GEMM 计算。
短阵相乘(GEMM)测试的结果。GEMM是DNN中的关键操作。 研究1:矩阵相乘(GEMM)测试 DNN 高度依赖矩阵相乘运算(GEMM)。常规DNN依赖FP32密集GEMM。 然而,较低精度、稀疏的新兴DNN 依赖低精度及/或稀疏的GEMM。英特尔团队评估了这些不同的GEMM。 FP32 密集GEMM:由于FP32密集GEMM已得到充分的研究,英特尔团队比较了FPGA和GPU数据表上的峰值数字。 低精度INT6 GEMM:为了表明FPGA在可定制性方面的优点,该团队将四个int6封装到一个DSP模块中,以研究FPGA的6位(Int6)GEMM。 该团队还在 GPU上测试了稀疏的 GEMM,但发现性能不如在GPU上执行密集的GEMM(矩阵一样大小)。
下面,本教程将带领大家一步步优化矩阵乘法GEMM。无需手工撸代码,编写繁杂冗长的底层汇编代码,只需十几行简洁的调度代码。 ? 在详细讲解优化步骤前,我们先谈谈优化的本质。 用Halide语言直接描述GEMM的计算过程。 Var x,y; RDom k(0, K); Func gemm("gemm"); gemm(x, y) += A(k, y) * B(x, k); 计算M=N=K=640的矩阵乘法。 gemm(x, y) += A(k, y) * B(x, k); gemm.update() .tile(x, y, xo, yo, xi, yi, 16, 8) .reorder 总的代码如下 Func prod; prod(x, y) += A(k, y) * B(x, k); gemm(x, y) = prod(x, y); gemm.tile(x, y, xi, yi
integration for releasing pypi package pip install pnnx on windows/linux/macos 3(medium) optimize loongarch gemm 优化 loongarch gemm write gemm_loongarch.cpp with lsx extension test with qemu 利用lsx扩展优化实现gemm_loongarch.cpp ,使用qemu测试 goals: support any packed layout good benchmark result 4(medium-hard) optimize risc-v gemm 优化 risc-v gemm write gemm_riscv.cpp with risc-v vector extension and zfh(fp16) test with qemu 利用risc-v vector和zfh(fp16)扩展优化实现gemm_riscv.cpp,使用qemu测试 goals: support any packed layout and fp32/fp16 good benchmark
基于高性能的 CUTLASS [5] grouped GEMM, 分成两个 gemm kernel 实现,并把 add_bias, softmax 等操作 fused 到 GEMM kernel 中。 1.CUTLASS grouped GEMM NVIDIA 开发的 grouped GEMM 可以在一个 kernel 中完成多个独立矩阵乘问题的计算,利用这个性质可以实现 Attention 中的 padding 每个矩阵乘子问题,把问题大小传入到 grouped GEMM,其中 seqlen 传递真实的 valid seqlen 即可。 grouped GEMM 原理示意图。 每个子问题拆解为不同数量的块,再对这些块均匀分配,高效地实现单个 kernel 计算多个独立 GEMM 问题 使用 grouped GEMM 实现 attention 时,由于子问题的数量 batch_size
线程块专业化将通信和计算工作负载隔离,避免相互干扰,GEMM 线程块采用默认实现,通信线程块负责数据读写。 (通用矩阵乘法)操作 ;图下方 “Consumer” 代表消费者部分,包含 GEMM 操作以及 TopK-reduce + All2all/ReduceScatter 操作。 生产者先通过 All2all/AllGather 操作收集数据,再进行 GEMM 计算;消费者接收数据后,先进行 GEMM 计算,然后进行 TopK-reduce + All2all/ReduceScatter 计算重调度:右侧展示了重新调度(Reschedule)的过程,根据源 token(by source tokens)对分解后的子张量进行排序(sort),然后分别送入不同的 GEMM 计算单元(GEMM #0、GEMM #1、GEMM #2)进行分组 GEMM 计算(GroupGEMM),最终得到计算序列(compute sequence) 。