- 博客(113)
- 资源 (3)
- 收藏
- 关注
原创 大模型中的各种并行:TP DP EP PP
TP并行技术摘要 Tensor Parallel(TP)是一种用于大模型训练和推理的并行技术,主要解决单GPU无法承载大模型计算和显存压力的问题。TP通过矩阵切分将计算分布到多GPU上,与按层切分的Pipeline Parallel不同,TP是在同一层内部切分矩阵计算。 TP主要采用两种切分方式:Column Parallel按输出维度切分权重列,每卡计算部分输出;Row Parallel按输入维度切分权重行,每卡计算部分结果后通过AllReduce求和。典型应用场景包括: 降低单卡显存占用和计算压力 支持
2026-07-04 09:51:24
286
原创 3 层 Decoder-only Transformer(以 Qwen3 为原型 )怎么计算的
阶段操作数据 shape 变化EmbeddingPrefill3 层 transformer[1,256,4096] → [1,256,4096](每层不变)存储 K/V每层: [1,32,256,128]Decode t=0输入 last hidden → 3 层 → logits重复上述 decode 步骤输入始终 [1,1,4096],cache 长度 +1阶段操作数据 shape 变化EmbeddingPrefill3 层 transformer。
2025-11-19 18:27:49
1040
原创 vllm-qwen3.py介绍
│├── lm_head (Linear) → 最后映射到词表│││││推理时实际调用路径(vLLM)如果你有具体的qwen3.py代码(比如来自 vLLM GitHub 的某版本),我可以进一步对照源码做逐行分析。希望这个梳理对你理解 Qwen3 在 vLLM 中的执行逻辑有所帮助!
2025-11-18 15:46:04
1151
原创 对比 DeepSeek(MLA)、Qwen 和 Llama 系列大模型在 Attention 架构/算法层面的核心设计及理解它们的本质区别。
主要大模型Attention架构对比显示:DeepSeek-V2/V3采用自研MLA技术,通过低维隐空间压缩KVCache,显存效率提升5-10倍,适合长文本;Qwen全系列使用GQA分组查询注意力,平衡速度与效果;Llama系列小模型用MHA,大模型(70B+)转向GQA。MLA为DeepSeek独有创新,Qwen和Llama均未采用,三者在显存效率、实现复杂度等方面各具优势。(149字)
2025-11-12 10:01:10
1499
原创 深入解析U-Net
U-Net 是一种“U形”编码-解码网络,靠保留细节。最初用于医学图像分割,现在广泛用于生成模型(如 Stable Diffusion 的核心去噪模块)。在 SD 中,U-Net 接收带噪 latent 和文本提示,预测噪声,反复迭代得到干净 latent。它不是大模型本身,但却是大模型(如 SD)能高效高质量生成的关键组件。参考深入解析U-Net:图像分割的经典架构与应用 - 知乎。
2025-11-06 15:40:45
804
原创 能例介绍ffn的作用吗
Attention 告诉模型“该关注什么”,FFN 告诉模型“这意味着什么”。收集信息(“Paris 和 France 一起出现”)理解信息(“所以 Paris 是法国的首都”)
2025-11-06 09:23:51
397
原创 至目前2025年,深度学习框架、训练框架、推理框架
来自Qwen3 Plus截至 2025 年,通用深度学习框架(General-Purpose Deep Learning Frameworks)是指能够支持多种任务.如需具体部署示例(如:用 vLLM 部署 Qwen2-7B,或用 TGI 启动 LLaMA-3),欢迎继续提问截至 2025 年,大模型(如 LLM、多模态模型)的训练对分布式能力、显存优化、通信效率、易用性提出了极高要求。当前主流的训练框架已从单一深度学习库演变为**“基础框架 + 分布式加速库 + 微调工具链”** 的组合生态。以下是当前主
2025-11-04 16:15:28
1354
原创 真实的大模型中,embedding映射的高维矩阵维度和 attention矩阵运算的规模尺寸?
本文详细分析了主流大语言模型(LLaMA-3、GPT-4等)中关键矩阵的规模及其计算复杂度。重点包括:1)Embedding矩阵可达[128K,8K]规模,占据约1.5%模型参数;2)注意力分数矩阵呈二次方增长,如8K序列长度会产生67M元素/头的庞大矩阵,32头叠加显存占用达4.3GB(fp16);3)KV缓存推理时每层约67MB,32层共2.15GB。文章指出,正是这些高维矩阵催生了FlashAttention等优化技术,通过避免显式存储注意力矩阵来降低显存消耗,使大模型训练成为可能。核心结论:注意力分
2025-11-04 13:51:01
1457
原创 传统卷积神经网络中的核心运算是卷积或者矩阵乘,请问transformer模型架构主要的计算
特性说明主导运算矩阵乘法(GEMM)是绝对主导,占总 FLOPs 的 90% 以上并行性高度并行,适合 GPU/TPU 加速内存瓶颈注意力矩阵(Q @ K.T)和 KV Cache 是显存瓶颈计算密度FFN 通常比 Attention 更“重”(更多 FLOPs)优化技术FlashAttention、PagedAttention、量化、稀疏化等用于提升效率🔹Transformer 的主要计算是:自注意力机制中的矩阵乘法Q@K.T, attn@V)前馈网络中的大规模矩阵乘法(W_1, W_2)🔹。
2025-10-31 10:49:25
825
原创 cuda 编译和运行操作背后发生了什么?
CUDA编译执行流程摘要(143字): 在CUDA编程中,编译器(nvcc)不直接调用Driver API。编译时,设备代码被编译为PTX中间码或直接生成SASS机器码。运行时存在两种执行路径:1)JIT编译路径:驱动程序将PTX动态编译为当前GPU架构的SASS;2)AOT路径:直接加载预编译的SASS。Driver API由Runtime API在运行时调用,负责加载模块/启动内核。PTX到SASS的转换由NVIDIA驱动完成,通过cuModuleLoad触发JIT编译,最终SM单元执行SASS指令。这
2025-10-20 16:10:34
869
原创 bash脚本 CMakeLists.txt Makefile学习记录
现在的软件构建和测试,基本依赖如下几种类型,这些应该是熟练的基本功,之前不重视的技能,陆陆续续积累起来。每个部分常用的接口符合二八定律,真正使用的没有多少。
2025-06-25 15:04:36
239
原创 CUDA 实践:隐式 GEMM 卷积 | CUDA
隐式 GEMM 卷积则可以直接从原始 feature map 和 weight 中取值,避免产生占用巨大内存/显存的中间结果矩阵
2025-05-30 09:55:52
438
原创 Attention-> flashAttention材料参考
一文看懂 Attention(本质原理+3大优点+5大类型)_attention结构-优快云博客22。另外,在语音识别、目标物体检测等领域,注意力机制同样取得了很好的效果。
2025-05-30 09:52:13
370
转载 cute swizzle
作者:水木皇工仔链接:https://www.zhihu.com/question/600927104/answer/4791506558来源:知乎著作权归作者所有。商业转载请联系作者获得授权,非商业转载请注明出处。 这篇文章已经讲得非常清楚,跟着里面的 demo 例程敲完,就基本可以理解了。本质就是上图,在 shared memory 中 每个 thread 会 access 到 8 个 2 bytes (fp16)的数据,然后去到 registers 中,warp 之间会进行分配,使得每个 th
2025-03-06 09:52:09
322
转载 cute 之 Layout
计算机中的内存是一维的线性地址空间,而数学计算问题所要处理的空间经常是高维的。如GEMM(General Matrix Multiplication)问题的数学计算体系是二维计算空间,Deep Learning计算体系是三维以上的计算空间(batch, height, width, channel, etc.)。如何高效的表达高维计算空间,如何高效便捷的将计算所要求的高维空间映射到一维空间变得越来越重要。
2025-03-05 10:12:49
439
转载 cute 之 Swizzle
本文介绍了shared memory的bank结构和特征,以及通过行列异或实现bank的交错的方法。继而介绍了cute中Swizzle的抽象,最后介绍了优化L2命中率的Thread Block Swizzle方法。至此矩阵乘法优化的理论性部分已经全部介绍完毕,接下来的文章我们会基于这些方法完成高效的矩阵乘法。
2025-03-04 11:23:14
1038
转载 cute 之 Copy抽象
cute提供了Copy能力来完成实现D = STensor的搬运,其通过对指令、适配层、原子能力、块状能力和线程级别的抽象分别形成了数据结构Copy_AtomTiledCopyThrCopy和cute::copy,在这些抽象的帮助下,我们可以实现逻辑Tensor由一个存储单元向另一个存储单元的拷贝,而不用过多的关注指令细节。这种能力和MMA一道共同构成cute下矩阵乘法的基础。在Copy和MMA抽象的帮助下我们可以在逻辑层次构造数据的搬运D = S和矩阵乘法。
2024-12-17 08:41:20
664
转载 理解 Tensor Core
在Tensor Core的MACC操作中,矩阵A的fragment包括8个FP16*2元素(即16个FP16元素),矩阵B的fragment包括另一8个FP16*2元素,以及针对FP16累加的4个FP16*2元素,或针对FP32累加的8个FP32元素。更具体地说,当threadgroup 0执行Set 0的HMMA指令(如下所示)时,将包含矩阵A的前4行和前4列的子块与包含矩阵B的前4行和前8列的子块相乘,其乘积与矩阵C的4×8子块累加,得到的和保存在矩阵D的4×8子块中,即上图第一行计算过程。
2024-12-03 12:41:08
860
原创 TMA swizzing
Swizzling 通过调整数据的存储布局来优化内存访问,避免了多个线程同时访问同一个内存bank。具体而言,它通过交错访问模式使得线程间的内存访问更加均匀,提升了并行计算的效率。在 CUDA 编程中,利用 Swizzling 技术可以显著减少bank冲突,提升共享内存的带宽利用率,从而提高整体计算性能。
2024-12-02 14:12:37
1622
转载 CUDA 编程入门(2):CUDA 调度模型
Block 对应的物理硬件概念是 SM,也就是说 SM 负责 block 中线程的执行,SM 会为每个 block 分配需求的资源,比如寄存器,共享内存等,由于 SM 自身资源有限,因此它被分配到的 block 数量也是有限的,这取决于 block 中线程的资源需求。每个 SM 中有 4 个处理块,每个处理块都有一个独立的 Warp Scheduler,Warp Scheduler 负责调度分配到当前处理块上的 warp,所有这些 warp 构成一个叫做 Active Warps 的集合,调度器在每个。
2024-09-09 16:47:31
399
原创 理一理Latency、Bandwidth、Throughput、Response Time概念的区别
在性能测试术语中,“在给定的时间段内从一个位置成功移动到另一个位置的数据量,通常以每秒比特数(bps)来度量,或每秒兆比特数(Mbps)、每秒千比特数(Gbps)”。带宽(Bandwidth):它显示了管道(通信通道)的容量。假设ISDN的带宽是64K,则我们可以增加一个64K的通道,所以总带宽是128K。经常,一些性能测试人员对延时(Latency)、吞吐(Throughput)、带宽(Bandwidth)和响应时间(Response Time)感到迷惑,今天,就用一些简单的例子来说明它们之间的区别。
2024-09-09 16:45:42
1939
转载 GPU 峰值算力及性能优化参考资料汇总
另外,这里用的是F32浮点峰值做例子,如果你的任务不需要浮点运算或是精度不是F32,这个值就意义不大,需要转换成你需要的那个操作。其实不看Tensor core的话,满血版一般有:F64:F32:F16=1:2:4,正好与占用的GPR成反比,这个其实是与GPR的带宽有很大的关联的,一般满血版的卡的功能单元配比就会尽量按极限的GPR带宽来设计。一般说来,实际应用中,较大尺寸的矩阵乘法(GEMM)是难得的能接近峰值性能的程序,有些实现能到98%峰值的效率。指令,这是很苛刻的条件。都是一个指令一个flop,
2024-09-09 13:44:12
1391
原创 simd vs simt
SIMD 架构是指在同一时间内对多个数据执行相同的操作,适用于向量化运算。例如,对于一个包含多个元素的数组,SIMD 架构可以同时对所有元素执行相同的操作,从而提高计算效率。在实际应用中,根据具体的计算需求和硬件环境选择合适的架构可以提高计算性能。SIMT 架构是指在同一时间内执行多个线程,每个线程可以执行不同的指令,但是这些线程通常会执行相同的程序。版权声明:本文为博主原创文章,遵循 CC 4.0 BY-SA 版权协议,转载请附上原文出处链接和本声明。常见的 SIMD 架构包括。
2024-09-09 13:42:02
1193
转载 CUDA编程:warp
因为一个warp中的线程必须执行相同的指令,不能一部分执行if中的指令,一部分执行else中的。由两个线程块组成的线程结构。,这些资源都被thread瓜分了,由于资源是有限的,所以,如果thread数量比较多,那么每个thread占用资源就比较少,反之如果thread数量较少,每个thread占用资源就较多,这需要根据自己的需求作出一个平衡。SM中的warp调度器会选择warp执行,被选中的成为selected,没选中但已经具备执行条件的成为eligible,没准备好的成为stalled。
2024-09-09 11:38:59
813
转载 CUDA编程方法论
GPGPU中一些问题的理解与思考(3)- 指令执行吞吐与指令集设计 gpu中的warp的divergence问题? 写CUDA到底难在哪? GPGPU中一些问题的理解与思考(2)- 各级存储的吞吐问题 NV GPU(SM50)的Register Bank Conflict和Reuse Cache的一些测试 GPGPU中一些问题的理解与思考(1)-上层架构 发布一个CUDA SASS
2024-09-09 11:24:51
278
转载 谈谈 cuda 线程束与内存模型
参考:1、2、Cuda thread 组成 block,block 组成 grid,形成嵌套网格状结构,目的是为了矩阵并行计算时,能对应矩阵位置。
2024-09-09 11:13:05
628
原创 Thread如何划分为Warp?
Thread Index和Thread ID之间有什么关系呢?(线程架构参考这里:CUDA C++ Programming Guide (nvidia.com)open in new window)1维的Thread Index,其Thread ID就是Thread Index2维的Thread Index,其Thread ID为3维的Thread Index,其Thread ID为由此再回到本文的问题:Thread如何划分为Warp?线程的组织形式对程序的性能影响是至关重要的,本篇博文主要以下面一种情况来
2024-09-08 23:21:18
1367
转载 CUDA编程4-执行模型
执行模型会提供一个操作视图,说明如何在特定的计算架构上执行指令。CUDA执行模型解释了GPU并行架构的抽象视图,有助于编写指令吞吐量和内存访问高效的代码。GPU架构围绕一个流式多处理器(SM)的可扩展阵列搭建,通过复制这种架构的构件来实现硬件并行。每个GPU通常有多个SM,每个SM都能支持数百个线程并发执行。当启动一个内核网格时,它的线程被分布在可用的SM上执行。线程块一旦被调度到一个SM上,其中的线程只会在那个指定的SM上并发执行;多个线程块可能被分配到同一个SM上,且根据SM资源的可用性进行调度;
2024-09-08 18:11:25
752
转载 GPU软件抽象:Grid,Block,Thread,Warp定义说明与硬件的映射执行细节 学生教室楼层的例子
一个Warp中的线程必然在同一个block中,如果block所含线程数目不是Warp大小的整数倍,那么多出的那些thread所在的Warp中,会剩余一些inactive的thread,也就是说,即使凑不够Warp整数倍的thread,硬件也会为Warp凑足,只不过那些thread是inactive状态,需要注意的是,即使这部分thread是inactive的,也会消耗SM资源。在执行功能上是没有问题的,那些不应该被执行分支的线程会被禁用,因此始终获得正确的结果,但是如果双方都很长,那么性能损失就很重要。
2024-09-08 18:07:44
752
转载 NVIDIA GPU以SIMT
NVIDIA GPU以SIMT(Single Instruction, Multiple Thread)的方式执行称为warp。许多 CUDA 程序通过利用 warp 执行来获得高性能。在这个博客中,我们将展示如何使用 CUDA 中引入的原语,使您的 warp 级编程安全有效。Warp-level 原语NVIDIA GPUs 和 CUDA 编程模型采用一种称为 SIMT(Single Instruction, Multiple Thread)的执行模型。SIMT 扩展了计算机体系结构的弗林分类学。
2024-09-08 15:36:40
829
转载 CUDA02 - 访存优化和Unified Memory
CUDA02 - 访存优化和Unified Memory - 猫猫子 - 博客园 (cnblogs.com)CUDA02 - 的内存调度与优化前面一篇(传送门)简单介绍了CUDA的底层架构和一些线程调度方面的问题,但这只是整个CUDA的第一步,下一个问题在于数据的访存:包括数据以何种形式在CPU/GPU之间进行通信、迁移,以及在GPU内部进行存储、访问。1 global 、shared 、constant、local通常来讲,待计算的数据都存放在内存或者硬盘(外部存储设备)中,由CPU
2024-09-07 23:13:25
1430
转载 CUDA01 - 硬件架构、warp调度、指令流水线和cuda并发流
超级棒这一部分打算从头记录一下CUDA的编程方法和一些物理架构上的特点;从硬件入手,写一下包括线程束的划分、流水线的调度等等微结构的问题,以及这些物理设备是如何与软件对应的。下一部分会写一下cuda中的几种内存划分,进行数据同步,以及优化cuda运行效率的几种方法。
2024-09-07 23:01:46
1287
转载 Warp Scheduler,合理块和线程的数量对GPU利用率非常重要
GPU算子GPU利用率的因素包括,算子线程数量设置合不合理、访存模式合不合理、数据传输次数是否过多、硬件特性是否被充分利用
2024-09-07 17:24:19
764
转载 英伟达GPU的硬件架构——显存带宽和算力
CUDA 是建立在英伟达GPU硬件架构之上的平台和编程模型。因此,就像要写好高性能的 CPU 程序,必须对 CPU 的硬件架构有一定的理解一样,我们写 CUDA 程序也必须对 GPU 有最基本的了解。
2024-09-07 16:02:24
4638
转载 CUDA 调度模型
Block 对应的物理硬件概念是 SM,也就是说 SM 负责 block 中线程的执行,SM 会为每个 block 分配需求的资源,比如寄存器,共享内存等,由于 SM 自身资源有限,因此它被分配到的 block 数量也是有限的,这取决于 block 中线程的资源需求。每个 SM 中有 4 个处理块,每个处理块都有一个独立的 Warp Scheduler,Warp Scheduler 负责调度分配到当前处理块上的 warp,所有这些 warp 构成一个叫做 Active Warps 的集合,调度器在每个。
2024-09-06 20:12:14
520
转载 CUDA 编程入门_CUDA 编程模型
1CUDA 的全称是 Compute Unified Device Architecture,它是建立在硬件架构之上的并行计算平台和编程模型。
2024-09-06 12:10:22
359
空空如也
TA创建的收藏夹 TA关注的收藏夹
TA关注的人
RSS订阅