大光叔叔-CSDN博客

原创 Conformer Encoder GPU 加速策略较全面汇总

最近一直在做WeNet conformer encoder模型的GPU TensorRT加速，也有幸参加了NVIDIA Hackathon 2022 加速 Wenet 的比赛和阅读了NVIDIA 内部团队关于 WeNet TensorRT加速的代码。学习到了很多东西，抛砖引玉进行汇总，欢迎各位大佬拍砖。以下加速策略描述以TensorRT为例进行简单描述。PS：阅读前需要非常了解conformer encoder的模型结构，并比较熟悉TensorRT or CUDA。

2022-11-02 17:31:51 1231 2

原创 libtorch 上线之路-学习篇

最近在尝试将组里几个pytorch的模型使用libtorch 做nn inference，折腾了几个周，一路踩坑，终于把整个流程都打通了。因为libtorch还比较新，网上经验性的资料还比较少，现在写下来总结一下经验，分享一下。0 环境python 3.6pytorch 1.3.1g++ 5.3.1 # 要高于4.8，否则会报释放内存的错误CentOS Linux release 7.2...

2019-12-31 11:55:26 5374 6

原创 CUDA进阶之路：开篇

合抱之木，生于毫末；九层之台，起于垒土；千里之行，始于足下。《老子·道德经·第六十四章》从大二开始，接触CUDA有六个年头了。从最初的实验室做CUDA项目，混技术群讨论问题，到研究生期间的GPU体系结构研究，再到现在的深度学习加速工作，都在跟GPU打交道。感觉自己怎么都算一个老司机了，在工作中却时常能碰见一些让人眼前一亮的CUDA知识技巧，或者是耐人琢磨的问题，每次弄明白都觉得贼舒服，...

2018-07-11 23:27:09 1917 2

原创 CUDA进阶第八篇-CUDA性能优化指南总结

本文为《CUDA C Programming Guide》chapter 5 Performance Guidelines章节的总结。精力有限，可能没做到覆盖所有的点，期望指正。该章的主要内容是指导如何进行性能优化。从三个大方面入手：最大化利用率;最大化内存吞吐量;和最大化指令吞吐量。最大化利用率应用层次：使用异步函数和流，最大化主机端（host）任务、设备端（device）任...

2018-05-18 17:35:30 8197

原创 CUDA进阶第七篇-如何获得核函数最佳线程块配置

最近在和几个小伙伴们一起翻译《cuda-c-programming-guide》，我选了chapter 5，收获良多。翻译到5.2.3.1 Occupancy Calculator的时候，正好跟现在手里的任务有些重合，就做了些实验，分享给大家。有说的不对的地方欢迎拍砖。这一小节的内容很简单，如何获得核函数最佳线程块大小配置，可以用几个API，根据kernel内寄存器和共享内存的使用量，计算出b...

2018-05-03 17:41:21 3972

原创 CUDA进阶第六篇-GPU资源（显存、句柄等）管理

最近在加速一个图像算法，符合《CUDA进阶第五篇-如何估算出程序的GPU加速比》中的第二种情况，程序由核函数和GPU API实现，但是资源管理特别差，显存和句柄在程序中使用时才申请。每次函数执行都要申请和释放一遍，非常耗费时间。优化方案一：C++重构我想到的第一个方案，就是C++重构，将所有的显存和句柄定义为成员变量，在构造函数中申请，析构函数中释放。难度比较低，但是工作量比较大，因...

2018-02-16 21:42:32 4222

原创 CUDA进阶第五篇-如何估算出程序的GPU加速比

当对一个程序进行加速的时候，很多时候需要预估出程序使用GPU加速后的加速比（比如你老板不懂GPU，或者甲方会问你预估加速比等等）。从大二接触GPU加速，到现在大概有6年时间，大大小小的项目也做了十几个，很多时候都需要事先回答加速比会有多少这个问题。这里简单的说一下自己的经验，欢迎各位大神指点。文中的经验基于目前主流的显卡，比如GTX1080，最低也得是GTX9**系列的。1.阿姆达尔定律...

2018-02-16 20:54:49 5523

原创 CUDA进阶补充篇：详析各种CUDA函数计时函数

