CUDA Warp Shuffle

束内洗牌函数（Warp Shuffle Functions）是设备端的内置函数，与束内表决函数一样，主要用于进行一些 warp 内数据交换操作。但有两个主要区别：没有数据的规约处理功能、交换的数据通常大于 1 bit。

一、函数描述

内部函数 __shfl_sync() 允许 warp 中的线程之间交换变量，而无需使用共享内存。交换同时发生在 warp 中的所有活动线程（使用 mask 指定），根据数据类型移动每个线程 或 个字节的数据。

warp 中的线程称为通道（lanes），并且每个通道具有介于 0 和 warpSize-1（包括）之间的索引，称之为通道 ID。当前支持四种源通道（source-lane）寻址模式：

  • __shfl_sync()：从索引通道直接复制。
  • __shfl_up_sync()：从相对于调用者 ID 较低的通道复制。
  • __shfl_down_sync()：从相对于调用者 ID 较高的通道复制。
  • __shfl_xor_sync()：基于自身通道 ID 的按位异或（XOR）从通道复制。

函数原型如下：

1
2
3
4
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);

mask 是参与的 lane 位掩码，几乎总传 0xffffffff（全 32 lane）。 关键好处：少一次 shared mem round trip，省 __syncthreads。

1.1 __shfl_sync

这个函数主要用来把某一个或几个线程的值广播到其它线程。

1
T __shfl_sync(unsigned mask, T var, int srcLane, int width=warpSize);
  • mask：参与线程的 mask，要参与的线程对应的位置所在的比特为 1，否则为 0
  • var：当前线程的var变量的值
  • width：32 个线程中没 width 个线程为一组，同步操作发生在组内

下面看一个例子：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
__global__ static void test(int id) {
    int tid = threadIdx.x;
    int value = tid;

    int ret = __shfl_sync(0xFFFFFFFF, value, 5, 32);
    printf("thread id = %d, before shuffle value = %d, after shuffle value=%d\n", tid, value, ret);
}

int main() {
    test<<<1, 32>>>(0);
    cudaDeviceSynchronize();

    return 0;
}

结果如下：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
thread id = 0, before shuffle value = 0, after shuffle value=5
thread id = 1, before shuffle value = 1, after shuffle value=5
thread id = 2, before shuffle value = 2, after shuffle value=5
thread id = 3, before shuffle value = 3, after shuffle value=5
thread id = 4, before shuffle value = 4, after shuffle value=5
thread id = 5, before shuffle value = 5, after shuffle value=5
thread id = 6, before shuffle value = 6, after shuffle value=5
thread id = 7, before shuffle value = 7, after shuffle value=5
thread id = 8, before shuffle value = 8, after shuffle value=5
thread id = 9, before shuffle value = 9, after shuffle value=5
thread id = 10, before shuffle value = 10, after shuffle value=5
thread id = 11, before shuffle value = 11, after shuffle value=5
thread id = 12, before shuffle value = 12, after shuffle value=5
thread id = 13, before shuffle value = 13, after shuffle value=5
thread id = 14, before shuffle value = 14, after shuffle value=5
thread id = 15, before shuffle value = 15, after shuffle value=5
thread id = 16, before shuffle value = 16, after shuffle value=5
thread id = 17, before shuffle value = 17, after shuffle value=5
thread id = 18, before shuffle value = 18, after shuffle value=5
thread id = 19, before shuffle value = 19, after shuffle value=5
thread id = 20, before shuffle value = 20, after shuffle value=5
thread id = 21, before shuffle value = 21, after shuffle value=5
thread id = 22, before shuffle value = 22, after shuffle value=5
thread id = 23, before shuffle value = 23, after shuffle value=5
thread id = 24, before shuffle value = 24, after shuffle value=5
thread id = 25, before shuffle value = 25, after shuffle value=5
thread id = 26, before shuffle value = 26, after shuffle value=5
thread id = 27, before shuffle value = 27, after shuffle value=5
thread id = 28, before shuffle value = 28, after shuffle value=5
thread id = 29, before shuffle value = 29, after shuffle value=5
thread id = 30, before shuffle value = 30, after shuffle value=5
thread id = 31, before shuffle value = 31, after shuffle value=5

然后我们再来看一下 width 这个参数。这个参数其实是控制我们的一个 warp 里面 32 个线程如何分组。默认值为 32，也就是把 32 个线程当成一组。

