QT + YOLO26 + TensorRT + CUDA 预处理完整方案（含完整 warpaffine_kernel / preprocess 代码）

一、背景

在基于 QT + YOLO26 + TensorRT 的工业级目标检测部署中，模型推理经过 TensorRT 加速后通常已经很快，但整体 FPS 往往仍然上不去。
真正的瓶颈很多时候不在推理，而在 预处理：

• resize
• copyMakeBorder
• cvtColor
• convertTo

其中 convertTo(float_img, CV_32F, 1/255.0) 在高分辨率场景下尤其耗时。

因此，更合理的思路不是继续在 CPU 上抠细节，而是：

把预处理整体搬到 GPU，一次 kernel 做完。

---

二、整体链路

QT读取图像/视频
    ↓
CUDA预处理
    1. Letterbox
    2. 双线性插值
    3. BGR → RGB
    4. /255.0
    5. HWC → CHW
    ↓
TensorRT推理
    ↓
后处理（逆仿射还原）
    ↓
QT渲染结果

---

三、为什么框会漂移

框漂移通常不是单纯画框函数的问题，而是：

• GPU 预处理几何变换与 CPU 版不一致
• preprocess / postprocess 没用同一套仿射关系
• 像素中心没有对齐 OpenCV INTER_LINEAR

正确原则

阶段	使用矩阵
preprocess 构建	s2d
kernel 采样	d2s
postprocess 坐标还原	同一个 d2s

---

四、仿射矩阵定义

struct AffineMatrix {
    float value[6];
};

---

五、完整 CUDA 预处理代码

下面这段代码是完整可用版本，包含：

• warpaffine_kernel
• preprocess

特点：

• 补全仿射矩阵平移项
• 使用半像素中心对齐，贴近 OpenCV INTER_LINEAR
• 支持 BGR → RGB
• 支持 /255.0
• 输出 CHW

1. 完整 warpaffine_kernel

__global__ void warpaffine_kernel(
    uint8_t* src, int src_line_size, int src_width,
    int src_height, float* dst, int dst_width,
    int dst_height, uint8_t const_value_st,
    AffineMatrix d2s, int edge)
{
    int position = blockIdx.x * blockDim.x + threadIdx.x;
    if (position >= edge) return;

    int dx = position % dst_width;
    int dy = position / dst_width;

    // d2s = [a b c; d e f]
    float m_x1 = d2s.value[0];
    float m_y1 = d2s.value[1];
    float m_z1 = d2s.value[2];
    float m_x2 = d2s.value[3];
    float m_y2 = d2s.value[4];
    float m_z2 = d2s.value[5];

    // 更贴近 OpenCV 的半像素中心对齐
    float src_x = m_x1 * (dx + 0.5f) + m_y1 * (dy + 0.5f) + m_z1 - 0.5f;
    float src_y = m_x2 * (dx + 0.5f) + m_y2 * (dy + 0.5f) + m_z2 - 0.5f;

    float c0 = const_value_st;
    float c1 = const_value_st;
    float c2 = const_value_st;

    if (!(src_x < -1.0f || src_x > src_width || src_y < -1.0f || src_y > src_height)) {
        int x_low = floorf(src_x);
        int y_low = floorf(src_y);
        int x_high = x_low + 1;
        int y_high = y_low + 1;

        float lx = src_x - x_low;
        float ly = src_y - y_low;
        float hx = 1.0f - lx;
        float hy = 1.0f - ly;

        uint8_t v1[3] = { const_value_st, const_value_st, const_value_st };
        uint8_t v2[3] = { const_value_st, const_value_st, const_value_st };
        uint8_t v3[3] = { const_value_st, const_value_st, const_value_st };
        uint8_t v4[3] = { const_value_st, const_value_st, const_value_st };

        if (y_low >= 0) {
            if (x_low >= 0 && x_low < src_width)
                memcpy(v1, src + y_low * src_line_size + x_low * 3, 3);
            if (x_high >= 0 && x_high < src_width)
                memcpy(v2, src + y_low * src_line_size + x_high * 3, 3);
        }

        if (y_high >= 0 && y_high < src_height) {
            if (x_low >= 0 && x_low < src_width)
                memcpy(v3, src + y_high * src_line_size + x_low * 3, 3);
            if (x_high >= 0 && x_high < src_width)
                memcpy(v4, src + y_high * src_line_size + x_high * 3, 3);
        }

        c0 = hy * hx * v1[0] + hy * lx * v2[0] + ly * hx * v3[0] + ly * lx * v4[0];
        c1 = hy * hx * v1[1] + hy * lx * v2[1] + ly * hx * v3[1] + ly * lx * v4[1];
        c2 = hy * hx * v1[2] + hy * lx * v2[2] + ly * hx * v3[2] + ly * lx * v4[2];
    }

    // BGR -> RGB
    float r = c2 / 255.0f;
    float g = c1 / 255.0f;
    float b = c0 / 255.0f;

    // HWC -> CHW
    int area = dst_height * dst_width;
    float* pdst_r = dst + dy * dst_width + dx;
    float* pdst_g = pdst_r + area;
    float* pdst_b = pdst_g + area;

    *pdst_r = r;
    *pdst_g = g;
    *pdst_b = b;
}

