Не могли бы вы рассказать мне, почему
a =(b>0)?1:0
лучше, чем
if (b>0)a=1; else a =0;
версия в CUDA? Пожалуйста, сообщите подробности. Большое спасибо.
Yik
Почему "a = (b> 0)? 1: 0" лучше, чем версия "if-else" в CUDA?
Ответ 1
В общем, я бы рекомендовал написать код CUDA в естественном стиле и позволить компилятору беспокоиться о локальном ветвлении. Помимо предикатов, оборудование GPU также реализует инструкции типа "select". Используя рамку talonmies и придерживающуюся исходный код плаката, я обнаружил, что один и тот же машинный код создается для обеих версий с компилятором CUDA 4.0 для sm_20. Я использовал -keep для хранения промежуточных файлов и утилиты cuobjdump для создания разборки. И тройной оператор, и оператор if преобразуются в инструкцию FCMP, которая является инструкцией "select".
Случай образца, рассмотренный talonmies, на самом деле является частным случаем. Компилятор распознает некоторые общие идиомы исходного кода, такие как конкретное тернарное выражение, часто используемое для выражения операций max() и min(), и соответственно генерирует код. Эквивалентный if-statement не распознается как идиома.
__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;
....................................
Ответ 2
Было время, когда компилятор NVIDIA использовал тестирование идиомы для создания более эффективного кода для тернарного оператора, чем если бы /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;
Таким образом, казалось бы, что, по крайней мере для GPU Fermi с CUDA 4.0 с такой конструкцией, тройной оператор производит меньше инструкций, эквивалентных if/then/else. Есть ли разница в производительности между ними, сводится к данным микрообнаружения, которых у меня нет.
Ответ 3
В общем, вам нужно избегать ветвей в коде CUDA, иначе вы можете получить разницу в деформации, которая может привести к большому результату. if/else, как правило, возникают ветки, основанные на проверке выражения. Один из способов устранения ветвей - использовать выражение, которое может быть реализовано без ветвей, если компилятор достаточно умен, - таким образом, все потоки в warp следуют одному и тому же пути кода.
Ответ 4
В обоих случаях компилятор попытается сделать то же самое, он будет стремиться использовать предикатное выполнение. Дополнительную информацию вы можете найти в Руководстве по программированию CUDA C (доступно через веб-сайт), а также на Wikipedia. По сути, для коротких ветвей, таких как это, оборудование может выдавать инструкции для обеих сторон ветки и использовать предикат, чтобы указать, какие потоки должны фактически выполнять инструкции.
Другими словами, минимальная разница в производительности. С более старыми компиляторами иногда помогал третичный оператор, но в настоящее время они эквивалентны.
Ответ 5
Не знаю, для CUDA, но в С++ и C99, используя первый, вы можете инициализировать константную переменную.
int const a = (b>0) ? 1 : 0;
В то время как с последним вы не можете сделать переменную const a const, поскольку вы должны объявить ее перед if.
Обратите внимание, что он может быть написан еще короче:
int const a = (b>0);
И вы даже можете удалить скобки... но IMHO не улучшает чтение.
Ответ 6
Мне легче читать. Сразу видно, что цель всего оператора - установить значение a.
Цель состоит в том, чтобы присвоить a одно из двух значений, а синтаксис термального условного оператора позволяет вам иметь только один a = в вашем заявлении.
Я думаю, что стандартный if/else все в одной строке уродлив (независимо от того, для чего он использовался).