cuda-samples 实战解析：使用 CUB DeviceSegmentedScan 实现分段扫描（ExclusiveSegmentedSum 与 InclusiveSegmentedScan）

cuda-samples 实战解析使用 CUB DeviceSegmentedScan 实现分段扫描ExclusiveSegmentedSum 与 InclusiveSegmentedScan【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samplescubDeviceSegmentedScan是 cuda-samples 仓库中 cpp/4_CUDA_Libraries/cubDeviceSegmentedScan 目录下的官方示例用于演示 CCCL 3.3 新增的cub::DeviceSegmentedScan设备端分段扫描算法。与普通前缀和Prefix Sum不同分段扫描允许在一次设备端调用中对多个连续数据段分别执行独立的扫描非常适合稀疏分段数据处理场景。读完本文你将掌握分段扫描与全局扫描的区别、ExclusiveSegmentedSum与InclusiveSegmentedScan两种操作的正确调用模式含临时存储两遍式分配、如何通过偏移量数组描述任意分段布局以及如何在 cuda-samples 中配置 CCCL 依赖并构建运行该示例。示例概览一次设备调用多个独立段的并行扫描分段扫描Segmented Scan的核心诉求是数据被划分为若干连续且互不重叠的段我们希望为每一段独立计算扫描scan结果段与段之间互不影响。如果用全局扫描再手工切分需要额外处理段边界的前缀传播问题而cub::DeviceSegmentedScan在单次内核调用内完成所有段的工作段间自动隔离。根据 README.md 的说明本示例展示两种操作操作语义使用的算子cub::DeviceSegmentedScan::ExclusiveSegmentedSum对每个段执行排他式exclusive前缀和即输出位置 i 的值等于段内当前位置之前所有元素之和不含自身内置加法cub::DeviceSegmentedScan::InclusiveSegmentedScan对每个段执行包含式inclusive扫描输出包含当前元素自身且算子可自定义自定义二元算子cuda::maximum段内运行最大值示例的 Key Concepts 明确归纳为CUB Device Algorithms、Segmented Scan、Prefix Sum 三个概念见 README.md说明该示例正是学习“设备端分段算法”的最佳入门代码。分段布局如何用偏移量数组描述任意分段分段扫描不需要为每一段单独启动内核而是通过一个**分段偏移量数组offsets**描述段边界。在 cubDeviceSegmentedScan.cu 中示例构造了如下布局thrust::device_vectorint d_in {1, 2, 3, 4, 5, 6, 7, 8}; thrust::device_vectorsize_t d_offsets {0, 3, 5, 8};offsets {0, 3, 5, 8}表示输入被切成 3 段[1,2,3]、[4,5]、[6,7,8]。段的数量由相邻偏移量对决定源码中的计算方式是const auto num_segments d_offsets.size() - 1; auto begin_offsets d_offsets.begin(); auto end_offsets d_offsets.begin() 1;这里begin_offsets与end_offsets构成一个“一对偏移量”的范围描述第 s 段的起始位置为offsets[s]结束位置为offsets[s1]段长即两者之差。只要 offsets 严格递增就可以表达任意长度、任意数量的连续分段这正是分段扫描相比多次调用普通DeviceScan的优势所在——一次 API 调用、一次内核启动覆盖全部段。两遍式调用模式先查临时内存再真正执行CUB 设备端算法普遍采用“两遍式”调用约定本示例的ExclusiveSegmentedSum是标准范本见 cubDeviceSegmentedScan.cusize_t temp_bytes 0; checkCudaErrors(cub::DeviceSegmentedScan::ExclusiveSegmentedSum( nullptr, temp_bytes, d_in.begin(), d_out.begin(), begin_offsets, end_offsets, num_segments)); thrust::device_vectorchar temp(temp_bytes); checkCudaErrors(cub::DeviceSegmentedScan::ExclusiveSegmentedSum(thrust::raw_pointer_cast(temp.data()), temp_bytes, d_in.begin(), d_out.begin(), begin_offsets, end_offsets, num_segments));第一步传入nullptr作为临时存储指针temp_bytes以引用方式返回算法所需的临时存储大小单位字节第二步按该大小分配临时缓冲区这里使用thrust::device_vectorchar再携带真实指针执行扫描。参数顺序依次为d_temp_storage——临时存储指针第一步为nullptrtemp_storage_bytes——临时存储大小输入/输出d_in——输入迭代器d_out——输出迭代器d_begin_offsets——段起始偏移量迭代器d_end_offsets——段结束偏移量迭代器num_segments——分段数量。调用之后紧跟cudaDeviceSynchronize()等待内核完成cubDeviceSegmentedScan.cu。代码中通过checkCudaErrors宏定义于 Common/helper_cuda.h包装所有 CUDA 与 CUB 调用一旦出错会打印文件名与行号并终止属于 cuda-samples 的标准错误处理范式。对于本示例数据{1,2,3,4,5,6,7,8}三个段的排他前缀和期望结果为段 1[1,2,3]→[0,1,3]段 2[4,5]→[0,4]段 3[6,7,8]→[0,6,13]自定义二元算子InclusiveSegmentedScan 与 cuda::maximumInclusiveSegmentedScan允许传入任意满足结合律的二元算子示例用其计算段内运行最大值running maximum。输入为{3,1,4,5,2,9,7,8}分段布局不变见 cubDeviceSegmentedScan.cu。自定义算子通过一个__host__ __device__Lambda 包装cuda::maximum来自 CCCL libcu 的 cuda/functional 头文件README 的 CUDA APIs 一节同样列出了cuda::maximumauto max_op [] __host__ __device__(int a, int b) - int { return cuda::maximum{}(a, b); };由于算子需要同时被主机端用于编译与设备端用于内核调用Lambda 必须标注__host__ __device__并在编译时开启--extended-lambda选项——这一点在 CMakeLists.txt 中有明确对应target_compile_options(cubDeviceSegmentedScan PRIVATE $$COMPILE_LANGUAGE:CUDA:--extended-lambda)随后调用InclusiveSegmentedScan相比ExclusiveSegmentedSum仅多出最后一个max_op算子参数cubDeviceSegmentedScan.cucheckCudaErrors(cub::DeviceSegmentedScan::InclusiveSegmentedScan( nullptr, temp_bytes, d_in.begin(), d_out.begin(), begin_offsets, end_offsets, num_segments, max_op)); // ... 分配 temp 后再次调用并传入 max_op同一套两遍式模式对自定义算子同样适用。若希望换成cuda::minimum求段内运行最小值或换成求最小公倍数、逻辑与等任意结合算子只需替换max_op即可无需改动调用骨架。主机端参考实现与结果校验为保证示例结果的正确性源码为两个操作各提供了一份主机端参考实现host reference这是 cuda-samples 中“GPU 结果 vs CPU 黄金参考”验证模式的典型体现host_exclusive_segmented_sum遍历每个段维护running累加器先写当前累加值再累加当前元素排他语义见 cubDeviceSegmentedScan.cuhost_inclusive_segmented_maxrunning初始化为std::numeric_limitsint::min()对每个元素先取std::max(running, input[i])再写出包含语义见 cubDeviceSegmentedScan.cu。设备端结果通过thrust::device_vector拷回主机std::vector与参考结果逐元素比较got expected并打印OK或FAIL见 cubDeviceSegmentedScan.cu。main中两个用例任一失败都会导致进程以EXIT_FAILURE退出cubDeviceSegmentedScan.cu因此运行示例本身即是一次自动化的正确性验证。示例启动时会调用findCudaDevice与cudaGetDeviceProperties同样来自 helper_cuda.h 与 CUDA Runtime API打印设备名称与计算能力int devID findCudaDevice(argc, (const char **)argv); cudaDeviceProp props; checkCudaErrors(cudaGetDeviceProperties(props, devID)); printf(Device: %s (Compute Capability %d.%d)\n\n, props.name, props.major, props.minor);构建与运行CCCL 依赖的自动获取与覆盖该示例依赖 CCCL 3.3README 的 Dependencies 一节明确要求。CCCLCUDA C Core Libraries是包含 CUB、libcu 和 Thrust 的统一发行包cub::DeviceSegmentedScan正是 CCCL 3.3 中新增的 API。仓库通过 CPMcmake/CPM.cmake在配置阶段自动拉取并固定到v3.3.3标签set(CCCL_SAMPLES_CCCL_TAG v3.3.3 CACHE STRING Tag/branch of NVIDIA/cccl to fetch for the CCCL samples)若网络环境不便或希望使用本地 CCCL 源码可通过-DCCCL_SOURCE_DIR覆盖拉取行为CMakeLists.txtif(DEFINED CCCL_SOURCE_DIR AND NOT CCCL_SOURCE_DIR STREQUAL ) CPMAddPackage(NAME CCCL SOURCE_DIR ${CCCL_SOURCE_DIR}) else() CPMAddPackage( NAME CCCL GIT_REPOSITORY https://github.com/NVIDIA/cccl GIT_TAG ${CCCL_SAMPLES_CCCL_TAG} ) endif()其余关键构建配置CMakeLists.txt还包括目标架构列表CMAKE_CUDA_ARCHITECTURES 75 80 86 87 89 90 100 110 120覆盖 SM 7.5 到 SM 12.0语言标准cxx_std_17与cuda_std_17Lambda 捕获等特性依赖 C17开启CUDA_SEPARABLE_COMPILATION支持设备端可分离编译链接CUDA::cudart与CCCL::CCCL。该示例已通过add_subdirectory(cubDeviceSegmentedScan)挂入 cpp/4_CUDA_Libraries/CMakeLists.txt因此既可整仓构建也可按根目录 README.md 的方式单独构建。Linux 下建议的命令序列为mkdir build cd build cmake .. make -j$(nproc)构建产物cubDeviceSegmentedScan位于 build 目录对应位置直接运行即可看到两个用例的输入、偏移量、GPU 结果、期望结果与 OK/FAIL 判定。支持范围与运行前提根据 README 的说明本示例的支持范围如下SM 架构SM 7.0 / 7.5 / 8.0 / 8.6 / 8.9 / 9.0 / 10.0 / 11.0 / 12.0操作系统Linux、WindowsCPU 架构x86_64、aarch64。运行前需安装与平台匹配的 CUDA Toolkit当前仓库版本对应 CUDA Toolkit 13.3见根目录 README.md并确保 CCCL 3.3 可用默认由 CPM 自动拉取无需手工安装。若以-DCCCL_SOURCE_DIR/path/to/cccl指定本地 CCCL 检出则要求该检出不低于 v3.3.3否则DeviceSegmentedScanAPI 可能不存在。小结cubDeviceSegmentedScan是一个麻雀虽小、五脏俱全的分段扫描教学示例它用最小化的代码量覆盖了分段扫描的段布局描述offsets、两遍式临时存储分配、排他/包含两种语义以及自定义二元算子四大关键知识点并以主机参考实现自动校验正确性。对于需要处理变长分段数据如稀疏矩阵行压缩、词频统计、时间序列分桶的开发者这套DeviceSegmentedScan调用模式可以直接迁移到自己的项目中。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考