四哥的专栏

转载 cute swizzle

作者：水木皇工仔链接：https://www.zhihu.com/question/600927104/answer/4791506558来源：知乎著作权归作者所有。商业转载请联系作者获得授权，非商业转载请注明出处。 这篇文章已经讲得非常清楚，跟着里面的 demo 例程敲完，就基本可以理解了。本质就是上图，在 shared memory 中 每个 thread 会 access 到 8 个 2 bytes （fp16）的数据，然后去到 registers 中，warp 之间会进行分配，使得每个 th

转载 cute 之 Layout

计算机中的内存是一维的线性地址空间，而数学计算问题所要处理的空间经常是高维的。如GEMM（General Matrix Multiplication）问题的数学计算体系是二维计算空间，Deep Learning计算体系是三维以上的计算空间(batch, height, width, channel, etc.)。如何高效的表达高维计算空间，如何高效便捷的将计算所要求的高维空间映射到一维空间变得越来越重要。​

转载 cute 之 Swizzle

本文介绍了shared memory的bank结构和特征，以及通过行列异或实现bank的交错的方法。继而介绍了cute中Swizzle的抽象，最后介绍了优化L2命中率的Thread Block Swizzle方法。至此矩阵乘法优化的理论性部分已经全部介绍完毕，接下来的文章我们会基于这些方法完成高效的矩阵乘法。

转载 cute 之 Copy抽象

cute提供了Copy能力来完成实现D = STensor的搬运，其通过对指令、适配层、原子能力、块状能力和线程级别的抽象分别形成了数据结构Copy_AtomTiledCopyThrCopy和cute::copy,在这些抽象的帮助下，我们可以实现逻辑Tensor由一个存储单元向另一个存储单元的拷贝，而不用过多的关注指令细节。这种能力和MMA一道共同构成cute下矩阵乘法的基础。在Copy和MMA抽象的帮助下我们可以在逻辑层次构造数据的搬运D = S和矩阵乘法。

转载 理解 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子块中，即上图第一行计算过程。

原创 TMA swizzing

Swizzling 通过调整数据的存储布局来优化内存访问，避免了多个线程同时访问同一个内存bank。具体而言，它通过交错访问模式使得线程间的内存访问更加均匀，提升了并行计算的效率。在 CUDA 编程中，利用 Swizzling 技术可以显著减少bank冲突，提升共享内存的带宽利用率，从而提高整体计算性能。

原创 CUDA 参考文章

CUDA：NVCC编译过程和兼容性详解_nvcc把cuda代码转换成什么-CSDN博客。

转载 CUDA 编程入门（2）：CUDA 调度模型

Block 对应的物理硬件概念是 SM，也就是说 SM 负责 block 中线程的执行，SM 会为每个 block 分配需求的资源，比如寄存器，共享内存等，由于 SM 自身资源有限，因此它被分配到的 block 数量也是有限的，这取决于 block 中线程的资源需求。每个 SM 中有 4 个处理块，每个处理块都有一个独立的 Warp Scheduler，Warp Scheduler 负责调度分配到当前处理块上的 warp，所有这些 warp 构成一个叫做 Active Warps 的集合，调度器在每个。

原创 理一理Latency、Bandwidth、Throughput、Response Time概念的区别

在性能测试术语中，“在给定的时间段内从一个位置成功移动到另一个位置的数据量，通常以每秒比特数(bps)来度量，或每秒兆比特数(Mbps)、每秒千比特数(Gbps)”。带宽（Bandwidth）：它显示了管道（通信通道）的容量。假设ISDN的带宽是64K，则我们可以增加一个64K的通道，所以总带宽是128K。经常，一些性能测试人员对延时（Latency）、吞吐（Throughput）、带宽（Bandwidth）和响应时间（Response Time）感到迷惑，今天，就用一些简单的例子来说明它们之间的区别。

转载 GPU 峰值算力及性能优化参考资料汇总

