Почему «a = (b> 0)? 1: 0» лучше, чем «if-else» версия в CUDA?

Question

Не могли бы вы сказать, почему

a =(b>0)?1:0

лучше, чем

if (b>0)a=1; else a =0;

версия в CUDA?Пожалуйста, дайте детали.Большое спасибо.

Yik

Answer 1

Было время, когда компилятор NVIDIA использовал тестирование идиом, чтобы сгенерировать более эффективный код для троичного оператора, чем конструкции if / then / else.Это результаты небольшого теста, чтобы выяснить, так ли это до сих пор:

__global__ void branchTest0(float *a, float *b, float *d)
{
        unsigned int tidx = threadIdx.x + blockDim.x*blockIdx.x;
        float aval = a[tidx], bval = b[tidx];
        float z0 = (aval > bval) ? aval : bval;

        d[tidx] = z0;
}

__global__ void branchTest1(float *a, float *b, float *d)
{
        unsigned int tidx = threadIdx.x + blockDim.x*blockIdx.x;
        float aval = a[tidx], bval = b[tidx];
        float z0;

        if (aval > bval) {
            z0 = aval;
        } else {
            z0 = bval;
        }
        d[tidx] = z0;
}

Компилируя эти два ядра для вычислительных возможностей 2.0 с помощью компилятора выпуска CUDA 4.0, секция сравнения выдает следующее:

branchTest0:
max.f32         %f3, %f1, %f2;

и

branchTest1:
setp.gt.f32     %p1, %f1, %f2;
selp.f32        %f3, %f1, %f2, %p1;

Тернарный оператор компилируется в одну инструкцию с плавающей запятой, тогда как if / then / else компилируется в две инструкции, после сравнения следует выбор.Оба кода выполняются условно - ни один не производит ветвления.Машинный код, испускаемый ассемблером для них, также отличается и точно повторяет PTX:

branchTest0:
    /*0070*/     /*0x00201c00081e0000*/     FMNMX R0, R2, R0, !pt;

и

branchTest1:
    /*0070*/     /*0x0021dc00220e0000*/     FSETP.GT.AND P0, pt, R2, R0, pt;
    /*0078*/     /*0x00201c0420000000*/     SEL R0, R2, R0, P0;

Так что, по крайней мере, для графических процессоров Fermi с CUDA4.0 с такой конструкцией, троичный оператор действительно производит меньше инструкций, чем эквивалентно if / then / else.Существует ли разница в производительности между ними, сводится к данным микробенчмаркинга, которых у меня нет.

Answer 2

В общем, я бы рекомендовал писать код CUDA в естественном стиле и позволить компилятору беспокоиться о локальном ветвлении.Помимо предикации, аппаратное обеспечение графического процессора также реализует инструкции типа «выбор».Используя каркас talonmies и придерживаясь исходного кода автора, я обнаружил, что один и тот же машинный код создается для обеих версий с помощью компилятора CUDA 4.0 для sm_20.Я использовал -keep для сохранения промежуточных файлов, а утилиту cuobjdump - для разборки.И троичный оператор, и оператор if транслируются в инструкцию FCMP, которая является инструкцией «выбора».

Пример случая, рассматриваемый талонмиями, на самом деле является особым случаем.Компилятор распознает некоторые общие идиомы исходного кода, такие как конкретное троичное выражение, часто используемое для выражения операций max () и min (), и генерирует код соответствующим образом.Эквивалентное выражение if не распознается как идиома.

__global__ void branchTest0(float *bp, float *d) 
{         
    unsigned int tidx = threadIdx.x + blockDim.x*blockIdx.x;
    float b = bp[tidx];
    float a = (b>0)?1:0;
    d[tidx] = a;
} 

__global__ void branchTest1(float *bp, float *d)
{
    unsigned int tidx = threadIdx.x + blockDim.x*blockIdx.x;
    float b = bp[tidx];
    float a;
    if (b>0)a=1; else a =0;
    d[tidx] = a;
}

code for sm_20
        Function : _Z11branchTest1PfS_
