- 博客(153)
- 收藏
- 关注
RSS订阅
原创 CMake实现跨平台高性能算子测试框架
之前使用pybind11作为转换器，将kernel编译为.so文件遇到过一个问题，那就是.cu文件的调用函数接口接受的参数是类似于torch::Tensor或者是numpy，这个编译过程比较慢，而且对于cuda代码来说，python提供了CUDAExtension的接口作为转换，但是对于CPU以及国产芯片比如说MLU就不知道如何替换，为此这里我们使用CMake来实现不同平台的kernel编译。代码架构如下所示：其中src文件夹存放的是不同算子的源文件，包括不同平台的kernel，test文件夹存放的是不
2024-11-13 16:35:13
754
原创 PDE的数值解法（有限元，有限差分法）综合介绍
以下内容均可参考本人知乎文章添加链接描述和添加链接描述有限差分法finite difference(FD)是求解微分方程的最为容易理解的方法，下面将针对几类常见的PDE来做一些具体的介绍。由于本人知识有限，关于误差分析和收敛性证明都不会介绍.一维例子我们以一个一维PDE的求解来介绍有限差分算法，这里给出精确解u^=x(1−x)exp(x)\hat{u} = x(1-x)\exp(x)u^=x(1−x)exp(x).{−uxx=f,x∈(0,1),f=exp(x)(−x2−x+1),u(0)=0,u(
2022-11-27 20:26:00
13071
2
原创 Linux安装anaconda和pytorch
anaconda安装这里我是使用课题组服务器，在服务器的/home/账户里安装anaconda。本人的安装环境：本人在登录节点安装，最终的结果送到计算 节点运行。1：首先要从网上下载anaconda.sh，这里大家登录anaconda官网寻找对应的版本链接：https://pan.baidu.com/s/1D0ooJFHvmz5ZIzqkIoXt2w提取码：hz3s上面是本人使用的sh文件，网盘自己取，这个生成的python版本是3.9，但是接下来我们要使用的pytorch对应的是3.8，读者自
2021-12-25 23:32:18
2870
原创 read me
关于该博客一开始是在2018年本人大三的时候，开始接触一些数值代数，微分方程数值解和学习matlab，那时候本人经常上网搜各种代码和算法，放到大作业里面，那时候经常接触优快云博客，发现很多人的博客对于代码讳莫如深，不肯把自己完整的代码放出来供读者下载使用学习，当然了，这样有助于读者自己学习摸索，只不过这个过程有点太复杂，不同的人学习代码有不同的根基和领悟能力，像博主这样的人偏偏就是领悟能力差的，所以对这种行为非常不喜大三保研结束以后，人生轨迹有了新的方向，本来一直打算进军数学方向，纯粹是因为那时候对代
2021-12-15 11:34:45
2253
1
原创 github合并多个commit message以及rebase解决文件冲突
为了安全起见，我不修改任何提交信息，直接使用Esc，以及:wq保存并且退出，打印出下面这个界面，下面这个git rebase --continue是上一步的，不需要额外再用。issue/18从main分支切出来，但是main分支被别人做了修改，导致issue/18目前和main存在文件冲突，现在使用rebase来解决冲突。比如说现在我要合并issue/108分支的提交记录，使用git log --oneline查看提交记录一共有6次，千万不要理会，千万不要使用git pull，直接强制推送即可，
2025-03-31 17:44:37
542
原创 寒武纪使用cnnl库函数实现卷积算子
完整代码参考寒武纪cnnl实现卷积的库函数说明参考cnnl调用卷积的函数这里我们使用的是cnnlConvolutionForward，这个函数的使用只是最后的计算过程，但是在计算之前，还有许多准备工作。下图展示了一个完整的cnnl调库实现卷积的全过程，事实上，cnnl调用其他库函数的过程也于此类似：1：申明一个descriptor，比如说这里申明的是卷积的描述器convDesc2：申明好相关的descriptor以后，调用create函数创建这个描述器。
2025-01-21 17:09:06
1210
原创 gather算子的CUDA编程和算子测试
知乎介绍参考添加链接描述完整测试框架参考本人仓库添加链接描述gather算子的onnx定义参考添加链接描述，该算子的主要变换参考下图：这里我们不妨以input = [A, dimsize, D]， indices = [B,C], axis = 1举例子，此时对应的output形状是[A,B,C,D]，并且根据gather算子定义，我们知道output[i,j,k,s] = input[i,indices[j,k], s]，下面我们引入一些定义：outputSize = ABCD, inputS
2025-01-19 22:38:19
117
原创 CUDA算子测试框架搭建-pybind11的使用
这里需要注意环境配置，之前使用本地计算机搭建linux服务器的时候使用的是WSL+Ubuntu，许多人使用默认选项安装的nvcc版本可能过低，比如说11.5，但是目前pytorch官网提供的针对CUDA的版本最低都是11.8，因此在安装nvcc的时候，也要升级到对应11.8以上才可以。nvcc11.8的安装可以参考链接添加链接描述根据官网提示，使用下面的命令可以把相关的文件下载好，wget https://developer.download.nvidia.com/compute/cuda/rep
2024-11-11 16:16:55
92
原创 pytorch调用手写CUDA算子和pybind11的使用
这段代码的主要作用是从cpu_input和cpu_output这两个py::array_t类型的NumPy数组中获取底层的浮点数数据指针，以便在C++代码中直接操作这些数据。通过request()方法获取数组的缓冲区信息，然后使用.ptr成员获取指向数据的指针，并通过类型转换确保指针类型与数组元素类型匹配。这为后续的数据处理（尽管在这段代码中没有展示）提供了基础。'softmaxCuda', # 扩展模块名称，后面import softmaxCuda['softmax.cu'], # CUDA源文件。
2024-10-23 21:15:05
487
原创 tensor core实现矩阵乘法的详细解读
之前关于tensor core的介绍可以参考链接基础的tensor core实现C=AB的代码可以参考下面这段内容：上面代码的几个注意事项：首先是加载mma.h头文件，这个是包含wmma模板类的头文件。其次是设置的WMMA_M=16,WMMA_N=16,WMMA_K=8，这三个参数的表示的意思是，对于一个线程块内的一个warp来说，这个线程簇warp一次能处理的是[16,8]@[8,16]这样小矩阵乘法。
2024-08-28 21:59:24
500
原创 大模型推理常见采样策略：Top-k, Top-p, Temperature以及rotary embedding的寒武纪编程
这个random sample本质上是一种采样策略，可以参考链接https://blog.youkuaiyun.com/qq_43243579/article/details/136331123给定input = {source, topp, topk, temperature}，输出output = {index}source是一个一维的长度为voc的向量，数据类型为float或者halftopp是一个位于(0,1)之间的float类型的常数topk是一个正整数。
2024-08-22 17:08:36
497
原创 CUDA实现矩阵乘法的性能优化策略
本人主要参考了https://zhuanlan.zhihu.com/p/435908830，https://zhuanlan.zhihu.com/p/410278370，https://zhuanlan.zhihu.com/p/518857175 ，下面的代码均是本人实现矩阵乘法的easy实现-V1C=AB,A∈RM×K,B∈RK×NC=AB,A∈R^{M×K},B∈R^{K×N}C=AB,A∈RM×K,B∈RK×Nci,j=∑0K−1ai,kbk,jc_{i,j}=\sum_{0}^{K-1} a_
2024-07-29 21:27:00
324
原创 flash attention的CUDA编程流水并行加速-V6
之前关于flash attention的介绍可以继续参考链接添加链接描述矩阵乘法的优化参考添加链接描述，我们发现矩阵乘法的最优配置为：BLOCK_DIM_x=BLOCK_DIM_y=16，同时每个线程处理一个8×8的子矩阵。线程网格设置如下所示：const int Rq = 8;const int Rv = 8; // 必须是4的倍数const int Br = 16;const int Bc = 16;const int Bk = 8; // 必须是4的倍数const int Bd =
2024-07-24 17:17:48
292
原创 寒武纪实现高维向量的softmax进阶优化和库函数对比
关于寒武纪编程可以参考本人之前的文章实验证明，axis=0和axis=-1的时候，手写softmax速度可以和库函数媲美，甚至于更甚一筹。
2024-07-01 16:05:39
191
原创 寒武纪显卡实现rmsNorm算子
rmsNorm算子的数学定义下面这个是针对一维向量的rmsNorm算子，对于高维向量，比如说形状为[A,B,C,D]的4维向量x，此时权重向量w长度默认为D，也就是说变换主要是针对axis=-1这个维度进行操作。关于寒武纪显卡的编程可以参考本人之前编写的文章添加链接描述和添加链接描述rmsNorm算子的实现思路由于寒武纪编程比较特殊，这里我们稍微复习一下寒武纪编程的特点：一张寒武纪显卡往往可以有8个或者4个cluster，其中一个cluster包含4个core，也就是说，对于4个cluster的显卡
2024-06-05 15:46:10
475
原创 flash attention的CUDA编程（一个线程处理多个元素）-V5
本文接着之前的文章介绍添加链接描述单线程处理多元素之前文章提到的算法都有一个特点，即一个(threadIdx.x,threadIdx.y)只计算输出O矩阵的一个元素，在内部循环计算matmul的时候也是如此，循环体内部的两个矩阵乘法也严格遵循了一个(threadIdx.x,threadIdx.y)计算输出矩阵的一个元素，但是在里面的matmul里面，为了计算一个输出，需要访问A,B矩阵的一行一列，这样就导致计算访存比非常小，为了加速内部的矩阵乘法，我们希望能够做出修改，一个线程同时计算多个元素，具体操作
2024-05-18 20:31:38
358
1
原创 flash attention的CUDA实现探讨-V4
之前关于flash attention的实现参考lash attention的数学变换：给定三个矩阵Q,K,V，形状都是[N,d]，计算S=QK.T，然后针对dim=1做softmax，然后和V继续做矩阵乘法得到形状为[N,d]的输出矩阵O，即O=softmax(QK.T,dim=1)V。下面本人的实现以下图为基础，从下图也可以看出attention算法的关键在于对Q,K,V的分块，因此我们需要首先考虑清楚编程过程中一个线程块到底处理哪些数据。
2024-05-13 17:31:02
511
原创 C++搭建深度学习的推理框架
回顾pytorch里面搭建的全连接神经网络，我们以最简单的全连接神经网络举例子，这种神经网络涉及向量的传输，矩阵乘法计算，以及激活函数定义，其中的矩阵计算以及激活函数就可以理解为神经网络里面的算子。之前在上一篇文章添加链接描述我们介绍了CUDA版本的矩阵乘法，这里我们不妨以OpenMP为工具搭建一个CPU版本的矩阵乘法算子。
2024-04-07 17:06:56
418
原创 CMakeLists.txt编写简单介绍：CMakeLists.txt同时编译.cpp和.cu以及寒武纪kernel
关于CMakeLists.txt的相关介绍，这里不赘诉，本人的出发点是借助于CMakeLists.txt掌握基本的C++构建项目流程，下面是本人根据网络资料以及个人实践掌握的资料。
2024-04-02 18:16:20
1102
原创 NCCL实现分布式矩阵乘法的CUDA代码
我们以矩阵乘法C=AB，其中A形状为[M,K],B形状为[K,N]，C形状为[M,N]举例子，下面的分布式算法我们默认以MPI来切分数据，其中每个进程之前的数据都是私有，进程之间的数据交互使用通信来完成。A=(a0,0a0,1a0,2⋯a0,K−1a1,0a1,1a1,2⋯a1,K−1⋮⋮⋱⋮aM−2,0⋯aM−2,K−3aM−2,K−2aM−2,K−1aM−1,0aM−1,1⋯aM−1,K−2aM−1,K−1)\begin{equation} A=\left(\begin{array}{ccccc} a
2024-03-28 21:57:52
228
原创 寒武纪显卡实现高维向量的softmax并行优化
关于寒武纪编程可以参考本人之前的文章添加链接描述，添加链接描述，添加链接描述高维向量的softmax实现更加复杂，回忆之前在英伟达平台上实现高维向量的softmax函数，比如说我们以形状为[1,2,3,4,5,6]的6维向量举例，变换维度假设axis=2，之前英伟达平台的实现，我们计算出变换维度的长度dimsize=3，其他维度的乘积othersize=1×2×4×5×6 = 240，步长stride= 1×6×5×4 = 120，使用othersize=240个线程块，其中每个线程块处理对应一份数据，计算
2024-02-06 15:35:24
342
原创 寒武纪显卡实现softmax的pingpong流水并行
在上一篇文章中我们介绍了寒武纪显卡实现基本的softmax代码，这里我们借助于寒武纪的流水并行来编写进一步的策略。
2024-01-17 17:35:08
379
原创 寒武纪显卡实现softmax算子
寒武纪实现softmax包括下面5个步骤，我们也采取5个kernel来实现softmax：unionMaxKernel(float* middle, float* source1, int num)，这个kernel使用的任务类型是union1，其中middle的长度为taskDim，达到的目的是middle对应元素存储的是不同taskId处理的那部分数据的局部max。
2024-01-10 17:49:23
780
原创 寒武纪bang的基础向量除法和规约编程
Queue是用来管理并行操作的一种方式。在同一个Queue中所下发的操作顺序是依次按照下发顺序执行的，在不同Queue中的执行操作则是乱序的。Queue所支持的操作不仅包括Kernel（在设备上的执行程序），还包括Notifier、内存拷贝、设备与主机的数据传输等一系列操作，所有需要按序执行的操作都可以下发至同一个Queue，Queue在当前操作满足执行条件后会依次按照进入顺序执行操作；而不同Queue的操作则不保证顺序，这些操作的顺序是无法预测的。
2024-01-10 16:54:01
839
原创 flash attention的CUDA编程进阶-V3
前面博客的实现逻辑是：以blockIdx.x作为矩阵Q的第i行输入，blockIdx.y × blockDim.y + threadIdx.y作为矩阵V的第j行输入，中间遍历矩阵K,V时使用for循环串行遍历。在for循环里面相当于每次只能获得QK.T的一个元素，我们的优化策略是以threadIdx.y来加速QK.T的矩阵乘法。这种做法的缺点是：如果N很大，d很小，那么会造成大量的线程浪费闲置，尤其是当d=1的时候，此时算法接近串行。
2023-11-19 10:44:13
789
原创 cublas,tensorcore矩阵乘法基本介绍
CαopAopBβCCαopAopB))βC，这里的运算符op决定是否转置。这里首先说明，cublasSgemm专门处理float类型的二维矩阵乘法，针对其他类型的矩阵，cublas有类似的函数接口cublasDgemm：处理double类型矩阵cublasCgemm：处理cuComplex类型矩阵cublasZgemm：处理cuDoubleComplex类型矩阵cublasHgemm：处理__half类型的矩阵。
2023-11-14 16:31:24
655
原创 softmax的高效CUDA编程和oneflow实现初步解析
本文参考了，其中oneflow实现softmax的CUDA编程源代码参考链接关于softmax的解读以及CUDA代码实现可以参考本人之前编写的几篇文章。
2023-10-31 09:19:27
619
原创 flash attention的重点解读和CUDA编程实现-V2
代码22行这里的i=blockIdx.x表示的输出矩阵的第i行，也是输入矩阵Q的第i行，代码23行的phd表示输出矩阵O的第phd列，也是输入矩阵V的第phd列。这两个索引在全过程不参与运算，只参与标记元素的位置。现在还剩下threadIdx.x没有派上用场，我们将利用threadIdx.x尝试遍历矩阵K的每一行，但是BLOCK_DIM_x是固定的，因此我们需要借助一个额外的循环来遍历K，因此phNumN就是总共要遍历的次数，这里一定要向上取整。
2023-09-26 16:28:52
2002
原创 flash attention的CUDA编程和二维线程块实现softmax-V1
本文参考了链接添加链接描述flash attention介绍flash attention的介绍可以参考论文：FlashAttention: Fast and Memory-Efficient Exact Attentionwith IO-Awareness，具体的数学公式参考下面这个图片：其中注意关于矩阵S有两个维度，softmax的操作维度是dim=1，用pytorch表示就是torch.softmax(S, dim=1)对于flash attention来说，里面有两次矩阵乘法，对于这样的二
2023-09-20 18:17:45
793
2
原创 softmax的cuda编程详细解读——算子融合
在上一篇博客我们介绍了softmax的内容以及相关的编程实现，总结一下softmax的特点如下：（没有特殊说明的情况下，考虑的仍然只是1D向量x1：获得向量的全局最大值M2：针对向量xexpxi−Mi求和得到S3：输出向量yxS从常理来看，我们至少需要发动3次kernel才能获得数值结果，但是这样非常耗时，有没有什么办法可以发动一次kernel就获得结果呢，下面不妨看看我们的分析过程。
2023-09-03 21:16:45
1270
原创 expand，where和softmax算子的cuda编程
expand和where介绍当谈到 Torch 中的 expand 函数时，我们实际上是指 PyTorch（Torch 的 Python 接口）中的 expand 方法。下面是对 expand 方法和 where 函数的介绍，包括它们的输入和输出：expand 方法：torch.Tensor.expand() 是 PyTorch 中 Tensor 类的一个方法，用于扩展张量的维度。输入：input 是要扩展的张量，size 是一个元组，指定了要扩展的每个维度的大小。输出：返回一个新的张量，形状是
2023-08-31 20:10:46
652
原创 数学系硕士研究生的科研过程——PDE约束下含参优化控制问题的深度学习算法
笔者今天上午收到了之前北大课题组老板的通知，得知研究生期间和学长合作的论文终于被siam接收，终于为自己研究生涯画上了一个句号。这里打算分享一下个人的科研过程以及这篇论文的工作，即将读研或者打算读研的同学或许可以从中获得益处。论文今天被接收，距离见刊尚有一段时间，因此这里只能提供arxiv上的论文链接。
2023-08-28 18:01:13
1012
8
原创 多重网格算法的cuda编程
这里写自定义目录标题多重网格算法介绍问题描述——五点差分法求解二维泊松方程五点差分法Gauss迭代算法限制算子介绍提升算子二重网格算法多重网格算法多重网格cuda代码编写串行代码mg.c两重网格cuda并行代码jacobi迭代的cuda编程device_jacobiMakefilecuda_mg.cucuda_mg.hmain.c多重网格算法介绍问题描述——五点差分法求解二维泊松方程{−Δu=f, in Ω,u=g, on ∂Ω,\left\{\begin{al
2023-08-08 16:58:42
526
原创 torch.nn.DataParallel实现深度学习求解PDE数据并行
泊松方程{−Δu=f, in Ω,u=g, on ∂Ω,\left\{\begin{aligned} -\Delta u &=f, \text { in } \Omega, \\ u &= g ,\text { on }\partial \Omega,\\ \end{aligned}\right.{−Δuu=f, in Ω,=g, on ∂Ω,这里我们选取
2023-08-02 13:31:47
754
原创 深度学习求解稀疏最优控制问题的并行化算法
稀疏最优控制问题问题改编自论文An FE-Inexact Heterogeneous ADMM for Elliptic Optimal Control Problems with L1-Control Cost{miny(μ),u(μ)J(y(μ),u(μ))=12∥y(μ)−yd(μ)∥L2(Ω)2+α2∥u(μ)∥L2(Ω)2+μ∥u(μ)∥L1, subject to {−Δy(μ)=u(μ)+f(μ), in Ω(μ),y(μ)=0,&nbs
2023-05-29 15:03:57
835
原创 mpi4py结合pytorch求解稀疏椭圆优化控制问题
稀疏椭圆优化控制问题{miny(μ),u(μ)J(y(μ),u(μ);μ):=12∥y(μ)−yd∥L2(Ω)2+α2∥u(μ)∥L2(Ω)2+μ∥u(μ)∥L1(Ω), subject to {−Δy(μ)+y(μ)3=u(μ) in Ω,y(μ)=0 on ∂Ω,andua≤u(μ)≤ub a.e. in Ω.\left\{\begin{aligned} &\min_{y(\bold
2023-05-29 15:03:29
1220
原创 pytorch实现导热率优化控制问题的并行算法
问题来源于论文Reduced Basis Methods—An Application to Variational Discretization of Parametrized Elliptic Optimal Control Problems问题描述{miny(μ),u(μ)J(y(μ),u(μ))=12∥y(μ)−yd∥L2(Ω)2+α2∥u(μ)∥L2(Ω)2, subject to {−∇⋅(a(μ)⋅∇y(μ))=u(μ), in Ω,y(
2023-05-29 15:01:42
838
原创 mpi4py结合pytorch求解多重参数最优控制问题
多重参数最优控制问题问题改编自论文Fast PDE-constrained optimization via self-supervised operator learningmin12∥y(μ)−yd(μ)∥L2(Ω)2+α(μ)2∥u(μ)∥L2(Ω)2+β(μ)∥u(μ)∥L1\min \frac{1}{2}\left\|y(\boldsymbol{\mu})- y_{d}(\boldsymbol{\mu})\right\|_{L_{2}\left(\Omega\right)}^{2}+\fra
2023-05-29 15:01:08
759
原创 mpi4py和pytorch求解含参优化控制问题-几何参数
含参优化控制问题数学模型下面三行代码分别表示：代码运行,–para 2表示选择第二种并行策略查看运行情况杀死命令nohup mpiexec -n 8 python hpann.py --para 2 >> hmpi.log 2>&1 &ps -ef | grep pythonps -ef | grep python | awk '{print $2}' | xargs -n 1 kill -9OCP(μ):{min(y(x,μ),u(x,μ))∈Y×UQ
2023-05-29 15:00:20
702
空空如也
空空如也
TA创建的收藏夹 TA关注的收藏夹
TA关注的人