使用条件分配避免 CUDA 内核中的线程发散

Question

我正在尝试找到避免CUDA 内核中线程发散（分支或扭曲发散）的方法。

例如，我有以下条件赋值（a和b是char值，x和y是unsigned int值）：

if (a == b) { ++x; }
else        { ++y; }

或者，或者：

if (a == b) { ++x; }
if (a != b) { ++y; }

如何重写上述操作以避免分支？

我查看了类型转换内在函数，但没有从boolto转换可用int。我在想可能有一些技巧可以使用min,max和绝对值（例如，__sad）来获得适当的整数结果来为每种情况添加（即，1，0或0，1）。

似乎没有常规的int绝对值函数，但我看到的是：

计算| X ？是 | + z ，绝对差之和。

__device__ ? unsigned int __sad ( int  x, int  y, unsigned int  z )

我想我可以提供一个z = 0论点，以获得正常的绝对值。也许是这样的：

const unsigned int mu = __sad(a, b, 1);
const unsigned int mv = __sad(a, b, 0);
const int u = __nv_min(1, mu);
const int v = __nv_min(1, mv);
x += u;
y += v;

但是，没有min功能（请参阅相关问题）。

Answer 1

tl;dr：考虑避免这种所谓的微优化。

让我们看看我们是否可以确定与问题中建议的原始公式存在哪些差异（如果有）：

if (a == b) { ++x; }
else        { ++y; }

以及另一个答案中建议的公式：

x += (a == b);
y += (a != b);

我们将使用这个测试代码：

$ cat t1513.cu
__global__ void k(char a, char b, unsigned int *dx, unsigned int *dy){

    unsigned int x=*dx;
    unsigned int y=*dy;
#ifndef USE_OPT
    if (a == b)
{
    ++x;
} else {
    ++y;
}
#else
x += (a == b);
y += (a != b);
#endif

    *dy = y;
    *dx = x;
}


$ nvcc -c t1513.cu
$ cuobjdump -sass t1513.o >out1.sass
$ nvcc -c t1513.cu -DUSE_OPT
$ cuobjdump -sass t1513.o >out2.sass
$ diff out1.sass out2.sass
28,29c28,29
<         /*0078*/                   BFE R7, R7, 0x1000;          /* 0x7000c0400071dc23 */
<                                                                 /* 0x22e04283f2828287 */
---
>         /*0078*/                   BFE R9, R7, 0x1000;          /* 0x7000c04000725c23 */
>                                                                 /* 0x22804283f2804287 */
31,41c31,41
<         /*0090*/                   ISET.EQ.AND R7, R8, R7, PT;  /* 0x110e00001c81dc23 */
<         /*0098*/                   LOP32I.AND R7, R7, 0x1;      /* 0x380000000471dc02 */
<         /*00a0*/                   LOP32I.XOR R8, R7, 0x1;      /* 0x3800000004721c82 */
<         /*00a8*/                   IADD R8, R6, R8;             /* 0x4800000020621c03 */
<         /*00b0*/                   IADD R7, R0, R7;             /* 0x480000001c01dc03 */
<         /*00b8*/                   ST.E [R4], R8;               /* 0x9400000000421c85 */
<                                                                 /* 0x200000000002f047 */
<         /*00c8*/                   ST.E [R2], R7;               /* 0x940000000021dc85 */
<         /*00d0*/                   EXIT;                        /* 0x8000000000001de7 */
<         /*00d8*/                   BRA 0xd8;                    /* 0x4003ffffe0001de7 */
<         /*00e0*/                   NOP;                         /* 0x4000000000001de4 */
---
>         /*0090*/                   ISET.NE.AND R7, R8, R9, PT;  /* 0x128e00002481dc23 */
>         /*0098*/                   ISET.EQ.AND R8, R8, R9, PT;  /* 0x110e000024821c23 */
>         /*00a0*/                   LOP32I.AND R7, R7, 0x1;      /* 0x380000000471dc02 */
>         /*00a8*/                   IADD R7, R6, R7;             /* 0x480000001c61dc03 */
>         /*00b0*/                   LOP32I.AND R6, R8, 0x1;      /* 0x3800000004819c02 */
>         /*00b8*/                   ST.E [R4], R7;               /* 0x940000000041dc85 */
>                                                                 /* 0x2000000002f04287 */
>         /*00c8*/                   IADD R6, R0, R6;             /* 0x4800000018019c03 */
>         /*00d0*/                   ST.E [R2], R6;               /* 0x9400000000219c85 */
>         /*00d8*/                   EXIT;                        /* 0x8000000000001de7 */
>         /*00e0*/                   BRA 0xe0;                    /* 0x4003ffffe0001de7 */
$

研究上面的差异输出，我们看到：

1. 在这两种实现中都没有分支（实际上甚至没有任何预测）。
2. 所谓的“优化”情况几乎相同，只是它比 if/else 情况长 1 条指令。

是的，我知道这不是“您的代码”。我只能处理所呈现的内容。

这给了我这些类型转换的直觉：

1. 需要努力（可能浪费时间）
2. 可能不会产生任何性能改进
3. 可能混淆代码，使维护更加困难

当然，随心所欲。