另外，这里用的是F32浮点峰值做例子，如果你的任务不需要浮点运算或是精度不是F32，这个值就意义不大，需要转换成你需要的那个操作。其实不看Tensor core的话，满血版一般有：F64:F32:F16=1:2:4，正好与占用的GPR成反比，这个其实是与GPR的带宽有很大的关联的，一般满血版的卡的功能单元配比就会尽量按极限的GPR带宽来设计。一般说来，实际应用中，较大尺寸的矩阵乘法（GEMM）是难得的能接近峰值性能的程序，有些实现能到98%峰值的效率。指令，这是很苛刻的条件。都是一个指令一个flop，

原创 simd vs simt

SIMD 架构是指在同一时间内对多个数据执行相同的操作，适用于向量化运算。例如，对于一个包含多个元素的数组，SIMD 架构可以同时对所有元素执行相同的操作，从而提高计算效率。在实际应用中，根据具体的计算需求和硬件环境选择合适的架构可以提高计算性能。SIMT 架构是指在同一时间内执行多个线程，每个线程可以执行不同的指令，但是这些线程通常会执行相同的程序。版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。常见的 SIMD 架构包括。

转载 CUDA编程：warp

因为一个warp中的线程必须执行相同的指令，不能一部分执行if中的指令，一部分执行else中的。由两个线程块组成的线程结构。，这些资源都被thread瓜分了，由于资源是有限的，所以，如果thread数量比较多，那么每个thread占用资源就比较少，反之如果thread数量较少，每个thread占用资源就较多，这需要根据自己的需求作出一个平衡。SM中的warp调度器会选择warp执行，被选中的成为selected，没选中但已经具备执行条件的成为eligible，没准备好的成为stalled。

转载 CUDA编程方法论

GPGPU中一些问题的理解与思考（3）- 指令执行吞吐与指令集设计 gpu中的warp的divergence问题? 写CUDA到底难在哪？ GPGPU中一些问题的理解与思考（2）- 各级存储的吞吐问题 NV GPU（SM50）的Register Bank Conflict和Reuse Cache的一些测试 GPGPU中一些问题的理解与思考（1）-上层架构 发布一个CUDA SASS

转载 谈谈 cuda 线程束与内存模型

参考：1、2、Cuda thread 组成 block，block 组成 grid，形成嵌套网格状结构，目的是为了矩阵并行计算时，能对应矩阵位置。

原创 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?线程的组织形式对程序的性能影响是至关重要的，本篇博文主要以下面一种情况来

转载 CUDA编程4-执行模型

执行模型会提供一个操作视图，说明如何在特定的计算架构上执行指令。CUDA执行模型解释了GPU并行架构的抽象视图，有助于编写指令吞吐量和内存访问高效的代码。GPU架构围绕一个流式多处理器（SM）的可扩展阵列搭建，通过复制这种架构的构件来实现硬件并行。每个GPU通常有多个SM，每个SM都能支持数百个线程并发执行。当启动一个内核网格时，它的线程被分布在可用的SM上执行。线程块一旦被调度到一个SM上，其中的线程只会在那个指定的SM上并发执行；多个线程块可能被分配到同一个SM上，且根据SM资源的可用性进行调度；

转载 GPU软件抽象：Grid，Block，Thread，Warp定义说明与硬件的映射执行细节 学生教室楼层的例子

一个Warp中的线程必然在同一个block中，如果block所含线程数目不是Warp大小的整数倍，那么多出的那些thread所在的Warp中，会剩余一些inactive的thread，也就是说，即使凑不够Warp整数倍的thread，硬件也会为Warp凑足，只不过那些thread是inactive状态，需要注意的是，即使这部分thread是inactive的，也会消耗SM资源。在执行功能上是没有问题的，那些不应该被执行分支的线程会被禁用，因此始终获得正确的结果，但是如果双方都很长，那么性能损失就很重要。

转载 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 扩展了计算机体系结构的弗林分类学。

转载 CUDA02 - 访存优化和Unified Memory

