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 时，两者的时间几乎相同。谢谢克里斯·多德。