CUDA 投票内在指令实战：基于 cuda-samples simpleVoteIntrinsics 解析 `__any_sync` 与 `__all_sync`

CUDA 投票内在指令实战基于 cuda-samples simpleVoteIntrinsics 解析__any_sync与__all_sync【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples导读本文以 CUDA Samples 仓库当前版本对应 CUDA Toolkit 13.3中的 simpleVoteIntrinsics 示例 为核心深入讲解 CUDA 投票Vote内在指令__any_sync与__all_sync的语义、用法与验证方法。你将了解到如何在 kernel 内让一个 warp 中的线程对某个谓词条件进行跨线程投票如何通过VOTE_DATA_GROUP分组构造可预期的测试数据以及如何用 CPU 端校验函数验证 GPU 投票结果。读完本文你可以直接在自有 kernel 中安全地使用同步投票指令实现分支收敛判断、数据有效性聚合等逻辑。示例概览它演示了什么官方 README 对该示例的定义非常简洁Simple program which demonstrates how to use the Vote (__any_sync,__all_sync) intrinsic instruction in a CUDA kernel.即用最小的 kernel 集合演示 CUDA 投票内在指令的两种核心操作。示例的关键概念Key Concepts只有一项——Vote Intrinsics。虽然 README 简短但其源码实现simpleVoteIntrinsics.cu 与 simpleVote_kernel.cuh提供了完整的三组测试 kernel覆盖了__any_sync、__all_sync在整 warp、半 warp 边界等场景下的行为是理解投票指令语义的最佳入门教材。从源码结构看该示例由两部分组成文件角色simpleVoteIntrinsics.cu主机端设备选择、测试数据生成、内存管理、三次 kernel 启动与结果校验simpleVote_kernel.cuh设备端三个__global__投票测试 kernelCMakeLists.txt独立可构建的 CMake 工程配置投票内在指令的语义投票指令允许一个 warp 内的所有线程对同一个谓词表达式求值然后通过硬件将结果广播回 warp 中的每一个线程。示例中的两个核心指令见 simpleVote_kernel.cuh__any_sync(mask, predicate)若 warp 内任意一个线程的谓词条件为真非零则所有线程都返回非零值1否则返回 0。__all_sync(mask, predicate)若 warp 内所有线程的谓词条件都为真则所有线程返回非零值1否则返回 0。其中mask是参与投票的线程掩码warp 内线程的位掩码0xffffffff表示整个 warp 的 32 个线程都参与。示例中所有 kernel 均使用完整掩码0xffffffff因为启动的线程数都是 warp 大小的整数倍不存在掩码收敛问题。需要强调的是_sync后缀意味着该指令同时是一个warp 级同步点所有掩码内线程必须都执行到该指令否则行为未定义。这是老版本无后缀指令如__any、__all现已废弃所不具备的约束也是 CUDA 9 之后投票/洗牌类指令的推荐写法。三个测试 kernel 的源码级解读Kernel #1VoteAnyKernel1—— 测试__any_sync__global__ void VoteAnyKernel1(unsigned int *input, unsigned int *result, int size) { int tx threadIdx.x; int mask 0xffffffff; result[tx] __any_sync(mask, input[tx]); }该 kernel 将每个线程的输入值input[tx]作为谓词调用__any_sync后把投票结果写回result[tx]。由于是any语义只要 warp 内存在一个非零输入该 warp 的所有 32 个线程都会得到非零结果。Kernel #2VoteAllKernel2—— 测试__all_sync__global__ void VoteAllKernel2(unsigned int *input, unsigned int *result, int size) { int tx threadIdx.x; int mask 0xffffffff; result[tx] __all_sync(mask, input[tx]); }结构与 Kernel #1 完全对称仅将指令换成__all_sync只有 warp 内全部输入非零时该 warp 的所有线程才得到非零结果。Kernel #3VoteAnyKernel3—— 跨 warp 与半 warp 的定向测试__global__ void VoteAnyKernel3(bool *info, int warp_size) { int tx threadIdx.x; unsigned int mask 0xffffffff; bool *offs info (tx * 3); // The following should hold true for the second and third warp *offs __any_sync(mask, (tx (warp_size * 3) / 2)); // The following should hold true for the upper half of the second warp, // and all of the third warp *(offs 1) (tx (warp_size * 3) / 2 ? true : false); // The following should hold true for the third warp only if (__all_sync(mask, (tx (warp_size * 3) / 2))) { *(offs 2) true; } }该 kernel 以 3 个 warp共 96 个线程见主机端VoteAnyKernel31, warp_size * 3验证跨 warp 边界的投票行为并输出每个线程的 3 个 bool 标志offs[0]__any_sync对谓词tx warp_size * 3 / 2的结果。预期第 2、3 个 warp上半部分为真第 1 个 warp 为假offs[1]同一谓词不经投票的普通布尔求值作为对照预期第 2 个 warp 的上半区 整个第 3 个 warp为真offs[2]__all_sync的结果仅当整个 warp 的谓词都为真即第 3 个 warp才写入true。这样设计可以精确检验any语义下只要有线程为真整个 warp 为真all语义下必须全部为真以及二者在 warp 边界第 32、64 号线程附近的正确性。测试数据的构造艺术四段式genVoteTestPattern投票指令的正确性依赖可预期的输入。主机端函数 genVoteTestPattern 把长度为VOTE_DATA_GROUP * warp_size默认VOTE_DATA_GROUP4即 128 个unsigned int的输入分成四段区间填充值针对的语义[0, size/4)全0x00000000any所有线程都为假结果应为 0[size/4, size/2)奇数索引填i偶数索引填0any半个 warp 为真结果应为 1[size/2, 3*size/4)奇数索引填0偶数索引填iall一半线程为假结果应为 0[3*size/4, size)全0xffffffffall所有线程都为真结果应为 1源码注释清晰标明了每段的意图前两段为VOTE.Any设计后两段为VOTE.All设计。对应的校验函数checkErrors1 与 checkErrors2分别检查期望结果为全 0sum 0即失败与期望结果为全 1sum ! warp_size即失败。以 Kernel #1 的校验 checkResultsVoteAnyKernel1 为例四段依次期望[0, 32)any结果全 0[32, 64)、[64, 96)any结果全 1因为每段内部都有非零输入[96, 128)any结果全 1。Kernel #2 的校验 checkResultsVoteAllKernel2 则相反前三段期望全 0all语义下只要有一个 0 就全 0最后一段期望全 1。Kernel #3 的校验 checkResultsVoteAnyKernel3 则按i % 3分派三种预期逐一核对hinfo[i]。最终main返回EXIT_SUCCESS仅当三个测试的错误计数全部为 0simpleVoteIntrinsics.cu因此该示例本身也是一份可自动判定的自检程序。主机端执行流程与 CUDA Runtime APImain的流程simpleVoteIntrinsics.cu如下设备选择调用 findCudaDevice位于 helper_cuda.h自动挑选最优 CUDA 设备并通过cudaGetDeviceProperties打印 Multi-Processors 数量与计算能力SM major.minor。内存准备主机端用malloc分配h_input/h_result设备端用cudaMalloc分配d_input/d_result随后cudaMemcpyHostToDevice上传测试模式。三次 kernel 启动Test 1/3VoteAnyKernel1dim3(1,1), dim3(VOTE_DATA_GROUP * warp_size, 1)即 1 个 block、128 线程Test 2/3VoteAllKernel2dim3(1,1), dim3(VOTE_DATA_GROUP * warp_size, 1)Test 3/3VoteAnyKernel31, warp_size * 3即 1 个 block、96 线程3 个 warp。 每次启动前后用cudaDeviceSynchronize保证同步并用getLastCudaError捕获启动错误。结果回拷与校验cudaMemcpyDeviceToHost后执行对应校验函数累加error_count[3]。资源释放cudaFree设备内存、free主机内存最后输出Shutting down...并按错误计数决定退出码。README 列出的 CUDA Runtime API 全部得到使用cudaMemcpy、cudaFree、cudaDeviceSynchronize、cudaMalloc、cudaGetDeviceProperties。所有 API 调用都包裹在checkCudaErrors宏中——该宏定义于 helper_cuda.h出错时打印文件、行号与错误字符串并退出是 CUDA Samples 统一采用的错误处理惯例。支持环境与硬件要求根据 README支持的 SM 架构SM 5.0、5.2、5.3、6.0、6.1、7.0、7.2、7.5、8.0、8.6、8.7、8.9、9.0覆盖 Maxwell 到 Hopper/Blackwell 的主流计算能力。支持的操作系统Linux、Windows。支持的 CPU 架构x86_64、armv7l、aarch64含 Tegra 等 ARM 平台。投票指令本身自 SM 1.2 起即受硬件支持见 simpleVote_kernel.cuh 中对 CUDA Programming Guide 的引用README 所列架构是其官方测试覆盖范围。示例的 CMake 工程CMakeLists.txt默认编译CMAKE_CUDA_ARCHITECTURES为75 80 86 87 89 90 100 110 120并启用CUDA_SEPARABLE_COMPILATION与 C17/CUDA 17 标准。如何构建与运行前置条件按官方说明安装与平台匹配的 CUDA Toolkit并确保 CMake 3.20 或更高版本。方式一从仓库根目录构建全部示例mkdir build cd build cmake .. make -j$(nproc)方式二单独构建本示例README 说明可从任意子目录或单个示例构建例如mkdir build cd build cmake ../cpp/0_Introduction/simpleVoteIntrinsics make -j$(nproc)Linux 下若需启用 cuda-gdb 设备端调试可在配置时追加-DENABLE_CUDA_DEBUGTrue这会为 nvcc 增加-G选项可能显著影响性能。Windows 下可直接用 Visual Studio2019 16.5导入该子目录或在x64 Native Tools Command Prompt中用cmake .. -G Visual Studio 16 2019 -A x64生成解决方案后构建。运行编译产物simpleVoteIntrinsics预期输出包含设备信息、[VOTE Kernel Test 1/3]、[VOTE Kernel Test 2/3]、[VOTE Kernel Test 3/3]三段测试及各自的OK结果最后以Shutting down...结束三个测试全部通过时进程退出码为 0。小结投票指令的适用场景__any_sync/__all_sync最常见的用途包括分支收敛判断判断一个 warp 是否全部满足某退出条件如迭代算法中的收敛检测此时__all_sync一次调用即可替代规约 广播两步数据有效性聚合判断数据块中是否存在非法值、是否全部有效__any_sync/__all_sync协作式提前退出结合__syncthreads与投票结果让整块线程在无有效工作时分叉退出。需要注意的是_sync版本要求掩码内线程必须全部到达该指令因此在使用时务必保证 warp 内无分歧或使用正确的收敛掩码。simpleVoteIntrinsics 通过四段式测试数据 三段定向 kernel CPU 端逐段校验为这两种指令提供了可直接对照的完整参考实现读者可在此基础上继续查阅 simpleVote_kernel.cuh 与 simpleVoteIntrinsics.cu 自行扩展实验。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考