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-samples

cubDeviceSegmentedScan是 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 的说明，本示例展示两种操作：

示例的 Key Concepts 明确归纳为：CUB Device Algorithms、Segmented Scan、Prefix Sum 三个概念（见 README.md），说明该示例正是学习“设备端分段算法”的最佳入门代码。

分段布局：如何用偏移量数组描述任意分段

分段扫描不需要为每一段单独启动内核，而是通过一个**分段偏移量数组（offsets）**描述段边界。在 cubDeviceSegmentedScan.cu 中，示例构造了如下布局：

thrust::device_vector<int> d_in = {1, 2, 3, 4, 5, 6, 7, 8}; thrust::device_vector<size_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[s+1]，段长即两者之差。只要 offsets 严格递增，就可以表达任意长度、任意数量的连续分段，这正是分段扫描相比多次调用普通DeviceScan的优势所在——一次 API 调用、一次内核启动覆盖全部段。

两遍式调用模式：先查临时内存，再真正执行

CUB 设备端算法普遍采用“两遍式”调用约定，本示例的ExclusiveSegmentedSum是标准范本（见 cubDeviceSegmentedScan.cu）：

size_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_vector<char> 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_vector<char>），再携带真实指针执行扫描。参数顺序依次为：

1. d_temp_storage——临时存储指针（第一步为nullptr）；
2. temp_storage_bytes——临时存储大小（输入/输出）；
3. d_in——输入迭代器；
4. d_out——输出迭代器；
5. d_begin_offsets——段起始偏移量迭代器；
6. d_end_offsets——段结束偏移量迭代器；
7. 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::maximum

InclusiveSegmentedScan允许传入任意满足结合律的二元算子，示例用其计算段内运行最大值（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::maximum）：

auto 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.cu）：

checkCudaErrors(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.cu；
• host_inclusive_segmented_max：running初始化为std::numeric_limits<int>::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.3+（README 的 Dependencies 一节明确要求）。CCCL（CUDA C++ Core Libraries）是包含 CUB、libcu++ 和 Thrust 的统一发行包，cub::DeviceSegmentedScan正是 CCCL 3.3 中新增的 API。仓库通过 CPM（cmake/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.txt）：

if(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_17（Lambda 捕获等特性依赖 C++17）；
• 开启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、Windows；
• CPU 架构：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