- 博客(216)
- 收藏
- 关注
RSS订阅
原创 End-to-End Object Detection with Transformers[DETR]
End-to-End Object Detection with Transformers[DETR]背景概述相关技术背景最近在做机器翻译的优化，接触的模型就是transformer， 为了提升性能，在cpu和GPU两个平台c++重新写了整个模型，所以对于机器翻译中transformer的原理细节还是有一定的理解，同时以前做文档图片检索对于图像领域的目标检测也研究颇深，看到最近各大公众号都在推送这篇文章就简单的看了一下，感觉还是蛮有新意的，由于该论文开源，所以直接就跟着代码来解读整篇论文。概述整体
2020-06-08 18:06:43
9806
18
原创 cuBLAS矩阵乘法
cuBLAS是cuda封装好的一个数学库，头文件为<cublas_v2.h>#define cublasSgemm cublasSgemm_v2CUBLASAPI cublasStatus_t CUBLASWINAPI cublasSgemm_v2( cublasHandle_t handle, cublasOperation_t transa, cublas...
2020-04-03 21:05:55
2575
2
原创 6. 设计模式之《单例模式》
单例模式，顾名思义就是只有一个对象，考虑一下什么时候只需要一个对象？常见比如日志对象，这种最好就是一个项目中一个对象。class Log{public: Log* getLog() { if(nullptr == m_log) log = new Log(); return m_log; }private: st...
2020-03-30 23:19:20
239
原创 24_gemm_grouped
最朴素的思想是一个gemm起一个kernel，然后循环就可以n次就行，但是有些gemm尺寸比较小，这样的会sm无法都利用，为此用多stream来启动，这样就可以保证多个kernel可以一块执行，但是效果不是很好，可能是因为多次启动的原因。为此需要继续优化思路，这里利用。
2025-03-20 18:26:47
133
原创 moe原理和优化
可以发现这三个矩阵的A矩阵尺寸不一样，不能用batched gemm，所以这里需要用grouped gemm来做，接下来可以看看为什么grouped gemm比batch gemm效果好。最近deepseek比较火，导致moe也比较热门，这里简单看看原理以及如何优化。1、首先每个token经过线性层和softmax选择自己的专家id。3、最后再根据不同专家的权重做一个合并。2、然后根据id选择不同的矩阵参数乘。
2025-03-18 14:16:41
180
原创 vllm入门
最近deepseek属实又被炒了一波，作为做AI工程的从业者，有必要撸一遍vllm看看，熟悉一下目前llm的推理部署情况，核心看看page attention以及分布式的实现。
2025-02-08 13:58:56
126
原创 Iterator
A矩阵举例，A是row major, 这里kContiguous =K，kStrided =M, 确实是的，连续数据的长度是k, 对于B矩阵，连续的值长度就是N，这里整明白了后可以看看其他DefaultMmaCore实例化的结果，验证发现是这个意思。这里A矩阵和B矩阵都是row-major，但是我们看到传入PitchLinearShape的尺寸和我们传统的行列顺序不一样，先传入的是列，然后才是行，为啥？
2025-01-21 17:58:35
356
原创 08_turing_tensorop_gemm
对于tensorcore的具体使用，需要配合文档才能使用，因为这里的mma是warp为概念的，对于每个线程只要做搬运数据搬运工作就行，程序员的作用就是根据文档，让每个线程搬运指定位置的数据到寄存器就行，得到的数据再放到指定位置就行，nv这块文档给的不详细，后来推荐用wmma的api，想搞细节可以看看tensorcore,不想搞的话用用wmma也行。和之前相比，区别在于做warpgemm的时候，上面的小绿色块是整个warp去计算得到的，之前是一个线程做。
2025-01-17 15:33:25
304
原创 00_basic_gemm
可以看到，编译期时候，程序员必须要定下输入矩阵的layout和数据类型。事实上真的是这样吗?我们来深究一下这个cutlass::gemm::device::Gemm，从这个名字就可以看出来，cutlass实现了一个gemm，有device, threadblock, warp, thread几个级别gemm，这个sample里面用的是device级别, 所谓的device级别就是在cpu端的代码可以调用的，这个其实和cub中的逻辑是一样的。这里研究的cutlass版本是3.5。
2025-01-07 18:34:02
385
原创 第二章、数值
无符号和有符号整数二者static_cast转换时候，底层二进制值不变，只是应用层解释不一样扩展时，无符号高位补0, 有符号高位补符号位，举个3bit的int a=-2(110=-4+2=-2)，扩展为4bit时候(1110=-8+4+2=-2)
2024-12-11 22:13:01
232
原创 右值引用使用说明
可以看到唯一的区别就是，一个调用了移动构造函数（这个性能会好很多），一个调用了copy构造函数，对于BaseClass而言，基本都是如下实现, 其实核心就是交接堆区域的内存，由于左值引用只能针对已经存在的对象，对于没有持久化地址、通常是临时对象、字面量、或在表达式中求值之后不再需要的对象无法使用，右值引用就是做这个事情的。比如一个函数的返回值是Type&& foo()这就代表了返回的是一个右值。比如一个函数foo(Type&& par)
2024-12-10 15:48:56
177
原创 cuda中的两种默认stream
本质上per-thread是为了多线程设计的流模式，当面对多线程的进程应用的时候，如果每个线程都使用的是legacy stream的时候，那么多线程的意义其实就失去了，因为在gpu上还是串行的。进行编译nvcc --default-stream per-thread ./pthread_test.cu -o pthreads_per_thread 可以得到如下结果，可以看到legacy就一个stream, 但是依然把整个进程中的所有流给同步了。还是上面的代码，我们先看看如何启动per-thread。
2024-10-14 15:12:14
1392
原创 多线程中的一些概念
可以理解为一个通行令牌，当有一个线程A拿到后，其他线程只有等待，等到A通行结束后，会把令牌归还，其余线程自己去抢。c++20之前是没有这个玩意的，下面可以用互斥量和条件变量去是现实一个。和二元信号量的区别是，互斥量是线程A获取，必须A去释放，而二元信号量是进程里所有的其他线程都可以去抢夺释放。这里使用场景是一定可以保证函数B先执行，所以可以看出二元信号量可以用作线程之间通信或者同步使用。可以理解为多个通行令牌，理解和上面一样，无非是一次可以有多个线程通行。
2024-10-11 14:26:27
266
原创 union不能被初始化由于有 non-trivial构造函数
如果类的成员是具有上述非平凡构造函数、析构函数或赋值操作符的类，则该成员也被认为是非平凡的。换句话说，如果一个类包含非平凡成员对象，那么这个类的构造函数、析构函数或赋值操作符也将是非平凡的。其实在union中的成员函数nt， 虽然是非平凡成员，但是构造函数和析构函数都有，为啥union不能直接调用相关函数生成一个默认的构造函数和析构函数呢？赋值操作符用于将一个对象的值赋给另一个对象。如果一个类从虚基类继承，那么这个类的构造函数、析构函数和赋值操作符也会变成非平凡的，因为需要处理虚基类的初始化和清理。
2024-08-27 17:44:26
1050
原创 C++对象初始化
以后要养成用大括号初始化的习惯，不然一堆坑！（这里推荐A a{1}, 而不是A a={1}, 两者还是有一些区别的，后续遇到问题再补充吧）
2024-08-22 17:50:27
304
原创 clang 编译cuda原理
最近在看一门julia的语言，里面是原生支持cuda的，不过在国产卡上却无法适配，为了开展工作有必要了解非常清楚整个编译的机制。此外在研究过程中发现openai 的triton，以及tvm等一些ai框架对nvidia的支持原理都及其类似，所以了解原理更加有必要。
2024-07-31 17:43:35
718
原创 cuda中的cooperative_groups
以前block内部的同步是用syncthreads(), block之间没用提供同步的接口，因为也合理，毕竟如果block太多的话，block_n要等block_0算完退出后才会进入sm, 但是block_0为了同步又要等block_n,这样就是锁死了，因为gpu的逻辑和cpu不一样，gpu单个block寄存器的值不会暂存到显存里来切换block_0,那合格api咋用。最近看到一个代码cooperative_groups.this_grid().sync()很好奇，这里好好梳理一下。
2024-07-17 11:16:30
368
原创 第四章、处理器体系结构
为什么要学习这一章？作为高性能计算的从业人员，对硬件知识如数家珍是基本的职业素养，而且硬件设计中有很多不错的优化思想值得借鉴，比如pipeline思想，总之，学习该章百利无一害，何乐不为，不过对于非电子行业看起来可能比较难，因为有数字逻辑电路和一些电路的名词，比如上跳沿，晶振这种，不过这章都是综诉性质的，不会深入晶体管PN节啥的，总体可学。与x86_64的汇编语言类似，没有特别难得地方。
2024-06-25 14:01:02
298
原创 cufftPlanMany参数说明
最近在看cufft这个库，传统的cufftPlan3d()这种plan接口逐渐被nvidia舍弃了，说是要用最新的cufftPlanMany，这个函数呢又依赖一个什么Advanced Data Layout()，最终把这个api搞得乌烟瘴气很难理解，为了理解自己写了一些测试来验证各个参数的意思，这里简单做一下总结。下面是函数声明以及对应的参数解释，看不懂的话可以结合后面的例子琢磨琢磨。图示如下，这里没有画输出，意思和输入一样，接下来做个实验试试。
2024-05-08 17:44:37
877
原创 第三章、汇编2
chat-gpt解释如下：其实jmp分为相对短跳转（short jump）和相对长跳转（near jump），二者的汇编虽然都是用jmp这三个字母，但是汇编器会根据输入立即数的大小翻译成不同的二进制, 指令分别被翻译EB(短跳转) 和 E9（长跳转）， 二者立即数范围分别是1字节（-128~+127）和4字节 （-2,147,483,648 ~ +2,147,483,647）。这个好处是可以利用一些简单的加法和乘法，减少指令个数。除法和乘法差不多，不过只有一个操作数，分子是128bit,分母是64bit。
2024-04-23 23:40:13
287
原创 第三章、汇编1
g：这是另一个编译选项，用于生成调试信息。使用 -g 选项编译程序会在生成的可执行文件中包含调试信息，包括源代码中的行号、变量名等，以便在调试器中进行调试时能够准确定位到源代码的位置。-g 选项通常与其他优化选项一起使用，以便在优化后的代码中进行调试。，以尽量保持生成的代码易读易调试。但是，它并不会进行像 -O1、-O2 或 -O3 这样的更加激进的优化。因此，当你希望在调试时仍然能够方便地查看源代码和变量名时，可以使用 -Og。-Og 的含义是“优化级别为 g”，其中的 “g” 代表了"g优化"。
2024-04-23 22:14:17
1681
原创 cutlass之基础类型
矩阵的layout, 其实就是把coord包装一下，主要看看两者的区别, 其实layout中主要就是包含上面说的一个stride, 也确实，一旦固定是行优先还是列优先后，只要定了stride后就固定了格式。W], 至于为啥是这个顺序我也不知道，估计是约定俗成的。理解步长其实就是在内存中，N0距离N1的长度是C。不是很明白为啥RowMajor是class, 而其他类型是struct？真是让我强迫症犯了。
2024-04-07 17:49:39
443
1
原创 cutlass序言
最近自己的服务器终于搞定了，虽然显卡是3060Ti, 但是基本目前nv主流的库和工具都支持，所以后续开始稳定研究cuda了，nv的库用过很多，目前玩过cub, thrust, cudnn，tensorrt等，但是只有少部分开源，比如cub和cutlass, cub是开源的但是都是一些经典并行算法，自己也移植过一个sort算法收获颇丰, cutlass更多是一些AI相关的算法，为了紧跟大模型时代，所以后面准备做个博客系列来分享学习经历，一方面记录，一方面看看能不能给国产GPU的高性能工程师一点优化参考，毕竟这
2024-03-27 14:19:50
251
原创 unique_ptr使用说明
指针问题一直是一个比较麻烦的事情，比如很多人说要用智能指针完全替换掉裸指针，有人说要用unique_ptr, 有人建议shared_ptr,可是实际看各种经典框架，发现一个框架什么指针都有，使用的方法也是无法八门，这里简单说一下unique_ptr的使用场景和注意事项。unique_ptr核心目的是：1.所有权唯一。2.避免所有人忘记delete。但凡符合这两个要求就可以建议使用unique_ptr。
2024-03-12 15:22:30
437
原创 CMakeLists.txt简单demo
这里需要主要一个-O3 -DNDEBUG, 我们知道-O3是编译选项，但是-DNDEBUG是宏定义，为啥都扯到FLAGS里面了，其实是我们用了set(CMAKE_BUILD_TYPE “Release”), cmake直接把“-O3 -DNDEBUG”这个字符串打包给了target_compile_options, 理论也不该这么写，但是-O3和-DNDEBUG一般同时出现，所以软件为了方便打包放一起了。这样做的目的可能是为了确保在链接阶段也使用相同的优化选项和宏定义，以保持整个编译过程的一致性。
2024-02-22 14:11:18
946
原创 第三章、汇编语言
大部分人对数值信息都了解七七八八，但汇编了解的人却不多，虽然我学过王爽的《汇编语言》，但那个时候对计算机理解不深刻等于没学，目前在工作中遇到很多相关问题每次查起来很麻烦，这次正好借csapp再系统的梳理一遍，顺便利用拆弹作业做个巩固。
2024-01-09 13:57:13
405
原创 第二章、信息的表示和处理
在实际编程中，依然对有些数值的处理和变换比较模糊，在看csapp的时候发现里面的论述很详细，常规问题不在此赘述，这里主要是记录大部分人的知识点盲区。参考：https://blog.youkuaiyun.com/feng__shuai/article/details/79111084。
2024-01-07 15:24:52
470
原创 CMake项目管理
假如项目名称叫project， 一般可以按照下面的方式组织代码，这里可以看到include里面又补了一个项目名字，这个目的是啥？目的是在一个project避免不同模块中里有相同的头文件，比如下面的A.h。目前看到很过很多框架，很好奇大家如何从头搭建一个C++的库，这里简单介绍一个基本模板。
2023-12-21 22:57:43
451
原创 find_package 和 find_library的区别
经常看CMakeLists.txt中有find_package和find_library，有时候没留意以为都一样，其实二者差距比较大，下面简单记录一下。find_package(NAME), 这段代码的本质就是在找一个NAME.cmake这个文件，一般在安装库的时候，会随行带一个文件NAME.cmake安装在系统的cmake文件里。https://zhuanlan.zhihu.com/p/631259689
2023-12-14 12:21:46
979
原创 realname，soname和linkname
其实soname设置为A.so也是可以的，这样的话，如果我更新为11的大版本后，理论上厂家A还是可以运行可执行文件，但是我版本变动太大，A在实际生产中会报错，为了把风险控制在可执行文件启动前，所以可执行文件的soname最好有版本号，这样程序会启动失败，提醒你却是版本10的so。当我发布一个动态库的时候，比如版本是maj.min.patch（10.2.1）的格式，当我改了小版本的号的时候（10.3.1），如果厂家A用的是我写的库，他需要做什么？
2023-11-26 23:16:28
524
空空如也
空空如也
TA创建的收藏夹 TA关注的收藏夹
TA关注的人