为什么"a =(b> 0)？1:0"比CUDA中的"if-else"版本好？

Question

你能告诉我为什么吗？

a =(b>0)?1:0

比...更好

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

CUDA中的版本？请详细说明.非常感谢.

奕

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;
}

使用CUDA 4.0发行版编译器为计算能力2.0编译这两个内核,比较部分产生:

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

和

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

三元运算符被编译成单个浮点最大指令,而if/then/else被编译成两个指令,一个比较后跟一个select.这两个代码都是有条件执行的 - 都不会产生分支.汇编程序为这些代码发出的机器代码也不同,并且与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;

所以看起来,至少对于具有这种构造的CUDA 4.0的Fermi GPU,三元运算符确实产生的指令数量等于if/then/else.它们之间是否存在性能差异归结为我没有的微基准测试数据.

Answer 2

一般来说,我建议以自然风格编写CUDA代码,让编译器担心本地分支.除了预测之外,GPU硬件还实现"选择"类型指令.使用talonmies的框架并坚持原始海报的代码,我发现使用用于sm_20的CUDA 4.0编译器为两个版本生成相同的机器代码.我使用-keep来保留中间文件,使用cuobjdump实用程序来生成反汇编.三元运算符和if语句都被转换为FCMP指令,这是一个"选择"指令.

由talonmies检查的样本案例实际上是一个特例.编译器识别一些常见的源代码习语,例如经常用于表达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;
        ....................................