企业官网建设流程全解析

cuda-samples 之 reductionMultiBlockCG：用 Multi-Block Cooperative Groups 实现单内核多块协同归约

【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples

导读

reductionMultiBlockCG是 NVIDIA cuda-samples 仓库中位于 cpp/2_Concepts_and_Techniques/reductionMultiBlockCG 的 CUDA 示例，核心演示了如何借助 Multi-Block Cooperative Groups（MBCG，多块协作线程组）在一次内核调用内完成对大规模数组的归约（reduction），即所谓的"单趟（single pass）归约"。读完本文，你将掌握 Cooperative Groups 的grid_group跨块同步原理、cudaLaunchCooperativeKernel的启动方式、基于 occupancy API 的网格配置方法，以及如何在 GPU 上与 CPU 参考实现（Kahan 求和）进行结果校验和带宽基准测试。

示例概览：从多趟归约到单趟归约

归约（Reduction）是并行算法中最常见的计算模式：只要需要用一个二元结合运算（如加法、min()、max()）把一组值合并成单个值，就可以用归约。典型场景包括均值、标准差等统计计算，以及求图像总亮度等图像处理任务。

本示例的输入数据假设长度为 2 的幂。与仓库中 reduction 示例（提供 7 个不同优化级别的内核，需要多趟内核启动完成）不同，reductionMultiBlockCG使用 Cooperative Groups 的跨线程块同步能力，将"每个块先归约、再把各块的部分和做最终归约"这两个阶段合并到同一个内核中完成。内核内部通过cg::grid_group的grid.sync()实现全网格同步，这是普通 CUDA 内核中不允许（也不可用）的。

硬件与软件前置条件

计算能力与设备要求

README 明确要求：

• 设备计算能力（compute capability）6.0 及以上，且支持计算抢占（compute preemption）；
• 官方列出的支持架构为 SM 6.0 / 6.1 / 7.0 / 7.2 / 7.5 / 8.0 / 8.6 / 8.7 / 8.9 / 9.0（即 Pascal 及之后的主要架构代次）；
• 支持操作系统：Linux、Windows；支持 CPU 架构：x86_64、aarch64。

仓库根目录 README.md 对 MBCG 特性的说明与此一致："Multi Block Cooperative Groups(MBCG) extends Cooperative Groups and the CUDA programming model to express inter-thread-block synchronization. MBCG is available on GPUs with Pascal and higher architecture."，即 MBCG 扩展了 CUDA 编程模型以表达线程块之间的同步，仅在 Pascal 及更高架构的 GPU 上可用。

运行期设备检查

即便硬件满足架构要求，程序在运行期仍会主动检查设备是否支持协作内核启动。见 reductionMultiBlockCG.cu 的main()：

dev = findCudaDevice(argc, (const char **)argv); checkCudaErrors(cudaGetDeviceProperties(&deviceProp, dev)); if (!deviceProp.cooperativeLaunch) { printf("\nSelected GPU (%d) does not support Cooperative Kernel Launch, " "Waiving the run\n", dev); exit(EXIT_WAIVED); }

若所选 GPU 不支持协作内核启动（deviceProp.cooperativeLaunch为 false），程序将以EXIT_WAIVED退出并跳过运行，而不是崩溃报错。依赖项方面，构建/运行还需要 CUDA Toolkit 中的 MBCG 支持与 C++11 CUDA 编译能力（仓库构建配置实际使用 C++17，见下文）。

内核实现剖析

块内归约：reduceBlock

reduceBlock（reductionMultiBlockCG.cu）负责把单个线程块的共享内存局部和归约成块内单一值，实现方式是用cg::tiled_partition<32>把块划分为 32 线程的 warp 级 tile，再调用cg::reduce进行硬件 shuffle 归约：

__device__ void reduceBlock(double *sdata, const cg::thread_block &cta) { const unsigned int tid = cta.thread_rank(); cg::thread_block_tile<32> tile32 = cg::tiled_partition<32>(cta); sdata[tid] = cg::reduce(tile32, sdata[tid], cg::plus<double>()); cg::sync(cta); double beta = 0.0; if (cta.thread_rank() == 0) { beta = 0; for (int i = 0; i < blockDim.x; i += tile32.size()) { beta += sdata[i]; } sdata[0] = beta; } cg::sync(cta); }

这一阶段对应经典的"共享内存归约"：先让每个 warp 通过cg::reduce（底层基于 shuffle 指令）得到 tile 内部分和，再让 0 号线程把各 tile 的部分和串行累加进sdata[0]。共享内存归约的复杂度为 O(n) 工作量、O(log n) 步数，适合 2 的幂大小的数组；注释中亦指出这是 Brent's Theorem 优化思想的体现（每个线程顺序累加多个元素以摊薄归约成本）。

单趟归约内核：reduceSinglePassMultiBlockCG

核心内核（reductionMultiBlockCG.cu）是单趟归约的关键，完整代码如下：

