企业官网建设流程全解析

CUB 库入门：GPU 并行计算的"瑞士军刀"

引言：为什么需要 CUB？

如果说 Thrust 是 GPU 编程的"高铁"——快速、舒适、不用操心细节，那 CUB 就是"赛车"——需要你亲自调校，但能跑出极限速度。

CUB（CUDA UnBound）是 NVIDIA 官方提供的底层 CUDA 并行原语库，专为追求极致性能的开发者设计。

CUB 是什么？

CUB 提供了一套可复用的、高性能的 CUDA 并行原语，覆盖三个层级：

┌─────────────────────────────────────────┐ │ Device Level（设备级） │ ← 整个 GPU │ 如：DeviceReduce、DeviceSort │ ├─────────────────────────────────────────┤ │ Block Level（线程块级） │ ← 一个 Block 内 │ 如：BlockReduce、BlockScan │ ├─────────────────────────────────────────┤ │ Warp Level（线程束级） │ ← 一个 Warp（32线程）内 │ 如：WarpReduce、WarpScan │ └─────────────────────────────────────────┘

这三个层级对应 CUDA 的三个并行粒度，CUB 在每个层级都提供了经过深度优化的实现。

CUB 与 Thrust 的关系

你的代码 ↓ Thrust（高层接口，易用） ↓ CUB（底层实现，高性能） ← CUB 就在这里 ↓ CUDA PTX 指令

Thrust 内部大量调用 CUB。当你调用thrust::sort时，背后很可能就是 CUB 的DeviceRadixSort在工作。

CUB 的核心模块

1. Device Level：整个 GPU 的并行操作

这是最常用的层级，一行代码启动整个 GPU 的并行计算。

DeviceReduce（归约）

#include<cub/cub.cuh>intN=1000000;float*d_in;// GPU 输入数组float*d_out;// GPU 输出（单个值）// 第一步：查询需要多少临时空间void*d_temp=nullptr;size_t temp_bytes=0;cub::DeviceReduce::Sum(d_temp,temp_bytes,d_in,d_out,N);// 第二步：分配临时空间cudaMalloc(&d_temp,temp_bytes);// 第三步：执行归约cub::DeviceReduce::Sum(d_temp,temp_bytes,d_in,d_out,N);

为什么要两次调用？

CUB 的设计哲学：让调用者控制内存。第一次调用只是"询问"需要多少临时空间，第二次才真正执行。这样你可以复用内存、精确控制显存使用。

DeviceReduce 支持的操作：

cub::DeviceReduce::Sum(...)// 求和cub::DeviceReduce::Min(...)// 求最小值cub::DeviceReduce::Max(...)// 求最大值cub::DeviceReduce::ArgMin(...)// 求最小值的下标cub::DeviceReduce::ArgMax(...)// 求最大值的下标cub::DeviceReduce::Reduce(...)// 自定义归约操作

DeviceSort（排序）

// 基数排序（整数类型极快）cub::DeviceRadixSort::SortKeys(d_temp,temp_bytes,d_keys_in,d_keys_out,// 输入、输出N);// 按键值对排序cub::DeviceRadixSort::SortPairs(d_temp,temp_bytes,d_keys_in,d_keys_out,d_values_in,d_values_out,N);

DeviceScan（前缀扫描）

intinput[]={1,2,3,4,5};// 包含扫描（inclusive scan）// 结果：{1, 3, 6, 10, 15}cub::DeviceScan::InclusiveSum(d_temp,temp_bytes,d_in,d_out,N);// 排除扫描（exclusive scan）// 结果：{0, 1, 3, 6, 10}cub::DeviceScan::ExclusiveSum(d_temp,temp_bytes,d_in,d_out,N);

DeviceSelect（筛选）

// 筛选出满足条件的元素autoselect_positive=[]__device__(floatx){returnx>0;};int*d_num_selected;// 输出：筛选出了多少个cub::DeviceSelect::If(d_temp,temp_bytes,d_in,d_out,d_num_selected,N,select_positive);

2. Block Level：线程块内的协作

Block Level 的原语在__global__kernel 内部使用，让同一个 Block 内的线程高效协作。

BlockReduce（块内归约）

