寒武纪显卡实现rmsNorm算子

已于 2024-09-27 13:55:34 修改
阅读量498 收藏
点赞数
分类专栏： 高性能计算 国产芯片编程 文章标签： 分布式
于 2024-06-05 15:46:10 首次发布
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.youkuaiyun.com/forrestguang/article/details/139462077
版权
本文详细介绍了在寒武纪显卡上实现rmsNorm算子的思路和步骤，包括针对不同dimsize与maxNum关系的处理、三级流水并行优化以及数据prefetch策略。此外，还提到了CMakeLists.txt构建项目的具体步骤，以帮助读者理解和实现该算子。

摘要生成于 C知道 ，由 DeepSeek-R1 满血版支持， 前往体验 >

rmsNorm算子的数学定义

下面这个是针对一维向量的rmsNorm算子，对于高维向量，比如说形状为[A,B,C,D]的4维向量x，此时权重向量w长度默认为D，也就是说变换主要是针对axis=-1这个维度进行操作。关于寒武纪显卡的编程可以参考本人之前编写的文章添加链接描述和添加链接描述
在这里插入图片描述

rmsNorm算子的实现思路

由于寒武纪编程比较特殊，这里我们稍微复习一下寒武纪编程的特点：一张寒武纪显卡往往可以有8个或者4个cluster，其中一个cluster包含4个core，也就是说，对于4个cluster的显卡来说，一张寒武纪显卡的并行度最多就是16（相比于CUDA显卡成千上万的thread就相形见绌）。但是寒武纪显卡的core非常特殊

