- 博客(103)
- 收藏
- 关注
RSS订阅原创 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
1063
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
780
原创 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
265
原创 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
236
原创 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
737
翻译 CUTLASS: Implicit GEMM Convolution
Implicit GEMM 是将卷积操作表述为 GEMM (广义矩阵-矩阵积)。卷积接受激活张量并对其应用滑动滤波器以产生输出张量。二维卷积可以映射到矩阵乘：组建一个包含激活张量元素的卷积矩阵，然后由滤波张量形成的矩阵乘以该矩阵。该算法的最早形式通过通常称为 im2col 的操作显式构造卷积矩阵。生成的矩阵按照滤波器大小复制每个激活元素，消耗额外的存储容量和内存带宽。隐式 GEMM 算法是 CUDA 中分块、分层 GEMM 计算的一种变体：当数据从全局内存加载到共享内存时，通过周密地更新指针和谓词，它
2023-08-15 19:21:06
1717
翻译 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
784
原创 OneFlow 中的 Softmax
Softmax 是深度学习模型中的常见算子。PyTorch 的 Softmax 算子直接调用 cuDNN 的接口。而 OneFlow 内部针对输入数据的类别数量，采用3个 kernel 来分别处理，在多数情况下都可以获得比 cuDNN 更优的性能表现。下面对其实现进行介绍。OneFlow 的静态分层结构如下图所示：
2023-08-10 19:18:13
345
原创 torchvision 中的 deform_conv2d
如 DCNv1 和 DCNv2 论文所述，DeformConv 相比常规卷积的参数量和计算量增加不多，但对网络的提升很大。然而，DeformConv 的计算模式并不利于高效实现，给网络带来的开销比纸面数值大：常规卷积可以采用 Implicit GEMM 的形式，非常高效；DeformConv 需要离散访存和插值，增加了 IO 量和内存占用。在 Torchvision 以及其他框架中，DeformConv2d 采用 Explicit GEMM 的方式实现。具体步骤为：deformable_im
2023-07-07 17:53:40
1781
原创 DCN v2阅读笔记
是研究的续作，发表在 CVPR 2019上。作者对的自适应行为进行研究，观察到虽然其神经特征的空间支持比常规的卷积神经网络更符合物体结构，但这种支持可能远远超出感兴趣区域，导致特征受到不相关图像内容的影响。为此，作者提出了改进版的，通过增加建模能力和更强的训练来提高其聚焦于相关图像区域的能力。
2023-07-05 12:51:51
767
原创 DCN v1阅读笔记
视觉识别（例如对象检测和语义分割）中的一个关键挑战是如何适应物体尺度、姿态、视角和零件变形中的几何变化或模型几何变换。以往通过扩充现有数据样本，构建具有足够所需变化的训练数据集来缓解。这两种方法的指导思想为在模块中增加额外偏移量的空间采样位置，并从目标任务中学习偏移量，而无需额外监督。新模块可以很容易地替换现有 CNN 中的普通模块，并且可以通过标准反向传播很容易地进行端到端训练，从而产生。
2023-06-26 20:14:14
542
原创 TensorFlow 中的 BatchToSpaceOp
Atrous Convolution 是中提出的卷积运算。不仅能够明确控制在深度卷积神经网络中计算特征响应的分辨率，还可以在不增加参数数量或计算量的情况下，有效地扩大滤波器的视场以纳入更大的上下文。作者通过在 Caffe 框架中的 im2col 层添加对底层特征图进行稀疏采样的选项来实现。而到了发布时，该操作已加入到 TensorFlow 官方支持中，即和。rr2r×r通过将空洞卷积化简为规则卷积，能够使用现有的高度优化的卷积程序。实现原理可参考。和会调用。fill:#333;
2023-04-23 19:34:48
637
1
原创 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
1086
原创 TensorFlow 中的 LRNOp
TensorFlow 中的 LRNOp 与 Caffe 的差异：* 直接使用平方和而不是像论文中一样使用平方和的均值，因此算子的推荐参数有所不同；* 仅支持 NHWC 格式的输入；* CPU 后端有多种实现： MKL、SingleThreadedLRN 和带状矩阵乘法实现；* GPU 后端仅有 cuDNN 实现，在前后插入转置操作。
2023-01-30 14:24:16
451
1
原创 Caffe 中的 LRNLayer
Caffe 中的 LRNLayer 支持两种模式： CrossChannel：为人们所熟知的局部响应归一化，在 AlexNet 中提出，在一些早期网络中使用； WithinChannel： Caffe 中独有的实现，未见网络中应用。本文略过。ReLU 具有不需要输入归一化以防止其饱和的理想特性。但 AlexNet 论文中发现 LRN 有助于提高泛化性。LRN CrossChannel 模式公式如下：
2023-01-17 11:54:38
397
翻译 Understanding Memory Formats
大多数计算都是关于数据的：分析数据、调整数据、读取和存储数据、生成数据等。DNN 领域也不例外。图像、权重/过滤器、声音和文本需要在计算机内存中高效表示，从而以最方便的方式快速执行操作。本文致力于数据格式的一种数据表示形式，它描述了多维数组（nD）如何存储在线性（1D）内存地址空间中，以及为什么这对 oneDNN 很重要。
2022-11-27 20:48:44
240
翻译 PyTorch Design Philosophy
本文档旨在帮助贡献者和模块维护者理解 PyTorch 演化出的高层设计原则。这些规则并不是硬性规定的，而是用来作为指南，以帮助权衡不同的关注点，解决开发 PyTorch 时可能出现的分歧。有关贡献、模块维护以及如何将分歧升级至核心维护者的更多信息，请参见 [ PyTorch Governance](https://pytorch.org/docs/master/community/governance.html)。
2022-11-23 16:58:02
129
原创 TensorFlow 中的 Conv2DOp
TensorFlow 中的2D 卷积主要依赖外部库，如 cuDNN、cuBLAS、ROCm 和 hfp/libxsmm，仅 DeepConv2D 为源码实现。Conv2DOpBinaryOpInitConv2DParameters 从 OpKernelConstruction 中读取设置到 Conv2DParameters 并进行检查。CudnnUseAutotune 标识是否开启自动调优。......
2022-08-03 10:35:49
799
原创 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
1146
原创 TensorFlow 中的 CTCGreedyDecoder
CTCGreedyDecoderOp 对输入中给出的 logits 执行贪婪解码（最佳路径）。
2022-04-30 21:47:25
1840
原创 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
929
原创 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
916
原创 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
1540
原创 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
906
原创 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
2086
原创 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
855
原创 TNN MatConverter Resize
TNN 的 resize 虽然分通道提供了多个接口，但底层是一起的。整个实现对于灰度图优化非常有限，而3通道或4通道的图像会有加速。缩放的映射关系较为简单，主要分为三步：
2021-05-09 16:53:29
850
1
原创 TNN MatConverter WarpAffine
TNN 的仿射变换形态介于 OpenCV 和 ncnn 之间。其处理流程与 OpenCV 较为相似并做了一些优化，不同的地方在于数据处理宽度为4，比较小。在性能表现方面中规中矩，小图上不及 ncnn。MatUtils::WarpAffine#mermaid-svg-FNwIOkXOm8kxHfXI .label{font-family:'trebuchet ms', verdana, arial;font-family:var(--mermaid-font-family);fill:#333;color
2021-05-09 16:42:07
474
3
原创 ncnn 中的 warpAffine
ncnn 的仿射变换对于深度学习的预处理即小图变换进行了优化，速度可达到 OpenCV 的两倍。详细请参考 [opencv ncnn warpaffine 性能测试](https://zhuanlan.zhihu.com/p/355147243)。在具体实现方面，优点是简洁明快，双线性插值采用10bit 量化，比 OpenCV 精度高；缺点是边界填充仅支持常量值。下面从 ncnn 的测试代码入手进行分析。
2021-05-09 16:26:02
1151
4
原创 OpenCV 中的 remap 函数
上一篇文章中提到 warpAffine 会分块处理，将坐标映射和插值系数分别存储下来，然后借助 remap 来实现最终的映射。而 remap 会根据映射关系取源像素并加权计算出目的像素值。其最核心的计算为 RemapVec_8u。cv::remap#mermaid-svg-lLtejyb6R6MyZRLO .label{font-family:'trebuchet ms', verdana, arial;font-family:var(--mermaid-font-family);fill:#333;c
2021-04-18 11:43:59
6612
6
原创 OpenCV 中的 warpAffine
warpAffine 是图像处理中比较常见的一种变换，可以将图像校正或对齐。对于线性插值方式，OpenCV 首先将坐标映射保存成两张图，然后调用 remap 函数。第二步是比较耗时的部分，并且 warpPerspective 亦采用此处理。remap 通过构建查找表来存储系数乘积，这样减少了乘法运算次数。由于篇幅过长，将文章分成 warpAffine 和 remap 两部分。
2021-04-18 11:12:02
5149
原创 CMake 编译 rkmedia
在瑞芯微的使用手册中，rkmedia 库使用 Buildroot 编译。然而由于配置文件众多，不易定位编译过程中的问题，所以本文以 CMake 进行构建。目标平台为 RV1109/1126。编译 libdrm-rockchiprkmedia 中的组件支持选项配置，但 drm 是必需的。所以第一步需要先编译该库。首先在电脑上按照编译过程中会用到的库。sudo apt-get install xutils-devsudo apt-get install libpciaccess-dev获取 lib
2021-03-07 20:01:58
2992
7
原创 Flatbuffers 的 C++使用示例
FlatBuffers 使用.来指定嵌套的名称空间/包。root_type声明序列化数据的根表（或结构）。除了基本类型外，Monster中包含一个表和一个结构体。表是在 FlatBuffers 中定义对象的主要方式，由名称（此处为Monster）和字段列表组成。每个字段都有一个名称、类型和可选的默认值（如果省略，则默认为0或NULL）。表中每个字段都是可选的：它不必出现在线路表示中，您可以选择忽略每个对象的字段。因此，您可以灵活地添加字段，而不必担心数据膨胀。
2021-02-27 19:37:21
5622
原创 TNN MatConverter CvtColor NV21TOBGR
OpenCV 中的 carotene 对于 armv7优化较好，而 armv8下则是 NEON 实现。TNN 提供了一套图像预处理接口并且进行了汇编优化。下面以 NV21TOBGR 为例进行介绍。MatUtils无成员变量，全部为静态函数。public: //copy cpu <-> device, cpu<->cpu, device<->device, src and dst dims must be equal. static Status Co
2021-02-21 21:38:45
599
原创 OpenCV cvtColor BGR2YUV420
OpenCV 中的 cvtColor 将图像从一种颜色空间转换为另一种颜色空间。虽然 OpenCV 可以支持从各种 YUV 格式转换到 BGR，但反向到 YUV420 却仅能输出 I420。尽管内部有 cvtBGRtoTwoPlaneYUV 函数，但是对外没有提供。Carotene 库并无该功能，所以 cvtColorBGR2ThreePlaneYUV 使用统一向量指令（universal intrinsics）来加速，维护成本低效率同样不高。在图像大于320x240时可能会启用线程。YCbCr 色彩空
2021-02-13 16:35:49
6008
1
翻译 6.47.2 Extended Asm - Assembler Instructions with C Expression Operands
使用扩展asm，您可以从汇编程序读取和写入 C 变量，并执行从汇编代码到 C 标号的跳转。扩展asm语法使用冒号（“:”）在汇编程序模板之后分隔操作数参数：asm asm-qualifiers ( AssemblerTemplate : OutputOperands [ : InputOperands [ : Clobbers ] ])asm asm-qualifiers ( AssemblerT
2021-02-10 21:19:29
499
原创 rkmedia 中的行人检测
RKMedia 是瑞芯微提供的媒体处理方案，可支持应用软件快速开发。rkmedia为 C 接口，其实现均在easymedia中。后者提供 C++ 接口。rkmedia中的行人检测示例，执行的操作为： 配置两个 RTSP 会话，初始化媒体处理平台； GetMediaBuffer 线程获取缓存并调用 rockx 进行检测跟踪； MainStream 线程绘制结果。
2020-12-26 11:20:21
5492
10
原创 OpenCV copyMakeBorder
copyMakeBorder 将源图像复制到目标图像的中间并在图像周围形成边框。当src已经在dst中间时，该函数不会复制src本身，而只是构造边框。在执行时函数会尝试使用 ROI 之外的像素来形成边界。若要禁用此功能并始终执行外推，就像src不是 ROI 一样，需要设置borderType | BORDER_ISOLATED。copyMakeBorder 是 OpenCV 中较为简单的一个函数，然而由于其浅拷贝和 ROI 机制的存在，实现也并不简单。
2020-11-29 16:10:43
7356
原创 MNN Interpreter and Session
MNN 中 Interpreter 和 Session 之间的关系如下图所示：Interpreter 为模型解释器，亦为会话管理器，负责从文件加载模型，创建并执行会话；一个 Interpreter 对应到一个 Net，管理基于该网络的多个任务；Session 表示推理任务，由 Pipeline 数组和输入输出张量字典组成；Session 根据 Schedule::schedule 生成的 Schedule::ScheduleInfo 创建；归属于同一 Net 的不同 Session 间不能并发调用
2020-09-12 17:56:16
2147
原创 MNN classficationTopkEval
MNN classficationTopkEval.cpp 输入模型和配置文件，测试模型在 ImageNet 数据集上的分类精度。程序分为3部分：ImageProcess 将图像转换为适当格式的 Tensor；Interpreter 由模型文件创建 Net 和 Session 并执行会话；computeTopkAcc 计算分类准确率。main#mermaid-svg-DIh1q4eQfe20BU3X .label{font-family:'trebuchet ms', verdana, ar
2020-08-30 14:22:48
372
空空如也
空空如也
TA创建的收藏夹 TA关注的收藏夹
TA关注的人