Использование регистра CUDA, чтобы избежать глобального доступа, вызывает замедление

Question

Я новичок в программировании CUDA, и у меня есть ядро, которое читает / записывает одно и то же место в глобальной памяти несколько раз. Итак, я решил создать локальные переменные, чтобы использовать эти значения в регистрах. Однако скорость GFLOP / s снижена до 38 с 42. Корректность программы сохраняется, но я ожидал ускорения после использования регистров.

Это первая версия ядра:

__global__ void compute1(unsigned char* image, float* diff_coef, float* std_dev, int width, int height,
                        float* north, float* south, float* east, float* west)
{
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int index = row * width + col;

    if((row < height - 1) && (col < width - 1) && (row > 0 && col > 0))
    {
        north[index] = image[index - width] - image[index];
        south[index] = image[index + width] - image[index];
        west[index] = image[index - 1] - image[index];
        east[index] = image[index + 1] - image[index];
        float gradient_square = ( north[index] * north[index] 
                                + south[index] * south[index] 
                                + west[index]  * west[index] 
                                + east[index]  * east[index] ) / (image[index] * image[index]);
        float laplacian = (north[index] + south[index] + west[index] + east[index]) / image[index];
        float num = (0.5 * gradient_square) - ((1.0 / 16.0) * (laplacian * laplacian));
        float den = 1 + (.25 * laplacian); 
        float std_dev2 = num / (den * den); 
        den = (std_dev2 - std_dev[0]) / (std_dev[0] * (1 + std_dev[0])); 
        diff_coef[index] = 1.0 / (1.0 + den); 
        if (diff_coef[index] < 0) {
            diff_coef[index] = 0;
        } else if (diff_coef[index] > 1){
            diff_coef[index] = 1;
        }
    }
}

Это вторая версия, которая использует регистры:

__global__ void compute1(unsigned char* image, float* diff_coef, float* std_dev, int width, int height,
                            float* north, float* south, float* east, float* west)
{
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int index = row * width + col;

    if(col < width - 1 && row < height - 1 && row > 0 && col > 0)
    {

        float image_k = image[index];
        float north_k = image[index - width] - image_k;
        float south_k = image[index + width] - image_k;
        float west_k = image[index - 1] - image_k;
        float east_k = image[index + 1] - image_k;

        float deviation = std_dev[0];

        float gradient_square = ( north_k * north_k 
                                + south_k * south_k
                                + west_k  * west_k 
                                + east_k  * east_k ) / (image_k * image_k);
        float laplacian = (north_k + south_k + west_k + east_k) / image_k;
        float num = (0.5 * gradient_square) - ((1.0 / 16.0) * (laplacian * laplacian));
        float den = 1 + (.25 * laplacian); 
        float std_dev2 = num / (den * den); 
        den = (std_dev2 - deviation) / (deviation * (1 + deviation)); 
        float diff_coef_k = 1.0 / (1.0 + den);

        north[index] = north_k;
        south[index] = south_k;
        east[index]  = east_k;
        west[index]  = west_k;

        if (diff_coef_k < 0) {
            diff_coef[index] = 0;
        } else if (diff_coef_k > 1){
            diff_coef[index] = 1;
        } else {
            diff_coef[index] = diff_coef_k;
        }
    }

}

Остальная часть программы такая же, поэтому в этом ядре происходит замедление. Я не вижу здесь логической ошибки.

Я использую: TESLA K20

И мой размер блока: 16 * 16