CUDA与OpenCV(CPU)矩阵加法性能阈值成因探究

我通过连续增加矩阵元素数量并测量运行时间，对比了简单CPU函数、CUDA与**OpenCV(CPU)**的矩阵加法性能，按数据类型（CUCV_8U=unsigned char、CUCV_16U=unsigned short、CUCV32F=float、CUCV64F=double）分别绘图。实验发现：当矩阵元素量未达到约2^12时，OpenCV与CUDA的运行时间保持恒定；超过该阈值后，运行时间开始分化（采用对数刻度）。

针对CUDA的情况，我原本推测是因为GTX 960的1024个CUDA核心，当元素量超过核心数时线程无法完全并行只能并发，但OpenCV(CPU)也呈现相同趋势，因此需要解释该阈值的成因。

一、CUDA侧的核心成因 #

• 线程调度与硬件warp机制：GTX 960的流多处理器（SM）以warp为调度单位，每个warp包含32个线程。当元素量小于2^12（4096）时，总线程数仅需少量warp即可覆盖，GPU能一次性将所有warp调度到SM中，此时线程启动开销占主导，计算时间被掩盖，因此运行时间恒定。当元素量超过4096后，所需warp数量超出SM的并发承载上限，GPU需要分批调度warp，计算时间开始随元素量线性增长。
• 内存缓存效应：小矩阵（<4096元素）的所有数据可完全放入GPU的L1/L2缓存，内存延迟被消除，计算速度仅受启动开销限制。当矩阵超过缓存容量后，需频繁从全局内存读写，内存延迟开始主导运行时间，导致时间随元素量上升。

二、OpenCV(CPU)侧的成因 #

• CPU缓存容量限制：现代CPU的L1/L2缓存容量通常在几KB到几十KB级别，以float类型为例，4096个元素的矩阵总大小为16KB，刚好能完全放入L2缓存。此时数据访问无缓存缺失，计算仅受CPU线程启动开销限制，运行时间恒定。当矩阵超过缓存容量后，需从主存加载数据，缓存缺失率飙升，内存延迟成为性能瓶颈，运行时间随元素量增加而上升。
• OpenCV内部优化策略：OpenCV的CPU矩阵加法针对小矩阵做了特殊优化，比如直接使用寄存器或栈存储数据，避免内存分配与拷贝开销。当矩阵超过阈值后，才会启用多线程并行或内存池分配，此时计算与内存开销开始显现，导致运行时间分化。

三、统一阈值的本质 #

2^12（4096）这个阈值本质上是缓存容量与硬件调度单元的共同临界点：无论是GPU还是CPU，当数据量小于这个规模时，硬件可将所有数据放入高速缓存，且计算单元的启动开销掩盖了实际计算时间；当数据量超过这个规模后，内存延迟和分批调度的开销开始主导性能，运行时间随数据量线性增长。

CUDA核代码 #

// CUDA Kernel
template <typename T>
__global__ void cuCV::kernel::add(DeviceCuMat<T> OUT, const DeviceCuMat<T> A, const DeviceCuMat<T> B) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int  ch = blockIdx.z * blockDim.z + threadIdx.z;

    int index = row * A.getWidth() + col + (A.getWidth()*A.getHeight()) * ch;  // linearisation of index

    if (col < A.getWidth() && row < A.getHeight() && ch < A.getNChannels())
        OUT.getDataPtr()[index] = A.getDataPtr()[index] + B.getDataPtr()[index];
}

内容的提问来源于stack exchange，提问作者Vinc37