cuda-samples 之 mergeSort 示例深度解析基于排序网络的 GPU 归并排序实现【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples本篇文章以 NVIDIA CUDA Samples 仓库中的 mergeSort 示例cpp/0_Introduction/mergeSort/README.md为核心深入讲解其算法原理、源码结构与实战运行方法。mergeSort 实现了一种被称为 Batchers sortBatcher 归并排序的排序网络算法特别适合在 GPU 上批量排序短到中等长度的key, value键值对数组。读完本文你将掌握该示例的三阶段归并流水线、底层共享内存排序内核、CPU 参考实现与双重验证机制并能够独立编译运行它。示例定位与适用场景在 CUDA Samples 的目录体系中mergeSort 归属于 cpp/0_Introduction入门级示例其核心概念被归类为Data-Parallel Algorithms数据并行算法。根据示例文档与 cpp/0_Introduction/README.md 的描述该示例实现的是归并排序merge sort它属于排序网络sorting networks这一类算法。文档同时给出了一个重要的工程判断在大规模序列排序上排序网络类算法相比具有更好渐进复杂度的算法如经典归并排序或基数排序通常效率较低但是当需要批量排序一组短到中等长度的key, value键值对数组时排序网络类算法往往是更合适的选择。这个定位解释了为何该示例被放在0_Introduction目录它不是要替代 Thrust、CUB 等库中的大规模排序方案而是展示如何在共享内存等片上资源受限的场景下以紧凑、确定性的比较交换网络完成小批量排序。说明示例文档提及了 H. W. Lang 关于排序网络的经典教程作为进一步学习材料本文不再赘述外部资料直接基于仓库源码展开分析。源码文件一览示例目录 cpp/0_Introduction/mergeSort 由 7 个文件组成职责划分清晰文件职责mergeSort.cuGPU 端完整实现共享内存排序内核、三阶段合并内核、顶层mergeSort入口bitonic.cu双调排序bitonic sort内核作为底层小规模排序的备选/辅助实现mergeSort_common.h公共头文件类型别名、关键常量、函数声明mergeSort_host.cppCPU 端仿真实现mergeSortHost用于对照验证mergeSort_validate.cpp排序结果验证键数组完整性/顺序检查、值数组正确性与稳定性检查main.cpp测试驱动程序数据生成、计时、调用与校验CMakeLists.txt该示例独立的 CMake 构建脚本从 mergeSort_common.h 可以看到两个决定算法规模的关键常量typedef unsigned int uint; #define SHARED_SIZE_LIMIT 1024U // 共享内存单次可容纳的元素数量 #define SAMPLE_STRIDE 128 // 采样步长即每个基本区间的最大长度所有排序数据均为uint32 位无符号整数这简化了共享内存布局与二进制搜索的实现。构建与运行该示例使用 CMake 构建构建方式与整个仓库一致。根据仓库根目录 README.md 的说明Linux 下的典型流程如下# 进入仓库根目录创建构建目录 mkdir build cd build # 配置项目CMake 3.20 或更高版本 cmake .. # 编译全部示例也可在 build 下的子目录单独构建 mergeSort make -j$(nproc)构建完成后可执行文件位于 build 目录下对应的子目录中直接运行即可./build/cpp/0_Introduction/mergeSort/mergeSort从 mergeSort/CMakeLists.txt 可以看到该示例的构建细节通过find_package(CUDAToolkit REQUIRED)依赖 CUDA Toolkit目标mergeSort由mergeSort.cu main.cpp mergeSort_host.cpp mergeSort_validate.cpp四个源文件编译而成启用CUDA_SEPARABLE_COMPILATION可分离编译使mergeSort.cu与bitonic.cu间的跨编译单元设备代码链接成为可能默认CMAKE_CUDA_ARCHITECTURES为75 80 86 87 89 90 100 110 120覆盖 Volta 到 Hopper/Blackwell 等主流架构编译标准为 C17 / CUDA 17并追加-lineinfo启用ENABLE_CUDA_DEBUG时替换为-G以支持 cuda-gdb 调试。该示例同样支持从cpp/0_Introduction/CMakeLists.txt其中包含add_subdirectory(mergeSort)随仓库整体构建。示例本身不读取外部数据文件运行时自动生成随机测试数据因此开箱即用。算法总览分层归并策略从 mergeSort.cu 的mergeSort()顶层入口第 471-537 行可以看出整体排序被划分为两个层次底层排序bottom-level sort调用mergeSortShared()以SHARED_SIZE_LIMIT 1024个元素为一块将整个数组切分为若干块每块由一个线程块在共享内存内部完成排序高层合并merge对排序好的 1024 元素块按stride 1024, 2048, 4096, ...逐轮两两归并每轮执行三个内核直到整个数组有序。合并轮数由 第 480-483 行 的循环计算得出并通过奇偶轮次切换输入/输出缓冲区ikey/ival与okey/oval互换从而在合并过程中避免额外的数据搬移uint stageCount 0; for (uint stride SHARED_SIZE_LIMIT; stride N; stride 1, stageCount) ;该实现还通过assert约束了输入规模N SAMPLE_STRIDE * MAX_SAMPLE_COUNT即128 * 32768 4,194,304个元素且N % SHARED_SIZE_LIMIT 0第 452-453 行 中的MAX_SAMPLE_COUNT 32768与initMergeSort()中预先分配的四个设备端辅助数组d_RanksA/d_RanksB/d_LimitsA/d_LimitsB相呼应。底层排序共享内存内核底层排序存在两种内核实现二者都是典型的排序网络mergeSortSharedKernel归并排序网络mergeSortSharedKernel 将每个 1024 元素块加载到共享内存数组s_key/s_val中每个线程负责 2 个元素然后以stride 1, 2, 4, ...递增的方式反复合并相邻子段。合并操作没有使用比较器交换而是采用基于二分搜索的插入对子段 A 中的元素keyA用binarySearchExclusive在子段 B 中计算其应插入位置排开等于keyA的 B 元素对子段 B 中的元素keyB用binarySearchInclusive在子段 A 中计算其应插入位置包含等于keyB的 A 元素。inclusive / exclusive 两种二分搜索的搭配保证了排序的稳定性相等键值下A 段的元素始终排在 B 段元素之前。两个模板函数分别定义在 第 67-103 行它们利用nextPowerOfTwo将搜索范围补齐为 2 的幂从而把二分查找写成无分支的固定迭代次数循环天然适配 GPU 的 SIMT 执行模型。内核还使用了 Cooperative Groups 的cg::thread_block与cg::sync(cta)第 112 行、第 130 行做块内同步替代传统的__syncthreads()。bitonicSortSharedKernel双调排序网络bitonic.cu 提供了另一种底层排序内核经典的 Batcher 双调排序网络。它通过Comparator比较交换单元第 36-48 行逐级构建双调序列并完成归并dir由线程号与size/2的按位与决定最后一级按sortDir统一方向。该内核通过extern C导出bitonicSortShared与bitonicMergeElementaryIntervalsmergeSort.cu 第 434-450 行 的声明与mergeSort.cu形成互补供上层按需选择。高层合并三阶段采样归并当stride SHARED_SIZE_LIMIT后两个待合并子段已超出单个线程块的共享内存容量无法整体载入。此时 mergeSort.cu 采用文献中经典的采样sample-based并行归并方案每一轮归并拆分为三个内核第 504-536 行阶段一生成样本秩 generateSampleRanksKernelgenerateSampleRanksKernel 以SAMPLE_STRIDE 128为步长从每个子段中抽取样本元素每 128 个元素取 1 个并用二分搜索在另一个子段中计算该样本的秩rank分别写入d_RanksA与d_RanksB。这样就把两个长度为 stride 的子段归并问题缩小为样本点之间的小区间归并问题。阶段二合并秩与索引 mergeRanksAndIndicesKernelmergeRanksAndIndicesKernel 对两个样本数组自身再做一次归并此时样本数量很小直接对秩数组做 inclusive/exclusive 二分搜索得到一组边界limits即每个样本点合并后在输出中的位置。结果存入d_LimitsA/d_LimitsB。这些边界把两个子段切分成若干互不重叠的基本区间elementary intervals每个区间的元素数量不超过SAMPLE_STRIDE 128。阶段三合并基本区间 mergeElementaryIntervalsKernelmergeElementaryIntervalsKernel 是流水线的最后一环每个线程块SAMPLE_STRIDE个线程负责一个基本区间将区间内来自 A、B 两段的元素加载进共享内存容量恰好为2 * SAMPLE_STRIDE再调用设备端merge()辅助函数完成归并并写回全局内存。merge()第 286-324 行延续了与底层内核一致的策略A 元素用 exclusive 二分搜索、B 元素用 inclusive 二分搜索随后一次性写回各自目标位置。尾部透传与缓冲区轮换每轮合并结束时第 516-527 行 处理了最后一个不完整段的特殊情况当lastSegmentElements stride时最后一段本身已有序直接通过cudaMemcpycudaMemcpyDeviceToDevice透传即可无需归并。随后ikey/okey互换进入下一轮。主程序与测试驱动main.cpp 是完整的测试驱动其参数均为硬编码常量第 46-48 行const uint N 4 * 1048576; // 待排序元素总数 4,194,304 const uint DIR 1; // 排序方向1 为升序0 为降序 const uint numValues 65536; // 键的取值范围 [0, 65536)测试流程如下数据生成srand(2009)固定随机种子h_SrcKey[i] rand() % numValues生成大量重复键fillValues将值数组填充为h_SrcVal[i] imergeSort_validate.cpp 第 102-106 行。用递增的值数组配合重复键恰好可以检验排序的稳定性。设备内存准备通过cudaMalloc分配d_SrcKey/d_SrcVal/d_DstKey/d_DstVal/d_BufKey/d_BufVal六块设备内存用cudaMemcpycudaMemcpyHostToDevice上传数据。初始化与排序调用initMergeSort()分配四个辅助缓冲区用sdkStartTimer/sdkStopTimer计时并执行mergeSort(...)随后cudaDeviceSynchronize()保证内核完成。结果回读与验证cudaMemcpy将结果拷回主机依次执行validateSortedKeys与validateSortedValues。退出清理closeMergeSort()释放辅助缓冲区cudaFree释放各设备数组最终以验证标志决定进程退出码。其中涉及的核心 CUDA Runtime API 与文档声明完全一致cudaMalloc、cudaMemcpy、cudaDeviceSynchronize、cudaFree。设备选择通过 Common/helper_cuda.h 中的findCudaDevice完成。正确性与稳定性验证验证逻辑集中在 mergeSort_validate.cpp这是该示例区别于一般教学示例的重要工程实践键数组验证 validateSortedKeysvalidateSortedKeys 从三个维度检查排序结果范围检查所有键必须落在[0, numValues)内多重集合不变性分别统计源数组与结果数组的键直方图srcHist/resHist并逐一比对——排序不允许丢失或新增任何元素有序性检查根据sortDir确认结果数组单调不减升序或单调不增降序。键值对正确性与稳定性 validateSortedValuesvalidateSortedValues 进一步利用val[i] i的构造正确性对每个位置j检查resKey[j] srcKey[resVal[j]]即值数组必须精确指向源数组中对应键的原始位置稳定性当相邻键相等resKey[j] resKey[j1]时要求resVal[j] resVal[j1]即相等键的相对顺序保持不变。程序会输出...stability property: stable!或NOT stable。配合 mergeSort_host.cpp 中的 CPU 参考实现mergeSortHost它逐行复刻了 GPU 的三阶段流程仅以bubbleSort代替底层内核、以memcpy代替cudaMemcpy开发者可以清晰对比 GPU 与 CPU 两条实现路径验证设备端算法的正确性。mergeSortHost内部同样包含checkOrder断言与区间边界assert如endPosA - startPosA SAMPLE_STRIDE从侧面印证了每个基本区间大小受SAMPLE_STRIDE约束的设计。支持的平台与 CUDA API根据示例 READMEmergeSort 的兼容性如下支持的 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 一代支持的操作系统Linux、Windows支持的 CPU 架构x86_64、armv7l依赖的 CUDA Runtime APIcudaMalloc、cudaDeviceSynchronize、cudaMemcpy、cudaFree。需要注意的是示例目录内的 CMakeLists.txt 默认仅针对75 80 86 87 89 90 100 110 120这些较新的架构生成设备代码若你的 GPU 架构更老如 SM 5.x/6.x需要自行调整CMAKE_CUDA_ARCHITECTURES后重新配置构建。小结cuda-samples 的 mergeSort 示例以排序网络为算法内核展示了在 GPU 上组织数据并行排序的完整工程思路共享内存内的确定性排序网络保证底层小块有序采样二分搜索的三阶段归并流水线解决大块跨线程块合并CPU 参考实现与双重验证函数兜底正确性而 inclusive/exclusive 两种二分搜索的设计细节则兼顾了性能与稳定性。对于需要在 GPU 上批量处理短到中等长度键值对数组的开发者这份示例既是可复用的参考实现也是理解 GPU 排序算法演进的绝佳起点。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考