#include<cub/cub.cuh>__global__voidsum_kernel(float*data,float*result,intN){// 声明 BlockReduce 类型（256 个线程的块）usingBlockReduce=cub::BlockReduce<float,256>;// 在共享内存中分配临时空间__shared__ BlockReduce::TempStorage temp_storage;intidx=blockIdx.x*blockDim.x+threadIdx.x;floatval=(idx<N)?data[idx]:0.0f;// 块内所有线程协作求和floatblock_sum=BlockReduce(temp_storage).Sum(val);// 只有线程 0 写结果if(threadIdx.x==0){atomicAdd(result,block_sum);}}

关键点：TempStorage放在共享内存里，这是 Block Level 原语高效的秘密——共享内存比全局内存快几十倍。

BlockScan（块内前缀扫描）

__global__voidprefix_sum_kernel(int*data,int*output,intN){usingBlockScan=cub::BlockScan<int,256>;__shared__ BlockScan::TempStorage temp_storage;intidx=blockIdx.x*blockDim.x+threadIdx.x;intval=(idx<N)?data[idx]:0;intprefix_sum;BlockScan(temp_storage).InclusiveSum(val,prefix_sum);if(idx<N)output[idx]=prefix_sum;}

BlockLoad / BlockStore（协作式内存加载）

这是 CUB 独有的特色功能——让整个 Block 协作加载数据，充分利用内存带宽：