CUDA02 - 访存优化和Unified Memory - 猫猫子 - 博客园 (cnblogs.com)CUDA02 - 的内存调度与优化前面一篇（传送门）简单介绍了CUDA的底层架构和一些线程调度方面的问题，但这只是整个CUDA的第一步，下一个问题在于数据的访存：包括数据以何种形式在CPU/GPU之间进行通信、迁移，以及在GPU内部进行存储、访问。1 global 、shared 、constant、local通常来讲，待计算的数据都存放在内存或者硬盘（外部存储设备）中，由CPU

转载 CUDA01 - 硬件架构、warp调度、指令流水线和cuda并发流

超级棒这一部分打算从头记录一下CUDA的编程方法和一些物理架构上的特点；从硬件入手，写一下包括线程束的划分、流水线的调度等等微结构的问题，以及这些物理设备是如何与软件对应的。下一部分会写一下cuda中的几种内存划分，进行数据同步，以及优化cuda运行效率的几种方法。

转载 Warp Scheduler，合理块和线程的数量对GPU利用率非常重要

​GPU算子GPU利用率的因素包括，算子线程数量设置合不合理、访存模式合不合理、数据传输次数是否过多、硬件特性是否被充分利用​

转载 英伟达GPU的硬件架构——显存带宽和算力

CUDA 是建立在英伟达GPU硬件架构之上的平台和编程模型。因此，就像要写好高性能的 CPU 程序，必须对 CPU 的硬件架构有一定的理解一样，我们写 CUDA 程序也必须对 GPU 有最基本的了解。

转载 ​CUDA 调度模型​

Block 对应的物理硬件概念是 SM，也就是说 SM 负责 block 中线程的执行，SM 会为每个 block 分配需求的资源，比如寄存器，共享内存等，由于 SM 自身资源有限，因此它被分配到的 block 数量也是有限的，这取决于 block 中线程的资源需求。每个 SM 中有 4 个处理块，每个处理块都有一个独立的 Warp Scheduler，Warp Scheduler 负责调度分配到当前处理块上的 warp，所有这些 warp 构成一个叫做 Active Warps 的集合，调度器在每个。

转载 CUDA 编程入门_CUDA 编程模型

1CUDA 的全称是 Compute Unified Device Architecture，它是建立在硬件架构之上的并行计算平台和编程模型。

转载 40岁职场选择

同时，40岁的职场人因为长期的职场打磨，对于工作方式、行为习惯有了长期一致稳定的SOP，放大看着就叫企业文化，是一种行为方式的偏好，十几二十年下来形成的习惯，很难改变，这个时候再跳槽企业也还会认为你不是白纸甚至是写满了各种各样内容不一的大杂烩，要进入新企业融入新的文化，很难，基于这一点，在雇佣40岁的职场人时，是有减分的。其他还包括，本身职场打磨长时间后，自身在生活成本、原薪资水平、职场观念上的影响，40多岁的职场人要价是不低的，同比与年轻人要高不少。

转载 单元流水线的延迟与吞吐量

当指令指针距离你还有 4KB 远的时候（这大约是 1500 条指令），你被 CPU 从内存取到指令缓存中。虽然从内存加载进入指令缓存需要一段时间，但是现在距离你被执行的时刻还很远，你有足够的时间。这个预取的过程属于流水线的第一级。你进入后发现你刚好站在你前面进来指令的后面，即使执行中的顺序可能已经不同，但你们退出的顺序继续保持一致。黑盒子会处理指令指针指向的指令，当处理完之后，会在内存里找到处理的结果。当指令指针离你越来越近，距离你还有 24 条指令的时候，你和你旁边的 5 个指令会被放到指令队列里面。

转载 DNN库中matmul分析和设计

matmul函数是深度学习中最基本的函数，表示两个张量矩阵相乘，底层主要是通过GEMM和GEMV实现，GEMM即为矩阵乘矩阵，GEMV为矩阵乘向量。GEMM在深度学习中是十分重要的，全连接层以及卷积层基本上都是通过GEMM来实现的，而网络中大约90%的运算都是在这两层中。而一个良好的GEMM的实现可以充分利用系统的多级存储结构和程序执行的局部性来充分加速运算。参数说明order存储方式：CblasColMajor或者CblasRowMajortransA矩阵A是否转置transB。