这里我们尝试把一个 warp 的线程分为 2 组，每组都通过广播得到该组里面 id 为 5 的线程的值：

1
int ret = __shfl_sync(0xFFFFFFFF, value, 5, 16);

结果如下：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
thread id = 0, before shuffle value = 0, after shuffle value=5
thread id = 1, before shuffle value = 1, after shuffle value=5
thread id = 2, before shuffle value = 2, after shuffle value=5
thread id = 3, before shuffle value = 3, after shuffle value=5
thread id = 4, before shuffle value = 4, after shuffle value=5
thread id = 5, before shuffle value = 5, after shuffle value=5
thread id = 6, before shuffle value = 6, after shuffle value=5
thread id = 7, before shuffle value = 7, after shuffle value=5
thread id = 8, before shuffle value = 8, after shuffle value=5
thread id = 9, before shuffle value = 9, after shuffle value=5
thread id = 10, before shuffle value = 10, after shuffle value=5
thread id = 11, before shuffle value = 11, after shuffle value=5
thread id = 12, before shuffle value = 12, after shuffle value=5
thread id = 13, before shuffle value = 13, after shuffle value=5
thread id = 14, before shuffle value = 14, after shuffle value=5
thread id = 15, before shuffle value = 15, after shuffle value=5
thread id = 16, before shuffle value = 16, after shuffle value=21
thread id = 17, before shuffle value = 17, after shuffle value=21
thread id = 18, before shuffle value = 18, after shuffle value=21
thread id = 19, before shuffle value = 19, after shuffle value=21
thread id = 20, before shuffle value = 20, after shuffle value=21
thread id = 21, before shuffle value = 21, after shuffle value=21
thread id = 22, before shuffle value = 22, after shuffle value=21
thread id = 23, before shuffle value = 23, after shuffle value=21
thread id = 24, before shuffle value = 24, after shuffle value=21
thread id = 25, before shuffle value = 25, after shuffle value=21
thread id = 26, before shuffle value = 26, after shuffle value=21
thread id = 27, before shuffle value = 27, after shuffle value=21
thread id = 28, before shuffle value = 28, after shuffle value=21
thread id = 29, before shuffle value = 29, after shuffle value=21
thread id = 30, before shuffle value = 30, after shuffle value=21
thread id = 31, before shuffle value = 31, after shuffle value=21

可以看到，因为 width=16，也就是 16 个线程为一组，那么一个 warp 里面 32 个线程就有两组。每组都被广播得到各自组里面的 id 为 5，也就是第 6 个线程(因为 id 从 0 开始)的值，所以 0 到 15 号线程得到的值都是 5，而 16 到 32 号线程得到值为 21，也就是第二组里面第六个线程的值。

那如果我们指定的 srcLane 比 width 大的话如何处理？比如

1
int ret = __shfl_sync(0xFFFFFFFF, value, 17, 16);

实际上，真实的lane是srcLane % width。

我们运行一下上面的代码，可以得到输出：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
thread id = 0, before shuffle value = 0, after shuffle value=1
thread id = 1, before shuffle value = 1, after shuffle value=1
thread id = 2, before shuffle value = 2, after shuffle value=1
thread id = 3, before shuffle value = 3, after shuffle value=1
thread id = 4, before shuffle value = 4, after shuffle value=1
thread id = 5, before shuffle value = 5, after shuffle value=1
thread id = 6, before shuffle value = 6, after shuffle value=1
thread id = 7, before shuffle value = 7, after shuffle value=1
thread id = 8, before shuffle value = 8, after shuffle value=1
thread id = 9, before shuffle value = 9, after shuffle value=1
thread id = 10, before shuffle value = 10, after shuffle value=1
thread id = 11, before shuffle value = 11, after shuffle value=1
thread id = 12, before shuffle value = 12, after shuffle value=1
thread id = 13, before shuffle value = 13, after shuffle value=1
thread id = 14, before shuffle value = 14, after shuffle value=1
thread id = 15, before shuffle value = 15, after shuffle value=1
thread id = 16, before shuffle value = 16, after shuffle value=17
thread id = 17, before shuffle value = 17, after shuffle value=17
thread id = 18, before shuffle value = 18, after shuffle value=17
thread id = 19, before shuffle value = 19, after shuffle value=17
thread id = 20, before shuffle value = 20, after shuffle value=17
thread id = 21, before shuffle value = 21, after shuffle value=17
thread id = 22, before shuffle value = 22, after shuffle value=17
thread id = 23, before shuffle value = 23, after shuffle value=17
thread id = 24, before shuffle value = 24, after shuffle value=17
thread id = 25, before shuffle value = 25, after shuffle value=17
thread id = 26, before shuffle value = 26, after shuffle value=17
thread id = 27, before shuffle value = 27, after shuffle value=17
thread id = 28, before shuffle value = 28, after shuffle value=17
thread id = 29, before shuffle value = 29, after shuffle value=17
thread id = 30, before shuffle value = 30, after shuffle value=17
thread id = 31, before shuffle value = 31, after shuffle value=17