/*0000*/     /*0x00005de428004404*/     MOV R1, c [0x1] [0x100];
/*0008*/     /*0x84009c042c000000*/     S2R R2, SR_Tid_X;
/*0010*/     /*0x94001c042c000000*/     S2R R0, SR_CTAid_X;
/*0018*/     /*0x10019de218000000*/     MOV32I R6, 0x4;
/*0020*/     /*0x20009ca320044000*/     IMAD R2, R0, c [0x0] [0x8], R2;
/*0028*/     /*0x1020dc435000c000*/     IMUL.U32.U32.HI R3, R2, 0x4;
/*0030*/     /*0x80211c03200d8000*/     IMAD.U32.U32 R4.CC, R2, R6, c [0x0] [0x20];
/*0038*/     /*0x90315c4348004000*/     IADD.X R5, R3, c [0x0] [0x24];
/*0040*/     /*0xa0209c03200d8000*/     IMAD.U32.U32 R2.CC, R2, R6, c [0x0] [0x28];
/*0048*/     /*0x00401c8584000000*/     LD.E R0, [R4];
/*0050*/     /*0xb030dc4348004000*/     IADD.X R3, R3, c [0x0] [0x2c];
/*0058*/     /*0x03f01c003d80cfe0*/     FCMP.LEU R0, RZ, 0x3f800, R0;
/*0060*/     /*0x00201c8594000000*/     ST.E [R2], R0;
/*0068*/     /*0x00001de780000000*/     EXIT;
        ....................................


        Function : _Z11branchTest0PfS_
/*0000*/     /*0x00005de428004404*/     MOV R1, c [0x1] [0x100];
/*0008*/     /*0x84009c042c000000*/     S2R R2, SR_Tid_X;
/*0010*/     /*0x94001c042c000000*/     S2R R0, SR_CTAid_X;
/*0018*/     /*0x10019de218000000*/     MOV32I R6, 0x4;
/*0020*/     /*0x20009ca320044000*/     IMAD R2, R0, c [0x0] [0x8], R2;
/*0028*/     /*0x1020dc435000c000*/     IMUL.U32.U32.HI R3, R2, 0x4;
/*0030*/     /*0x80211c03200d8000*/     IMAD.U32.U32 R4.CC, R2, R6, c [0x0] [0x20];
/*0038*/     /*0x90315c4348004000*/     IADD.X R5, R3, c [0x0] [0x24];
/*0040*/     /*0xa0209c03200d8000*/     IMAD.U32.U32 R2.CC, R2, R6, c [0x0] [0x28];
/*0048*/     /*0x00401c8584000000*/     LD.E R0, [R4];
/*0050*/     /*0xb030dc4348004000*/     IADD.X R3, R3, c [0x0] [0x2c];
/*0058*/     /*0x03f01c003d80cfe0*/     FCMP.LEU R0, RZ, 0x3f800, R0;
/*0060*/     /*0x00201c8594000000*/     ST.E [R2], R0;
/*0068*/     /*0x00001de780000000*/     EXIT;
        ....................................

Answer 3

В общем, вам нужно избегать ветвей в коде CUDA, в противном случае вы можете получить расхождение деформации , что может привести к значительному снижению производительности. Предложения if / else обычно приводят к ответвлениям, основанным на проверке выражения. Одним из способов устранения ветвей является использование выражения, которое может быть реализовано без ветвей, если компилятор достаточно умен - таким образом, все потоки в деформации следуют по одному и тому же пути кода.

Answer 4

В обоих случаях компилятор попытается сделать то же самое, он будет стремиться использовать предикатное выполнение. Вы можете найти более подробную информацию в Руководстве по программированию CUDA C (доступно на веб-сайте ), а также в Wikipedia . По существу, для таких коротких ветвей аппаратное обеспечение способно передавать инструкции для обеих сторон ветви и использовать предикат, чтобы указать, какие потоки должны фактически выполнять инструкции.

Другими словами, разница в производительности будет минимальной. В старых компиляторах иногда помогал третичный оператор, но в настоящее время они эквивалентны.

Answer 5

Мне легче читать. Сразу видно, что цель всего оператора - установить значение a.

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

Я думаю, что стандартный if / else все в одной строке некрасив (независимо от того, для чего он используется).

Answer 6

Не знаю для CUDA, но в C ++ и C99, используя первый, вы можете инициализировать переменную const.

int const a = (b>0) ? 1 : 0;

Принимая во внимание, что с последним вы не можете сделать вашу переменную a const, как вы должны объявить ее до if.

Обратите внимание, что это может быть написано еще короче:

int const a = (b>0);

И вы могли бы даже удалить скобки ... но ИМХО это не улучшает чтение.