_lecture5-CUDA性能优化课件课件.ppt

? David Kirk/NVIDIA and Wen-mei W. Hwu Taiwan, June 30 – July 2, 2008 2013 HPC 课程HPC课程:GPU编程之CUDA 性能优化彭博 * 目标 将CUDA的性能优化知识用于实际工作 一些貌似有效的策略可能不一定能提高性能 在不同情况的应用中，不同的约束可能占主要作用 通过实例的学习帮助建立并行编程的思路，习惯和方法 算法模式可导致更好的效率以及更好的硬件利用率 * 并行归约 给定一个数组的数据，以一种并行的模式将它们减少到一个单独的数值 例子 求和归约: 一个数组中所有值的和 最大值归约: 一个数组中的最大值 典型的并行实现: 递归的减半线程，每个线程加两个值 N个元素需要log(N)步骤, 需要 N/2 线程 * 矢量归约例子-求和 假定用全局内存进行数据区域内的归约 原始的矢量存储于设备(device)的全局内存 使用全局内存存储一个部分的和矢量 每一次迭代会将部分和进一步相加，逐步达到最终的结果 最终的结果将会位于元素0 * 简单实现-用全局内存 假定数组partialSum中存有16个元素 unsigned int t = threadIdx.x; for (unsigned int stride = 1; stride < blockDim.x; stride *= 2) { __syncthreads(); if (t % (2*stride) == 0) partialSum[t] += partialSum[t+stride]; } * 具有分支发散特性的矢量归约 0 1 2 3 4 5 7 6 10 9 8 11 0+1 2+3 4+5 6+7 10+11 8+9 0...3 4..7 8..11 0..7 8..15 1 2 3 数组元素 迭代 Thread 0 Thread 8 Thread 2 Thread 4 Thread 6 Thread 10 * 较复杂实现-用共享内存 假定已经将数组导入到共享内存 __shared__ float partialSum[] unsigned int t = threadIdx.x; for (unsigned int stride = 1; stride < blockDim.x; stride *= 2) { __syncthreads(); if (t % (2*stride) == 0) partialSum[t] += partialSum[t+stride]; } * 具有分支发散特性的矢量归约 0 1 2 3 4 5 7 6 10 9 8 11 0+1 2+3 4+5 6+7 10+11 8+9 0...3 4..7 8..11 0..7 8..15 1 2 3 数组元素 迭代 Thread 0 Thread 8 Thread 2 Thread 4 Thread 6 Thread 10 * 一些观察 在每次迭代中，对于每个Warp，两个控制流路径将会串行的遍历执行 一些线程执行加操作，而另外一些线程就不会做加操作 这些不执行加操作的线程可能导致过多的时钟周期开销（取决于发散的实现过程） 在任何时刻，不超过一般的线程将会被执行 所有奇数索引线程从一开始就被禁用! 在一个Warp执行时间内，平均来说，少于1/4的线程被激活。 经过5次迭代后，每个Block中的整个Warp将会终止执行，造成低资源利用率，但是没有发散. * 该实现的缺点 假定已经将数组导入到共享内存 __shared__ float partialSum[] unsigned int t = threadIdx.x; for (unsigned int stride = 1; stride < blockDim.x; stride *= 2) { __syncthreads(); if (t % (2*stride) == 0) partialSum[t] += partialSum[t+stride]; } 缺点: 归因为间隔的分支预测，会产生分支发散 * 更好的实现 假定已经将数组导入到共享内存 __shared__ float partialSum[] unsigned int t = threadIdx.x; for (unsigned int stride = blockDim.x; stride > 1; stride >> 1) { __syncthreads(); if (t < stride) partialSum[t] += partialSum[t+strid