1.2 __shfl_up_sync & __shfl_down_sync

1
2
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);

函数作用：首先，按照width把一个warp的线程进行分组，每组有width个线程。然后，在每一组内部，0到delta-1的线程值保持不变，lane 为 delta 到 width-1 的线程的值为 lane - delta 的线程的值。

换句话说，一个组内所有的值都往右移动delta位置，低delta个线程的值不变，大于等于delta的位置的值用新值覆盖。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
__global__ static void test(int lid)
{
    int tid = threadIdx.x;
    int value = tid;

    int ret = __shfl_up_sync(0xFFFFFFFF, value, 5, 16);

    printf("thread id = %d, before shuffle value = %d, after shuffle value=%d\n", tid, value, ret);
}

int main()
{
    test<<<1, 32>>>(0);
    cudaDeviceSynchronize();
    return 0;
}

结果：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
thread id = 0, before shuffle value = 0, after shuffle value=0
thread id = 1, before shuffle value = 1, after shuffle value=1
thread id = 2, before shuffle value = 2, after shuffle value=2
thread id = 3, before shuffle value = 3, after shuffle value=3
thread id = 4, before shuffle value = 4, after shuffle value=4
thread id = 5, before shuffle value = 5, after shuffle value=0
thread id = 6, before shuffle value = 6, after shuffle value=1
thread id = 7, before shuffle value = 7, after shuffle value=2
thread id = 8, before shuffle value = 8, after shuffle value=3
thread id = 9, before shuffle value = 9, after shuffle value=4
thread id = 10, before shuffle value = 10, after shuffle value=5
thread id = 11, before shuffle value = 11, after shuffle value=6
thread id = 12, before shuffle value = 12, after shuffle value=7
thread id = 13, before shuffle value = 13, after shuffle value=8
thread id = 14, before shuffle value = 14, after shuffle value=9
thread id = 15, before shuffle value = 15, after shuffle value=10
thread id = 16, before shuffle value = 16, after shuffle value=16
thread id = 17, before shuffle value = 17, after shuffle value=17
thread id = 18, before shuffle value = 18, after shuffle value=18
thread id = 19, before shuffle value = 19, after shuffle value=19
thread id = 20, before shuffle value = 20, after shuffle value=20
thread id = 21, before shuffle value = 21, after shuffle value=16
thread id = 22, before shuffle value = 22, after shuffle value=17
thread id = 23, before shuffle value = 23, after shuffle value=18
thread id = 24, before shuffle value = 24, after shuffle value=19
thread id = 25, before shuffle value = 25, after shuffle value=20
thread id = 26, before shuffle value = 26, after shuffle value=21
thread id = 27, before shuffle value = 27, after shuffle value=22
thread id = 28, before shuffle value = 28, after shuffle value=23
thread id = 29, before shuffle value = 29, after shuffle value=24
thread id = 30, before shuffle value = 30, after shuffle value=25
thread id = 31, before shuffle value = 31, after shuffle value=26

__shfl_down_sync 函数做的事情恰好相反，把所有位置的值都往低位移动 delta 个位置，高位的 delta 个值保持不变。

1
2
3
T __shfl_down_sync(unsigned mask, T var, unsigned int delta, int width=warpSize);

int ret = __shfl_down_sync(0xFFFFFFFF, value, 5, 16);