写在前面：之前写了一篇CUDA进阶第三篇：CUDA计时方式，列出了几种当时遇到的CUDA计时方式，只是个教程式的东西，没有太多技术含量，也不太全面。前几天在CUDA Professional（45157483）群里和大佬们讨论到CUDA官方event函数在计时cpu和cpu混合代码时有问题，虎躯一震，tm这么多年要是一直用的都是错的就瞎了。今天特花时间实验探究一番。有不足之处还望各位前辈指点。

2017-08-25 20:28:33 6485 3

原创 cublasGemmEx函数应用-探究8bit矩阵乘

介绍cublasGemmEx 是CUDA8.0中cuBLAS新出的函数，是cublasgemm()类函数的扩展，也是目前来看功能最强大的矩阵乘函数了。该函数另一强大之处在于支持多种计算模式（compute type），其中就包括CUDA 8.0新出的FP16和INT8。

2017-05-22 21:11:48 8356 3

原创 Doxygen学习

最近在阅读源码的时候学习了一个神器-Doxygen，安利给大家。顺便放出学习过程中查到的一些好资料。介绍：Doxygen是为C ++源代码生成文档的的强大工具，但它也支持其他流行的编程语言，如C，Objective-C，C＃，PHP，Java，Python，IDL（Corba，Microsoft和UNO / OpenOffice flavor ），Fortran，VHDL，Tcl，部分支持D语言.

2017-03-16 15:00:13 1042 1

原创安装NCCL问题：nvcc fatal : Value 'gnu++0x' is not defined for option 'std'

今天安装NCCL的时候碰见下面这个问题。本博客为解决方案。环境： CentOS Linux release 7.2.1511 CUDA 7.5 gcc version 4.8.5 20150623报错：nvcc fatal : Value 'gnu++0x' is not defined for option 'std'在Ubuntu 14.04上安装成功，调试Makefile后发现Ubu

2017-03-01 16:35:47 3098

原创 Fast Multi-GPU collectives with NCCL-翻译

Fast Multi-GPU collectives with NCCL-翻译本文是对NCCL博客的翻译，方便大家学习交流。因水平有限，如有问题欢迎更正。目前许多服务器包含8个或更多的GPU。理论上，将应用程序从一个GPU扩展到多个GPU应该提供巨大的性能提升。然而在实践中，性能提升往往很难达到理论值。这个问题往往是由两个原因导致。第一是分配任务时没有充分利用的并行性，从而导致

2017-02-28 20:48:38 5675

原创 CUDA进阶第四篇-如何实时获得GPU各种资源利用率（NVML）

CUDA进阶第四篇-如何实时获得GPU各种资源利用率前几天在玩digits的时候发现在训练的过程中网页上能够动态显示显卡利用率，显存利用率和温度。突然想起来前几天有人问过这个问题并查了一下目前好像并没有类似的中文博客，所以就探究了一下这个问题。首先我去翻了一下digits的源码，发现是用python调用了NVML(nvidia-management-library)这么个库。原来NVI

2016-08-11 12:45:51 18921

原创写CUDA必备sublime text插件：CUDA Snippets

CUDA自NVIDIA提出以来，经过多年的发展，俨然已经成为高性能计算的首选。编辑器不够成熟，代码调试工具上手难度较高，出现bug千奇百怪，CUDA开发一直让并行程序员们“痛不欲生”。本人写了几年CUDA，也是深受其害。既然等不到NVIDIA出IDE，那就自己挖坑了。写CUDA一直用的编辑器是sublime text，目前仅有的一个跟CUDA相关的插件是CUDA C++，但已经很长时间没有更新了

2016-02-15 22:03:02 4563

原创 CUDA进阶第三篇：CUDA计时方式

写CUDA，追求的就是加速比，想要得到准确的时间，计时函数就是必不可少计时通常分为两种情况，（1）直接得到接口函数的时间，一般用于得到加速比；（2）获得接口函数内核函数、内存拷贝函数等所耗时间，一般用于优化代码时。情况（1）方法有两种，CPU计时函数和GPU计时函数。情况（2）有三种工具nsight，nvvp，nvprof本博客会详细介绍情况（1）的两种方法；情况（2），nsight不会用，简单介绍一下nv

2016-02-04 00:02:32 14441 3

原创 CUDA进阶第二篇：巧用PTX

