不会吧，不会吧，都AI大模型时代了，还有人不会写CUDA？

多处理器以一组 32 个并行线程（称为 warp）为单位创建、管理、调度和执行线程。组成一个线程的各个线程从相同的程序地址开始，但它们有自己的指令地址计数器和寄存器状态，因此可以自由分支和独立执行。（CUDA C++ Programming Guide Release 12.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);
• T __shfl_xor_sync(unsigned mask, T var, unsigned int delta, int width=warpSize);

在本题中，我们可以简单将warp理解为一个长为 $$$32$$$ 的数组 $$$a[0],\cdots,a[31]$$$，这些指令可以快速地对指定的下标集合进行「数据传输」。

$$$\text{mask}$$$ 用于确定参与数据传输的下标集合 $$$o$$$。设 $$$f(x)$$$ 为 $$$x$$$ 二进制表示中所有为 $$$1$$$ 的位的集合，如果 $$$f(\text{mask})\subseteq f(p)$$$，那么下标 $$$p$$$ 将参与此次的数据传输。

例如若 $$$\text{mask}=(10)_{10}=(01010)_2$$$，则 $$$f(10)=\{1,3\}$$$。$$$(8)_{10}=(01000)_2,f(8)=\{3\}$$$，则 $$$8$$$ 不符合条件；而 $$$(31)_{10}=(11111)_2,f(31)=\{0,1,2,3,4\}$$$，则 $$$31$$$ 符合条件。

$$$\text{delta}$$$ 和函数名共同确定传输的所有「下标对」。例如对于 __shfl_up_sync来说，对于任意下标 $$$x\in o$$$，若存在下标 $$$y$$$ 满足 $$$y\in o$$$ 且 $$$y=x-\text{delta}$$$，则将 $$$a[x]$$$ 传输到 $$$a[y]$$$，记为 $$$x\to y$$$。若不存在这样的 $$$y$$$，则无事发生。

注意：每次数据传输对于所有下标来说是同时发生的，没有先后次序。因此若存在 $$$z\to y,y\to x$$$，则 $$$a[y]$$$ 接收到的是数据传输前的 $$$a[z]$$$，同理 $$$a[x]$$$ 接收到的是数据传输前的 $$$a[y]$$$。

在本题中，「数据传输」有以下三种方式：

• 0: __shfl_up_add(mask, delta)：从相对于调用者下标较低的下标复制并累加，即 a[p] += a[p - delta]（若 $$$p\in o,p-\text{delta}\in o$$$）
• 1: __shfl_down_add(mask, delta)：从相对于调用者下标较高的下标复制并累加，即 a[p] += a[p + delta]（若 $$$p\in o,p+\text{delta}\in o$$$）
• 2: __shfl_xor_add(mask, delta)：从调用者下标的异或的下标复制并累加，即 a[p] += a[p ^ delta]（若 $$$p\in o,p\oplus\text{delta}\in o$$$，其中 $$$\oplus$$$ 为异或）

你的任务就是模拟这些指令，得到最终的warp。

提示：所谓异或，即按二进制逐位异或。异或的规则为：$$$1 \oplus 1=0,0 \oplus 0=0,1 \oplus 0=1,0 \oplus 1=1$$$，例如$$$(001)_2\oplus(101)_2=(100)_2$$$，即$$$(1)_{10}\oplus (5)_{10}=(4)_{10}$$$。