温州平阳县营销型网站建设,给别人做网站收8000贵不贵,广西注册公司网站,怎么上百度推广产品本篇博文转载于https://www.cnblogs.com/1024incn/tag/CUDA/#xff0c;仅用于学习。 Exposing Parallelism
这部分主要介绍并行分析#xff0c;涉及掌握nvprof的几个metric参数#xff0c;具体的这些调节为什么会影响性能会在后续博文解释。 代码准备
下面是我们的kernel…本篇博文转载于https://www.cnblogs.com/1024incn/tag/CUDA/仅用于学习。 Exposing Parallelism
这部分主要介绍并行分析涉及掌握nvprof的几个metric参数具体的这些调节为什么会影响性能会在后续博文解释。 代码准备
下面是我们的kernel函数sumMatrixOnGPUD
__global__ void sumMatrixOnGPU2D(float *A, float *B, float *C, int NX, int NY) {unsigned int ix blockIdx.x * blockDim.x threadIdx.x;unsigned int iy blockIdx.y * blockDim.y threadIdx.y;unsigned int idx iy * NX ix;if (ix NX iy NY) {C[idx] A[idx] B[idx];}
}
我们指定一个比较大的数据矩阵包含16384个元素
int nx 114;
int ny 114;
下面的代码用来配置main函数的参数也就是block的维度配置
if (argc 2) {dimx atoi(argv[1]);dimy atoi(argv[2]);
}
dim3 block(dimx, dimy);
dim3 grid((nx block.x - 1) / block.x, (ny block.y - 1) / block.y);
编译
$ nvcc -O3 -archsm_20 sumMatrix.cu -o sumMatrix
Checking Active Warps with nvprof
在做各项数据比较的时候需要有个基准这里使用四个block配置的时间消耗作为基准观察分别为323232,1616,32和16,16本文开始时有提到第一个参数是x象限维度第二个参数是y象限维度。
下面是几种配置的时间消耗输出结果
$ ./sumMatrix 32 32
sumMatrixOnGPU2D (512,512), (32,32) elapsed 60 ms
$ ./sumMatrix 32 16
sumMatrixOnGPU2D (512,1024), (32,16) elapsed 38 ms
$ ./sumMatrix 16 32
sumMatrixOnGPU2D (1024,512), (16,32) elapsed 51 ms
$ ./sumMatrix 16 16
sumMatrixOnGPU2D (1024,1024),(16,16) elapsed 46 ms
比较这几个结果不难发现最慢的是第一个32,32最快的是第二个32,16这里可以猜测的到的是拥有更多的block并行性更好。这个猜测可以使用nvprof 的achieved_occupancy这个metric参数来验证。该参数的定义公式在上一篇博文有介绍实际上就是指每个SM在每个cycle能够达到的最大active warp数目占总warp的比例。下面是使用该参数后得到的结果
$ nvprof --metrics achieved_occupancy ./sumMatrix 32 32
sumMatrixOnGPU2D (512,512), (32,32) Achieved Occupancy 0.501071
$ nvprof --metrics achieved_occupancy ./sumMatrix 32 16
sumMatrixOnGPU2D (512,1024), (32,16) Achieved Occupancy 0.736900
$ nvprof --metrics achieved_occupancy ./sumMatrix 16 32
sumMatrixOnGPU2D (1024,512), (16,32) Achieved Occupancy 0.766037
$ nvprof --metrics achieved_occupancy ./sumMatrix 16 16
sumMatrixOnGPU2D (1024,1024),(16,16) Achieved Occupancy 0.810691
从上面的输出可以得知两件事儿
由于第二个配置比第一个有更多的blockdevice就会达到更多active warp跟鸡蛋放在多个篮子的道理差不多。也就是第二个性能优于第一个的原因。第四个的achieved Occupancy最高但是却不是最快的因此较高的achieved Occupancy并不一定就意味着更好的性能也就是说还有更多的因素影响着GPU的性能。 checking memory operations with nvprof
对于C[idx] A[idx] B[idx]来说共有三个memory操作两个memory load和一个memory store。要查看这些操作的效率可以使用nvprof的两个metric参数如果想要查看memory的throughput则可使用gld_throughput
$ nvprof --metrics gld_throughput./sumMatrix 32 32
sumMatrixOnGPU2D (512,512), (32,32) Global Load Throughput 35.908GB/s
$ nvprof --metrics gld_throughput./sumMatrix 32 16
sumMatrixOnGPU2D (512,1024), (32,16) Global Load Throughput 56.478GB/s
$ nvprof --metrics gld_throughput./sumMatrix 16 32
sumMatrixOnGPU2D (1024,512), (16,32) Global Load Throughput 85.195GB/s
$ nvprof --metrics gld_throughput./sumMatrix 16 16
sumMatrixOnGPU2D (1024,1024),(16,16) Global Load Throughput 94.708GB/s
不难看到第四个拥有最高的load throughput但是却比第二个慢第二个也就是第四个的一半所以较高的load throughput也不一定就有较高的性能。之后讲到memory transaction时会具体分析这种现象的原因简单说就是高load throughput有可能是一种假象如果需要的数据在memory中存储格式未对齐不连续会导致许多额外的不必要的load操作所以本文中的efficiency会这么低。
然后我们可以使用nvprof的gld_efficiency来度量load efficiency该metric参数是指我们确切需要的global load throughput与实际得到global load memory的比值。这个metric参数可以让我们知道APP的load操作利用device memory bandwidth的程度
$ nvprof --metrics gld_efficiency ./sumMatrix 32 32
sumMatrixOnGPU2D (512,512), (32,32) Global Memory Load Efficiency 100.00%
$ nvprof --metrics gld_efficiency ./sumMatrix 32 16
sumMatrixOnGPU2D (512,1024), (32,16) Global Memory Load Efficiency 100.00%
$ nvprof --metrics gld_efficiency ./sumMatrix 16 32
sumMatrixOnGPU2D (1024,512), (16,32) Global Memory Load Efficiency 49.96%
$ nvprof --metrics gld_efficiency ./sumMatrix 16 16
sumMatrixOnGPU2D (1024,1024),(16,16) Global Memory Load Efficiency 49.80%
从上述结果可知最后两个的load efficiency只是前两个的一半。这也可以解释为什么较高的throughput和较高的Occupancy却没有产生较好的性能。尽管最后两个的load操作数目要多不少因为二者throughput较高但是他们的load effecitiveness却低不少由于efficiency较低。
观察最后两个可以发现他们block的x象限配置是warp的一半由前文推测知该象限应该保持为warp大小的整数倍。关于其具体原因将在后续博文详细解释。 Exposing More Parallelism
我们现在可以得出一个结论就是blockDim.x应该是warp大小的整数倍。这样做是很容易就提升了load efficiency。现在我们可能还有其他疑惑比如
继续调整blockDim.x是否会继续增加load throughput还有其他方法能增大并行性吗
现在我们重新整一个基准数据出来这两个问题可以从这个基准分析个大概
$ ./sumMatrix 64 2
sumMatrixOnGPU2D (256,8192), (64,2) elapsed 0.033567 sec
$ ./sumMatrix 64 4
sumMatrixOnGPU2D (256,4096), (64,4) elapsed 0.034908 sec
$ ./sumMatrix 64 8
sumMatrixOnGPU2D (256,2048), (64,8) elapsed 0.036651 sec
$ ./sumMatrix 128 2
sumMatrixOnGPU2D (128,8192), (128,2) elapsed 0.032688 sec
$ ./sumMatrix 128 4
sumMatrixOnGPU2D (128,4096), (128,4) elapsed 0.034786 sec
$ ./sumMatrix 128 8
sumMatrixOnGPU2D (128,2048), (128,8) elapsed 0.046157 sec
$ ./sumMatrix 256 2
sumMatrixOnGPU2D (64,8192), (256,2) elapsed 0.032793 sec
$ ./sumMatrix 256 4
sumMatrixOnGPU2D (64,4096), (256,4) elapsed 0.038092 sec
$ ./sumMatrix 256 8
sumMatrixOnGPU2D (64,2048), (256,8) elapsed 0.000173 sec
Error: sumMatrix.cu:163, code:9, reason: invalid configuration argument
从上面数据我们能够分析出来的是
最后一个配置256,8不可行block中总共的thread数目超过了1024这是GPU的硬件限制。最好的结果是第四个1282。第一个启动了最多的block但不是最快的。因为第二个与第四个在一个block中拥有相同数目的thread本应猜测二者有相同的表现但是实际却是第二个略逊色所以blockDim.x的大小是很关键的。剩下的相对第四个都有较少的block数目所以并行规模也是影响性能的关键因素。
现在我们又得猜测了拥有block最少的应该会有一个最低的achieved Occupancy吧而拥有最多block的应该会达到最高的achieved Occupancy吧为了验证这些想法我们再看一组数据
$ nvprof --metrics achieved_occupancy ./sumMatrix 64 2
sumMatrixOnGPU2D (256,8192), (64,2) Achieved Occupancy 0.554556
$ nvprof --metrics achieved_occupancy ./sumMatrix 64 4
sumMatrixOnGPU2D (256,4096), (64,4) Achieved Occupancy 0.798622
$ nvprof --metrics achieved_occupancy ./sumMatrix 64 8
sumMatrixOnGPU2D (256,2048), (64,8) Achieved Occupancy 0.753532
$ nvprof --metrics achieved_occupancy ./sumMatrix 128 2
sumMatrixOnGPU2D (128,8192), (128,2) Achieved Occupancy 0.802598
$ nvprof --metrics achieved_occupancy ./sumMatrix 128 4
sumMatrixOnGPU2D (128,4096), (128,4) Achieved Occupancy 0.746367
$ nvprof --metrics achieved_occupancy ./sumMatrix 128 8
sumMatrixOnGPU2D (128,2048), (128,8) Achieved Occupancy 0.573449
$ nvprof --metrics achieved_occupancy ./sumMatrix 256 2
sumMatrixOnGPU2D (64,8192), (256,2) Achieved Occupancy 0.760901
$ nvprof --metrics achieved_occupancy ./sumMatrix 256 4
sumMatrixOnGPU2D (64,4096), (256,4) Achieved Occupancy 0.595197
看到了吧64,2的achieved Occupancy竟然是最低的尽管他有最多的block高中做物理题也是这感觉它达到了硬件对block数量的限制。
第四个128,2和第七个256,2拥有拥有差不多的achieved Occupancy。我们对这两个再做一个试验再次增大将blockDim.y设置为1这也减少了block的大小
$ ./sumMatrix 128 1
sumMatrixOnGPU2D (128,16384),(128,1) elapsed 0.032602 sec
$ ./sumMatrix 256 1
sumMatrixOnGPU2D (64,16384), (256,1) elapsed 0.030959 sec 这次配置产生了最佳的性能特别是256,1要比128,1要更好再次检查achieved Occupancyload throughput和load efficiency
$ nvprof --metrics achieved_occupancy ./sumMatrix 256 1
$ nvprof --metrics gld_throughput ./sumMatrix 256 1
$ nvprof --metrics gld_efficiency ./sumMatrix 256 1 输出
Achieved Occupancy 0.808622
Global Load Throughput 69.762GB/s
Global Memory Load Efficiency 100.00%
现在可以看出最佳配置既不是拥有最高achieved Occupancy也不是最高load throughput的。所以不存在唯一metric来优化计算性能我么需要从众多metric中寻求一个平衡。
总结
在大多数情形下并不存在唯一的metric可以精确的优化性能。哪个metric或者event对性能的影响大是由kernel具体的代码决定的。在众多相关的metric和event中寻求一个平衡。Grid/blcok heuristics启发 为调节性能提供了不错的切入点。