并行线程执行（Parallel Thread eXecution，PTX）代码是编译后的GPU代码的一种中间形式，它可以再次编译为原生的GPU微码。利用PTX来进行试验，我们可以解决一些在写代码时遇到的不确定问题。

2016-01-12 18:33:39 19513 7

原创 CUDA进阶第一篇：CUDA调试

0. 写在前面 “初学CUDA，好不容易自己写完一段cuda代码，一运行，满屏的语法bug，语法bug还好说，竟然还有逻辑bug，逻辑bug怎么改啊，wtf！！” “从别人手里接到一段CUDA代码，WTF，为什么还有bug！！还没有注释！！没有文档！！写代码的人怎么不去死啊！！” 同事的代码出bug了，找你调bug，内心独白：“tmd这写的都是啥”，“tmd这不是我写...

2016-01-05 15:02:01 34275 1

原创 comparing app performance to maximize theoretical limits

占坑+2

2015-12-15 10:57:34 497

原创 GPU体系结构-the miss status holding register（MSHR）原理简介

最近在看论文的时候里面有一段介绍MSHR的原理，以后有可能用做这方面的东西，先写出来。英文原文On a cache hit, a request will be served by sending data to the register ﬁle immediately. On a cache miss, the miss handling logic will ﬁrst check the mi

2015-12-08 22:36:14 6383

原创 GPGPU-Sim学习（代码解读）-如何实现Bypass L1D cache

GPGPU-Sim bypass

2015-12-05 11:28:24 3547 6

原创 OpenCUDA-基于CUDA的图像并行算法开源程序库

OpenCUDACUDA(Compute Unified Device Architecture)，是显卡厂商NVIDIA推出的运算平台。随着GPU的发展，CUDA使用人数也越来越多。但关于CUDA的库基本都是不开源的，目前手里有上一个项目遗留下的一套图像处理代码，准备整理整理，一步一步的开源出来。想分享出来和大家一起学习。也希望各位CUDA大牛能够奉献自己的源码。源码地址：https

2015-12-02 09:35:47 4921 3

原创举办cuda竞赛-有感

今年中秋和国庆时间较近，群里就举办了一个CUDA算法加速比赛-欢庆中秋喜迎国庆，2015双庆杯CUDA算法竞赛，作为群里唯二的管理员（另一个是风大有木有！），自然就成为组织者之一。经过一番折腾，比赛终于在昨天落下帷幕。故写篇有感博客，写一下自己的经验和感想比赛说明欢庆中秋喜迎国庆，2015双庆杯CUDA算法竞赛参与方式，以小组为单位，最少2人，最多3人，组员必须是本群群友。不管你是大咖

2015-11-26 18:44:14 1359

原创 GPGPU-Sim（番外）-如何加快GPGPU-Sim的运行速度

序言GPGPU-Simer都应该体验过Sim在虚拟机中奇慢无比的速度，一个在真实环境下几ms就跑完的benchmark，在Sim下通常需要几十分钟到几小时。跑一次实验benchmark多了基本就得一天，比较浪费时间。等待总是无聊的，本篇博客将给出我自己总结的几种加快Sim运行速度的解决方案。方案一：堆硬件ssd，加内存，上好CPU，使劲堆，有钱任性方案二：减少Sim输出输出信息打印出来是很耗时间

2015-11-24 15:20:24 2617

原创 GPGPU-Sim（番外）-GPGPU-Sim的模拟机制

这是我本科毕设内容的一部分，内容是一篇论文中的一部分，仅供学习，欢迎拍砖如图，我们展示了CUDA应用程序在GPU和GPGPU-Sim下的编译运行过程。 CUDA应用程序分为device代码和host代码。当在GPU上编译CUDA应用程序时，nvcc编译器会将device代码和host代码区分开。device代码会被编译成ptx，再由ptxas编译成cubin.bin文件，最后再和host代码、l

2015-11-24 14:59:09 2144

原创 OpenCV如何实现透明（alpha channel）图像的读取和写入

http://docs.opencv.org/modules/highgui/doc/reading_and_writing_images_and_video.html http://blog.csdn.net/wangyaninglm/article/details/41685607

2015-05-02 10:49:16 22281

原创拷贝global memory,cudaMemcpyToSymbol 和cudaMemcpy函数是否有区别

