Компилятор 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времена для обоих стали практически одинаковыми. Спасибо Крис Додд.