2. 完整 preprocess

void preprocess(
    uint8_t* src, const int& src_width, const int& src_height,
    float* dst, const int& dst_width, const int& dst_height,
    cudaStream_t stream, float& scale)
{
    AffineMatrix s2d, d2s;

    scale = std::min(dst_width / (float)src_width,
                     dst_height / (float)src_height);

    float new_w = src_width * scale;
    float new_h = src_height * scale;

    float tx = (dst_width - new_w) * 0.5f;
    float ty = (dst_height - new_h) * 0.5f;

    // src -> dst
    s2d.value[0] = scale;
    s2d.value[1] = 0.0f;
    s2d.value[2] = tx;
    s2d.value[3] = 0.0f;
    s2d.value[4] = scale;
    s2d.value[5] = ty;

    cv::Mat m2x3_s2d(2, 3, CV_32F, s2d.value);
    cv::Mat m2x3_d2s(2, 3, CV_32F, d2s.value);
    cv::invertAffineTransform(m2x3_s2d, m2x3_d2s);
    memcpy(d2s.value, m2x3_d2s.ptr<float>(0), sizeof(d2s.value));

    int jobs = dst_height * dst_width;
    int threads = 256;
    int blocks = (jobs + threads - 1) / threads;

    warpaffine_kernel<<<blocks, threads, 0, stream>>>(
        src, src_width * 3, src_width, src_height,
        dst, dst_width, dst_height,
        114, d2s, jobs);
}

---

六、这段代码相对旧版修正了什么

1. 补上仿射矩阵平移项

错误写法常见是：

float src_x = m_x1 * dx + m_y1 * dy;
float src_y = m_x2 * dx + m_y2 * dy;

这样会丢掉 tx / ty，直接导致 Letterbox 偏移丢失。

修正后：

float src_x = m_x1 * (dx + 0.5f) + m_y1 * (dy + 0.5f) + m_z1 - 0.5f;
float src_y = m_x2 * (dx + 0.5f) + m_y2 * (dy + 0.5f) + m_z2 - 0.5f;

---

2. 增加半像素中心对齐

这一步是为了尽量贴近 OpenCV INTER_LINEAR 的采样行为。
如果少了这一步，图像整体会产生系统性偏差，检测框容易漂移。

---

3. 保持颜色和布局与 CPU 版一致

输出仍然是：

• 输入：BGR uint8
• 输出：RGB float
• 范围：[0,1]
• 排列：CHW

这和 CPU 版 cvtColor + convertTo + HWC2CHW 保持一致。

---

七、后处理必须怎么配套

既然预处理时 kernel 使用的是 d2s 做逆仿射采样，那么后处理时，检测框恢复到原图坐标也必须使用同一套逆仿射关系。

正确做法

如果模型输出是：

x1, y1, x2, y2, conf, cls

那么应当做：

float x1_ori = d2s.value[0] * x1 + d2s.value[1] * y1 + d2s.value[2];
float y1_ori = d2s.value[3] * x1 + d2s.value[4] * y1 + d2s.value[5];

float x2_ori = d2s.value[0] * x2 + d2s.value[1] * y2 + d2s.value[2];
float y2_ori = d2s.value[3] * x2 + d2s.value[4] * y2 + d2s.value[5];

不推荐再做

(x - pad) / scale