extern "C" __global__ void reduceSinglePassMultiBlockCG(const float *g_idata, float *g_odata, unsigned int n) { // Handle to thread block group cg::thread_block block = cg::this_thread_block(); cg::grid_group grid = cg::this_grid(); extern double __shared__ sdata[]; // Stride over grid and add the values to a shared memory buffer sdata[block.thread_rank()] = 0; for (int i = grid.thread_rank(); i < n; i += grid.size()) { sdata[block.thread_rank()] += g_idata[i]; } cg::sync(block); // Reduce each block (called once per block) reduceBlock(sdata, block); // Write out the result to global memory if (block.thread_rank() == 0) { g_odata[blockIdx.x] = sdata[0]; } cg::sync(grid); if (grid.thread_rank() == 0) { for (int block = 1; block < gridDim.x; block++) { g_odata[0] += g_odata[block]; } } }

算法流程分四步：

1. 网格级 stride 归约（grid stride loop）：grid.thread_rank()给出整个网格内所有线程的全局编号，grid.size()给出网格内线程总数。每个线程以grid.size()为步长遍历输入数组g_idata，把命中的元素累加到自己的共享内存槽sdata[block.thread_rank()]。这一步让每个线程顺序累加多个元素，显著减少了全局内存访问开销。
2. 块内归约：cg::sync(block)保证共享内存写入完成，随后调用reduceBlock把块内threads个局部和归约为单值。
3. 写出块部分和：每个块的 0 号线程把本块结果写入g_odata[blockIdx.x]。
4. 网格级最终归约：cg::sync(grid)是单趟归约的关键——它保证所有块的部分和都已写入g_odata；随后全局 0 号线程串行累加g_odata[1..gridDim.x-1]，得到最终结果写入g_odata[0]。

第 2 步与第 4 步之间依赖的就是grid.sync()这一 MBCG 提供的跨块同步原语：它使得"全网格数据就绪 → 单一线程做最终累加"可以在同一个内核里安全完成，而不必像传统做法那样启动第二个内核或在 CPU 侧做尾段求和。这是本示例与 reduction 示例（其内核 6 需要把各块部分和读回主机做最终累加，或额外内核完成）最本质的区别。

内核启动：cudaLaunchCooperativeKernel与占用率计算

使用协作启动 API

普通内核用<<<grid, block>>>语法启动，而 MBCG 内核必须通过运行时 APIcudaLaunchCooperativeKernel启动，见 reductionMultiBlockCG.cu 的封装函数：

void call_reduceSinglePassMultiBlockCG(int size, int threads, int numBlocks, float *d_idata, float *d_odata) { int smemSize = threads * sizeof(double); void *kernelArgs[] = { (void *)&d_idata, (void *)&d_odata, (void *)&size, }; dim3 dimBlock(threads, 1, 1); dim3 dimGrid(numBlocks, 1, 1); cudaLaunchCooperativeKernel((void *)reduceSinglePassMultiBlockCG, dimGrid, dimBlock, kernelArgs, smemSize, NULL); getLastCudaError("Kernel execution failed"); }

要点：

• 内核参数以void *数组形式显式传入（d_idata、d_odata、size），这是cudaLaunchCooperativeKernel的签名要求，与<<<>>>的隐式参数传递不同；
• 动态共享内存大小smemSize = threads * sizeof(double)，注意内核里共享内存声明为double数组（extern double __shared__ sdata[]），即使输入输出是float，块内归约阶段也使用double精度累加，以减少浮点截断误差；
• 流参数传NULL表示默认流，dimGrid与dimBlock必须满足协作启动的占用率约束。

通过 occupancy API 计算网格规模

MBCG 内核要求网格中所有线程块必须同时驻留在 GPU 上，因此网格规模不能超过设备实际能容纳的块数。示例用两个 occupancy API 计算启动配置（reductionMultiBlockCG.cu 与 L326-L341）：

void getNumBlocksAndThreads(int n, int maxBlocks, int maxThreads, int &blocks, int &threads) { if (n == 1) { threads = 1; blocks = 1; } else { checkCudaErrors(cudaOccupancyMaxPotentialBlockSize(&blocks, &threads, reduceSinglePassMultiBlockCG)); } blocks = min(maxBlocks, blocks); }

随后通过cudaOccupancyMaxActiveBlocksPerMultiprocessor计算每 SM 实际可同时驻留的块数，并把网格规模裁剪到numBlocksPerSm * numSms以内：

int numBlocksPerSm = 0; checkCudaErrors(cudaOccupancyMaxActiveBlocksPerMultiprocessor( &numBlocksPerSm, reduceSinglePassMultiBlockCG, numThreads, numThreads * sizeof(double))); int numSms = prop.multiProcessorCount; if (numBlocks > numBlocksPerSm * numSms) { numBlocks = numBlocksPerSm * numSms; } printf("numThreads: %d\n", numThreads); printf("numBlocks: %d\n", numBlocks);

这里cudaOccupancyMaxPotentialBlockSize负责给出"最可能达到最高占用率"的线程数/块数组合，cudaOccupancyMaxActiveBlocksPerMultiprocessor负责确认真实可驻留块数——两者共同保证cudaLaunchCooperativeKernel启动的网格能完全驻留，从而grid.sync()不会死锁。这也是协作内核编程中最容易出错、必须用 API 而非猜测来确定网格大小的原因。

命令行参数与运行方式

示例支持三个命令行参数（reductionMultiBlockCG.cu）：

其中maxblocks的默认值意味着"在保持高占用率的前提下尽可能铺满所有 SM"。块数最终还会被numBlocksPerSm * numSms二次裁剪（见上文），因此实际启动的块数可能小于--maxblocks指定值。运行时程序会打印实际采用的numThreads与numBlocks，便于观察。

典型的运行方式（在仓库根目录下构建后）：

# 使用默认 33554432 个元素 ./reductionMultiBlockCG # 指定元素个数与每块线程数 ./reductionMultiBlockCG --n=16777216 --threads=256 # 限制最大块数 ./reductionMultiBlockCG --n=33554432 --maxblocks=64

程序会先随机生成输入数据（(rand() & 0xFF) / (float)RAND_MAX，保持数值较小以避免求和截断误差），再执行 100 次内核调用取平均耗时，输出平均时间与带宽，最后与 CPU 参考结果比对并给出 PASS/FAIL 判定。

正确性校验与性能基准

CPU 参考实现：Kahan 求和

reduceCPU（reductionMultiBlockCG.cu）采用Kahan 求和算法计算 CPU 参考值，这是为提高大规模数组浮点求和的精度而设计的补偿求和法——它维护一个补偿项c来记录每一步累加中被丢弃的低位误差：

template <class T> T reduceCPU(T *data, int size) { T sum = data[0]; T c = (T)0.0; for (int i = 1; i < size; i++) { T y = data[i] - c; T t = sum + y; c = (t - sum) - y; sum = t; } return sum; }

判定阈值与带宽统计

基准流程位于benchmarkReduce与runTest（reductionMultiBlockCG.cu）：

• 重复执行 100 次（testIterations = 100）内核调用，用sdkStartTimer/sdkStopTimer计时，最后以sdkGetAverageTimerValue输出平均耗时；
• 带宽按(size * sizeof(int)) / (reduceTime * 1.0e6)GB/s 计算（输入按 int 大小计字节）；
• GPU 结果通过cudaMemcpy（DeviceToHost）从d_odata[0]读回，与 CPU 的 Kahan 结果比较，判定条件为diff < 1e-8 * size，即允许相对输入规模放大的微小浮点误差。

double threshold = 1e-8 * size; double diff = abs((double)gpu_result - (double)cpu_result); bTestPassed = (diff < threshold);

程序最终打印GPU result、CPU result、Average time与Bandwidth，并以进程退出码反映测试是否通过。

构建方式

示例使用 CMake 构建，配置文件为 cpp/2_Concepts_and_Techniques/reductionMultiBlockCG/CMakeLists.txt：

• find_package(CUDAToolkit REQUIRED)依赖 CUDA Toolkit；
• CMAKE_CUDA_ARCHITECTURES设置为75 80 86 87 89 90 100 110 120，对应 SM 7.5 至 SM 12.0 代次；
• 编译选项包含--extended-lambda与-lineinfo（调试构建时切换为-G以便 cuda-gdb 调试）；
• 启用CUDA_SEPARABLE_COMPILATION（可分离编译），并编译语言标准为cxx_std_17/cuda_std_17；
• 头文件路径包含仓库的 Common 目录（依赖helper_cuda.h、helper_functions.h等工具头文件）。

典型构建命令（在仓库根目录）：

mkdir build && cd build cmake ../cpp/2_Concepts_and_Techniques/reductionMultiBlockCG make

构建前需先安装与设备匹配的 CUDA Toolkit，并确保设备满足前文"硬件与软件前置条件"中的 MBCG 依赖（README.md 中 "Multi-block Cooperative Groups" 一节列出的特性）。

与其他归约示例的对比

从源码结构看，reductionMultiBlockCG的共享内存归约部分（reduceBlock）与reduction示例中 reduce6 等内核的共享内存归约思想一脉相承（reduction_kernel.cu），差异在于用 MBCG 把跨块同步内联进了内核，从而用一次内核启动换掉了传统方案的多次启动/回读开销。值得一提的是，reduceBlock内部采用double精度在共享内存中做块内归约，体现了示例对大规模求和数值稳定性的额外考虑。

总结

reductionMultiBlockCG是理解 Cooperative Groups 与 MBCG 编程模型的理想入口：从cg::this_grid()获取grid_group句柄、用grid.thread_rank()/grid.size()组织网格级 stride 归约、用grid.sync()完成跨块同步，再到用cudaLaunchCooperativeKernel与两个 occupancy API 安全地确定可完全驻留的网格规模，整个示例完整覆盖了协作内核从设计、启动到验证的实战链路。对于需要在单次内核内完成全局规约（如全局统计量、全局极值等）的 CUDA 应用，本示例提供了可直接参考的范式。运行前请务必确认目标 GPU 满足 SM 6.0+、支持计算抢占，并且所选设备的cooperativeLaunch属性为真。

【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples