我记得几年前我被告知在GLSL着色器中做得更好
a = condition ? statementX : statementY;
过度
if(condition) a = statementX;
else a = statementY;
因为在后一种情况下,对于不满足条件的每个片段,执行将停止,而对于满足条件的片段执行statementX;然后对这些片段执行会等到对其他片段执行statementY;而在前一种情况下,所有statementX和statementY将并行执行相应的片段。 (我想这对Workgroups等来说有点复杂,但我认为这就是它的主旨)。事实上,即使对于我曾经看过的多个语句:
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;
}
这仍然是这样吗?或改进架构或编译器?这是不值得追求的过早优化吗?还是很有意义? (对于不同类型的着色器,它是否相同?片段,顶点,计算等)。
在这两种情况下,您通常都会有一个分支,几乎可以肯定两者都会导致同一个程序集。
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
你看到的@!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