对于3060这种桌面级的GPU来说，像yolov8s动不动就数十、上百帧的模型还是太小了，瓶颈很容易卡在模型以外的其他地方，本文为了更加明确地反映不同的操作对于推理速度的影响，将使用最大的yolov8x模型。
加速推理有两种途径，一种是降低推理过程中模型的精度，用精度换速度，另一种是通过CUDA并行加速一些预处理、后处理操作，至于替换激活函数这种更加狠的活本文暂且不涉及。
降低精度
FP32
TensorRT默认是使用TF32精度，可以在终端中配置如下环境变量来禁用TF32，强制使用FP32精度构建engine
export NVIDIA_TF32_OVERRIDE=0推理速度如下：
yolov8x : preprocess:6.69934ms; inference: 46.2639ms; postprocess:2.25514ms; fps:18.1099TF32
全名TensorFloat32，这是英伟达在Ampere架构中提出的一种数据类型，他比FP32小，却能表示和FP32同样范围的数据，简而言之就是加速版本的FP32。
在build.cu中配置
// 设置精度
config->setFlag(nvinfer1::BuilderFlag::kTF32);推理速度如下：
yolov8x : preprocess:6.3109ms; inference: 42.7947ms; postprocess:2.33598ms; fps:19.4395FP16
半精度浮点可显著提高模型推理的速度。
// 设置精度
config->setFlag(nvinfer1::BuilderFlag::kFP16);推理速度如下：
yolov8x : preprocess:6.11168ms; inference: 11.1036ms; postprocess:2.17895ms; fps:51.5618INT8
这里先要吐槽一下NV，官方的sample写得绕来绕去的，根本不是给初学者看的
从上述文档可以得知，如果要用INT8量化的话，需要构造一个类，继承IInt8EntropyCalibrator2、IInt8MinMaxCalibrator这种量化算法，类里面需要包含下面4个函数，同时配置config
getBatchSize()
getBatch(void* bindings[], const char* names[], int nbBindings)
readCalibrationCache(size_t& length)
writeCalibrationCache(const void* cache, size_t length)
config->setFlag(nvinfer1::BuilderFlag::kINT8);
config->setInt8Calibrator(calibrator.get());这部分代码我参考的tensorrtx
class Int8Calibrator : public nvinfer1::IInt8MinMaxCalibrator
{
public:
Int8Calibrator(int batchsize, int input_w, int input_h, const char* img_dir, const char* calib_table_name, const char* input_blob_name, bool read_cache = false);
virtual ~Int8Calibrator();
int getBatchSize() const noexcept override;
bool getBatch(void* bindings[], const char* names[], int nbBindings) noexcept override;
const void* readCalibrationCache(size_t& length) noexcept override;
void writeCalibrationCache(const void* cache, size_t length) noexcept override;
private:
int batchsize_;
int input_w_;
int input_h_;
int img_idx_;
std::string img_dir_;
std::vector<std::string> img_files_;
size_t input_count_;
std::string calib_table_name_;
const char* input_blob_name_;
bool read_cache_;
void* device_input_;
std::vector<char> calib_cache_;
};
static inline int read_files_in_dir(const char *p_dir_name, std::vector<std::string> &file_names) {
DIR *p_dir = opendir(p_dir_name);
if (p_dir == nullptr) {
return -1;
}
struct dirent* p_file = nullptr;
while ((p_file = readdir(p_dir)) != nullptr) {
if (strcmp(p_file->d_name, ".") != 0 &&
strcmp(p_file->d_name, "..") != 0) {
//std::string cur_file_name(p_dir_name);
//cur_file_name += "/";
//cur_file_name += p_file->d_name;
std::string cur_file_name(p_file->d_name);
file_names.push_back(cur_file_name);
}
}
closedir(p_dir);
return 0;
}
Int8Calibrator::Int8Calibrator(int batchsize, int input_w, int input_h, const char* img_dir, const char* calib_table_name,
const char* input_blob_name, bool read_cache)
: batchsize_(batchsize)
, input_w_(input_w)
, input_h_(input_h)
, img_idx_(0)
, img_dir_(img_dir)
, calib_table_name_(calib_table_name)
, input_blob_name_(input_blob_name)
, read_cache_(read_cache)
{
input_count_ = 3 * input_w * input_h * batchsize;
CUDA_CHECK(cudaMalloc(&device_input_, input_count_ * sizeof(float)));
read_files_in_dir(img_dir, img_files_);
}
Int8Calibrator::~Int8Calibrator()
{
CUDA_CHECK(cudaFree(device_input_));
}
int Int8Calibrator::getBatchSize() const noexcept
{
return batchsize_;
}
bool Int8Calibrator::getBatch(void* bindings[], const char* names[], int nbBindings) noexcept
{
if (img_idx_ + batchsize_ > (int)img_files_.size()) {
return false;
}
std::vector<cv::Mat> input_imgs_;
for (int i = img_idx_; i < img_idx_ + batchsize_; i++) {
std::cout << img_files_[i] << " " << i << std::endl;
cv::Mat temp = cv::imread(img_dir_ + img_files_[i]);
if (temp.empty()){
std::cerr << "Fatal error: image cannot open!" << std::endl;
return false;
}
cv::Mat LetterBoxImg;
cv::Vec4d params;
LetterBox(temp, LetterBoxImg, params, cv::Size(640, 640));
input_imgs_.push_back(LetterBoxImg);
}
img_idx_ += batchsize_;
cv::Mat blob = cv::dnn::blobFromImage(input_imgs_[0], 1 / 255.0, cv::Size(640, 640), cv::Scalar(0, 0, 0), true, false, CV_32F);
CUDA_CHECK(cudaMemcpy(device_input_, blob.ptr<float>(0), input_count_ * sizeof(float), cudaMemcpyHostToDevice));
// assert(!strcmp(names[0], input_blob_name_));
bindings[0] = device_input_;
return true;
}
const void* Int8Calibrator::readCalibrationCache(size_t& length) noexcept
{
std::cout << "reading calib cache: " << calib_table_name_ << std::endl;
calib_cache_.clear();
std::ifstream input(calib_table_name_, std::ios::binary);
input >> std::noskipws;
if (read_cache_ && input.good())
{
std::copy(std::istream_iterator<char>(input), std::istream_iterator<char>(), std::back_inserter(calib_cache_));
}
length = calib_cache_.size();
return length ? calib_cache_.data() : nullptr;
}
void Int8Calibrator::writeCalibrationCache(const void* cache, size_t length) noexcept
{
std::cout << "writing calib cache: " << calib_table_name_ << " size: " << length << std::endl;
std::ofstream output(calib_table_name_, std::ios::binary);
output.write(reinterpret_cast<const char*>(cache), length);
}IInt8EntropyCalibrator2
推理速度如下，加速了个寂寞
yolov8x : preprocess:6.21368ms; inference: 13.1919ms; postprocess:2.18218ms; fps:46.3226IInt8MinMaxCalibrator
改用MinMaxCalibrator后，加速效果显著
yolov8x : preprocess:6.52782ms; inference: 6.83574ms; postprocess:2.20803ms; fps:64.2195在构建engine时，trt给出了下面的输出：
大概意思就是有的层无法使用INT8，这部分使用的是FP32，可见下图的issue
可以试一试混合精度
FP16/INT8混合精度
修改build config，加入FP16精度
config->setFlag(nvinfer1::BuilderFlag::kINT8);
config->setFlag(nvinfer1::BuilderFlag::kFP16);
auto* calibrator = new Int8Calibrator(1, 640, 640, "./coco_calib/", "int8calib.table", "");
config->setInt8Calibrator(calibrator);推理速度如下：
yolov8x : preprocess:5.64279ms; inference: 6.26466ms; postprocess:2.20027ms; fps:70.8831CUDA加速后处理
目前的后处理流程是一个遍历+NMS，总共耗时大概是2.2ms，经过测量，遍历过程用掉了1.6ms以上的时间，现在尝试通过CUDA直接在GPU上处理这些数据。
快过年了，挖坑以后再埋