目录
背景
博主也在学习，根据网络上的资料+官方资料+ai分析+个人分析，所以有些观点可能错误，需要深度学习之后建立框架，在回头看现在的观点，可能大家的硬件不一样结果不一样，需要理论+实践进行合理分析，如有错误欢迎大家指正
上一章当中我们学到了一个访存效率的瓶颈，只能达到百分之10左右，核心原因在于L2缓存的不断换入换出，导致缓存命中率降低，从而不断的访问显存，降低了效率
本章将换一个实战场景，不再是简单的矩阵加法，从实战场景引出共享内存
并行归约
矩阵加法
C[i] = A[i] + B[i]中，A[i]和B[i]都只被读一次，算完就再也不会用了。
共享内存的优势是让一个线程块内的线程可以反复读取同一份数据。如果数据只被用一次，直接从全局内存读一次就够了，搬进共享内存反而多了一次搬运开销，毫无收益。并行归约就是提供一个数据复用的场景
“归约” 这个词，可以拆成“归”（递归迭代）和“约”（逐渐缩减）意思就是反复用同一种计算方式，把一堆数据一步步缩减成最终的一个结果。这完美契合加法或乘法的结合律和交换律。
任何具有结合性和交换性的计算，都可以使用归约
每次迭代计算方式都是相同的（归），从一组多个数据最后得到一个数（约）
int res=a+b+c+d+e+f+g+h; //那么可以先算a+b,c+d,e+f,g+h //然后再算(a+b+c+d)+(e+f+g+h) //最后由两部分的和得出结果 每一轮的计算方式完全相同，数据量却以对数速度减少（8 → 4 → 2 → 1）。 这种重复使用中间结果的特性，正是共享内存大显身手的地方。归约的方式基本包括如下几个步骤：
- 将输入向量划分到更小的数据块中
- 用一个线程块计算一个数据块的部分和
- 对每个数据块的部分和再求和得到最终的结果。
1.针对一个大型的数据，可能百万千万上亿级别的计算，一个block最多开1024线程数，不可能用一个block去处理很大的数据量，block和block又是独立的，没不能互相通信，不同的block是在不同的SM上并行的，为了适配SM并行，我们数据就要分块。
2.一个block里面有多个线程，那这些线程就可以通过协作的方式算一个数据块
3.针对block算出来的部分和，然后所有的block算出来的在加在一起，就是结果了
一般来说，算出来block的和之后，我们需要在cpu上加在一起才能得出结果
因为block之间无法通信，所以我们需要移到cpu上进行相加
一般来说得出来的部分和数据就会很小了，比如几百个，你就可以cpu上相加，但是如果数据量还是很大，那就建议在启动一个核函数即可
针对一个加法的场景
归约最常见的就是把数据分对，然后每个线程算一对，得到的结果作为输入继续分成对，迭代的进行，直到最后一个元素
相邻配对：元素与他们的相邻元素进行配对
交错配对：元素与一定距离的元素配对
以下代码是cpu上的一个交错配对的并行归约：
每一轮将数组的后半部分元素逐一加到前半部分对应元素上，结果存放在前半部分。
如果数组长度为奇数，单独把最后一个元素加到
data[0]。然后用新的长度
stride（约原长度的一半）递归调用。整个过程原地修改输入数组，最终结果保存在
data[0]。int recursiveReduceCpu(int *data, int const size) { if (size == 1) return data[0]; int const stride = size / 2; if (size % 2 == 1){ for (int i = 0; i < stride; i++) data[i] += data[i + stride]; data[0] += data[size - 1]; } else{ for (int i = 0; i < stride; i++) data[i] += data[i + stride]; } return recursiveReduce(data, stride); }下面我们就来通过不同的配对方式，不同的数据组织来看CUDA的执行效率。
之前提到过线程束分化，即使你不执行这个命令，你也要等到这轮结束
橙色圆圈就是每个线程做的操作，可见线程threadIdx.x=0 的线程进行了三次计算，下标奇数线程一致在陪跑，没做过任何计算，但是根据之前的介绍，这些线程虽然什么都不干，但是不可以执行别的指令，4号线程做了两步计算，2号和6号只做了一次计算。
图里的蓝色框，是每一轮被闲置、无有效数据的存储位置
整个朴素归约的过程，全程都是直接在全局内存（显存）里读写数据：
- 每一轮线程都从显存里读两个数，算完加法再写回显存，只会操作到部分的内存；
- 蓝色框标注的内存，就是只读了一次、后续迭代再也不会被访问的闲置位置 —— 它们还存在显存里，只是逻辑上不再使用。
这种朴素归约的瓶颈就在于你对同一个地方的显存反复利用，你没有加到缓存里面，而是不断的读写显存，性能会降低
//相邻分配 __global__ void reduceNeighbored(int * g_idata,int * g_odata,unsigned int n){ unsigned int tid = threadIdx.x; //块内索引 unsigned int gid=tid+blockIdx.x*blockDim.x; //全局索引 if (gid >= n) return; //越界则退出 int *idata = g_idata + blockIdx.x*blockDim.x; for (int stride = 1; stride < blockDim.x; stride *= 2) { if ((tid % (2 * stride)) == 0) { idata[tid] += idata[tid + stride]; } __syncthreads(); } if (tid == 0) g_odata[blockIdx.x] = idata[0]; }threadIdx.x就是这个块内的坐标
idata就是数据要进行分块，你要算这个块内的线程是算哪个数据块
blockIdx.x*blockDim.x就是进行分块用的
下面的stride就是进行步数，这里采用的是相邻分配
加完__syncthreads()保证同一轮内所有线程都完成读写后，再进入下一轮，避免数据竞争。
如果不加__syncthreads()就会出现读写脏数据，因为上一轮有些线程都没算完，你就开始用
这个数据块的结果放到这个块的0号处
接下来就是创建随机数，然后进行测试即可
#include <iostream> #include "../freshman.hpp" using namespace std; int recursiveReduceCpu(int *data, int const size){ if (size == 1) return data[0]; int const stride = size / 2; if (size % 2 == 1){ for (int i = 0; i < stride; i++) data[i] += data[i + stride]; data[0] += data[size - 1]; } else{ for (int i = 0; i < stride; i++) data[i] += data[i + stride]; } return recursiveReduceCpu(data, stride); } //预热函数 __global__ void warmUp(int * g_idata,int * g_odata,unsigned int n){ unsigned int tid = threadIdx.x; if (tid >= n) return; int *idata = g_idata + blockIdx.x*blockDim.x; for (int stride = 1; stride < blockDim.x; stride *= 2) { if ((tid % (2 * stride)) == 0) { idata[tid] += idata[tid + stride]; } __syncthreads(); } if (tid == 0) g_odata[blockIdx.x] = idata[0]; } //相邻分配 __global__ void reduceNeighbored(int * g_idata,int * g_odata,unsigned int n){ unsigned int tid = threadIdx.x; //块内索引 unsigned int gid=tid+blockIdx.x*blockDim.x; //全局索引 if (gid >= n) return; //越界则退出 int *idata = g_idata + blockIdx.x*blockDim.x; for (int stride = 1; stride < blockDim.x; stride *= 2) { if ((tid % (2 * stride)) == 0) { idata[tid] += idata[tid + stride]; } __syncthreads(); } if (tid == 0) g_odata[blockIdx.x] = idata[0]; } // 随机初始化整数数组 void initialData(int *ip, int size) { srand((unsigned)time(NULL)); // 设置随机种子 for (int i = 0; i < size; i++) { ip[i] = rand() & 0xFF; // 生成 0 ~ 255 之间的整数 } } int main(){ int dev=0; cudaSetDevice(dev); int num=1<<24; int blocksize=1024; dim3 block(blocksize); dim3 grid((num + block.x - 1) / block.x); //开辟内存 int* tmp=new int[num]; int* tmpCpu=new int[num]; int* tmpGpu=nullptr; int* resultGpu=nullptr; int* resultGpuToCpu=new int[grid.x]; //把gpu的结果传输给cpu printf("grid.x = %d\n", grid.x); CHECK(cudaMalloc((void**)&tmpGpu,sizeof(int)*num)); //开辟显存空间 CHECK(cudaMalloc((void**)&resultGpu,sizeof(int)*grid.x)); //开辟显存空间 //初始化数据 initialData(tmp,num); //启动cpu函数 memcpy(tmpCpu,tmp,sizeof(int)*num); double startCpu,endCpu,sumCpu=0; startCpu=efficiency(); sumCpu=recursiveReduceCpu(tmpCpu,num); endCpu=efficiency(); cout<<"cpu耗时："<<endCpu-startCpu<<endl; //先预热 CHECK(cudaMemcpy(tmpGpu,tmp,sizeof(int)*num,cudaMemcpyHostToDevice)); double startWarm,endWarm,sumWarm=0; startWarm=efficiency(); warmUp<<<grid,block>>>(tmpGpu,resultGpu,num); //等预热结束 cudaDeviceSynchronize(); endWarm=efficiency(); cout<<"warm耗时："<<endWarm-startWarm<<endl; //预热结束启动真正的核函数 double startGpu,endGpu,sumGpu=0; CHECK(cudaMemcpy(tmpGpu,tmp,sizeof(int)*num,cudaMemcpyHostToDevice)); startGpu=efficiency(); reduceNeighbored<<<grid,block>>>(tmpGpu,resultGpu,num); cudaDeviceSynchronize(); endGpu=efficiency(); //计算完之后拷贝回来 CHECK(cudaMemcpy(resultGpuToCpu,resultGpu,grid.x*sizeof(int), cudaMemcpyDeviceToHost)); //cpu最后进行计算总和 for(int i=0;i<grid.x;i++){ sumGpu+=resultGpuToCpu[i]; } cout<<"gpu耗时："<<endGpu-startGpu<<endl; cout<<"sumGpu："<<sumGpu<<endl; cout<<"sumCpu："<<sumCpu<<endl; if(sumGpu==sumCpu) cout<<"结果正确"<<endl; else{ cout<<"结果错误"<<endl; } return 0; }
可以看出gpu对于一些简单的运算，并行下还是非常快的
但是以上的朴素归约存在的问题：
1.线程束分化，对于一个加法，只会利用下标为偶数的线程，下标奇数的线程没有用到
2.访存，有些重复利用的空间是否可以加载到缓存当中，没必要每次都访问全局显存
一个block1024的线程，理论上来说需要10次归约才能算出结果，最后的读是2N，写是N，总的读写量达到3N
3.硬件是按照缓存行来读取的，一行是128个字节，即32个int大小，如果stride跨越非常大，到最后可能需要两个事务才能读取
之前提到过占有率，占有率=实际的warp数量/理论上SM可容纳的warp数量，无论这个warp里是一个线程再运行，其他31个全部空跑，它还是算一个，所以理论分析这个占有率应该不低，但是实际计算的效率是很低的
基于以上，我们使用ncu进行分析一下
一、线程束分化
指标
含义
smsp__thread_inst_executed_per_inst_executed.ratio每条指令平均有多少线程真正执行。理想值是 32（一个 Warp 全满）。数值越低，说明越多线程因分化而空转。
smsp__sass_branch_targets_threads_divergent.sum发生分化的分支指令总条数。这个数字越大，说明代码中导致 Warp 分裂的分支越多。
smsp__sass_branch_targets_threads_uniform.sum 分支目标一致的指令。
二、显存读写数据量
统计内核运行期间实际进出显存的总数据量，用于评估带宽压力和计算效率。
指标
含义
dram__bytes_read.sum内核执行期间从显存（DRAM）读取的总字节数。
dram__bytes_write.sum内核执行期间写入显存的总字节数。
gpu__compute_memory_throughput.avg.pct_of_peak_sustained_elapsed显存吞吐量占理论峰值的百分比。接近 100% 说明带宽被充分利用；较低则可能因访问模式不佳或缓存命中率高而未用满带宽。
三、访存效率 / 合并访问
衡量每次内存请求是否高效。理想情况下，一次请求就能拿到所需数据；非合并访问会导致多次请求，浪费带宽。
指标
含义
l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_ld.ratio每次全局加载请求平均访问的 L1 sector 数。理论最优为 1。数值越大，说明访存越不合并，一次请求被拆成了多次。
同理
..._st.ratio写操作的合并程度，含义同上，方向为存储。
四、缓存命中率
反映数据在 L1/L2 缓存中被找到的频率，直接影响数据访问延迟。
指标
含义
l1tex__t_sector_hit_rate.pctL1 缓存命中率（按 sector 计）。越高越好，意味着更多数据在最近的 L1 中找到，避免了更慢的 L2/DRAM 访问。
lts__t_sector_hit_rate.pctL2 缓存命中率。越高越好，可以减少对显存的访问。朴素归约中由于全局内存直接操作且无数据复用，L2 命中率通常很低。
五、Occupancy（占用率）
用于排除“Warp 数量不足”这个干扰因素。
指标
含义
sm__warps_active.avg.pct_of_peak_sustained_active平均活跃 Warp 占比。越高说明 SM 上有越多的 Warp 可供调度。
六、一把梭命令
ncu --set full -o report_xxx ./your_program会收集所有可用指标并生成一个报告文件（.ncu-rep），非常适合完整分析和截图存档。基于以上，我们进行测试
词根 含义 smStreaming Multiprocessor，流多处理器 smspSM Sub-Partition，SM 子分区 dramDynamic Random Access Memory，显存 l1texL1 Texture Cache，L1 纹理缓存 ltsL2 Cache，L2 缓存 gpu整个 GPU 层面 inst_executed已执行的指令 thread_inst_executed按线程数加权后的已执行指令 sassShader Assembly，GPU 的底层汇编指令 branch_targets分支跳转的目标地址 divergent分化的（线程跳转到不同目标） uniform统一的（所有线程跳转到同一个目标） avgAverage，平均值 pctPercentage，百分比 pct_of_peak_sustained_active持续活跃时间占理论峰值的百分比 per_inst_executed按每条已执行指令平均 ratio比率 sum总和 t_sectorTexture Sector，缓存的最小单元 hit_rate命中率 per_request按每次请求平均 pipe_lsuLoad/Store Unit Pipeline，加载/存储单元流水线 mem_global_op_ldGlobal Memory Load Operation，全局内存加载操作
活跃线程占有率是65.51：中等偏上，说明occupancy不是这个kernel的主要瓶颈，可以排除"是因为并行度不够"这个干扰项
访问显存存方面：显存带宽利用率只有24.43%，远没有跑满。这说明瓶颈 不是数据量太大、带宽不够用，而是访存模式本身效率低、没能把带宽用满。
访存合并效率： 理论最优值是1（warp内32个线程的访问刚好落在连续的、最少数量的sector里）。实际是2.67，说明 每次load/store请求平均要触碰2.67个sector ，访存严重不合并。这是本次实验最关键的发现：
idata[tid] += idata[tid + stride]这种写法，随着stride从1递增到512，同一个warp内活跃线程的访存地址间隔越来越大，导致访问越来越分散，无法合并成连续cache line，这直接拉低了显存带宽利用率（呼应上面24.43%这个低数字）。
缓存命中率：L1/L2命中率都不低，说明数据局部性本身还可以（归约过程中访问范围在不断收缩，旧数据容易留在缓存里）。 但这和"访存合并度低"是两件不同的事：命中率高只代表"数据多半能在缓存里找到"，不代表"访问请求本身打包得高效"——即使命中缓存，2.67 sector/request依然意味着同一个warp要发出更多次内存事务才能拿到所有线程需要的数据。
线程束分化：divergent占比 = 3,129,344 / (3,129,344+7,880,704) ≈ 28.4% 线程级有效利用率 = 28.39/32 ≈ 88.7%结论：存在中等程度的divergence。约28.4%的分支判断在warp内部意见不一致，主要集中在stride较小的早期迭代轮（
tid % (2*stride)==0这种写法导致同一warp内只有部分线程满足条件）；随着stride增大，越来越多warp整体满足/整体不满足该条件，分化逐渐减少。这一点divergence是性能损耗的一部分，但不是最大头，最大头应该是访存。取模分支判断 (tid % (2*stride)==0) │ ├──→ 28.4%分支divergent，线程级利用率降到88.7% ← 性能损耗来源① │ └──→ 活跃线程访存地址间隔不连续 (stride递增) │ └──→ sector/request = 2.67（理论最优1） ← 性能损耗来源②（更大头） │ └──→ 显存带宽利用率仅24.43%（远未跑满） （即使L1/L2命中率高达91.6%/84.4%， 依然无法弥补访问模式本身的低效）综合结论
两个理论瓶颈都被数据证实，但访存不合并的影响权重明显大于divergence：
- Divergence造成约11.3%的线程级浪费（28.39 vs 32）
- 访存不合并造成sector利用率仅37.5%（1/2.67），且最终反映为带宽利用率不到1/4
Occupancy 65.51%说明GPU没有闲着，有足够warp可调度，进一步印证瓶颈确实是"访存效率"和"指令利用率"这两个kernel内部的问题，而不是"并行度不够掩盖延迟"的问题。
优化reduceNeighboredLess
基于以上我们主要发现两个问题：
1.访存效率
2.线程束分化
太多的线程用不上，并且这些线程还必须等待这一轮指令的完成
- tid=0,2,4,6 工作；tid=1,3,5,7 摸鱼
- 工作的线程ID是离散的、间隔分布的：0, 2, 4, 6
优化：
还是采用相邻分配，但是这次我们让warp前面的线程工作，后面的线程不工作，把工作的集中到一块，也就是0123好线程工作，4567号线程不工作
__global__ void reduceNeighboredLess(int * g_idata,int *g_odata,unsigned int n) { unsigned int tid = threadIdx.x; unsigned idx = blockIdx.x*blockDim.x + threadIdx.x; int *idata = g_idata + blockIdx.x*blockDim.x; if (idx > n) return; for (int stride = 1; stride < blockDim.x; stride *= 2){ int index = 2 * stride *tid; if (index < blockDim.x){ idata[index] += idata[index + stride]; } __syncthreads(); } if (tid == 0) g_odata[blockIdx.x] = idata[0]; }工作的线程ID变成了连续的：tid=0,1,2,3；不工作的也是连续的：tid=4,5,6,7。
空闲线程数量没有减少。 真正变化的是空闲线程在warp内的排列方式。
那么放大来看，你的数据量可是1<<24这么大，一个线程束仍然32个线程，一个block里面有32个线程束，即1024个线程，那么这个block在执行的时候，我们把工作的线程集中在前面0-511，不工作的线程集中在后面512-1023，那么不工作的线程是连续的在同一个warp里面，那么是不是不会分化了
换句话说：divergence只会发生在边界所在的那一个warp，其余warp是uniform的。而原版几乎每一个warp（只要stride不太大）都会发生divergence，因为交错模式让每个warp内部都同时存在工作/摸鱼线程。
第一张图：快了将近一倍，但需要跑多次取平均值，这里博主只跑了几次，大概也是0.7，0.8倍左右
第二张图：
最主要的先看分支：完全符合理论预期。把交错分布的工作线程改成连续聚集分布后，divergence从"几乎每个warp都中招"降到"只有极少数边界warp"，线程级利用率从88.7%提升到99%。这是这个优化版本最干净、最直接的收益，证据非常扎实。（之前的分支大概是300多万条，这里仅仅有8万多条）
带宽利用率提升了约62%（相对值），这个提升不是来自访存变得更合并（马上会讲为什么），而主要来自：divergence减少后，SM的指令发射效率提高，单位时间内能发出更多有效访存请求，从而在相同时间窗口内推送更多数据，带宽利用率随之被"挤"上去了。
访存其实是一样的，sector的总量基本一致，主要是request数量不一样，这个比值是通过sector/request算出来的，优化版本的request更少，所以比值高
最重要的是要清楚第二个版本的一个warp的request更重，他要申请的sector更多两者最终带宽利用率：27.08% → 43.92%，净效果是提升的，说明在这个case里，"消除divergence带来的指令效率提升"完全压过了"单次request变重"带来的效益。
这是一个很好的论点：本次实验两个瓶颈不是独立可加的，优化一个可能会让另一个变差，最终结果取决于谁的影响权重更大。
reduceNeighboredLess通过将线程到数据的映射方式从"交错离散"改为"连续聚集"，几乎完全消除了warp divergence（divergent分支占比从28.4%降至0.74%，线程级有效利用率从88.7%提升至99.0%）。然而，这一映射方式同时带来单次的request变重总结
- 朴素归约有两个独立瓶颈：divergence（线程行为不一致）+ 访存效率（地址跨度大、不合并）
reduceNeighboredLess只改变了"哪些tid在工作"的分布方式（让活跃tid连续），没有改变地址的映射逻辑本身，所以：
- Divergence问题解决了（因为divergence看的是warp内行为是否一致，活跃tid连续=同一warp内行为一致）
- 访存效率问题不仅没解决，还加剧了单次request，最终等待时间可能会变久
- 综合带宽利用率提升，不是因为"访存变合并了"，而是因为"divergence消除后，请求发出的频率/效率提升了"，这是两个独立维度的此消彼长，本次实验中"频率提升"的收益压过了"单次效率下降"的损失
优化reduceInterLeaved
采用交错配对
//交错分配，线程uniform __global__ void reduceInterleaved(int *g_idata, int *g_odata, unsigned int n) { unsigned int tid = threadIdx.x; unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x; int *idata = g_idata + blockIdx.x * blockDim.x; if (idx >= n) return; for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) { if (tid < stride) { idata[tid] += idata[tid + stride]; } __syncthreads(); } if (tid == 0) g_odata[blockIdx.x] = idata[0]; }为什么这个可以合并？？？
首先要搞清楚
- A. 一个线程访问的两个地址之间的距离（即
idata[tid]和idata[tid+stride]之间隔多远）- B. 同一个warp内，32个线程访问的起始地址，相互之间是否连续
我们主要看得是B这种情况，而不是A，因为访问连续我们才可以合并，为什么上一轮合并效果不好，因为访问的是idata[0] [2] [4] [6]这样，效果怎么会好
而这一轮，
idata[tid]这部分，32个线程访问的是idata[0..31]——这32个地址是完全连续的！ 这就是一次标准的合并访问，32×4字节=128字节，对应4个sector，效率很高。
idata[tid+512]这部分，32个线程访问的是idata[512..543]——这也是连续的32个地址！ 同样是一次完美合并访问。
idata[tid] += idata[tid + stride];最重要的就是这条你每次相隔stride跨度，但是你每次的idata都是连续的128字节的地址
所以这样分配的好处可以访存合并
需要注意这个比值效率来说是：sector ÷ request
为什么第一次是2.67，因为它后面跨度非常大之后，导致一个线程束可能发出一次请求用一个sector就行，所以导致整体的ratio会变小，而后续测出来的3.58可以说明访存合并效果更好，一次访存需要差不多4个sector
把sector/request和gpu内存吞吐结合起来看，才是完整的故事：
- sector/request 衡量"单次请求的装载效率"（这次request有没有把sector用满），但它的分母会受divergence影响而失真，不能孤立地只看这一个数字判断好坏
- gpu__compute_memory_throughput 衡量"整体数据搬运的吞吐总量"，这个才是最终关心的"性能好不好"的直接体现
reduceInterleaved同时做到了：divergence很低（thread_inst_executed_per_inst_executed=31.67，接近满载32）+ sector/request合理（3.58，接近理论基准4，是健康的合并访问，不是因为divergence被稀释出来的假低值）+ 吞吐总量最高（42.17%，三个版本里仅次于reduceNeighboredLess的47.56%，但更健康）。虽然reduceNeighboredLess的47.56%高于reduceInterLeaved的42.17%，但是你搬运了很多垃圾值，而reduceInterLeaved虽然搬运的总量小，但是搬运的有用值多于reduceNeighboredLess，导致最后的时间上效果更好
总结
所有的实际情况来说，每个人的应用场景不同，硬件不同，测出来的效果都不一样
- 理论分析：从代码逻辑出发，预判可能存在的瓶颈（一开始就判断出divergence和访存效率两个候选瓶颈，这一步做得很好）
- 指标验证：用ncu具体测出每个理论瓶颈的严重程度，量化"到底有多差"
- 实测验证：用实际运行耗时（wall-clock time）确认优化是否真的让程序变快了，因为指标只能解释现象，不能替代真实的性能结果
这里主要注意第二个版本，减少线程束分化带来了带宽的提升，你搬运数据的总量得到了提升，但是你同时每次搬运sector/request垃圾值过多，但是这个实验来看，减少线程束分化带来的效益比较高，实际我们应该根据场景来判断
reduceInterleaved不是银弹，它只是这次实验范围内的阶段最优解，更进一步还有shared memory、warp shuffle等技术
通过学习，博主还发现一个结论，但是不明确是否正确，主要是第一个和第二个版本
其实分析来看都是相邻分配，sector总量应该来说是一样的，但是request不一样，因为第二个版本把工作线程压缩在一个warp里面，那就会导致发出的request更少
那么会不会存在一些特殊的情况导致第一个版本实际优于第二个版本呢？如果你的warp发出请求的时候，很慢，会阻塞，那就需要调度其他的warp上来，但是发现如果此时没有可用的，那就会等，所以效率就很慢
反观第一个版本，虽然你发出的request变多了，但是每次都很轻量，所以很快能返回，而且你是把每次都分摊在多个warp上面，所以如果有一个慢，那就调度别的上来运行
设想一种情况，如果你的数据量很大，那么你越到后面第二个版本的warp会很重，等很久又没有别的调度，但是第一个版本的warp很轻，所以会不会变快呢？总之就是用更多轻量的warp去提高并发度
但是还得看吞吐量，第二个版本减少了线程束分化还带来了吞吐量的提升
其实本质来说是上节学的隐藏延迟
上手实验：
这里不用管结果，可能是因为溢出还是啥吧，但是可以看出第二个版本比第一个版本慢太多了，这里的数据量是1<<30次方
当数据量足够大时，reduceNeighboredLess后期"重warp、少并发"的劣势会被放大，最终反超reduceNeighbored，变成更慢的版本。
等待的时间变长了，导致你的带宽效率都下降了，时间线拉长再加上ncu，效率变低了
所以
- 小数据规模(
1<<24)下：divergence是主导瓶颈，reduceNeighboredLess和reduceInterleaved都显著优于reduceNeighbored，吞吐量和实测耗时排序基本一致- 大数据规模(
1<<30)下：reduceNeighbored反而实测最快，reduceNeighboredLess明显最慢，说明divergence优化的收益和"warp级并行掩盖延迟的能力"之间存在权衡，数据规模是影响这个权衡天平倾向哪一边的关键变量- 吞吐量百分比和实测耗时的排序在大数据量下出现了不一致，说明在grid规模极大、调度压力本身成为显著因素时，ncu的吞吐量指标不能单独作为效率判断依据，必须结合实测耗时
所以没有任何一种技术是银弹，必须结合实际
433
扫一扫
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
