SASS 中的扭曲发散如何体现？

Question

当一个 warp 中的不同线程执行不同的代码时，不同的分支被序列化，非活动 warp 是 "disabled."

如果发散路径包含少量指令，例如使用分支谓词，那么 "disabled" 的意思就很清楚了（线程被谓词转向 on/off），而且它也是在 sass 转储中清晰可见。

如果不同的执行路径包含大量指令（确切数量取决于 some compiler heuristics），则插入分支指令以可能跳过一个执行路径或另一个。这是有道理的：如果一个长分支很少被采用，或者没有被某个 warp 中的任何线程采用，那么允许 warp 跳过这些指令是有利的（而不是在所有情况下都被迫执行两条路径作为谓词）。

我的问题是：非活动线程如何 "disabled" 在与分支发散的情况下？第 2 页上 this presentation 左下角的幻灯片似乎表明分支是根据条件进行的，不参与的线程通过分支目标处的指令附带的谓词关闭。但是，这不是我在 SASS.

中观察到的行为

这是一个最小的可编译示例：

#include <stdio.h>

__global__ void nonpredicated( int* a, int iter )
{
  if( a[threadIdx.x] == 0 )
    // Make the number of divergent instructions unknown at 
    // compile time so the compiler is forced to create branches
    for( int i = 0; i < iter; i++ )
    {
      a[threadIdx.x] += 5;
      a[threadIdx.x] *= 5;
    }
  else
    for( int i = 0; i < iter; i++ )
    {
      a[threadIdx.x] += 2;
      a[threadIdx.x] *= 2;
    }
}

int main(){}

这里的 SASS 转储显示分支指令是断言的，但分支目标处的代码不是断言的。在执行这些分支目标期间，未采用分支的线程是否以某种方式在 SASS 中不直接可见而隐式关闭？我经常在各种 Cuda 文档中看到像 "active mask" 这样的术语，但我想知道这在 SASS 中是如何体现的，如果它是一个独立于谓词的机制。

