Низкопроизводительное ядро ​​

Question

У меня есть ядро ​​CUDA, где есть много операций и несколько ветвей. Похоже

__global__ 
void kernel(Real *randomValues, Real mu, Real sigma) 
{ 
    int row = blockDim.y * blockIdx.y + threadIdx.y; 
    int col = blockDim.x * blockIdx.x + threadIdx.x; 

    if (row >= cnTimeSteps || col >= cnPaths) return; 

    Real alphaLevel = randomValues[row*cnPaths+col]; 
    Real q = 0.0; 
    Real x = 0.0; 

    if (alphaLevel < p_low) 
    { 
     q = sqrt(-2*log(alphaLevel)); 
     x = (((((c1*q+c2)*q+c3)*q+c4)*q+c5)*q+c6)/((((d1*q+d2)*q+d3)*q+d4)*q+1); 
    } 
    else if (alphaLevel < p_high) 
    { 
     q = alphaLevel-0.5; 
     Real r = q*q; 
     x= (((((a1*r+a2)*r+a3)*r+a4)*r+a5)*r+a6)*q/(((((b1*r+b2)*r+b3)*r+b4)*r+b5)*r+1); 
    } 
    else 
    { 
     q = sqrt(-2*log(1.0-alphaLevel)); 
     x = -(((((c1*q+c2)*q+c3)*q+c4)*q+c5)*q+c6)/((((d1*q+d2)*q+d3)*q+d4)*q+1); 
    } 

    randomValues[row*cnPaths+col] = sigma * x + mu; 
} 

где все a 'ы, b' ы, c 's и d' s являются постоянными величинами (в устройстве постоянной памяти)

static __device__ __constant__ Real a1 = 1.73687; 
static __device__ __constant__ Real a2 = 1.12321100; 

и так далее.

После профилирования ядра я обнаружил, что теоретическое занятие составляет 100%, но я получаю не более 60%.

Я прошел через this и this GTC говорит, чтобы попытаться оптимизировать мое ядро.

С одной стороны у меня есть, что IPC сообщает в среднем 1,32 выданных инструкций и 0,62 выполненных. Сериализация инструкций составляет около 50%, но активность SM составляет почти 100%. С другой стороны, существует около 38 активных перекопов, но 8 имеют право выполнять следующую инструкцию, но при эффективности эмиссионного анализа я получаю, что около 70% циклов не имеют подходящей основы. Причины срыва указаны как «Прочие», которые, как я полагаю, связаны с вычислением log и sqrt.

1. Как может активность SM быть 99,82%, если в большинстве циклов нет подходящей основы?
2. Как я могу уменьшить стойло?
3. Поскольку потоки в деформации не могут попасть в одну ветвь, запросы на постоянную память, вероятно, серализованы, это правда? Должен ли я помещать эти константы в глобальную память (возможно, также использовать разделяемую память)?

Впервые я использую Nsight Visual Studio, поэтому я пытаюсь понять смысл всего анализа производительности. Кстати, моя карточка Quadro K4000.

Answer 1

1) Как может активность SM быть 99,82%, если в большинстве циклов нет подходящих оснований?

Варп активен, если регистры и слот деформации выделены для основы. SM активен, если на SM активна не менее 1 warp.

Активность SM не должна путаться с эффективностью.

2) Как я могу уменьшить стойло?

В случае кода, находящегося выше перекоса, задерживается, ожидая, что будут доступны блоки исполнения двойной точности. Quadro K4000 имеет пропускную способность 8 потоков/циклов для операций двойной точности.

Средства правовой защиты для этой проблемы: a. Уменьшите количество операций двойной точности. Например, перемещение последовательных операций в float может значительно повысить производительность, поскольку пропускная способность с одинарной точностью с плавающей запятой - это пропускная способность с двойной точностью 24x. b. Выполните ядро ​​на GK110, который имеет 8X пропускную способность с двойной точностью GK10x.

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

Эксперимент «Достигнутый FLOP» может использоваться для подтверждения того, связана ли производительность ядра с двойной пропускной способностью.

3) Поскольку потоки в деформации не могут попасть в одну ветвь, запросы на постоянную память, вероятно, являются серализованными, это правда? Должен ли я помещать эти константы в глобальную память (возможно, также использовать разделяемую память)?

Код не имеет расхождения адресов памяти в нагрузках постоянной памяти. Расхождение потока управления деформацией просто означает, что по каждому запросу на часть потоков будет активна.

Начальная глобальная нагрузка не может быть объединена. Вы должны указать значение cnPaths для просмотра кем-то. Вы также можете посмотреть эксперименты с памятью или эксперименты с корреляцией источника.

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

Answer 2

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

if (alphaLevel < p_low) 
{ 
    q = sqrt(-2*log(alphaLevel)); 
    x = (((((c1*q+c2)*q+c3)*q+c4)*q+c5)*q+c6)/((((d1*q+d2)*q+d3)*q+d4)*q+1); 
} 
else if (alphaLevel < p_high) 
{ 
    q = alphaLevel-0.5; 
    Real r = q*q; 
    x= (((((a1*r+a2)*r+a3)*r+a4)*r+a5)*r+a6)*q/(((((b1*r+b2)*r+b3)*r+b4)*r+b5)*r+1); 
} 
else 
{ 
    q = sqrt(-2*log(1.0-alphaLevel)); 
    x = -(((((c1*q+c2)*q+c3)*q+c4)*q+c5)*q+c6)/((((d1*q+d2)*q+d3)*q+d4)*q+1); 
} 

к:

if (alphaLevel >= p_low && alphaLevel < p_high) 
{ 
    q = alphaLevel-0.5; 
    Real r = q*q; 
    x= (((((a1*r+a2)*r+a3)*r+a4)*r+a5)*r+a6)*q/(((((b1*r+b2)*r+b3)*r+b4)*r+b5)*r+1); 
} 
else 
{ 
    alphaLevel = alphaLevel >= p_low ? 1.0-alphaLevel : alphaLevel; 
    q = sqrt(-2*log(alphaLevel)); 
    x = -(((((c1*q+c2)*q+c3)*q+c4)*q+c5)*q+c6)/((((d1*q+d2)*q+d3)*q+d4)*q+1); 
} 

Answer 3

Я полагаю, ваш Real тип данных является ЬурейеЕ поплавка. Вы можете попробовать добавить суффикс f к постоянным значениям, которые используются, чтобы компилятор не добавлял ненужные приведения.

E.g.

q = alphaLevel-0.5; 

Постоянная 0.5 - двойное значение, альфа-уровень - значение real = float. alphaLevel будет передано в двойное. q имеет тип float. Результат от вычитания должен быть уменьшен до плавающего значения.

Если Real - это typedef dobule, все ваши вычисления смешиваются в два раза и плавают, что приводит к тому же литье вверх и вниз.