fp16 量化 如果显卡支持 fp16 运算加速，那么使用 fp16 量化能显著提升模型的推理速度，且由于 fp16 的表示范围为 (-65504, +65504)，一般都能包含模型的权重数值范围，所以直接截断对精度的影响很小。因此在 TensorRT 中实现 fp16 量化的方法相当简单，只需要在构建 engine 时添加一行配置即可 config = builder.create_builder_config() config.set_flag(trt.BuilderFlag.FP16) int8 量化 int8 量化能进一步压缩模型体积以及提升推理速度，但是 8bit 整数能表示的范围只有 -127~128，模型权重值一般都超过了这个范围。因此需要对权重进行 scale...

安装 TensorRT 的注意事项 TensorRT 的安装方式在官方安装导引 里已经有详细的说明了，这里提下需要注意的地方，首先是尽量保证 TensorRT 的安装方式 和 CUDA 的安装方式相同，否则可能会出现找不到 CUDA 的情况，比如在 Ubuntu 中，如果 CUDA 不是通过 deb 包安装的，后面用 deb 安装 TensorRT 就会报错找不到 CUDA 依赖。 从 PyTorch 导出 ONNX 模型 ONNX 是一种通用的神经网络模型交换格式，TensorRT 有专门的 ONNX 解析器，可以解析 ONNX 模型并创建 TensorRT 自有的网络结构并做后续的优化工作。因此，为了让 TensorRT 优化我们使用 PyTorch 训练的模型，可以先将 P...

算术强度的计算问题 前面我们在介绍算术强度的时候，非常理想化的把内存读取总量等价成了参与运算的数据总量。但实际上，由于寄存器数量限制，没法完全容纳全部的数据量，每个数据往往需要读很多次，因此实际的算术强度只可能比理论值更低。 图 1. 数据流动示意图 上图右边展示了实际的数据流动过程，每次加载只能读取一部分数据到寄存器，这部分数据算完之后，再读下一部分，由于算法的原因，两次计算之间有可能需要用到一些相同的数据，因此存在重复加载的现象。 共享内存的作用 通过上一节的分析我们可以看到，寄存器容量的限制会直接导致算法的实际算术强度达不到理论值，会对程序的性能造成不利影响。在没法无限制增加寄存器数量的情况下，一种折中的方案就是用一块速度比普通内存块，容量比寄存器大的特殊区域来做中间...

程序计时[1] 既然我们专注于高性能计算，那么程序的运行时间就是一个非常值得关注的指标，在 CUDA 中，提供了很有用的 API 来帮助我们对程序在 Device 端的运行时间进行统计。下面的代码展示了如何来使用这些工具： void vectorAdd(const float *a, const float *b, const int n, float* c) { float *d_a, *d_b, *d_c; cudaMalloc((void**)&d_a, n * sizeof(float)); cudaMalloc((void**)&...

项目结构 上一篇我们写了一个简单的小程序，使用一条命令就可以编译通过，但实际的项目结构要复杂得多，因此在本篇文章中，我们先介绍如何使用 VSCode 搭建一个 CMake 工程来进行 CUDA 开发，并集成 googletest 框架来做测试。 首先，我们的项目结构如下 |.vscode/ |--c_cpp_properties.json |include/ |--vector_add.h |src/ |--vector_add_kernel.cu |test/ |--CMakeLists.txt |--test.cpp |CMakeLists.txt 这里的 c_cpp_properties.json 主要是用来指引 VSCode，使其能够识别在源码中 include 的头...