__global__voidprocess_kernel(float*input,float*output,intN){usingBlockLoad=cub::BlockLoad<float,256,4>;// 每线程加载4个元素usingBlockStore=cub::BlockStore<float,256,4>;__shared__union{BlockLoad::TempStorage load;BlockStore::TempStorage store;}temp_storage;floatitems[4];// 每个线程持有4个元素// 协作式加载：256线程 × 4元素 = 一次加载1024个元素BlockLoad(temp_storage.load).Load(input,items);// 处理数据for(inti=0;i<4;i++)items[i]*=2.0f;// 协作式存储BlockStore(temp_storage.store).Store(output,items);}

为什么这比普通加载快？CUB 会自动选择最优的内存访问模式（合并访问、向量化加载等），最大化内存带宽利用率。

3. Warp Level：32 线程的极速协作

Warp 是 GPU 的最小调度单位（32 个线程），Warp Level 原语利用硬件的 warp shuffle 指令，无需共享内存，速度极快。

__global__voidwarp_reduce_kernel(float*data,float*result){usingWarpReduce=cub::WarpReduce<float>;__shared__ WarpReduce::TempStorage temp_storage[4];// 4个warpintwarp_id=threadIdx.x/32;floatval=data[threadIdx.x];// Warp 内归约，无需 __syncthreads()floatwarp_sum=WarpReduce(temp_storage[warp_id]).Sum(val);if(threadIdx.x%32==0){atomicAdd(result,warp_sum);}}

实战：用 CUB 实现高性能点积

任务：计算两个大向量的点积sum(a[i] * b[i])。

#include<cub/cub.cuh>#include<cuda_runtime.h>// 自定义归约操作：先乘后加structDotProductOp{float*a;float*b;__device__floatoperator()(inti)const{returna[i]*b[i];}};floatdot_product(float*d_a,float*d_b,intN){// 使用 DeviceReduce::Reduce 配合转换迭代器cub::CountingInputIterator<int>counting_iter(0);// 转换迭代器：访问第 i 个元素时，自动计算 a[i]*b[i]autotransform_iter=cub::TransformInputIterator<float,DotProductOp,cub::CountingInputIterator<int>>(counting_iter,{d_a,d_b});float*d_result;cudaMalloc(&d_result,sizeof(float));void*d_temp=nullptr;size_t temp_bytes=0;cub::DeviceReduce::Sum(d_temp,temp_bytes,transform_iter,d_result,N);cudaMalloc(&d_temp,temp_bytes);cub::DeviceReduce::Sum(d_temp,temp_bytes,transform_iter,d_result,N);floatresult;cudaMemcpy(&result,d_result,sizeof(float),cudaMemcpyDeviceToHost);cudaFree(d_temp);cudaFree(d_result);returnresult;}

CUB 的设计哲学

1. 显式临时存储

// CUB 的风格：你来管内存void*d_temp=nullptr;size_t temp_bytes=0;cub::DeviceReduce::Sum(d_temp,temp_bytes,...);// 查询cudaMalloc(&d_temp,temp_bytes);// 分配cub::DeviceReduce::Sum(d_temp,temp_bytes,...);// 执行

好处：

• 可以在多次调用间复用临时空间
• 精确控制显存使用
• 避免隐式的cudaMalloc（它会触发 GPU 同步，很慢）

2. 策略可配置

CUB 的很多操作支持通过模板参数调整内部策略：

// 自定义 BlockReduce 的算法usingBlockReduce=cub::BlockReduce<float,// 数据类型256,// 线程数cub::BLOCK_REDUCE_WARP_REDUCTIONS// 算法：基于warp归约>;

3. 迭代器抽象

CUB 支持各种迭代器，让你无需额外内存就能做数据变换：

// 常量迭代器：所有元素都是 1.0fcub::ConstantInputIterator<float>ones(1.0f);// 计数迭代器：0, 1, 2, 3, ...cub::CountingInputIterator<int>counter(0);// 变换迭代器：对每个元素应用函数autosquared=cub::TransformInputIterator<float,Square,float*>(d_data,Square{});

性能对比：CUB vs 手写 Kernel

以归约（求和）为例，在 A100 GPU 上对 1 亿个 float 求和：

CUB 之所以快，是因为它针对每种 GPU 架构（Volta、Ampere、Hopper 等）都有专门调优的实现，这些优化积累了 NVIDIA 工程师多年的经验。

什么时候用 CUB？

✅ 适合用 CUB

1. 性能是首要目标，需要榨干 GPU 的每一分算力
2. 需要精确控制内存，避免隐式分配
3. 在自定义 kernel 内部需要高效的块级/束级原语
4. Thrust 性能不够，需要更底层的控制

⚠️ 可以先用 Thrust

1. 快速原型开发，性能要求不极致
2. 代码可读性优先
3. 算法逻辑复杂，不想被内存管理分心

快速上手

安装

CUB 随 CUDA Toolkit 自带（CUDA 11.0+），也可以单独从 GitHub 获取：

# 随 CUDA 自带，直接包含头文件即可#include <cub/cub.cuh>

编译

nvcc-arch=sm_80 my_program.cu-omy_program

一个完整的可运行示例

#include<cub/cub.cuh>#include<iostream>intmain(){constintN=10;inth_data[]={5,2,8,1,9,3,7,4,6,0};// 分配 GPU 内存int*d_in,*d_out;cudaMalloc(&d_in,N*sizeof(int));cudaMalloc(&d_out,sizeof(int));cudaMemcpy(d_in,h_data,N*sizeof(int),cudaMemcpyHostToDevice);// CUB 归约求和void*d_temp=nullptr;size_t temp_bytes=0;cub::DeviceReduce::Sum(d_temp,temp_bytes,d_in,d_out,N);cudaMalloc(&d_temp,temp_bytes);cub::DeviceReduce::Sum(d_temp,temp_bytes,d_in,d_out,N);intresult;cudaMemcpy(&result,d_out,sizeof(int),cudaMemcpyDeviceToHost);std::cout<<"Sum = "<<result<<std::endl;// Sum = 45cudaFree(d_in);cudaFree(d_out);cudaFree(d_temp);return0;}

总结

CUB 是什么？

CUB 是 NVIDIA 官方的底层 CUDA 并行原语库，提供设备级、块级、束级三个层次的高性能并行操作，是追求极致 GPU 性能的首选工具。

核心要点

1. 三个层级：Device（全 GPU）、Block（线程块内）、Warp（32线程内）
2. 显式内存管理：两步调用模式，精确控制临时空间
3. 极致性能：针对每种 GPU 架构深度优化，接近硬件理论峰值
4. Thrust 的底层：Thrust 的高性能来自于 CUB

学习路径建议

入门 GPU 编程 ↓ 学习 Thrust（高层，快速上手） ↓ 遇到性能瓶颈 ↓ 学习 CUB Device Level（最常用） ↓ 需要自定义 kernel 内部优化 ↓ 学习 CUB Block/Warp Level

一句话：如果 Thrust 是让你"能跑起来"，那 CUB 是让你"跑得飞快"。🚀

后记

2026年8月15日于上海，在claude opus 4.8辅助下完成。