结果：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
thread id = 0, before shuffle value = 0, after shuffle value=5
thread id = 1, before shuffle value = 1, after shuffle value=6
thread id = 2, before shuffle value = 2, after shuffle value=7
thread id = 3, before shuffle value = 3, after shuffle value=8
thread id = 4, before shuffle value = 4, after shuffle value=9
thread id = 5, before shuffle value = 5, after shuffle value=10
thread id = 6, before shuffle value = 6, after shuffle value=11
thread id = 7, before shuffle value = 7, after shuffle value=12
thread id = 8, before shuffle value = 8, after shuffle value=13
thread id = 9, before shuffle value = 9, after shuffle value=14
thread id = 10, before shuffle value = 10, after shuffle value=15
thread id = 11, before shuffle value = 11, after shuffle value=11
thread id = 12, before shuffle value = 12, after shuffle value=12
thread id = 13, before shuffle value = 13, after shuffle value=13
thread id = 14, before shuffle value = 14, after shuffle value=14
thread id = 15, before shuffle value = 15, after shuffle value=15
thread id = 16, before shuffle value = 16, after shuffle value=21
thread id = 17, before shuffle value = 17, after shuffle value=22
thread id = 18, before shuffle value = 18, after shuffle value=23
thread id = 19, before shuffle value = 19, after shuffle value=24
thread id = 20, before shuffle value = 20, after shuffle value=25
thread id = 21, before shuffle value = 21, after shuffle value=26
thread id = 22, before shuffle value = 22, after shuffle value=27
thread id = 23, before shuffle value = 23, after shuffle value=28
thread id = 24, before shuffle value = 24, after shuffle value=29
thread id = 25, before shuffle value = 25, after shuffle value=30
thread id = 26, before shuffle value = 26, after shuffle value=31
thread id = 27, before shuffle value = 27, after shuffle value=27
thread id = 28, before shuffle value = 28, after shuffle value=28
thread id = 29, before shuffle value = 29, after shuffle value=29
thread id = 30, before shuffle value = 30, after shuffle value=30
thread id = 31, before shuffle value = 31, after shuffle value=31

1.3 __shfl_xor_sync

这个函数是做 reduce 的关键。函数声明如下：

1
T __shfl_xor_sync(unsigned mask, T var, int laneMask, int width=warpSize);

同样的，一个 warp 里面 32 个线程根据 width 进行分组，每组有 width 个线程。一个组里面 id 为 lane 的线程的最终值是组里面原始线程 (lane^laneMask) 的值。

当 width < 32 的时候，那么连续的 width 个线程可以取比它们原始 id 小的线程的值，哪怕这些线程在其它的 group 里面。但是它们如果要取比它们原始id还要大的线程的值，那么它们的值保持不变。

下面看一个例子：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
__global__ static void test(int lid)
{
    int tid = threadIdx.x;
    int value = tid;

    int ret = __shfl_xor_sync(0xFFFFFFFF, value, 16, 32);

    printf("thread id = %d, before shuffle value = %d, after shuffle value=%d\n", tid, value, ret);
}

int main()
{
    test<<<1, 32>>>(0);
    cudaDeviceSynchronize();
    return 0;
}

也就是把一个 warp 里面的 32 个线程都当成一组，laneMask 设置为 16。这样做完异或之后，0 到 15 会变成 16 到 31，而 16 到 31 会变成 0 到 15。我们看一下输出：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
thread id = 0, before shuffle value = 0, after shuffle value=16
thread id = 1, before shuffle value = 1, after shuffle value=17
thread id = 2, before shuffle value = 2, after shuffle value=18
thread id = 3, before shuffle value = 3, after shuffle value=19
thread id = 4, before shuffle value = 4, after shuffle value=20
thread id = 5, before shuffle value = 5, after shuffle value=21
thread id = 6, before shuffle value = 6, after shuffle value=22
thread id = 7, before shuffle value = 7, after shuffle value=23
thread id = 8, before shuffle value = 8, after shuffle value=24
thread id = 9, before shuffle value = 9, after shuffle value=25
thread id = 10, before shuffle value = 10, after shuffle value=26
thread id = 11, before shuffle value = 11, after shuffle value=27
thread id = 12, before shuffle value = 12, after shuffle value=28
thread id = 13, before shuffle value = 13, after shuffle value=29
thread id = 14, before shuffle value = 14, after shuffle value=30
thread id = 15, before shuffle value = 15, after shuffle value=31
thread id = 16, before shuffle value = 16, after shuffle value=0
thread id = 17, before shuffle value = 17, after shuffle value=1
thread id = 18, before shuffle value = 18, after shuffle value=2
thread id = 19, before shuffle value = 19, after shuffle value=3
thread id = 20, before shuffle value = 20, after shuffle value=4
thread id = 21, before shuffle value = 21, after shuffle value=5
thread id = 22, before shuffle value = 22, after shuffle value=6
thread id = 23, before shuffle value = 23, after shuffle value=7
thread id = 24, before shuffle value = 24, after shuffle value=8
thread id = 25, before shuffle value = 25, after shuffle value=9
thread id = 26, before shuffle value = 26, after shuffle value=10
thread id = 27, before shuffle value = 27, after shuffle value=11
thread id = 28, before shuffle value = 28, after shuffle value=12
thread id = 29, before shuffle value = 29, after shuffle value=13
thread id = 30, before shuffle value = 30, after shuffle value=14
thread id = 31, before shuffle value = 31, after shuffle value=15