转载 性能峰值那点事--CPUFP分析

工欲善其事,必先利其器，在DNN库算子设计之前，首先要了解CPU的理论硬件性能。cpufp就是一款可以测出实际理论性能的工具。。

转载 浮点峰值那些事儿_知乎高洋

新的架构很有意思，高端版本支持一个周期发射两条AVX512版本的乘加指令，其中一条来自port0和port1的256位FMA的端口融合，另外一条来自port5。除了用到不同的指令以外，与SNB还有如下一些区别：fma指令的某个参数寄存器，既做输入也做输出，这样前后两个相邻循环间的同一条fma，就形成了RAW型寄存器依赖。浮点峰值的计算，一般是计算单位时间内，乘法和加法的最大总吞吐量，单位是GFLOPS或者TFLOPS，表示每秒钟计算乘法和加法的总次数。通过测试结果与理论上限的差距，评估算法的可能优化空间。

转载 {AI}System

本开源项目主要是跟大家一起探讨和学习人工智能、深度学习的系统设计，而整个系统是围绕着在 NVIDIA、ASCEND 等芯片厂商构建算力层面，所用到的、积累、梳理得到 AI 系统全栈的内容。希望跟所有关注 AI 开源项目的好朋友一起探讨研究，共同促进学习讨论。正在上传…重新上传取消。

转载 C++ STL 容器、迭代器、适配器(深入了解,一文学会)

原文链接：https://blog.csdn.net/qq_37529913/article/details/120052137。C++ STL 容器、迭代器、适配器(深入了解,一文学会)目录。6.1、reverse_iterator反向迭代器适配器。6.2、insert_iterator插入迭代器适配器。6.3、stream_iterator流迭代器。4.1、unordered_map容器。4.2、unordered_set容器。5.2、queue容器适配器详解。5.1、stack容器适配器。

转载 AI编译器TVM与MLIR框架分析

图5 常见ai compiler设计架构

转载 CUDA生态和ROCm生态对比分析

原创 GPU并行计算-参考资料收集list

1、2、3、4、

原创 先进编译实验室

循环优化（三）：循环分布 循环优化（四）：循环交换循环优化（六）：循环分段循环优化（七）：循环分块循环优化（八）：循环分裂循环优化（九）：循环倾斜

转载 C/C++ 性能优化背后的方法论：TMAM

开发过程中我们多少都会关注服务的性能，然而性能优化是相对比较困难，往往需要多轮优化、测试，属于费时费力，有时候还未必有好的效果。但是如果有较好的性能优化方法指导、工具辅助分析可以帮助我们快速发现性能瓶颈所在，针对性地进行优化，可以事半功倍。性能优化的难点在于找出关键的性能瓶颈点，如果不借助一些工具辅助定位这些瓶颈是非常困难的，例如：c++程序通常大家可能都会借助perf /bcc这些工具来寻找存在性能瓶颈的地方。性能出现瓶颈的原因很多比如 CPU、内存、磁盘、架构等。

转载 利用英特尔® VTune™ Profiler 分析、调整和提升应用性能的 8 种方法

8 Ways to Maximize Application Performance With VTune™ Profiler (intel.com)A present-day developer has a wide range of hardware and software tools options available when building and maintaining an application. At the software level, they can choose progra

转载 【无标题】vtune TMA中四个指标的画图解释

tuningguide-intelxeonprocessor-scalablefamily-2ndgen-181827.pdf

原创 Top-down Microarchitecture Analysis Method

1、英文链接：1. https://www.intel.com/content/www/us/en/docs/vtune-profiler/cookbook/2023-0/top-down-microarchitecture-analysis-method.html2. http://portal.nacad.ufrj.br/online/intel/vtune2017/help/GUID-02271361-CCD4-410C-8338-4B8158157EB6.htmlVTune Profiler aut

转载 Analyzing Open vSwitch* with DPDK Bottlenecks Using Intel® VTune™ Amplifier

【代码】Analyzing Open vSwitch* with DPDK Bottlenecks Using Intel® VTune™ Amplifier。