allocator类来分配不同deive的内存，使用buffer类来管理内存，use_external判断是否由buffer管理，
buffer是智能指针，可以自动释放内存，析构函数去释放ptr
- base_forward 调用每个算子的前向计算
- 在base_forward中 get input and weights
- seletct kernel, 根据设备的类型选择算子的实现,返回函数指针
- 传入inputs and weights ，调用kernel的计算函数， 将结果返回给到output中
### course3-张量的设计和实现 张量：多维度数组，在推理流程中管理、传递数据，并结合Buffer类来自动管理内存或者显存资源
tensor释放的时候，buffer也会释放
步长为后续所有维度的乘积
基本公式如下所示:
分CPU和GPU两种实现，CPU实现使用armadillo库，GPU实现使用cuda。
GPU：
最小的执行单元是thread
最小的调度单元是warp，硬件会一次性把一个wrap放在就绪的硬件执行单元中。
执行的时候将输入的数据打包成float4，减少内存的访问次数，提高内存带宽利用率。
使用blockreduce
在做规约的时候，以32个线程为单位进行， 假设一个block有128个线程， 那么将其分为4份，每一份计算完成之后
保存到shared memory中，最后对这四个进行相加， 避免数据竞争重复计算。
Andrej karpathy 提供的权重dump工具， int8 weight only , group weight
- 使用transformers库加载llama结构的模型
- 从模型的配置config.json中构造模型参数
- 根据配置信息创建一个导出的模型
- 为导出的模型配置权重，权重来自huggingface的预训练权重
- 开始导出权重
常规计算例子：
float sum = 0.0f;
for(int i=tid; i<size; i+=blockDim.x){
sum += in[i] * in[i];
}每次读取四个数据，充分利用带宽
向量化存取例子：
float4* in_pack = reinterpret_cast<float4*>(in);
for (int i = tid; i < pack_num; i += blockDim.x) {
float4 in_float4 = *(in_pack + i);
sum += in_float4.x * in_float4.x;
sum += in_float4.y * in_float4.y;
sum += in_float4.z * in_float4.z;
sum += in_float4.w * in_float4.w;
}- 这样会提升内存和L2的吞吐率
- gpu运算的指令减少
- L2 cache的命中率提升
Allocator类里面实现内存分配和释放接口 CUDA：
- 调用cudaMalloc等接口有一定的耗时，设计buffer避免重复分配和释放，减少内存分配和释放的开销
- 对于小块显存，先不用cudafree释放， 先保存起来，以后用到的时候直接返回该内存块
- 设定空闲内存块大于一定阈值，才开始释放
- 以字节为单位打开文件，使用的时候直接按字节数量strncpy即可
- 按需加载数据
- 减少数据拷贝，直接将文件映射到进程的地址空间
权重文件格式:
--
dim, hidden_dim, layer_num ... 前面的28个字节
group_size ... 量化参数信息 (optional)
--
floa权重
--
- 用MMAP打开权重文件之后
- 计算权重的数量，通过维度累乘
- 将这块权重指针赋值给Buffer(不管理内存，由mmap自动映射)
- 将buffer实例赋值给层的权重
权重(主存) - > 算子weight(GPU)
下图为模型权重的文件，黄色区域代表权重的位置
算子后端:
- 根据传入的device_type 返回对应的kernel函数指针
CPU
- 调用armadillo库
- 矩阵内存复用
- armadillo是列主序，需要转置
GPU
- 规约计算，每个block负责计算乘法计算中的一行，一个block有多个wrap组成.
将k×dim维度的矩阵query拆分为两部分
- 包含0~dim-1 行的query1矩阵，维度为 （dim-1） × dim
- 第二部分是仅仅包含第K行的query2矩阵，维度为1 × dim。在进行自回归计算时，只需要计算query2矩阵×key矩阵即可。
K cache
- 将K矩阵分为k1和k2, k1为前面k-1个计算步骤所得到的结果，k2是当前步骤中所得到的结果
- k1 也就是之前的计算结果，直接缓存到K cache中，当计算到第K步时，直接从K cache中取出即可。
- k2 = input_token3 * W_k ，计算得到K2, 然后qeury 与k1+k2计算
显存计算：
memory = K(步长 or token长度) × dim(V的维度) × N(transformer的层数) × sizeof(float)
# News：新课发布，《动手自制大模型推理框架》，全手写cuda算子，课程框架支持LLama2和3.x Hi，各位朋友们好！我是 KuiperInfer 的作者。KuiperInfer 作为一门开源课程，迄今已经在 GitHub 上已斩获 2.4k star。 如今在原课程的基础上，**我们全新推出了《动手自制大模型推理框架》， 新课程支持Llama系列大模型（包括最新的LLama3.2），同时支持 Cuda 加速和 Int8 量化，自推出以来便广受好评。**
https://l0kzvikuq0w.feishu.cn/docx/ZF2hd0xfAoaXqaxcpn2c5oHAnBc
如果你对大模型推理感兴趣，想要深入了解并掌握相关技术，那么这门课程绝对不容错过。快来加入我们，一起开启学习之旅吧！ 感兴趣的同学欢迎扫一扫课程下方二维码或者添加微信 lyrry1997 参加课程
LLama1.1b fp32模型，视频无加速，运行平台为Nvidia 3060 laptop，速度为60.34 token/s
一、项目整体架构和设计
学习架构思维，防止自己只会优化局部实现
- 环境的安装和课程简介
- 资源管理和内存管理类的设计与实现
- 张量类的设计与实现
- 算子类的设计与实现
- 算子的注册和管理
二、支持LLama2模型结构
本节将为大家补齐算法工程师思维，在算法层面讲解大模型和Transformer的原理之后，开始对LLama2进行支持
- LLama模型的分析
- MMap内存映射技术打开大模型的权重文件
- LLama模型文件的参数和权重载入
- LLama中各个层的初始化以及输入张量、权重张量的分配和申请
- 实现大模型中的KV Cache机制
三、模型的量化
为了减少显存的占用，我们开发了int8模型量化模块
- 量化模型权重的导出
- 量化系数和权重的加载
- 量化乘法算子的实现
四、Cuda基础和算子实现
带你学Cuda并在实战大模型算子的实现，为大模型推理赋能
- Cuda基础入门1 - 内容待定
- Cuda基础入门2 - 内容待定
- Cuda基础入门3 - 内容待定
- Cuda基础入门4 - 内容待定
- RMSNorm算子的Cuda实现
- Softmax算子的Cuda实现
- Add算子的Cuda实现
- Swiglu算子的Cuda实现
- GEMV算子的Cuda实现
- 多头注意力机制的Cuda实现
- 让框架增加Cuda设备的支持和管理
- 完成Cuda推理流程
五、用推理框架做点有趣的事情
- 文本生成
- 讲一段小故事
- 让大模型和你进行多轮对话
六、学习其他商用推理框架的实现，查漏补缺
-
LLama.cpp的设计和实现讲解
这里有多个小节
-
Miopen（AMD出品，对标CUDNN）的设计和实现讲解
这里有多个小节
-
总结
- google glog https://github.com/google/glog
- google gtest https://github.com/google/googletest
- sentencepiece https://github.com/google/sentencepiece
- armadillo + openblas https://arma.sourceforge.net/download.html
- Cuda Toolkit
openblas作为armadillo的后端数学库，加速矩阵乘法等操作，也可以选用Intel-MKL，这个库用于CPU上的推理计算
-
LLama2 https://pan.baidu.com/s/1PF5KqvIvNFR8yDIY1HmTYA?pwd=ma8r 或 https://huggingface.co/fushenshen/lession_model/tree/main
-
Tiny LLama
- TinyLLama模型 https://huggingface.co/karpathy/tinyllamas/tree/main
- TinyLLama分词器 https://huggingface.co/yahma/llama-7b-hf/blob/main/tokenizer.model
需要其他LLama结构的模型请看下一节模型导出
python export.py llama2_7b.bin --meta-llama path/to/llama/model/7B
# 使用--hf标签从hugging face中加载模型， 指定--version3可以导出量化模型
# 其他使用方法请看export.py中的命令行参数实例 mkdir build
cd build
# 需要安装上述的第三方依赖
cmake ..
# 或者开启 USE_CPM 选项，自动下载第三方依赖
cmake -DUSE_CPM=ON ..
make -j16./llama_infer llama2_7b.bin tokenizer.model