这样，通过下面的代码就可以实现一个warp里面每个线程都得到全部线程的累加值

1
2
for (int i = 16; i >= 1; i >>= 1)
       value += __shfl_xor_sync(0xffffffff, value, i, 32);

二、应用

2.1 warp 内广播单个变量的值

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
__global__ void bcast(int arg) {
    int laneId = threadIdx.x & 0x1f;
    int value;
    // Note unused variable for all threads except lane 0
    if (laneId == 0)        
        value = arg;       
    // Synchronize all threads in warp, and get "value" from lane 0
    value = __shfl_sync(0xffffffff, value, 0);   
    if (value != arg)
        printf("Thread %d failed.\n", threadIdx.x);
}

int main() {
    bcast<<< 1, 32 >>>(1234);
    cudaDeviceSynchronize();

    return 0;
}

首先在寄存器内定义一个变量 value，随后在 warp 内通道 ID 为 的线程中将 value 赋值为 arg，此时其他 个线程的 value 还未进行初始化。然后每个线程中调用 value = __shfl_sync(0xffffffff, value, 0)，相当于把 ID 为 的线程中的 value 的值广播到其他 个线程并赋值给 value。

2.2 束内规约

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
__global__ void warpReduce() {
    int laneId = threadIdx.x & 0x1f;
    // Seed starting value as inverse lane ID
    int value = 31 - laneId;

    // Use XOR mode to perform butterfly reduction
    for (int i=16; i>=1; i/=2)
        value += __shfl_xor_sync(0xffffffff, value, i, 32);

    // "value" now contains the sum across all threads
    printf("Thread %d final value = %d\n", threadIdx.x, value);
}

int main() {
    warpReduce<<< 1, 32 >>>();
    cudaDeviceSynchronize();

    return 0;
}

上述代码，基于折半规约的思想，调用 __shfl_xor_sync(0xffffffff, value, i, 32) 将 warp 内 value 的值规约求和到通道 ID 为 的线程中。

2.3 Prefix Sum (Scan)

定义：out[i] = in[0] + in[1] + ... + in[i]（inclusive）。 在 LLM 里用于：top-p 采样（按 cumulative prob 找截断点）、PagedAttention 的页索引计算、稀疏 softmax。

warp 内 inclusive scan（Hillis-Steele，5 步）：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
__device__ __forceinline__ float warp_inclusive_scan(float v) {
    #pragma unroll
    for (int o = 1; o < 32; o <<= 1) {
        float t = __shfl_up_sync(0xffffffff, v, o);
        if ((threadIdx.x & 31) >= o) v += t;
    }
    return v;
}

__global__ void scan_block(const float* in, float* out, int n) {
    int tid = threadIdx.x;
    int gid = blockIdx.x * blockDim.x + threadIdx.x;
    float v = (gid < n) ? in[gid] : 0;
    v = warp_inclusive_scan(v);

    __shared__ float warp_tails[32];
    int lane = tid & 31, wid = tid >> 5;
    if (lane == 31) warp_tails[wid] = v;
    __syncthreads();

    if (wid == 0) {
        float t = (tid < BLOCK / 32) ? warp_tails[lane] : 0;
        t = warp_inclusive_scan(t);
        if (tid < BLOCK / 32) warp_tails[tid] = t;
    }
    __syncthreads();

    if (wid > 0) v += warp_tails[wid - 1];
    if (gid < n) out[gid] = v;
}

block 内 scan：每 warp 算 partial，warp tail 收集到 shared，再用 1 个 warp 扫一次，最后把偏移加回去。

C++ > CUDA
#C++ #CUDA
CUDA Warp Shuffle
http://example.com/2026/07/22/CUDA-Warp-Shuffle/
作者
Yu xin
发布于
2026年7月22日
许可协议