企业官网建设流程全解析

网站网络推广教程,网上商城什么意思,wordpress+假用户插件,重庆森林经典台词1. Reduce算子优化入门#xff1a;从基础实现到性能翻倍 在GPU编程中#xff0c;Reduce算子是最基础也是最常用的操作之一。简单来说#xff0c;Reduce就是对数组中的元素进行归约计算#xff0c;比如求和#xff08;sum#xff09;、求最大值#xff08;max#xff0…1. Reduce算子优化入门从基础实现到性能翻倍在GPU编程中Reduce算子是最基础也是最常用的操作之一。简单来说Reduce就是对数组中的元素进行归约计算比如求和sum、求最大值max、求最小值min等。想象一下你有一长串数字需要计算它们的总和——这就是Reduce操作的典型场景。为什么Reduce在GPU上如此重要因为GPU擅长并行计算而Reduce操作虽然简单但在深度学习、科学计算等领域无处不在。一个优化良好的Reduce算子可以显著提升整体性能。我刚开始接触CUDA时第一个优化的算子就是Reduce当时性能提升了近200%这种成就感至今难忘。我们先来看最基础的实现版本Kernel 0__global__ void reduce_v0(float *g_idata, float *g_odata) { __shared__ float sdata[BLOCK_SIZE]; unsigned int tid threadIdx.x; unsigned int i blockIdx.x*blockDim.x threadIdx.x; sdata[tid] g_idata[i]; __syncthreads(); for(unsigned int s1; s blockDim.x; s * 2) { if (tid % (2*s) 0) { sdata[tid] sdata[tid s]; } __syncthreads(); } if (tid 0) g_odata[blockIdx.x] sdata[0]; }这个实现分为三个步骤每个线程加载一个数据到共享内存在共享内存中进行树形归约线程0将结果写回全局内存但实测下来这个基础版本在V100上的带宽利用率只有40.97%性能瓶颈非常明显。问题主要出在两个方面一是取模运算%性能很差二是条件判断导致warp divergence线程束分化。2. 性能优化第一战解决warp divergence和bank conflict2.1 消除warp divergence在Kernel 1中我们通过改变线程的工作方式来解决这两个问题__global__ void reduce_v1(float *g_idata, float *g_odata) { __shared__ float sdata[BLOCK_SIZE]; unsigned int tid threadIdx.x; unsigned int i blockIdx.x*blockDim.x threadIdx.x; sdata[tid] g_idata[i]; __syncthreads(); for(unsigned int s1; s blockDim.x; s * 2) { int index 2 * s * tid; if (index blockDim.x) { sdata[index] sdata[index s]; } __syncthreads(); } if (tid 0) g_odata[blockIdx.x] sdata[0]; }这个改进的关键在于不再是每个线程固定处理一个元素而是让活跃线程处理更多工作。这样在早期迭代中整个warp的线程要么都工作要么都空闲避免了warp divergence。实测性能从788.29us提升到502.43us加速比达到1.56倍。2.2 解决bank conflictKernel 1虽然解决了warp divergence但引入了新的问题——bank conflict。在共享内存中相邻地址位于不同的bank当线程访问间隔为2*s时可能导致多个线程访问同一个bank。这时候Kernel 2采用了顺序寻址的策略__global__ void reduce_v2(float *g_idata, float *g_odata) { __shared__ float sdata[BLOCK_SIZE]; unsigned int tid threadIdx.x; unsigned int i blockIdx.x*blockDim.x threadIdx.x; sdata[tid] g_idata[i]; __syncthreads(); for(unsigned int sblockDim.x/2; s0; s 1) { if (tid s) { sdata[tid] sdata[tid s]; } __syncthreads(); } if (tid 0) g_odata[blockIdx.x] sdata[0]; }这个版本让相邻线程访问相邻的内存地址避免了bank conflict。性能进一步提升到375.90us带宽利用率达到85.79%。从性能数据可以看出解决bank conflict带来的提升比解决warp divergence更显著。3. 高阶优化技巧充分利用硬件资源3.1 解决线程闲置问题观察前面的kernel会发现一个问题在归约过程中有一半线程会逐渐闲置。Kernel 3通过让每个线程处理更多数据来解决这个问题__global__ void reduce_v3(float *g_idata, float *g_odata) { __shared__ float sdata[BLOCK_SIZE]; unsigned int tid threadIdx.x; unsigned int i blockIdx.x*(blockDim.x*2) threadIdx.x; sdata[tid] g_idata[i] g_idata[i blockDim.x]; __syncthreads(); for(unsigned int sblockDim.x/2; s0; s 1) { if (tid s) { sdata[tid] sdata[tid s]; } __syncthreads(); } if (tid 0) g_odata[blockIdx.x] sdata[0]; }这个改动让每个线程在加载数据时就进行一次加法操作减少了闲置线程的数量。性能直接翻倍达到205.89us带宽利用率提升到81.72%。3.2 展开最后一个warp当活跃线程数小于等于32时即1个warp我们可以省略同步操作因为warp内的线程是天然同步的。Kernel 4实现了这个优化__device__ void warpReduce(volatile float* cache, unsigned int tid) { cache[tid] cache[tid32]; cache[tid] cache[tid16]; cache[tid] cache[tid8]; cache[tid] cache[tid4]; cache[tid] cache[tid2]; cache[tid] cache[tid1]; } __global__ void reduce_v4(float *g_idata, float *g_odata) { __shared__ float sdata[BLOCK_SIZE]; unsigned int tid threadIdx.x; unsigned int i blockIdx.x*(blockDim.x*2) threadIdx.x; sdata[tid] g_idata[i] g_idata[i blockDim.x]; __syncthreads(); for(unsigned int sblockDim.x/2; s32; s 1) { if (tid s) { sdata[tid] sdata[tid s]; } __syncthreads(); } if (tid 32) warpReduce(sdata, tid); if (tid 0) g_odata[blockIdx.x] sdata[0]; }这个版本性能提升到176.86us。需要注意的是对于计算能力7.0以上的GPU如V100需要使用__syncwarp()来保证正确性这就是Kernel 4.1的改进。4. 终极优化完全展开与向量化访存4.1 完全展开循环Kernel 5通过模板参数将循环完全展开让编译器生成更优化的指令template unsigned int blockSize __device__ void warpReduce(volatile float* cache, int tid) { if(blockSize 64) cache[tid] cache[tid32]; if(blockSize 32) cache[tid] cache[tid16]; if(blockSize 16) cache[tid] cache[tid8]; if(blockSize 8) cache[tid] cache[tid4]; if(blockSize 4) cache[tid] cache[tid2]; if(blockSize 2) cache[tid] cache[tid1]; } template unsigned blockSize __global__ void reduce_v5(float *g_idata, float *g_odata) { __shared__ float sdata[BLOCK_SIZE]; unsigned int tid threadIdx.x; unsigned int i blockIdx.x*(blockDim.x*2) threadIdx.x; sdata[tid] g_idata[i] g_idata[i blockDim.x]; __syncthreads(); if (blockSize 1024) { if (tid 512) sdata[tid] sdata[tid512]; __syncthreads(); } if (blockSize 512) { if (tid 256) sdata[tid] sdata[tid256]; __syncthreads(); } if (blockSize 256) { if(tid 128) sdata[tid] sdata[tid128]; __syncthreads(); } if (blockSize 128) { if (tid 64) sdata[tid] sdata[tid64]; __syncthreads(); } if (tid 32) warpReduceblockSize(sdata, tid); if (tid 0) g_odata[blockIdx.x] sdata[0]; }4.2 向量化访存最后的性能杀手锏是向量化访存一次性读取多个数据template typename T, int pack_size struct alignas(sizeof(T) * pack_size) Packed { T elem[pack_size]; __device__ void operator(PackedT, pack_size packA) { #pragma unroll for (int i 0; i pack_size; i) { elem[i] packA.elem[i]; } } }; __global__ void reduce_v8(float *g_idata, float *g_odata, unsigned int n) { __shared__ float warpLevelSums[kWarpSize]; unsigned int i blockDim.x * blockIdx.x threadIdx.x; Packedfloat, 4 sum_pack(0.0); const auto *pack_ptr reinterpret_castconst Packedfloat, 4*(g_idata); for (int linear_index i; linear_index n/4; linear_index blockDim.x*gridDim.x) { sum_pack pack_ptr[linear_index]; } float sum sum_pack.elem[0] sum_pack.elem[1] sum_pack.elem[2] sum_pack.elem[3]; // ... 后续reduce操作 }经过所有这些优化最终性能从最初的788.29us提升到162.21us加速比达到4.86倍带宽利用率提升到34.3%。虽然看起来带宽利用率不高但这已经接近Reduce这种低计算强度算子的理论极限了。