了解本专栏 订阅专栏 解锁全文
寒武纪显卡实现softmax算子
forrestguang的博客
01-10 797
寒武纪实现softmax包括下面5个步骤，我们也采取5个kernel来实现softmax：unionMaxKernel(float* middle, float* source1, int num)，这个kernel使用的任务类型是union1，其中middle的长度为taskDim，达到的目的是middle对应元素存储的是不同taskId处理的那部分数据的局部max。
极智AI | 寒武纪 pytorch-mlu 添加逐层算子方法
极智视界
11-24 2897
本教程分享了在寒武纪设备上 pytorch-mlu 中添加逐层算子的方法。
算子实现_「专利解密」寒武纪AI芯片 混合拼接算子
weixin_32848617的博客
01-12 521
【嘉德点评】寒武纪发明的AI芯片，借助于人工智能处理库中的常用算子，可以十分方便的完成卷积计算，同时，混合而成的拼接算子不仅可以提高运算效率，也能够节省开发成本。集微网消息，“寒武纪”是中国科学院计算技术研究所发布的全球首个能够“深度学习”的“神经网络”处理器芯片。它是给电脑创造出模仿人类大脑多层大规模人工神经网络的芯片。在深度学习处理器的运行当中，计算系统的运算能力提升是决定深度学习处理效率的关...
英伟达、寒武纪、GTI的对比
weixin_44563740的博客
01-21 7171
英伟达、寒武纪、GTI的比照 　　现在活泼在AI芯片范畴的，构成影响力的有Nvidia和华为麒麟970带红的寒武纪等我国新星，他们也有不少创新。可是大部分品牌能做到在终端产品内大容量本地核算，不必后台核算机支撑吗?功耗方面，他们和GTI同级别的产品显然有很大的差距。 　　GTI于本年9月推出的芯片Lightspeeur™光矛系列是全球首款可一同支撑图像与视频、语音与自然语言处理的智能神经网络专用处...
寒武纪思元370GPU性能及应用场景
2402_84466582的博客
06-03 2941
寒武纪思元370GPU性能及应用场景
在国产GPU寒武纪MLU上快速上手Pytorch使用指南
qq_39691492的博客
12-20 9104
在国产GPU寒武纪MLU上快速上手Pytorch使用指南
【寒武纪（16）】寒武纪和GPU的硬件架构的差异
Hali_Botebie的博客
11-21 2037
如果想更深入地了解GPU的设计细节、实现细节，可阅读GPU厂商定期发布的白皮书和各大高校、机构发布的论文。，虽然是多年前的视频，但比较系统、全面地讲解了GPU的机制和技术。
深度学习中的算子
xw555666的博客
04-04 3035
在深度学习中，层是模型的构建单元，它们通过各自定义的算子作用于函数空间，将输入数据映射到越来越抽象的特征表示，直至输出层给出最终的预测结果。模型由一系列算子的组合构成，模型的整体特点和其中的构成算子有关，模型的训练也和其中的算子有关。因此有必要仔细观察并讨论一下算子。
OpenCV中常用算子
Goodness2020的博客
08-03 1770
opencv中常见算子
RMSNorm算子的CUDA实现
最新发布
周同学的博客
02-19 1128
RMSNorm算子的CUDA实现
寒武纪如何查看mlu的运行状态
LDJKWX
10-20 4090
寒武纪（mlu）的运行状态查看问题描述： 问题描述： 查看mlu的运行状态： 在正确安装 neuware‑driver 包后，CNMon 即安装成功。查看 CNMon 安装后信息，在终端输入 cnmon 命令，显示内容如下所示： ...
AI处理器-寒武纪NPU芯片简介
热门推荐
耐心的小黑的博客
10-31 1万+
一、前言 当今时代，人工智能（AI）正被广泛运用于各式各样的应用上。人工智能的三大支撑是硬件、算法和数据，其中硬件指的是运行 AI 算法的芯片与相对应的计算平台。由于使用场景变多，所需处理的数据量变大，人们的需求也更高，这就使得AI算法必须能够高效的运行在硬件平台上。在硬件方面，目前主要是使用 GPU 并行计算神经网络，同时，还有 FPGA 和 ASIC 也具有未来异军突起的潜能。 GPU称为图形处理器，它是显卡的“心脏”，与 CPU 类似，只不过是一种专门进行图像运算工作的微处理器。GPU 在浮点运算、
寒武纪软件开发面试题C++
bolinzhiyi的博客
06-07 1924
寒武纪软件开发面试题1 虚函数2 C++纯虚函数3 抽象类4 虚函数与纯虚函数总结5 C++虚函数表剖析6 C++三种继承方式下的访问权限控制7 C++ 宏定义与内嵌函数 1 虚函数 定义一个函数为虚函数，不代表函数为不被实现的函数。（虚函数是可以被实现的） 定义他为虚函数是为了允许用基类的指针来调用子类的这个函数。（主要用于实现运行时多态） 定义一个函数为纯虚函数，才代表函数没有被实现。（纯虚函数是不被实现的） 定义纯虚函数是为了实现一个接口，起到一个规范的作用，规范继承这个类的程序员必须实现这个函数。
PETSC的安装
forrestguang的博客
10-21 5767
PETSC PETSC理解为一个软件，就像我们编译Python用的pycharm一样，在上面写代码，然后运行。 C或者C++ 这个软件的语言主要依赖于C和C++，注意，我们这个软件的语言只是语法类型和C很像，虽然差别很小，但是也不是C。这部分内容可以理解为高性能计算的基础。下面开始介绍安装过程： 1：虚拟机 我们首先要在电脑上安装虚拟机，我用的是oracle box,在这个软件上添加虚拟机。 2：PETSC.tar.gz 在假设有虚拟机以及会一点linux的基础语法基础上，我们开始安装PETSC。首先我们
Petsc入门基础篇
forrestguang的博客
02-04 4735
欧拉常数的计算 本人之前讲述过在虚拟机环境中安装Petsc，下面都默认在已经有虚拟机的情况运行。 e=∑n=01n!e=\sum_{n=0}\frac{1}{n!}e=∑n=0​n!1​ 下面我们先通过一个简单的代码来介绍一些基础知识： #include <petsc.h> int main(int argc,char **argv) { PetscErrorCode ierr; PetscMPIInt rank; PetscMPIInt i; PetscRea
基于PETSC的线性方程求解
forrestguang的博客
10-27 3318
使用KSP方法求解小型线性系统 需要注意的是，在PETSC中出现的语法和C语言基本相同，区别大部分都只存在于不同的表达方式 比如在C中，int表示整型，在PETSC中，PetscInt表示整型，这里大家百度用法，或者参考我上次的学习报告，我们这里重点介绍在PETSC中一些简单函数的定义调用，并且使用这种函数来求解实际问题（求解线性系统）。 我们慢慢来解释：第一句static char help[] = “Solve a 4x4 linear system using KSP.\n”;是为了对这个代码进行说明
区域分解与并行计算
forrestguang的博客
05-27 2220
区域分解算法 以下都以泊松方程来举例 问题引入 给定二维空间的泊松方程： −Δu=f,x∈Ω-\Delta u = f,x \in \Omega−Δu=f,x∈Ω u=g,x∈∂Ωu = g,x \in \partial \Omegau=g,x∈∂Ω 我们把区域分成若干部分分别求解，然后相邻区域做一个通信，这样就可以得到一个并行算法。下面我们先把区域分成两部分做一个介绍： 考虑区域Ω1,Ω2\Omega_1,\Omega_2Ω1​,Ω2​ 其中红色的是Ω1\Omega_1Ω1​，蓝色的是Ω2\Omega
flash attention的重点解读和CUDA编程实现-V2
forrestguang的博客
09-26 2084
代码22行这里的i=blockIdx.x表示的输出矩阵的第i行，也是输入矩阵Q的第i行，代码23行的phd表示输出矩阵O的第phd列，也是输入矩阵V的第phd列。这两个索引在全过程不参与运算，只参与标记元素的位置。现在还剩下threadIdx.x没有派上用场，我们将利用threadIdx.x尝试遍历矩阵K的每一行，但是BLOCK_DIM_x是固定的，因此我们需要借助一个额外的循环来遍历K，因此phNumN就是总共要遍历的次数，这里一定要向上取整。
MPI的基础编程代码
forrestguang的博客
10-11 1878
这个代码我们来观察，如果rank=0，那么这个0号进程就会把token发送给1号进程，如果rank=1，那么1号进程就会接收来自0号进程的token，同时把token发送回去给0号进程，只有当rank=0的时候才会打印结果。这个代码很简单，但是如果0号进程把token发送出去，1号进程没有接收到token，此时1号进程是否会启动自己的发送命令呢？
linux查看寒武纪显卡nvidia-smi
02-06
### 使用 nvidia-smi 命令检查寒武纪 GPU 状态 `nvidia-smi` 是专为 Nvidia 显卡设计的管理工具，主要用于监控和管理系统中的 Nvidia GPU 设备[^1]。然而，寒武纪（Cambricon）是一家专注于人工智能芯片研发的企业，其生产的 AI 加速器并不兼容 `nvidia-smi` 工具。 对于寒武纪系列的加速卡，官方提供了专门的 SDK 和工具集用于管理和监测设备的状态。具体操作如下： #### 安装寒武纪驱动程序和支持软件包 为了能够正常访问并控制寒武纪硬件，在 Linux 上需先安装对应的驱动以及支持库。通常这些资源可以从寒武纪官方网站获取到最新版本。 #### 查看寒武纪 GPU 信息的方法 由于寒武纪产品线多样，不同型号可能配备不同的诊断工具。一般情况下，可以通过以下方式查询设备详情： - **mluinfo**：这是寒武纪提供的一款命令行工具，用来展示 MLU 卡的相关配置参数。 ```bash mluinfo -v ``` 此命令将会输出有关已安装的寒武纪加速模块的具体规格、运行状况以及其他重要属性。 #### 验证安装成功与否 完成上述步骤之后，可通过简单的测试应用程序验证是否正确识别到了寒武纪设备。这部分内容建议参照随附文档或在线帮助指南来进行。
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包

打赏作者

谨慎付费（看不懂试读博客不要订阅）

你的鼓励将是我创作的最大动力

¥1 ¥2 ¥4 ¥6 ¥10 ¥20
扫码支付：¥1
获取中
扫码支付

您的余额不足，请更换扫码支付或充值

打赏作者

实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

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

余额充值