CUDA зарегистрировать давление

Question

У меня ядро ​​линейно подходит по методу наименьших квадратов. Оказывается, потоки используют слишком много регистров, поэтому их заполненность низкая. Вот ядро,

__global__
void strainAxialKernel(
    float* d_dis,
    float* d_str
){
    int i = threadIdx.x;
    float a = 0;
    float c = 0;
    float e = 0;
    float f = 0;
    int shift = (int)((float)(i*NEIGHBOURS)/(float)WINDOW_PER_LINE);
    int j;
    __shared__ float dis[WINDOW_PER_LINE];
    __shared__ float str[WINDOW_PER_LINE];

    // fetch data from global memory
    dis[i] = d_dis[blockIdx.x*WINDOW_PER_LINE+i];
    __syncthreads();

    // least square fit
    for (j=-shift; j<NEIGHBOURS-shift; j++)                                     
    {                                                                           
        a += j;                                                                 
        c += j*j;                                                               
        e += dis[i+j];                                                          
        f += (float(j))*dis[i+j];                                               
    }                                                                       
    str[i] = AMP*(a*e-NEIGHBOURS*f)/(a*a-NEIGHBOURS*c)/(float)BLOCK_SPACING;    

    // compensate attenuation
    if (COMPEN_EXP>0 && COMPEN_BASE>0)                                          
    {                                                                           
        str[i]                                                                  
        = (float)(str[i]*pow((float)i/(float)COMPEN_BASE+1.0f,COMPEN_EXP));     
    }   

    // write back to global memory
    if (!SIGN_PRESERVE && str[i]<0)                                             
    {                                                                           
        d_str[blockIdx.x*WINDOW_PER_LINE+i] = -str[i];                          
    }                                                                           
    else                                                                        
    {                                                                           
        d_str[blockIdx.x*WINDOW_PER_LINE+i] = str[i];                           
    }
}

У меня 32х404 блока с 96 нитями в каждом блоке. На GTS 250 SM должен обрабатывать 8 блоков. Тем не менее, визуальный профилировщик показывает, что у меня есть 11 регистров на поток, в результате занятость составляет 0,625 (5 блоков на SM). Кстати, общая память, используемая каждым блоком, составляет 792 В, поэтому проблема в регистре. Спектакль не конец света. Мне просто любопытно, могу ли я обойти это. Спасибо.

Answer 1

Занятость не является проблемой.

SM в GTS 250 (вычислительная возможность 1.1) может быть в состоянии хранить 8 блоков (8x96 потоков) одновременно в своих регистрах, но он имеет только 8 исполнительных блоков, что означаетчто только 8 из 8x96 (или, в вашем случае, 5x96) потоков будут продвигаться в любой момент времени.Пытаться втиснуть больше блоков в перегруженный SM очень мало смысла.

Фактически, вы можете попытаться поиграть с опцией -maxrregcount, чтобы УВЕЛИЧИТЬ количество регистров, что может оказать положительное влияние на производительность.

Answer 2

Всегда существует компромисс между быстрой, но ограниченной регистрацией / общей памятью и медленной, но большой глобальной памятью.Там нет никакого способа «обойти» этот компромисс.Если вы используете сокращение использования регистров с помощью глобальной памяти, вы должны получить более высокую загрузку, но более медленный доступ к памяти.

Тем не менее, вот некоторые идеи использовать меньшее количество регистров:быть предварительно вычисленным и сохраненным в постоянной памяти?Тогда каждому потоку просто нужно посмотреть shift [i].

Должны ли a и c быть числами с плавающей точкой? Или, a и c можно удалить из цикла и вычислить один раз?И, таким образом, удаляется полностью?

a вычисляется как простая арифметическая последовательность, поэтому уменьшите ее ... (примерно так)

a = ((NEIGHBORS-shift) - (-shift) + 1) * ((NEIGHBORS-shift) + (-shift)) / 2

или

a = (NEIGHBORS + 1) * ((NEIGHBORS - 2*shift)) / 2

поэтому вместо этого сделайте что-то вроде следующего (возможно, вы сможете еще больше сократить эти выражения):

str[i] = AMP*((NEIGHBORS + 1) * ((NEIGHBORS - 2*shift)) / 2*e-NEIGHBOURS*f)
str[i] /= ((NEIGHBORS + 1) * ((NEIGHBORS - 2*shift)) / 2*(NEIGHBORS + 1) * ((NEIGHBORS - 2*shift)) / 2-NEIGHBOURS*c)
str[i] /= (float)BLOCK_SPACING;

Answer 3

Вы можете использовать границы запуска, чтобы дать указание компилятору сгенерировать отображение регистров для максимального количества потоков и минимального количества блоков на многопроцессорный процессор.Это может уменьшить количество регистров, чтобы вы могли достичь желаемой занятости.

В вашем случае калькулятор занятости Nvidia показывает теоретическую пиковую занятость 63%, что, по-видимому, является тем, что вы достигаете.Как вы упомянули, это связано с количеством регистров, но также с количеством потоков в блоке.Увеличение числа потоков на блок до 128 и уменьшение числа регистров до 10 приводит к 100% теоретической пиковой загруженности.

Чтобы контролировать границы запуска для вашего ядра:

__global__ void
__launch_bounds__(128, 6)
MyKernel(...)
{
    ...
}

Затем просто запуститес размером блока 128 потоков и наслаждайтесь своей занятостью.Компилятор должен сгенерировать ваше ядро ​​так, чтобы оно использовало 10 или менее регистров.