- 博客(116)
- 收藏
- 关注
翻译 Bicubic interpolation
双三次插值是二维图像处理中常用的高阶插值方法,通过16个邻近像素(4×4)计算插值点,比双线性插值(2×2)和最近邻插值更平滑。其数学原理是在单位正方形四个角点匹配函数值、一阶和二阶导数,构建双三次多项式曲面。通过求解16个系数的线性方程组,可获得连续且导数连续的插值结果。该方法在图像重采样中能提供更高质量的视觉效果,但计算复杂度较高,且可能引入特定插值伪影。
2025-09-24 20:11:05
204
原创 Ascend DrivingSDK 中的 modulated_deform_conv2d(二)
本文介绍了modulated_deformable_conv2d的反向传播实现。主要内容包括:1) 反向计算需要处理4个输入的梯度,执行效率低于前向;2) 介绍了DrivingSDK中的modulated_deformable_conv2d_backward函数实现,包括参数检查、张量初始化和两种计算路径选择;3) 详细分析了DeformableConv2dGradKernel的处理流程,包含预处理和梯度计算过程。实现中考虑了分组和不分组情况,并对不同输入通道数进行了优化处理。
2025-08-14 12:43:34
693
原创 Ascend DrivingSDK 中的 modulated_deform_conv2d(一)
Ascend DrivingSDK是为昇腾NPU平台开发的自动驾驶算子和模型加速库,其modulated_deform_conv2d算子采用两种实现:v1版本缓存卷积窗口索引并复用计算结果,v2版本优化3x3卷积性能但功能更简陋。两者均存在参数顺序不一致、功能不完善等问题,且缺乏文档说明,使用体验较差。该SDK虽开源但代码质量未达工业级标准,主要亮点在于利用NPU硬件特性(如大容量L2缓存)提升性能,但仍有较大改进空间。
2025-08-11 19:13:16
766
原创 MARLIN: Mixed-Precision Auto-Regressive Parallel Inference on Large Language Models
MARLIN提出了一种高效的4位权重量化内核,在16-32批量下可实现接近理论4倍的加速效果。该技术通过优化量化布局、异步内存访问、任务调度流水线等创新方法,显著提升了GPU推理效率。作为通用混合精度实现方案,MARLIN已集成到vLLM框架,并开源了稠密和稀疏版本。研究显示,该方法在批量增大时仍保持显著加速优势,为LLM推理提供了重要优化手段。
2025-07-09 09:57:46
1980
原创 MMCV 中的 ModulatedDeformConv2d
MMCV 是一个面向计算机视觉的基础库。相比 torchvision, MMCV 的功能更丰富。而且除了 CPU 和 GPU 之外,还支持昇腾、寒武纪和摩尔线程的硬件。然而 MMCV 后端接入缺乏统一标准,DIOPI 早夭更是加剧了这一局面。MMCV 中 DeformConv2d 和 ModulatedDeformConv2d 分别对应 DCNv1 和 DCNv2 的实现。DeformConv2d 对外提供参数,而 ModulatedDeformConv2d 则没有。DeformConv2d 可能已
2025-06-30 18:56:13
1769
翻译 cuBLAS 3.1.5. Narrow Precision Data Types Usage
cuBLAS我们这里所说的最初是作为8位浮点数据类型(FP8)随 Ada 和 Hopper GPU(计算能力8.9及以上)引入的,旨在进一步加速矩阵乘法。:除非另有说明,FP8 指的是和两种数据类型。随着 Blackwell GPU(计算能力10.0及以上)的推出,cuBLAS 增加了对4位浮点数据类型 (FP4)的支持。E2和 M1分别表示2位指数和1位尾数。更多详细信息,请参阅。为了保持精度,窄精度数据在计算前需要进行缩放或反量化,并在计算后可能需要进行量化。
2025-03-27 19:12:23
714
原创 DaDianNao:一种无主存储器的多核加速器
DianNao 发现内存访问而非计算是深度学习模型执行的瓶颈,并研究了一种最小化内存传输的加速器架构。后续的 DaDianNao 和 ShiDianNao 将该思路贯彻到了极致: ShiDianNao 是一种 YX 切分、输出驻留的嵌入式 CNN 加速器; DaDianNao:一种分布式驻留权值的多核加速器。DaDianNao 本意是一种多芯片AI 超算系统,类似于 Groq LPU 的 Scale Out。然而数据在多芯片间的移动会成为新的瓶颈。以现在的标准来看,DaDianNao
2025-02-11 20:11:02
1393
原创 DianNao:一种支持 Conv、FC 和 Pooling 的单核加速器
DianNao: A Small-Footprint High-Throughput Accelerator for Ubiquitous Machine-Learning 提出了首个深度学习处理器架构,开创了深度学习处理器方向,并且获得当年的 ASPLOS 2014最佳论文奖。在此之后,CNN 模型日新月异,各类 AI 加速器层出不穷。 NVIDIA 的 NVDLA 即是采用了和 DianNao 类似的 KC 切分、权值驻留结构,而搭载了 NVDLA 的 Jetson Xavier 和 Jetson Or
2024-12-24 18:42:25
1379
翻译 CUDA C++ Programming Guide 7.28. Asynchronous Data Copies using cuda::pipeline
CUDA提供了cuda::pipeline同步对象来管理异步数据移动并将其与计算重叠。libcudacxx API 中提供了cuda::pipeline的 API 文档。 流水线对象是一个具有头部和尾部的双端 N 级队列,用于按先进先出(FIFO)顺序处理工作。流水线对象有以下成员函数,用于管理流水线的各个阶段。
2024-10-24 10:02:29
382
翻译 CUDA C++ Programming Guide 7.27. Asynchronous Data Copies
CUDA 11通过memcpy_async API 引入异步数据操作,允许设备代码显式管理数据的异步复制。memcpy_async特性使 CUDA 内核能够在数据移动的同时进行计算。
2024-10-24 09:59:08
541
翻译 CUDA C++ Programming Guide 7.26. Asynchronous Barrier
NVIDIA C++ 标准库引入了 std::barrier 的 GPU 实现。除了std::barrier的实现之外,该库还提供了扩展功能,允许用户指定屏障对象的作用范围。屏障 API 的作用范围在 Thread Scopes 中有详细说明。 计算能力为8.0或更高的设备为屏障操作提供了硬件加速,并将其与memcpy_async 功能集成。 在计算能力低于8.0但不低于7.0的设备上,这些屏障可以使用但没有硬件加速。
2024-09-11 09:30:16
1057
原创 CUTLASS 中的 47_ampere_gemm_universal_streamk 示例
前一篇文章介绍了 Stream-K: Work-centric Parallel Decomposition for Dense Matrix-Matrix Multiplication on the GPU 论文,下面对其代码实现进行分析。cutlass 的 examples/47_ampere_gemm_universal_streamk 展示了 GEMM Stream-K 算法在 Ampere 架构上的使用。对比了普通 Gemm 以及 Split-K 算法和 Stream-K 的性能
2024-08-21 19:36:09
2248
原创 Stream-K: Work-centric Parallel Decomposition for Dense Matrix-Matrix Multiplication on the GPU
在 NVIDIA GTC22 秋季会议上,CUTLASS: Python API, Enhancements, and NVIDIA Hopper 介绍了 CUTLASS~2.11 中引入的 Stream-K 分解:在这里插入图片描述几个月后公开的 Stream-K: Work-centric Parallel Decomposition for Dense Matrix-Matrix Multiplication on the GPU 论文对其进行了更详细的介绍。CUTLASS 的 GEMM 实现由三大
2024-05-31 18:18:45
3437
原创 onnxruntime 中的 Gather 算子
上一篇文章中介绍了 Division by Invariant Integers using Multiplication 的原理,很多框架均才用该算法优化除法运算。onnxruntime 是已知实现中最为简洁的,因此本文结合 onnxruntime 的 Gather 实现进行介绍。 Gather 算子是一个索引类算子,kernel 中每个线程计算偏移时使用 fast_divmod 避免除法运算。 注意:ONNX 中的 Gather 功能与 numpy.take 相同
2024-03-27 19:06:20
3718
2
原创 Division by Invariant Integers using Multiplication
表 1.1 比较了一些处理器上乘法和除法的时间。这张表展示了乘法和除法时间差距的增长趋势。因此,中提出了使用整数乘法进行任意非零整数常数和运行时不变量之间除法的算法。文档中记录了更广泛的处理指令性能,其中 Intel IceLake 处理器的乘除法指令延迟和吞吐倒数如下表所示:可以看出,在现代 CPU 处理器上除法开销大的情况并未发生改变。NVIDIA 和 AMD GPU 均不支持整数除法指令,CUDA C++ Programming Guide。
2024-03-19 20:00:37
1148
原创 CUTLASS 1.3.3中的 Volta884_h884gemm
CUTLASS 是 CUDA C++ 模板抽象的集合,用于在 CUDA 内的所有级别和规模上实现高性能矩阵-矩阵乘法 (GEMM) 和相关计算。它采用了类似于 cuBLAS 和 cuDNN 中实现的分层分解和数据移动策略。CUTLASS 最新版本为3.3,相比1.3.3变动较大。然而重温一下1.3.3仍然是有意义的。因为它更易于理解: 与 PROGRAMMING TENSOR CORES: NATIVE VOLTA TENSOR CORES WITH CUTLASS 中介绍的内容相匹配;
2023-11-22 19:45:13
826
原创 Programming Tensor Cores: NATIVE VOLTA TENSOR CORES WITH CUTLASS
PROGRAMMING TENSOR CORES: NATIVE VOLTA TENSOR CORES WITH CUTLASS 源自于 GTC Silicon Valley-2019: cuTENSOR: High-performance Tensor Operations in CUDA,介绍了 CUTLASS 1.3 中基于 Volta Tensor Core 实现高效矩阵乘法计算的策略。主要内容为以下三点: CUDA 10.1中mma.sync指令介绍; Global Memory
2023-11-22 17:18:50
1374
原创 Modeling Deep Learning Accelerator Enabled GPUs
Modeling Deep Learning Accelerator Enabled GPUs 发表在 ISPASS 2019 上。文章研究了 NVIDIA 的 Volta 和 Turing 架构中张量核的设计,并提出了 Volta 中张量核的架构模型。基于实现该模型,并且支持 CUTLASS 运行。发现其性能与硬件非常吻合,与 Titan V GPU 相比,获得了99.6%的 IPC 相关性。文中还展示了 Turing 架构中张量核的操作数矩阵元素到线程的映射,并发现它们与 Volta 张量核的行为不同。
2023-10-23 17:54:42
1432
翻译 CUTLASS: Implicit GEMM Convolution
Implicit GEMM 是将卷积操作表述为 GEMM (广义矩阵-矩阵积)。卷积接受激活张量并对其应用滑动滤波器以产生输出张量。二维卷积可以映射到矩阵乘:组建一个包含激活张量元素的卷积矩阵,然后由滤波张量形成的矩阵乘以该矩阵。该算法的最早形式通过通常称为 im2col 的操作显式构造卷积矩阵。生成的矩阵按照滤波器大小复制每个激活元素,消耗额外的存储容量和内存带宽。隐式 GEMM 算法是 CUDA 中分块、分层 GEMM 计算的一种变体:当数据从全局内存加载到共享内存时,通过周密地更新指针和谓词,它
2023-08-15 19:21:06
4916
翻译 CUTLASS: Efficient GEMM in CUDA
CUTLASS 实现了 CUTLASS: Fast Linear Algebra in CUDA C++ 和 CUTLASS GTC2018 talk 中描述的分层分块结构。基本的三重嵌套循环计算矩阵乘法可以应用分块和拼贴,以匹配硬件、内存局部性和并行编程模型中的并发性。CUTLASS 中 GEMM 映射到 NVIDIA GPU 的结构如以下嵌套循环所示。
2023-08-10 19:39:18
3437
原创 OneFlow 中的 Softmax
Softmax 是深度学习模型中的常见算子。PyTorch 的 Softmax 算子直接调用 cuDNN 的接口。而 OneFlow 内部针对输入数据的类别数量,采用3个 kernel 来分别处理,在多数情况下都可以获得比 cuDNN 更优的性能表现。下面对其实现进行介绍。OneFlow 的静态分层结构如下图所示:
2023-08-10 19:18:13
2966
原创 torchvision 中的 deform_conv2d
如 DCNv1 和 DCNv2 论文所述,DeformConv 相比常规卷积的参数量和计算量增加不多,但对网络的提升很大。然而,DeformConv 的计算模式并不利于高效实现,给网络带来的开销比纸面数值大:常规卷积可以采用 Implicit GEMM 的形式,非常高效;DeformConv 需要离散访存和插值,增加了 IO 量和内存占用。在 Torchvision 以及其他框架中,DeformConv2d 采用 Explicit GEMM 的方式实现。具体步骤为:deformable_im
2023-07-07 17:53:40
4632
4
原创 DCN v2阅读笔记
是研究的续作,发表在 CVPR 2019上。作者对的自适应行为进行研究,观察到虽然其神经特征的空间支持比常规的卷积神经网络更符合物体结构,但这种支持可能远远超出感兴趣区域,导致特征受到不相关图像内容的影响。为此,作者提出了改进版的,通过增加建模能力和更强的训练来提高其聚焦于相关图像区域的能力。
2023-07-05 12:51:51
1497
原创 DCN v1阅读笔记
视觉识别(例如对象检测和语义分割)中的一个关键挑战是如何适应物体尺度、姿态、视角和零件变形中的几何变化或模型几何变换。以往通过扩充现有数据样本,构建具有足够所需变化的训练数据集来缓解。这两种方法的指导思想为在模块中增加额外偏移量的空间采样位置,并从目标任务中学习偏移量,而无需额外监督。新模块可以很容易地替换现有 CNN 中的普通模块,并且可以通过标准反向传播很容易地进行端到端训练,从而产生。
2023-06-26 20:14:14
712
原创 TensorFlow 中的 BatchToSpaceOp
Atrous Convolution 是中提出的卷积运算。不仅能够明确控制在深度卷积神经网络中计算特征响应的分辨率,还可以在不增加参数数量或计算量的情况下,有效地扩大滤波器的视场以纳入更大的上下文。作者通过在 Caffe 框架中的 im2col 层添加对底层特征图进行稀疏采样的选项来实现。而到了发布时,该操作已加入到 TensorFlow 官方支持中,即和。rr2r×r通过将空洞卷积化简为规则卷积,能够使用现有的高度优化的卷积程序。实现原理可参考。和会调用。fill:#333;
2023-04-23 19:34:48
926
原创 cuDNN 的初始设计
cuDNN V1.0在2014年的发布,并集成到 Caffe、Paddle 等深度学习框架中。论文介绍了 NVIDIA 对于该库的设计和实现。近十年间,NVIDIA 迭代推出了8代架构,cuDNN 也更新到 8.9。硬件上引入了 Tensor Core,软件方面 cuDNN V8 中的 Graph API 相比之前变化较大。然而,作为深度学习领域影响广泛的软件基础设施,一窥 cuDNN 的初始设计仍是有意义的。cuDNN 类似于 cuBLAS,提供了深度学习原语的高效实现。
2023-04-22 20:07:02
1539
原创 TensorFlow 中的 LRNOp
TensorFlow 中的 LRNOp 与 Caffe 的差异:* 直接使用平方和而不是像论文中一样使用平方和的均值,因此算子的推荐参数有所不同;* 仅支持 NHWC 格式的输入;* CPU 后端有多种实现: MKL、SingleThreadedLRN 和带状矩阵乘法实现;* GPU 后端仅有 cuDNN 实现,在前后插入转置操作。
2023-01-30 14:24:16
600
1
原创 Caffe 中的 LRNLayer
Caffe 中的 LRNLayer 支持两种模式: CrossChannel:为人们所熟知的局部响应归一化,在 AlexNet 中提出,在一些早期网络中使用; WithinChannel: Caffe 中独有的实现,未见网络中应用。本文略过。ReLU 具有不需要输入归一化以防止其饱和的理想特性。但 AlexNet 论文中发现 LRN 有助于提高泛化性。LRN CrossChannel 模式公式如下:
2023-01-17 11:54:38
509
翻译 Understanding Memory Formats
大多数计算都是关于数据的:分析数据、调整数据、读取和存储数据、生成数据等。DNN 领域也不例外。图像、权重/过滤器、声音和文本需要在计算机内存中高效表示,从而以最方便的方式快速执行操作。本文致力于数据格式的一种数据表示形式,它描述了多维数组(nD)如何存储在线性(1D)内存地址空间中,以及为什么这对 oneDNN 很重要。
2022-11-27 20:48:44
435
翻译 PyTorch Design Philosophy
本文档旨在帮助贡献者和模块维护者理解 PyTorch 演化出的高层设计原则。这些规则并不是硬性规定的,而是用来作为指南,以帮助权衡不同的关注点,解决开发 PyTorch 时可能出现的分歧。有关贡献、模块维护以及如何将分歧升级至核心维护者的更多信息,请参见 [ PyTorch Governance](https://pytorch.org/docs/master/community/governance.html)。
2022-11-23 16:58:02
286
原创 TensorFlow 中的 Conv2DOp
TensorFlow 中的2D 卷积主要依赖外部库,如 cuDNN、cuBLAS、ROCm 和 hfp/libxsmm,仅 DeepConv2D 为源码实现。Conv2DOpBinaryOpInitConv2DParameters 从 OpKernelConstruction 中读取设置到 Conv2DParameters 并进行检查。CudnnUseAutotune 标识是否开启自动调优。......
2022-08-03 10:35:49
1098
原创 PaddlePaddle 中的 CTCGreedyDecoder
TensorFlow 中的 CTCGreedyDecoder 仅包含 CPU 实现。而 PaddlePaddle 框架则更贴近实际需求,可以在 GPU 上运行。简单来说,PaddlePaddle 内部通过拼接方式,先通过 topk 算子找到最大类别,然后通过 CTCAlignOp 完成后处理。ctc_greedy_decodercheck_variable_and_dtype 检查变量的类型以及数据类型。LayerHelper 主要是在各个 layers 函数之间共享代码。内部调用 topk 算子得
2022-05-04 22:31:01
1309
原创 TensorFlow 中的 CTCGreedyDecoder
CTCGreedyDecoderOp 对输入中给出的 logits 执行贪婪解码(最佳路径)。
2022-04-30 21:47:25
2150
原创 Gables: A Roofline Model for Mobile SoCs
为了帮助构建 SoC 思维并指导早期移动 SoC 设计,Gables: A Roofline Model for Mobile SoCs 提出了 Gables 模型,该模型改进和重新定位了 Roofline 模型(最初为多核芯片的性能和带宽限制而设计)来对 SoC 上的每个加速器进行建模,在不同的加速器之间并发的分配工作(由文中用例分析证明),并计算 SoC 性能上限。作者使用现有 SoC (Snapdragon 835)评估 Gables 模型并开发了多个扩展,使 Gables 能够为早期移动 SoC 设
2022-02-12 21:03:11
1373
原创 Applying the Roofline Model for Deep Learning performance optimizations
Applying the Roofline Model for Deep Learning performance optimizations 以 Intel Xeon 为例,介绍了一种为非统一内存访问( NonUnified Memory Access,NUMA[8])自动创建 Roofline 模型的方法,并对 Intel oneDNN 库中实现的高效深度学习原语进行了评估。2 Description of methodology所有实验均在禁用 Intel Turbo Boost 技术的 Inte
2022-01-15 11:49:02
1270
原创 Design and Implementation of a Highly Efficient DGEMM for 64-bit ARMv8 Multi-Core Processors
Design and Implementation of a Highly Efficient DGEMM for 64-bit ARMv8 Multi-Core Processors 针对64位 ARMv8八核处理器,设计并实现了一种基于 OpenBLAS 的高效 DGEMM。作者首先为此架构开发性能模型,然后根据理论指导用汇编语言系统地开发高度优化的 GEBP 内核。性能模型表明,优化 DGEMM 的峰值性能(效率)需要在内存层次结构的所有级别上最大化其计算内存访问比率。而提高 GEBP 的性能的主要
2022-01-01 11:20:01
2455
原创 Roofline Model Toolkit: A Practical Tool for Architectural and Program Analysis
Roofline Model Toolkit: A Practical Tool for Architectural and Program Analysis 描述了 Roofline Toolkit 的原型架构表征引擎。该引擎由一组使用消息传递接口(Message Passing Interface,MPI )以及用于表示线程级并行性的 OpenMP 实现的便携式设备化微基准组成,可量化多核、众核和加速系统的带宽和计算特性。这些微观测试侧重于在编译器和运行时环境以及线程级并行、指令级并行和显式 SIMD
2021-12-25 15:58:46
1665
原创 Roofline-on-NVIDIA-GPUs代码分析
Roofline 代码现状:CS Roofline Toolkit 为 Roofline Model Toolkit: A Practical Tool for Architectural and Program Analysis 的实现,uo-cdux/ert-mirror 为 github 上的一个镜像;cyanguwa/nersc-roofline 为 Hierarchical Roofline Analysis: How to Collect Data using Performance To
2021-12-04 10:28:22
2843
原创 Hierarchical Roofline Performance Analysis for Deep Learning Applications
Roofline 模型是劳伦斯伯克利国家实验室在2008年提出的一个性能模型,后续很多工作亦出自该实验室。考虑到分层 Roofline 这一概念已在先前的Hierarchical Roofline analysis for GPUs: Accelerating performance optimization for the NERSC-9 Perlmutter system 和 Hierarchical Roofline Analysis: How to Collect Data using Perfo
2021-11-28 21:37:31
1672
原创 TNN MatConverter Resize
TNN 的 resize 虽然分通道提供了多个接口,但底层是一起的。整个实现对于灰度图优化非常有限,而3通道或4通道的图像会有加速。缩放的映射关系较为简单,主要分为三步:
2021-05-09 16:53:29
1042
1
空空如也
空空如也
TA创建的收藏夹 TA关注的收藏夹
TA关注的人
RSS订阅