CUDA实现Patch Match性能优化问询：GPU与CPU性能持平问题

哇，GT730跑出来和i5多线程差不多甚至更慢，确实挺让人失望的——毕竟咱们对GPU的性能预期可不是这样。先帮你拆解下当前代码的核心问题，再给你几个实打实的优化方向：

一、当前代码的核心瓶颈 #

1. 线程调度效率极低：你设置每个线程块只有15个线程，而CUDA的调度单元是32线程的Warp，这意味着每个Warp里有17个线程完全空转，硬件利用率直接砍半都不止！
2. 全局内存重复访问：所有线程都在反复读取同一个D_SMALL矩阵，全局内存的延迟极高，这部分浪费了大量时间。
3. 分支与内存细节的小问题：虽然影响不如前两点大，但也拖慢了整体速度。

二、具体优化步骤 #

1. 先把线程配置改对（立竿见影） #

线程块大小必须是32的倍数（比如128、256、512），这是CUDA硬件调度的基本要求。重新计算线程块和线程数：

int numOfIndentations = 300 - 50;
int blockSize = 256; // 选32的倍数，256是比较通用的选择
int numBlock = (numOfIndentations + blockSize - 1) / blockSize; // 向上取整，覆盖所有250个任务
matCompare_cuda<<< numBlock, blockSize >>>(D_SUB, D_SMALL, D_RSLTS, step, numOfIndentations, SMALL_size);

这样每个Warp的32个线程都会被充分利用，硬件利用率直接拉满。

2. 用只读缓存优化小矩阵访问 #

D_SMALL是所有线程都要读取的只读数据，给它加上const __restrict__修饰，让CUDA编译器自动把它放到**只读缓存（Read-Only Cache）**里，避免重复访问全局内存：

__global__ void matCompare_cuda(uint8_t *D_SUB, const __restrict__ uint8_t *D_SMALL, float *D_RSLTS, unsigned int step, int numOfIndentations, int SMALL_size) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid >= numOfIndentations) return; // 提前退出，避免无效操作

    int success = 0, sumZero = 0;
    int LoopIndex = tid * step;
    for (int j = 0; j < SMALL_size; j++) {
        int i = j + LoopIndex;
        if (D_SUB[i] == 0) {
            sumZero++;
            if (D_SMALL[j] == 0) {
                success++;
            }
        }
    }

    // 给默认值，避免未初始化的垃圾数据
    D_RSLTS[tid] = (sumZero > 500 && success > 0) ? (100.0f * ((float)success / sumZero)) : 0.0f;
}

如果想用共享内存的话，注意GT730的每个线程块共享内存只有48KB，而你的小矩阵是50KB，刚好超了，所以只读缓存是更简单有效的选择。

3. 向量化内存访问，提升吞吐量 #

因为处理的是uint8_t类型，可以一次加载多个元素（比如4个），减少循环次数和内存访问次数。修改循环部分：

// 50*1000=50000，刚好是4的倍数，完美适配向量化加载
for (int j = 0; j < SMALL_size; j += 4) {
    // 一次加载4个uint8_t元素
    uint4 sub_vals = *reinterpret_cast<uint4*>(&D_SUB[LoopIndex + j]);
    uint4 small_vals = *reinterpret_cast<uint4*>(&D_SMALL[j]);

    // 逐个比较每个字节
    if (sub_vals.x == 0) {
        sumZero++;
        success += (small_vals.x == 0) ? 1 : 0;
    }
    if (sub_vals.y == 0) {
        sumZero++;
        success += (small_vals.y == 0) ? 1 : 0;
    }
    if (sub_vals.z == 0) {
        sumZero++;
        success += (small_vals.z == 0) ? 1 : 0;
    }
    if (sub_vals.w == 0) {
        sumZero++;
        success += (small_vals.w == 0) ? 1 : 0;
    }
}

cudaMalloc分配的内存是自动对齐的，所以不用担心对齐问题，这样能把内存访问效率提升4倍左右。

4. 检查内存拷贝开销 #

你有没有把CPU到GPU的数据拷贝时间算进去？如果每次运行都要做cudaMemcpy传输大矩阵和小矩阵，这部分时间可能占比不小。如果是多次运行任务，建议把数据留在GPU上，只拷贝一次，避免重复开销。

5. 关于GT730的硬件限制 #

最后提一句：GT730是入门级老GPU（Kepler架构，计算能力3.5），CUDA核心只有384个，内存带宽也不高，即使优化后可能也达不到“快一个数量级”的预期，但至少能比CPU快3-5倍，不会再和CPU持平了。

三、优化后的预期效果 #

调整线程配置+只读缓存+向量化访问后，你应该能看到GPU耗时降到10ms以内，明显超过CPU的30ms。如果还想进一步优化，可以考虑用CUDA Thrust库来简化代码，或者尝试把计算逻辑拆分成更细的任务，让GPU的并行性更充分发挥。

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