- 博客(234)
- 收藏
- 关注
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
9987
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
2734
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
272
原创 nano-vllm-8
摘要：Speculative Decoding是一种加速推理的方法，使用小模型（草稿模型）先预测多个token（如"我是个大好人"），再交由大模型（目标模型）验证这些token的概率值，决定是否采纳。该方法将原本需要多次大模型前向推理的过程，优化为少量小模型推理+一次大模型验证，显著提升了推理效率。（99字）
2025-10-31 14:20:46
165
原创 AWQ_3
这篇文章介绍了对神经网络权重进行量化缩放后的处理方法。代码展示了如何根据不同的计算版本（gemm/gemv/marlin等）选择合适的量化线性模块，将原始线性层转换为量化版本。关键步骤包括：对缩放因子和零点值进行转置处理，根据计算版本选择对应的量化模块类，最后用量化后的线性层替换原始层。该过程支持不同位宽和分组大小的量化，适用于各种模型架构。
2025-10-27 17:46:20
229
原创 AWQ_2
本文介绍了AWQ（Adaptive Weight Quantization）量化过程的代码实现。代码首先处理输入数据的位置编码和设备迁移，然后针对Transformer模型进行分层量化。主要步骤包括：提取线性模块、计算并应用缩放因子、应用剪裁（可选）以及量化权重。作者发现nano-vllm自定义的Qwen3模型与标准Transformer在rotary_emb结构上存在差异，并通过截图展示了两种模型结构的对比情况。代码中包含了版本兼容性处理，特别是针对Transformer 4.48.0及以上版本的特殊处理
2025-10-23 15:38:47
133
原创 AWQ_1
因为这批数据首先是为了确定如何放缩哪几列的weight）后, 然后开始获取每个矩阵的输入X，算法里叫做激活值，反正爱怎么叫就怎么叫吧。在拿到校准数据（其实我理解这玩意不叫。
2025-10-23 14:40:39
183
原创 AWQ量化
本文介绍了使用AWQ进行模型量化的实践指南。主要内容包括：1)环境配置，需安装特定版本transformer(4.51.3)并设置单卡运行；2)代码示例展示如何加载Qwen3-0.6B模型，配置4bit量化参数，并执行量化过程；3)注意事项指出目前AWQ支持的模型有限，且需处理版本兼容性问题。量化过程支持自定义校准数据集，最终生成量化模型可保存至指定路径。该方案适用于需要模型轻量化的应用场景，但需注意硬件和版本适配问题。
2025-10-22 16:36:28
359
原创 nano-vllm-6
文章摘要： CUDA Graph通过缓存kernel启动前的准备工作，显著减少重复执行的耗时。在nano-vllm中，针对decode阶段（输入长度固定但batch size可变）预先生成了多个不同batch size版本的graph。实际运行时通过padding到最近的batch size版本进行推理。系统为1-512的batch size准备了多组graph，运行时选择最接近的graph执行。虽然KV cache位置会变化，但由于输入指针固定且通过block_table管理显存访问，CUDA Graph
2025-10-16 16:55:55
330
原创 nano-vllm-1.5
这个框架核心就是启动多个进程，每个进程维护一个模型运行器（ModelRunner）, 主进程的调度器（Scheduler）统筹(BlockManager)好显存（kv_cache）和输入(Sequence)后, 然后让每个进程的模型运行器做推理。
2025-10-14 17:29:45
213
原创 nano-vllm-5
这篇主要介绍一下attention算子，因为前面反复涉及到flash-attention的api, 而且很多prepare_prefill和prepare_decode都是为这个api在服务，所以有必要了解.
2025-10-11 16:44:05
300
原创 nano-vllm-4
该代码实现了一个分布式模型运行器ModelRunner，主要用于多GPU环境下的大模型推理任务。主要功能包括：1) 通过共享内存和事件机制实现多进程通信；2) 提供模型预热(warmup)功能以优化显存使用；3) 管理KV缓存分配，支持多卡并行计算。核心方法包含初始化、退出处理、循环执行、内存读写以及KV缓存分配等。代码支持Qwen3等大模型，通过分布式处理提高推理效率，并利用pickle进行进程间数据序列化传输。特别值得注意的是其显存管理机制，包括峰值统计重置、缓存分配策略等，确保GPU资源的高效利用。
2025-10-10 17:36:56
200
原创 nano-vllm-2
以上就是所有的squence源码，本质就是给token ids配了一些简单的属性值， 接下来看看scheduler。sequence本质就是把原来的token序列附加一些额外的属性而已，仅此而已。
2025-09-26 17:31:06
254
原创 nano-vllm-1
本文分析了LLMEngine.py文件的核心功能。该文件定义了LLMEngine类，主要包含请求处理、推理执行和结果生成等功能。关键方法包括：exit()用于进程终止管理，add_request()处理输入请求，step()执行推理步骤，generate()实现完整生成流程。类通过调度器管理请求序列，调用模型执行推理，并支持多进程协同工作。文章特别说明了主进程通过共享内存和事件机制控制其他进程退出的细节。
2025-09-25 18:12:50
274
原创 nano-vllm-0
大模型训练环境配置摘要 本文介绍了大模型训练所需的环境配置要求。硬件方面建议至少配备两块NVIDIA显卡，每卡显存8GB以上，内存至少16GB。软件环境需要安装PyTorch和flash_attention库，其中flash_attn推荐直接下载兼容的whl包安装以避免编译问题，例如文中使用的flash_attn-2.8.3版本。该配置特别针对后续需要讲解的张量并行技术场景，强调显存和内存越大越好。
2025-09-24 18:04:35
404
原创 大模型推理并行
目前大模型的参数以及计算量越来越大，如果放在多卡上处理成为关键，这里简单记录一下每种并行策略的概念。目前大模型核心就是gemm、FFN(MLP)、attention, 所以下面的说明也以这三个算子作为说明。每个gpu上储存一份模型参数，通过切分batch来实现并行推理。
2025-08-21 17:31:35
340
原创 move-only 类型
摘要：C++中的move-only类型是指只能移动不能复制的类型，常见于管理唯一资源的场景。标准库中典型的move-only类型包括：std::unique_ptr（独占指针）、std::mutex（互斥锁）、std::unique_lock（锁管理）、std::thread（线程对象）、std::promise/std::future（异步通信）。开发者也可自定义move-only类型，通过禁用拷贝构造函数/赋值运算符，只保留移动操作来实现资源的安全转移。这些类型的设计确保了资源的独占所有权和线程安全性。
2025-07-30 13:59:01
351
原创 C++初始化相关
摘要：C++初始化方式主要有小括号()、大括号{}和=三种。小括号初始化可能引发most vexing parse问题，且允许隐式窄化转换；大括号初始化能避免解析歧义，对窄化转换会发出警告（仅表达式传参有效），但会优先匹配initializer_list构造函数。initializer_list支持显式声明、auto推导和函数传参三种使用方式。不同初始化方式各有优缺点，开发者需根据场景谨慎选择。(148字)
2025-07-30 11:04:13
269
原创 24_gemm_grouped
最朴素的思想是一个gemm起一个kernel，然后循环就可以n次就行，但是有些gemm尺寸比较小，这样的会sm无法都利用，为此用多stream来启动，这样就可以保证多个kernel可以一块执行，但是效果不是很好，可能是因为多次启动的原因。为此需要继续优化思路，这里利用。
2025-03-20 18:26:47
1038
原创 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
860
原创 nano-vllm入门
最近deepseek属实又被炒了一波，作为做AI工程的从业者，有必要撸一遍vllm看看，熟悉一下目前llm的推理部署情况，核心看看page attention以及分布式的实现。
2025-02-08 13:58:56
399
原创 Iterator
A矩阵举例，A是row major, 这里kContiguous =K，kStrided =M, 确实是的，连续数据的长度是k, 对于B矩阵，连续的值长度就是N，这里整明白了后可以看看其他DefaultMmaCore实例化的结果，验证发现是这个意思。这里A矩阵和B矩阵都是row-major，但是我们看到传入PitchLinearShape的尺寸和我们传统的行列顺序不一样，先传入的是列，然后才是行，为啥？
2025-01-21 17:58:35
611
原创 08_turing_tensorop_gemm
对于tensorcore的具体使用，需要配合文档才能使用，因为这里的mma是warp为概念的，对于每个线程只要做搬运数据搬运工作就行，程序员的作用就是根据文档，让每个线程搬运指定位置的数据到寄存器就行，得到的数据再放到指定位置就行，nv这块文档给的不详细，后来推荐用wmma的api，想搞细节可以看看tensorcore,不想搞的话用用wmma也行。和之前相比，区别在于做warpgemm的时候，上面的小绿色块是整个warp去计算得到的，之前是一个线程做。
2025-01-17 15:33:25
403
原创 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
664
原创 第二章、数值
无符号和有符号整数二者static_cast转换时候，底层二进制值不变，只是应用层解释不一样扩展时，无符号高位补0, 有符号高位补符号位，举个3bit的int a=-2(110=-4+2=-2)，扩展为4bit时候(1110=-8+4+2=-2)
2024-12-11 22:13:01
278
原创 右值引用使用说明
可以看到唯一的区别就是，一个调用了移动构造函数（这个性能会好很多），一个调用了copy构造函数，对于BaseClass而言，基本都是如下实现, 其实核心就是交接堆区域的内存，由于左值引用只能针对已经存在的对象，对于没有持久化地址、通常是临时对象、字面量、或在表达式中求值之后不再需要的对象无法使用，右值引用就是做这个事情的。比如一个函数的返回值是Type&& foo()这就代表了返回的是一个右值。比如一个函数foo(Type&& par)
2024-12-10 15:48:56
232
原创 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
2057
原创 多线程中的一些概念
可以理解为一个通行令牌，当有一个线程A拿到后，其他线程只有等待，等到A通行结束后，会把令牌归还，其余线程自己去抢。c++20之前是没有这个玩意的，下面可以用互斥量和条件变量去是现实一个。和二元信号量的区别是，互斥量是线程A获取，必须A去释放，而二元信号量是进程里所有的其他线程都可以去抢夺释放。这里使用场景是一定可以保证函数B先执行，所以可以看出二元信号量可以用作线程之间通信或者同步使用。可以理解为多个通行令牌，理解和上面一样，无非是一次可以有多个线程通行。
2024-10-11 14:26:27
320
原创 union不能被初始化由于有 non-trivial构造函数
如果类的成员是具有上述非平凡构造函数、析构函数或赋值操作符的类，则该成员也被认为是非平凡的。换句话说，如果一个类包含非平凡成员对象，那么这个类的构造函数、析构函数或赋值操作符也将是非平凡的。其实在union中的成员函数nt， 虽然是非平凡成员，但是构造函数和析构函数都有，为啥union不能直接调用相关函数生成一个默认的构造函数和析构函数呢？赋值操作符用于将一个对象的值赋给另一个对象。如果一个类从虚基类继承，那么这个类的构造函数、析构函数和赋值操作符也会变成非平凡的，因为需要处理虚基类的初始化和清理。
2024-08-27 17:44:26
1296
原创 C++对象初始化
以后要养成用大括号初始化的习惯，不然一堆坑！（这里推荐A a{1}, 而不是A a={1}, 两者还是有一些区别的，后续遇到问题再补充吧）
2024-08-22 17:50:27
349
原创 clang 编译cuda原理
最近在看一门julia的语言，里面是原生支持cuda的，不过在国产卡上却无法适配，为了开展工作有必要了解非常清楚整个编译的机制。此外在研究过程中发现openai 的triton，以及tvm等一些ai框架对nvidia的支持原理都及其类似，所以了解原理更加有必要。
2024-07-31 17:43:35
1117
原创 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
613
原创 第四章、处理器体系结构
为什么要学习这一章？作为高性能计算的从业人员，对硬件知识如数家珍是基本的职业素养，而且硬件设计中有很多不错的优化思想值得借鉴，比如pipeline思想，总之，学习该章百利无一害，何乐不为，不过对于非电子行业看起来可能比较难，因为有数字逻辑电路和一些电路的名词，比如上跳沿，晶振这种，不过这章都是综诉性质的，不会深入晶体管PN节啥的，总体可学。与x86_64的汇编语言类似，没有特别难得地方。
2024-06-25 14:01:02
338
空空如也
空空如也
TA创建的收藏夹 TA关注的收藏夹
TA关注的人