此外，对于 Volta 之前的架构，程序计数器是每个 warp 共享的，因此谓词分支指令的想法让我感到困惑。为什么要将每线程谓词附加到一条可能会更改 warp 中所有线程共享的某些内容（程序计数器）的指令？

    code for sm_20
            Function : _Z13nonpredicatedPii
    .headerflags    @"EF_CUDA_SM20 EF_CUDA_PTX_SM(EF_CUDA_SM20)"
    /*0000*/         MOV R1, c[0x1][0x100];                       /* 0x2800440400005de4 */
    /*0008*/         S2R R0, SR_TID.X;                            /* 0x2c00000084001c04 */
    /*0010*/         MOV32I R3, 0x4;                              /* 0x180000001000dde2 */
    /*0018*/         IMAD.U32.U32 R2.CC, R0, R3, c[0x0][0x20];    /* 0x2007800080009c03 */
    /*0020*/         IMAD.U32.U32.HI.X R3, R0, R3, c[0x0][0x24];  /* 0x208680009000dc43 */
    /*0028*/         LD.E R0, [R2];                               /* 0x8400000000201c85 */
    /*0030*/         ISETP.EQ.AND P0, PT, R0, RZ, PT;             /* 0x190e0000fc01dc23 */
    /*0038*/     @P0 BRA 0xd0;                                    /* 0x40000002400001e7 */
    /*0040*/         MOV R4, c[0x0][0x28];                        /* 0x28004000a0011de4 */
    /*0048*/         ISETP.LT.AND P0, PT, R4, 0x1, PT;            /* 0x188ec0000441dc23 */
    /*0050*/         MOV R4, RZ;                                  /* 0x28000000fc011de4 */
    /*0058*/     @P0 EXIT;                                        /* 0x80000000000001e7 */
    /*0060*/         NOP;                                         /* 0x4000000000001de4 */
    /*0068*/         NOP;                                         /* 0x4000000000001de4 */
    /*0070*/         NOP;                                         /* 0x4000000000001de4 */
    /*0078*/         NOP;                                         /* 0x4000000000001de4 */
    /*0080*/         IADD R4, R4, 0x1;                            /* 0x4800c00004411c03 */
    /*0088*/         IADD R0, R0, 0x2;                            /* 0x4800c00008001c03 */
    /*0090*/         ISETP.LT.AND P0, PT, R4, c[0x0][0x28], PT;   /* 0x188e4000a041dc23 */
    /*0098*/         SHL R0, R0, 0x1;                             /* 0x6000c00004001c03 */
    /*00a0*/     @P0 BRA 0x80;                                    /* 0x4003ffff600001e7 */
    /*00a8*/         ST.E [R2], R0;                               /* 0x9400000000201c85 */
    /*00b0*/         BRA 0x128;                                   /* 0x40000001c0001de7 */
    /*00b8*/         NOP;                                         /* 0x4000000000001de4 */
    /*00c0*/         NOP;                                         /* 0x4000000000001de4 */
    /*00c8*/         NOP;                                         /* 0x4000000000001de4 */
    /*00d0*/         MOV R0, c[0x0][0x28];                        /* 0x28004000a0001de4 */
    /*00d8*/         MOV R4, RZ;                                  /* 0x28000000fc011de4 */
    /*00e0*/         ISETP.LT.AND P0, PT, R0, 0x1, PT;            /* 0x188ec0000401dc23 */
    /*00e8*/         MOV R0, RZ;                                  /* 0x28000000fc001de4 */
    /*00f0*/     @P0 EXIT;                                        /* 0x80000000000001e7 */
    /*00f8*/         MOV32I R5, 0x19;                             /* 0x1800000064015de2 */
    /*0100*/         IADD R0, R0, 0x1;                            /* 0x4800c00004001c03 */
    /*0108*/         IMAD R4, R4, 0x5, R5;                        /* 0x200ac00014411ca3 */
    /*0110*/         ISETP.LT.AND P0, PT, R0, c[0x0][0x28], PT;   /* 0x188e4000a001dc23 */
    /*0118*/     @P0 BRA 0x100;                                   /* 0x4003ffff800001e7 */
    /*0120*/         ST.E [R2], R4;                               /* 0x9400000000211c85 */
    /*0128*/         EXIT;                                        /* 0x8000000000001de7 */
            .....................................

Answer 1

Are the threads that did not take the branch switched off implicitly during execution of those branch targets, in some way that is not directly visible in the SASS?

是的。

有一个扭曲执行或 "active" 掩码，它与 predication as defined in the PTX ISA manual.

的正式概念分开

谓词执行可以允许在逐条指令的基础上为特定线程执行（或不执行）指令。编译器还可以发出谓词指令来执行条件跳转或分支。

但是 GPU 还维护一个扭曲活动掩码。当机器观察到 warp 内的线程执行已经发散时（例如在谓词分支点，或者可能是任何谓词指令），它将相应地设置活动掩码。这个过程在 SASS 级别上并不是真正的 "visible"。 AFAIK 发散 warp 的低级执行过程（不是通过预测）没有很好地指定，所以关于 warp 保持发散多长时间和重新同步的确切机制的问题没有很好地指定，并且 AFAIK 可能受到影响编译器选择，在某些架构上。 This 是最近的一个讨论（特别注意@njuffa 的评论）。

Why would you attach a per-thread predicate to an instruction that might change something (the program counter) that is shared by all threads in the warp?

这就是执行条件跳转或分支的方式。由于所有执行都是锁步的，如果我们要执行特定指令（无论掩码状态或预测状态如何），PC 最好指向该指令。但是，GPU 可以在执行时根据需要执行指令重播来处理不同的情况。

一些其他注意事项：

提到 "active mask" 是 here：

The scheduler dispatches all 32 lanes of the warp to the execution units with an active mask. Non-active threads execute through the pipe.
活动遮罩的一些 NVIDIA 工具 allow for inspection。

SASS 中的扭曲发散如何体现？

How does warp divergence manifest in SASS?

cuda

gpu

simd