upupwords
0%

warp function

发表于 更新于 分类于
warp-level 函数的结果测试

以下是关于warp function的一些理解
主要参考资料是nVidia的官方文档

前提

warp的概念不加赘述，且建议在capability 7.x以上且cuda9.0以上的GPU上测试，有些函数在例如 any, all, and __ballot等在cuda9.0上已经被移除。

predicate表示线程的一种状况
以下所有的实例，使用一个block，block块的大小是32，所说的线程号就是laneID
mask中的第i个bit表示第i个线程

warp vote functions

  • 允许给定warp中的线程执行规约和广播操作
  • 所有线程的lane的mask必须一致
1
2
3
4
int __all_sync(unsigned mask, int predicate);
int __any_sync(unsigned mask, int predicate);
unsigned __ballot_sync(unsigned mask, int predicate);
unsigned __activemask()

__all_sync 评估mask中所有non-exited的线程（线程对应mask的位是1）的predicate，返回非零值当且仅当mask中对应线程的predicate都是非零的。

__any_sync 评估mask中所有non-exited的线程（线程对应mask的位是1）的predicate，返回非零值当且仅当mask中对应线程的predicate存在非零的。

_ballot_sync评估mask中所有non-exited线程的predicate，返回一个unsigned数，其第N位为1当且仅当mask中线程的predicate非零

__activemask() 返回unsigned数 mask，表示这个warp中所有的active状态的线程

1
2
3
4
5
6
7
8
__global__ void wall(){
    int laneId=threadIdx.x & 0x1f;
    int predicate= laneId%2;
    unsigned n;
    n=__all_sync(0x55555555,predicate);
    printf("Thread %d final n= %x\n", threadIdx.x, n);
    
}

0x55555555表示所有偶数位是1，奇数位是0。predicate是所有的偶数线程是0，奇数线程是1。在__all_sync下，所有线程的n是0，因为其只统计mask中non-exited的线程，0x55555555使得__all_sync只检查偶数位的线程，结果偶数位线程的predicate都是0。其余函数同理。

warp match functions

  • 执行warp内线程之间变量的广播和比较操作
1
2
3
unsigned int __match_any_sync(unsigned mask, T value);
unsigned int __match_all_sync(unsigned  mask, T value, int *pred)
//T可以是int/unsigned int/long/unsigned long/long long/unsigned long long/float/double

warp shuffle function

  • 交换warp内部的线程的值
  • 采用可选的width，必须是2的幂次，且不能大于warpsize
  • mask的值没有影响，不管mask是什么，结果都一致
    1
    2
    3
    4
    5
    6
    7
    T __shfl_sync(unsigned mask, T var, int srcLane, int width=warpSize);
    T __shfl_up_sync(unsigned mask, T var, unsigned int delta, int width=warpSize);
    T __shfl_down_sync(unsigned mask, T var, unsigned int delta, int width=warpSize);
    T __shfl_xor_sync(unsigned mask, T var, int laneMask, int width=warpSize);

    //T can be int, unsigned int, long, unsigned long, long long, unsigned long long, float or double. With the cuda_fp16.h header included, T can also be __halfor __half2. Similarly, with the cuda_bf16.h header included, T can also be __nv_bfloat16 or __nv_bfloat162.

    __shfl_sync()返回srclane号线程的var，width将warpsize的线程数进行划分，每个子段长度为width。每个width中的线程得到srclane所指示的var，注意这里srclane都是width中的相对位置。
    若是每个线程的var是它本身的线程值，那么

    __shfl_sync(mask,threadIdx.x,0,4)，得到的结果是 4x0, 4x4, 4x8, 4x12, 4x16, 4x20, 4x24, 4x28, 4x32
    __shfl_sync(mask,x,2,8), 得到的结果是2x8, 10x8, 18x8, 26x8

_shfl_up_sync() 返回向前偏移为 delta 的线程中的变量 var 的值，其余线程返回0。width将warpsize划分成warpsize/width个部分，每个部分返回的是当前的线程-delta的线程的value，若是减法结果为-，那么结果就不会变。注意这里是相对值，记得上面这个减法必须是要一个分组内。

例如n=__shfl_up_sync(0xffffffff,value,15,16);
结果是0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 0 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 16

__shfl_down_sync() 线程返回向后偏移为 delta 的线程中的变量 var 的值，其余线程返回0 。

调用 __shfl_down_sync(mask, x, 2, 16); 则标号为 0-13 的线程分别获得标号为 2-15 的线程中变量 x 的值；标号为 16 -29 的线程分别获得标号为 18 - 31 的线程中变量 x 的值。

__shfl_xor_sync()通过对调用者的通道ID与laneMask进行按位异或（XOR）运算来计算源通道ID。返回值为计算所得源通道中的var值。此模式实现了蝶形寻址模式。如果width小于warpsize，那么对于异或的结果，若是处于前面的group，那么可以获取异或的结果，若是处于后面的group，则会返回本身的var

例如n=__shfl_xor_sync(0,threadIdx.x,3,4);的结果是3 2 1 0 7 6 5 4 11 10 9 8 15 14 13 12 19 18 17 16 23 22 21 27 26 25 24 31 30 29 28
例如n=__shfl_xor_sync(0,threadIdx.x,3,2);的结果是0 1 1 0 4 5 5 4 8 9 9 8 12 13 13 12 16 17 17 16 20 21 21 20 24 25 25 24 28 29 29 28