Является ли производительность лучше использовать (несколько) условных тернарных операторов, чем оператор if в GLSL

Question

Я помню, как много лет назад мне сказали, что в шейдере GLSL лучше сделать

a = condition ? statementX : statementY;

более

if(condition) a = statementX;
else a = statementY;

, потому что в последнем случаедля каждого фрагмента, который не удовлетворяет условию, выполнение будет остановлено, в то время как StatementX был выполнен для фрагментов, которые выполнили условие;и затем выполнение этих фрагментов будет ждать, пока оператор Y будет выполнен для других фрагментов;тогда как в первом случае все StatementX и StatementY будут выполняться параллельно для соответствующих фрагментов.(Я думаю, что это немного сложнее с рабочими группами и т. Д., Но я думаю, в этом суть).На самом деле, даже для нескольких утверждений я видел это:

a0 = condition ? statementX0 : statementY0;
a1 = condition ? statementX1 : statementY1;
a2 = condition ? statementX2 : statementY2;

вместо

if(condition) {
    a0 = statementX0;
    a1 = statementX1;
    a2 = statementX1;
} else {
    a0 = statementY0;
    a1 = statementY1;
    a2 = statementY1;
}

Это все еще так?или архитектуры или компиляторы улучшились?Не стоит ли проводить эту преждевременную оптимизациюИли все еще очень актуально?(и одинаково ли для разных типов шейдеров? фрагмент, вершина, вычисление и т. д.).

Answer 1

В обоих случаях у вас обычно будет ветвь, и почти наверняка оба приведут к одной и той же сборке.

8 __global__ void simpleTest(int *in, int a, int b, int *out)
9 {
10     int value = *in;
11     int p = (value != 0) ? __sinf(a) : __cosf(b);
12     *out = p;
13 }   
14 
15 __global__ void simpleTest2(int *in, int a, int b, int *out)
16 {
17     int value = *in;
18     int p;
19     if (value != 0)
20     {
21         p = __sinf(a);
22     }   
23     else
24     {
25         p = __cosf(b);
26     }   
27     *out = p;
28 }   

Вот как SASS выглядит для обоих:

MOV R1, c[0x0][0x44]
MOV R2, c[0x0][0x140]
MOV R3, c[0x0][0x144]
LD.E R2, [R2]
MOV R5, c[0x0][0x154]
ISETP.EQ.AND P0, PT, R2, RZ, PT
@!P0 I2F.F32.S32 R0, c[0x0] [0x148]
@P0 I2F.F32.S32 R4, c[0x0] [0x14c]
@!P0 RRO.SINCOS R0, R0
@P0 RRO.SINCOS R4, R4
@!P0 MUFU.SIN R0, R0
@P0 MUFU.COS R0, R4
MOV R4, c[0x0][0x150]
F2I.S32.F32.TRUNC R0, R0
ST.E [R4], R0
EXIT
BRA 0x98

The @! P0 и @ P0, которые вы видите, являются предикатами.Каждый поток будет иметь свой собственный бит предиката, основанный на результате.В зависимости от бита, когда блок обработки проходит через код, будет решено, должна ли быть выполнена инструкция (может также означать, что результат был зафиксирован?).

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

8 __global__ void simpleTest(int *in, int a, int b, int *out)
9 {
10     int value = *in;
11     int p = (value != 0) ? a : b;
12     *out = p;
13 }
14 
15 __global__ void simpleTest2(int *in, int a, int b, int *out)
16 {
17     int value = *in;
18     int p;
19     if (value != 0)
20     {
21         p = a;
22     }
23     else
24     {
25         p = b;
26     }
27     *out = p;
28 }

И вот как SASS выглядит для обоих:

MOV R1, c[0x0][0x44]
MOV R2, c[0x0][0x140] ; load in pointer into R2
MOV R3, c[0x0][0x144]
LD.E R2, [R2] ; deref pointer
MOV R6, c[0x0][0x14c] ; load a. b is stored at c[0x0][0x148]
MOV R4, c[0x0][0x150] ; load out pointer into R4
MOV R5, c[0x0][0x154]
ICMP.EQ R0, R6, c[0x0][0x148], R2 ; Check R2 if zero and select source based on result. Result is put into R0.
ST.E [R4], R0
EXIT
BRA 0x60

Здесь нет ответвлений.Вы можете думать о результате как о линейной интерполяции A и B:

int cond = (*p != 0)
*out = (1-cond) * a + cond * b