如果你已经在后处理中使用了同一个 d2s，就不要再重复做一遍简化版 pad/scale 变换，否则会再次偏移。

---

八、QT .pro 配置（Windows + MSVC + CUDA）

下面是一份适合 Qt + MSVC + CUDA 11.8 的 .pro 配置片段。

CUDA_SOURCES += infer.cu
CUDA_DIR = C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v11.8
CUDA_ARCH = sm_86

win32 {
    CUDA_NVCC = $$quote($$CUDA_DIR/bin/nvcc.exe)
}

INCLUDEPATH += $$CUDA_DIR/include
LIBS += -L"$$CUDA_DIR/lib/x64" -lcudart

CUDA_INC =
for(path, INCLUDEPATH) {
    CUDA_INC += -I\"$$path\"
}

win32-msvc {
    QMAKE_CFLAGS_RELEASE -= /MT
    QMAKE_CXXFLAGS_RELEASE -= /MT
    QMAKE_CFLAGS_DEBUG -= /MTd
    QMAKE_CXXFLAGS_DEBUG -= /MTd

    QMAKE_CFLAGS_RELEASE += /MD
    QMAKE_CXXFLAGS_RELEASE += /MD
    QMAKE_CFLAGS_DEBUG += /MDd
    QMAKE_CXXFLAGS_DEBUG += /MDd

    CONFIG(debug, debug|release) {
        CUDA_RUNTIME_FLAG = -Xcompiler /MDd
    } else {
        CUDA_RUNTIME_FLAG = -Xcompiler /MD
    }
}

cuda.input = CUDA_SOURCES
cuda.output = ${QMAKE_FILE_BASE}.obj
cuda.variable_out = OBJECTS
cuda.commands = $$CUDA_NVCC -c -arch=$$CUDA_ARCH $$CUDA_RUNTIME_FLAG $$CUDA_INC ${QMAKE_FILE_NAME} -o ${QMAKE_FILE_OUT}

QMAKE_EXTRA_COMPILERS += cuda

---

九、性能收益

在实际工程里，使用 CUDA 预处理替代 OpenCV CPU 流程后，通常会得到明显收益：

分辨率	OpenCV CPU	CUDA 预处理
640×640	~8ms	~0.5ms
1080P	~35ms	~1.2ms

收益主要来自：

• 去掉 convertTo
• 去掉 CPU resize + copyMakeBorder
• 减少 Host ↔ Device 往返
• 让 TensorRT 直接吃显存输入

---

十、常见坑总结

坑 1：kernel 调用语法写错

错误写法：

warpaffine_kernel << <blocks, threads, 0, stream >> >(...)

正确写法：

warpaffine_kernel<<<blocks, threads, 0, stream>>>(...)

---

坑 2：后处理与预处理不一致

如果 preprocess 用了 d2s，postprocess 却又改回：

(x - pad) / scale

就会造成框偏移。

---

坑 3：MSVC 运行库不统一

如果 Qt 主工程用 /MT，而 .cu 编译时用了 /MD，就会出现：

LNK2038 RuntimeLibrary mismatch

必须统一运行库。

---

坑 4：.cu 文件编译了，但没有参与最终链接

一定要有：

cuda.variable_out = OBJECTS

否则容易出现：

LNK2019 unresolved external symbol

---

十一、总结

这套方案的核心不是“把 convertTo 改成 CUDA”这么简单，而是：

• 把预处理全部搬到 GPU
• 保证几何关系和 OpenCV 对齐
• 保证 preprocess / postprocess 使用同一套逆仿射逻辑
• 让 TensorRT 直接使用显存输入

最终效果通常体现在：

• 预处理耗时下降 90%+
• 整体 FPS 提升 2～3 倍
• 检测框不漂移
• QT + TensorRT + CUDA 的工程链路更加完整稳定

---

十二、适用范围

这套思路不仅适用于 YOLO26，也可迁移到：

• YOLOv5
• YOLOv8
• 其他 TensorRT 检测模型

尤其适合：

• 工业视觉
• 安防监控
• 实时检测
• QT 桌面部署场景

参考资料

1. https://blog.csdn.net/qq_34717531/article/details/135650724
2. https://blog.csdn.net/Sakuya__/article/details/141264954
3. https://blog.csdn.net/qq_59302895/article/details/157513019