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 GroupsMBCG多块协作线程组在一次内核调用内完成对大规模数组的归约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 capability6.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 支持与 C11 CUDA 编译能力仓库构建配置实际使用 C17见下文。内核实现剖析块内归约reduceBlockreduceBlockreductionMultiBlockCG.cu负责把单个线程块的共享内存局部和归约成块内单一值实现方式是用cg::tiled_partition32把块划分为 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_tile32 tile32 cg::tiled_partition32(cta); sdata[tid] cg::reduce(tile32, sdata[tid], cg::plusdouble()); 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 的幂大小的数组注释中亦指出这是 Brents 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]; } } }算法流程分四步网格级 stride 归约grid stride loopgrid.thread_rank()给出整个网格内所有线程的全局编号grid.size()给出网格内线程总数。每个线程以grid.size()为步长遍历输入数组g_idata把命中的元素累加到自己的共享内存槽sdata[block.thread_rank()]。这一步让每个线程顺序累加多个元素显著减少了全局内存访问开销。块内归约cg::sync(block)保证共享内存写入完成随后调用reduceBlock把块内threads个局部和归约为单值。写出块部分和每个块的 0 号线程把本块结果写入g_odata[blockIdx.x]。网格级最终归约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-L341void 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参数含义默认值--nN参与归约的元素个数33554432即1 25--threadsN每块线程数设备maxThreadsPerBlock--maxblocksN允许启动的最大线程块数multiProcessorCount * (maxThreadsPerMultiProcessor / maxThreadsPerBlock)其中maxblocks的默认值意味着在保持高占用率的前提下尽可能铺满所有 SM。块数最终还会被numBlocksPerSm * numSms二次裁剪见上文因此实际启动的块数可能小于--maxblocks指定值。运行时程序会打印实际采用的numThreads与numBlocks便于观察。典型的运行方式在仓库根目录下构建后# 使用默认 33554432 个元素 ./reductionMultiBlockCG # 指定元素个数与每块线程数 ./reductionMultiBlockCG --n16777216 --threads256 # 限制最大块数 ./reductionMultiBlockCG --n33554432 --maxblocks64程序会先随机生成输入数据(rand() 0xFF) / (float)RAND_MAX保持数值较小以避免求和截断误差再执行 100 次内核调用取平均耗时输出平均时间与带宽最后与 CPU 参考结果比对并给出 PASS/FAIL 判定。正确性校验与性能基准CPU 参考实现Kahan 求和reduceCPUreductionMultiBlockCG.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与runTestreductionMultiBlockCG.cu重复执行 100 次testIterations 100内核调用用sdkStartTimer/sdkStopTimer计时最后以sdkGetAverageTimerValue输出平均耗时带宽按(size * sizeof(int)) / (reduceTime * 1.0e6)GB/s 计算输入按 int 大小计字节GPU 结果通过cudaMemcpyDeviceToHost从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.txtfind_package(CUDAToolkit REQUIRED)依赖 CUDA ToolkitCMAKE_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 一节列出的特性。与其他归约示例的对比维度reductionreductionMultiBlockCG本文内核数量7 个优化内核reduce0~reduce61 个单趟内核归约趟数需要多趟内核启动完成最终归约单次内核调用内完成全部归约跨块同步不支持依赖多内核/CPU 尾段依赖 MBCG 的grid.sync()启动方式普通cudaLaunchCooperativeKernel硬件要求较低需 SM 6.0 且支持协作启动与计算抢占从源码结构看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创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 cuda-samples 之 reductionMultiBlockCG：用 Multi-Block Cooperative Groups 实现单内核多块协同归约
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 GroupsMBCG多块协作线程组在一次内核调用内完成对大规模数组的归约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 capability6.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 支持与 C11 CUDA 编译能力仓库构建配置实际使用 C17见下文。内核实现剖析块内归约reduceBlockreduceBlockreductionMultiBlockCG.cu负责把单个线程块的共享内存局部和归约成块内单一值实现方式是用cg::tiled_partition32把块划分为 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_tile32 tile32 cg::tiled_partition32(cta); sdata[tid] cg::reduce(tile32, sdata[tid], cg::plusdouble()); 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 的幂大小的数组注释中亦指出这是 Brents 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]; } } }算法流程分四步网格级 stride 归约grid stride loopgrid.thread_rank()给出整个网格内所有线程的全局编号grid.size()给出网格内线程总数。每个线程以grid.size()为步长遍历输入数组g_idata把命中的元素累加到自己的共享内存槽sdata[block.thread_rank()]。这一步让每个线程顺序累加多个元素显著减少了全局内存访问开销。块内归约cg::sync(block)保证共享内存写入完成随后调用reduceBlock把块内threads个局部和归约为单值。写出块部分和每个块的 0 号线程把本块结果写入g_odata[blockIdx.x]。网格级最终归约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-L341void 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参数含义默认值--nN参与归约的元素个数33554432即1 25--threadsN每块线程数设备maxThreadsPerBlock--maxblocksN允许启动的最大线程块数multiProcessorCount * (maxThreadsPerMultiProcessor / maxThreadsPerBlock)其中maxblocks的默认值意味着在保持高占用率的前提下尽可能铺满所有 SM。块数最终还会被numBlocksPerSm * numSms二次裁剪见上文因此实际启动的块数可能小于--maxblocks指定值。运行时程序会打印实际采用的numThreads与numBlocks便于观察。典型的运行方式在仓库根目录下构建后# 使用默认 33554432 个元素 ./reductionMultiBlockCG # 指定元素个数与每块线程数 ./reductionMultiBlockCG --n16777216 --threads256 # 限制最大块数 ./reductionMultiBlockCG --n33554432 --maxblocks64程序会先随机生成输入数据(rand() 0xFF) / (float)RAND_MAX保持数值较小以避免求和截断误差再执行 100 次内核调用取平均耗时输出平均时间与带宽最后与 CPU 参考结果比对并给出 PASS/FAIL 判定。正确性校验与性能基准CPU 参考实现Kahan 求和reduceCPUreductionMultiBlockCG.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与runTestreductionMultiBlockCG.cu重复执行 100 次testIterations 100内核调用用sdkStartTimer/sdkStopTimer计时最后以sdkGetAverageTimerValue输出平均耗时带宽按(size * sizeof(int)) / (reduceTime * 1.0e6)GB/s 计算输入按 int 大小计字节GPU 结果通过cudaMemcpyDeviceToHost从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.txtfind_package(CUDAToolkit REQUIRED)依赖 CUDA ToolkitCMAKE_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 一节列出的特性。与其他归约示例的对比维度reductionreductionMultiBlockCG本文内核数量7 个优化内核reduce0~reduce61 个单趟内核归约趟数需要多趟内核启动完成最终归约单次内核调用内完成全部归约跨块同步不支持依赖多内核/CPU 尾段依赖 MBCG 的grid.sync()启动方式普通cudaLaunchCooperativeKernel硬件要求较低需 SM 6.0 且支持协作启动与计算抢占从源码结构看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创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考