这是今天在群里有人问了这么一个问题 cudaMemcpyToSymbol可以将数据从host拷贝到global，cudaMemcpy也是从host到>global，这种情况下二个函数有什么区别吗？和各位大佬讨论一下后，和大家分享一下~首先，学到了cudaMemcpyToSymbol竟然还有将数据从host拷贝到global的功能，以前只用过这个函数拷贝constant memory。拷贝方式

2015-04-14 19:56:20 7684 1

原创乱七八糟-局域网内解决路由器无法支持ipv6问题……

情景：为了访问局域网资源，实验室内用路由器搭建了局域网。问题来了……外网支持ipv6，买的路由器不支持ipv6，ipv6一过路由器就没了。在网上找了很久，网上大部分的场景是在寝室为了省钱……目前这个场景下没找到解决方式。解决办法一换一个支持ipv6的路由器……这是最直接的解决方式解决办法二比较折中……现在好一点的路由器基本都支持vpn了。所以可以将电脑连接到外网上，正常使用ipv6，当需要使用内网

2015-04-03 21:11:16 9944

原创 GPGPU-Sim学习-如何禁用GPGPU-Sim内的L1 cache

最近在做一个观察实验，把L1 cache禁用以后，看看benchmark的IPC。在Sim 文档中并没有禁用L1 cache的描述，但是Sim的代码中却考虑到了这一因素，所以我们可以直接修改几行代码达到禁用L1 cache的效果。修改代码如下即可： gpu-cache.h class cache_config 将 bool m_disabled; 属性修改为 public。 shad

2015-03-29 10:33:00 1559 1

原创 GPGPU-Sim学习（二）(补充)gpgpu-sim官网所给虚拟机与自己所搭建虚拟机的使用感受

最近GPGPU-Sim的官网竟然给出了自己的虚拟机，这绝对是GPGPU-Sim学习者的最大福利了。因为安装sim的周期一般需要一个周，需要反复尝试很多次（20+），而且就目前我所遇见的simer，花这么多时间，50%安装失败了。好吧，一共只有六个人而已，可见这个仿真器在国内是多么的小众。偶然发现官网上的虚拟机，就赶紧下下来用了一下~下面是一些自己的使用感受类别官方所给虚拟机自己配

2015-03-24 21:35:11 4233 8

原创 GPGPU-Sim学习（二）搭建GPGPU-Sim环境（ubuntuServer 10.04 安装GPGPU-Sim）

最近发现在sim的官网那个已经给出了安装好gpgpu-sim的虚拟机，也是大家的福音。请戳官网链接http://gpgpu-sim.org/，我的写一篇《GPGPU-Sim学习（二）(补充)gpgpu-sim官网所给虚拟机与自己所搭建虚拟机的使用感受》本阶段的主要任务就是安装GPGPU-Sim，安装步骤是从带我那个学长写的一篇日志里，加上我自己的安装过程中的一些体会。学长日志链接：http:/

2015-03-18 21:47:45 4370

原创 GPGPU-Sim学习（一）-GPGPU-Sim介绍

本部分内容分为以下几个部分：概括性地介绍GPGPU-Sim、学习GPGPU-Sim的建议步骤、部分资料概括性地介绍GPGPU-SimGPGPU-Sim，通过单线程方式模拟出GPU多线程运行效果，并输出统计数据，是当前比较流行的GPGPU模拟器。目前国内外相当一部分关于GPGPU的研究都是使用的GPGPU-Sim。GPGPU-Sim是用C语言实现的，使用者可以通过合理修改GPGPU-S

2015-03-18 11:53:50 5731

原创 GPGPU-Sim学习（零）-写在前面

博客再一次搬家（好吧，其实是自己始终没坚持下去）。迁到CSDN上的原因一是这里有关于CUDA的东西（以前在科学网和博客园上都木有），还有一个就是可以用Markdown。加上新学期，又有三分热血写博客了。在做GPGPU-Sim的这段时间内，经历了很多坎坷，自己也慢慢脱离了菜鸟阶段。也意识到以前想写《 GPGPU-Sim研究》系列是多么的无知，所以改为《 GPGPU-Sim学习》。这段时间内也碰见了

2015-03-17 20:26:14 1708 1