Почему изменение размера блока и сетки оказывает такое большое влияние на время выполнения?

Question

Я работаю над некоторыми cuda tutorial конвертирует изображение RGBA в оттенки серого. Но я не мог понять, почему изменение blockSize и gridSize делает X33-временную модернизацию.

__global__ 
void rgba_to_greyscale(const uchar4* const rgbaImage, 
         unsigned char* const greyImage, 
         int numRows, int numCols) 
{ 
    int i = blockIdx.x*numCols + threadIdx.x; 
    float channelSum = .299f * rgbaImage[i].x + .587f * rgbaImage[i].y + .114f * rgbaImage[i].z; 
    greyImage[i]= channelSum; 
} 

void your_rgba_to_greyscale(const uchar4 * const h_rgbaImage, uchar4 * const d_rgbaImage, 
          unsigned char* const d_greyImage, size_t numRows, size_t numCols) 
{ 
    const dim3 blockSize(numCols, 1, 1); 
    const dim3 gridSize(numRows, 1 , 1); 
    rgba_to_greyscale<<<gridSize, blockSize>>>(d_rgbaImage, d_greyImage, numRows, numCols); 

    cudaDeviceSynchronize(); checkCudaErrors(cudaGetLastError()); 
} 

Когда я установил, как описано выше:

const dim3 blockSize(numCols, 1, 1); 
const dim3 gridSize(numRows, 1 , 1); 

Я получаю Your code executed in 0.030304 ms

Когда я установил:

const dim3 blockSize(1, 1, 1); 
const dim3 gridSize(numRows, numCols , 1); 

и обновление функции потоков для работы с новым индексом:

int i = blockIdx.x*numCols + blockIdx.y; 

Я получаю Your code executed in 0.995456 ms.

1. Я хотел бы ожидать, что это будет наоборот, так как графический процессор может рассчитать все пиксель отдельно на второй сетке расщеплению Связан ли это вопросам когерентности кэша? почему я получаю эти результаты?
2. Какова лучшая теория теории и размер блока для этой проблемы? можно ли вычислить его во время выполнения?

FYI:

numRows = 313 numCols =557 

---

Технические характеристики:

#uname -a && /usr/bin/nvidia-settings -v 
    Linux ip-10-16-23-92 3.2.0-39-virtual #62-Ubuntu SMP Thu Feb 28 00:48:27 UTC 2013 x86_64 x86_64 x86_64 GNU/Linux 

    nvidia-settings: version 304.54 (buildmeister@swio-display-x86-rhel47-11) 

Answer 1

Ни один из конфигурации сетки/блок рекомендуется. Первый из них не масштабируется, потому что количество потоков на один блок ограничено для графического процессора, поэтому в конечном итоге он не сможет увеличить размер изображения. Второй - плохой выбор, потому что в блоке есть только 1 поток, который не рекомендуется, так как занятость GPU будет очень низкой. Вы можете проверить его с помощью GPU Occupancy Calculator, входящего в комплект CUDA Toolkit. Рекомендуемый размер блока должен быть кратным размеру GPU (16 или 32) в зависимости от графического процессора.

Общий и масштаб способен подход для 2D сетки и размера блока в вашем случае будет что-то вроде этого:

const dim3 blockSize(16, 16, 1); 
const dim3 gridSize((numCols + blockSize.x - 1)/blockSize.x, (numRows + blockSize.y - 1)/blockSize.y , 1); 

Вы можете изменить размер блока от 16 х 16 до любого размера вы хотите, при условии, вы держитесь в пределах устройства. Максимально 512 потоков на блок разрешено для устройств с вычислительной мощностью от 1.0 до 1.3. Для устройства с возможностью вычисления 2.0 вперед этот предел составляет 1024 потока на блок.

В настоящее время, сетки и блока 2 одномерно, индексирование внутри ядра будет изменен следующим образом:

int i = blockIdx.x * blockDim.x + threadIdx.x; //Column 
int j = blockIdx.y * blockDim.y + threadIdx.y; //Row 

int idx = j * numCols + i; 

//Don't forget to perform bound checks 
if(i>=numCols || j>=numRows) return; 

float channelSum = .299f * rgbaImage[idx].x + .587f * rgbaImage[idx].y + .114f *  rgbaImage[idx].z; 
greyImage[idx]= channelSum;