c - nvcc 编译器未优化

Question

为什么编译器不做一些可以在内核中完成的简单优化？我有以下矩阵乘法代码:

__global__ void matrixMultiply(float * A, float * B, float * C,
        int numARows, int numAColumns,
        int numBRows, int numBColumns,
        int numCRows, int numCColumns) {

    int n=numAColumns;
    int Row=blockIdx.x*blockDim.x+threadIdx.x;
    int Col=blockIdx.y*blockDim.y+threadIdx.y;
    if((Row<numCRows) && (Col<numCColumns)){
        for(int k=0;k<n;++k){
                    C[Row*numCColumns+Col]+=
                        A[Row*numAColumns+k]*B[k*numBColumns+Col];
        }
    }   
}

如果我使用临时寄存器 Cvalue 来存储总和，该示例会运行得更快:

__global__ void matrixMultiply(float * A, float * B, float * C,
        int numARows, int numAColumns,
        int numBRows, int numBColumns,
        int numCRows, int numCColumns) {

    int n=numAColumns;
    int Row=blockIdx.x*blockDim.x+threadIdx.x;
    int Col=blockIdx.y*blockDim.y+threadIdx.y;
    if((Row<numCRows) && (Col<numCColumns)){
        float Cvalue=0;
        for(int k=0;k<n;++k){
            Cvalue+=A[Row*numAColumns+k]*B[k*numBColumns+Col];
        }
        C[Row*numCColumns+Col]=Cvalue;
    }   
}

在最后一种情况下，C 的全局内存仅被访问一次，而在第一种情况下，它在循环中被访问多次。这种优化通常不是由编译器完成的吗？在我的测试中，这两个代码的性能差异约为 30%，我正在执行 nvcc -O3 ...

Answer 1

自 C未声明为 __restrict__编译器无法知道是否 C与 A 相同的矩阵或B ，因此它无法执行提到的优化。当我改为使用float* __restrict__ C时，两者的时间几乎相同。谢谢克里斯·多德。