parallel reduction（并行规约）和unroll last warp（同步问题）

在 CUDA 编程中，提升 parallel reduction 类型程序性能的一种有效策略是采用 unroll last warp 技术，该方法在官方提供的示例代码中如 CUDA Radix Sort（Thrust 库）、CUDA Parallel Reduction 以及 scan 等均有体现，并且在 CUDA_sample 中也有所提及：

所包含的 RadixSort 类可用于对键值对（键为浮点数或无符号整数）进行排序，也可仅对键进行排序。本示例中的优化代码（同样适用于 reduction 和 scan 示例）采用了一种称为 warp-synchronous 编程的技术，其原理在于：在一个 warp 内运行的线程，在 CUDA GPU 上会同步执行指令。该技术被用于避免当 warp 内线程通过 shared 内存共享数据时所需的 __syncthreads() 操作。需要特别指出的是，为了确保所有 GPU 上的程序能够正确运行并避免出现竞态条件，这些 warp-synchronous 表达式中使用的共享内存必须被声明为 volatile。若未声明为 volatile，并且缺少 __syncthreads() 的情况下，编译器可能会自由地延迟将数据写入 shared 内存，并将数据保留在寄存器中（这是编译器优化的一部分），从而导致执行结果错误。因此，在这些示例