более низкая занятость - лучшая производительность

Question

Меня беспокоит следующее.

Запуск того же ядра с двумя разными устройствами, один с возможностью вычисления 1.3, а другой с возможностью вычисления 2.0, я получаю лучшую производительность с большим количеством потоков на блок (с высоким уровнем занятости) в 1.3, но наоборот в 2.0. Пик производительности для 2.0, по-видимому, составляет 16 потоков на блок, а занятость 17%. Что-то меньшее или что-то большее, чем эта точка, имеет худшую производительность.

Поскольку, скорее всего, причиной этого является природа самого ядра здесь.

__global__ void 
kernel_CalculateRFCH (int xstart, int ystart, int xsize, 
      int ysize, int imxsize, int imysize, int *test, int *dev_binIm, int *per_block_results) 
{ 
    int x2, y2, bin, bin2; 
    __shared__ int s_pixels[blockDim.x*blockDim.y]; //this wouldn't compile in reailty 

    int tx = threadIdx.x; 
    int ty = threadIdx.y; 
    int tidy = threadIdx.y + blockIdx.y * blockDim.y; 
    int tidx = threadIdx.x + blockIdx.x * blockDim.x; 

    if (xstart + xsize > imxsize) 
    xsize = imxsize - xstart; 
    if (ystart + ysize > imysize) 
    ysize = imysize - ystart; 

    s_pixels[tx * blockDim.y + ty] = 0; 

    if (tidy >= ystart && tidy < ysize + ystart && tidx >= xstart && tidx < xsize + xstart) 
{ 
     bin = dev_binIm[tidx + tidy * imxsize]; 

     if (bin >= 0) 
    { 
     x2 = tidx; 
     y2 = tidy; 

     while (y2 < ystart + ysize) 
      { 
      if (x2 >= xstart + xsize || x2 - tidx > 10) 
      { 
        x2 = xstart; 
        y2++; 
        if (tidx - x2 > 10) 
        x2 = tidx - 10; 
        if (y2 - tidy > 10) 
        { 
         y2 = ystart + ysize; 
         break; 
        } 
        if (y2 >= ystart + ysize) 
         break; 
       } 

      bin2 = dev_binIm[x2 + y2 * imxsize]; 

      if (bin2 >= 0) 
       { 
       test[(tidx + tidy * imxsize) * 221 + s_pixels[tx * blockDim.y + ty]] = bin + bin2 * 80; 
       s_pixels[tx * blockDim.y + ty]++; 
       } 
      x2++; 
     }   
    }   

    } 

    for (int offset = (blockDim.x * blockDim.y)/2; offset > 0; offset >>= 1) 
    { 
    if ((tx * blockDim.y + ty) < offset) 
     { 
     s_pixels[tx * blockDim.y + ty] += s_pixels[tx * blockDim.y + ty + offset]; 
     } 
     __syncthreads(); 
    } 

    if (tx * blockDim.y + ty == 0) 
    { 
     per_block_results[blockIdx.x * gridDim.y + blockIdx.y] = s_pixels[0]; 

    } 

} 

Я пользуюсь 2-D резьбой.

ptxas информация: Компиляция функции ввода '_Z20kernel_CalculateRFCHiiiiiiPiS_' для '' sm_10 ptxas информация: Используется 16 регистров, 128 байт SMEM, 8 байт cmem [1] .

16 регистров отображаются в каждом случае на каждом устройстве.

Любые идеи, почему это может происходить, были бы очень поучительными.

Answer 1

Помимо общих замечаний, сделанных выше, ваше ядро ​​- это особый случай, поскольку большинство нитей вообще не работают. Почему бы вам не добавить xstart и ystart в tidx и tidy сразу и выбрать меньшую сетку? Ваша лучшая производительность при меньших размерах блоков может быть просто артефактом того, как область интересов разбивается на блоки.

Это также объясняет, почему вы видите большую разницу между вычислительными возможностями 1.x устройств и устройствами CC 2.0+. Начиная с CC 2.0 графические процессоры Nvidia стали намного лучше работать с ядрами, где время выполнения существенно варьируется между блоками.
При возможности вычисления 1.x новая волна блоков запланирована только после того, как закончились все текущие запущенные блоки, а из CC 2.0 в новом блоке запускается, как только закончится любой старый блок.