CUDA学习笔记(七)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 = 1<<14;
int ny = 1<<14;

下面的代码用来配置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 -arch=sm_20 sumMatrix.cu -o sumMatrix

Checking Active Warps with nvprof

在做各项数据比较的时候需要有个基准,这里使用四个block配置的时间消耗作为基准观察,分别为(32,32)(32,16)(16,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

从上面的输出可以得知两件事儿:

  1. 由于第二个配置比第一个有更多的block,device就会达到更多active warp(跟鸡蛋放在多个篮子的道理差不多)。也就是第二个性能优于第一个的原因。
  2. 第四个的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的硬件限制。
  • 最好的结果是第四个(128,2)。
  • 第一个启动了最多的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 Occupancy,load 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(启发) 为调节性能提供了不错的切入点。

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处:http://www.mzph.cn/news/114167.shtml

如若内容造成侵权/违法违规/事实不符,请联系多彩编程网进行投诉反馈email:809451989@qq.com,一经查实,立即删除!

相关文章

类和对象介绍

一、类 1.类的声明 class Box{//类名private://私有成员double width,length,height;//长&#xff0c;宽&#xff0c;高 public://公有成员void init(double l,double w,double h){//初始化 lengthl;widthw;heighth;}double S(){//求表面积 return (height*lengthheight*width…

【Lua语法】字符串

Lua语言中的字符串是不可变值。不能像在C语言中那样直接改变某个字符串中的某个字符&#xff0c;但是可以通过创建一个新字符串的方式来达到修改的目的 print(add2(1 , 2 ,15,3))a "no one"b string.gsub(a , "no" , "on1111")print(a) print…

【Python】AttributeError: module lib has no attribute X509_V_FLAG_CB_ISSUER_CHECK

问题&#xff1a; 运行脚本报错&#xff1a;module lib has no attribute X509_V_FLAG_CB_ISSUER_CHECK 原因&#xff1a; pyOpenSSL版本与python版本不匹配 解决方案&#xff1a; window系统 重新安装pyOpenSSL 1、卸载当前版本pyOpenSSL pip uninstall pyOpenSSL2、重新…

flink问题 集合

1.flink 乱码 在配置文件flink-conf.yaml增加配置项&#xff1a;env.java.opts: -Dfile.encodingUTF-8

神经网络的不同类型的层

神经网络可以包含多个不同类型的层&#xff0c;每种层都具有不同的特点和应用场景。以下是常见的神经网络层的分类及其详细特点和应用场景&#xff1a; 1. 输入层&#xff08;Input Layer&#xff09;&#xff1a; 特点&#xff1a;输入层通常不包含任何权重或激活函数&#…

第87步 时间序列建模实战:LSTM回归建模

基于WIN10的64位系统演示 一、写在前面 这一期&#xff0c;我们介绍大名鼎鼎的LSTM回归。 同样&#xff0c;这里使用这个数据&#xff1a; 《PLoS One》2015年一篇题目为《Comparison of Two Hybrid Models for Forecasting the Incidence of Hemorrhagic Fever with Renal…

Compose Desktop 使用中的几个问题(分平台加载资源、编写Gradle 任务下载平台资源、桌面特有组件、鼠标键盘事件)

前言 在我之前的文章 Compose For Desktop 实践&#xff1a;使用 Compose-jb 做一个时间水印助手 中&#xff0c;我们使用 Compose For Desktop 写了一个用于读取照片 EXIF 中的拍摄日期参数并以文字水印的方式添加到照片上的桌面程序。 但是事实上&#xff0c;这个程序的名字…

从入门到进阶 之 ElasticSearch SpringData 继承篇

&#x1f339; 以上分享 从入门到进阶 之 ElasticSearch SpringData 继承篇&#xff0c;如有问题请指教写。&#x1f339;&#x1f339; 如你对技术也感兴趣&#xff0c;欢迎交流。&#x1f339;&#x1f339;&#x1f339; 如有需要&#xff0c;请&#x1f44d;点赞&#x1f…

PHP数据加密传输和存储问题

PHP数据加密的类型 md5()&#xff0c;sha1()&#xff0c;crypt() 双md5加密加盐

AAPCS:最新的ARM子程序调用规则

AAPCS是arm公司发布的ARM架构应用程序二进制&#xff08;ABI&#xff09;程序调用接口&#xff0c;该文档由多个版本&#xff0c;博主第一次ARM程序调用规则是在《ARM体系与结构编程》&#xff0c;但书中描述的是ATPCS&#xff0c;AAPCS是ATPCS的升级版。后面去ARM官网看到了AA…

自然语言处理基础——词表示

词表示 把自然语言中最基本的语言单元——词转换为机器能够理解的 词表示能完成以下两个能力 词相似度计算 词与词之间语义的关系 近义词&上位词 使用近义词或上位词表示的问题 遗漏差异 遗漏新的释义 带有主观性 数据吸收 需要大量人工构建 One-Hot Representation …

Kafka学习(最新版3.6.0)

文章目录 一、初识MQ1.1 什么是MQ1.2 同步和异步通讯1.1.1 同步通讯1.1.2 异步通讯 1.3 技术对比1.4 MQ的两种模式 二、初识Kafka2.1 Kafka的使用场景2.2 Kafka基本概念2.3 Topic与Partition 三、Kafka基本使用3.1 部署前的准备3.2 启动kafka服务器3.3 Kafka核心概念之Topic3.4…

【每日一题Day354】LC2316统计无向图中无法互相到达点对数 | 并查集

统计无向图中无法互相到达点对数【LC2316】 给你一个整数 n &#xff0c;表示一张 无向图 中有 n 个节点&#xff0c;编号为 0 到 n - 1 。同时给你一个二维整数数组 edges &#xff0c;其中 edges[i] [ai, bi] 表示节点 ai 和 bi 之间有一条 无向 边。 请你返回 无法互相到达…

059:mapboxGL监听键盘事件,通过eastTo控制左右旋转

第059个 点击查看专栏目录 本示例是介绍演示如何在vue+mapbox中监听键盘事件,通过eastTo控制左右旋转。 本例通过easeTo方法来加减一定数值的bearing角度,通过.addEventListener的方法来监听键盘的按键动作。这里一定要设置interactive: false, 否则展现不出来旋转效果。 直…

机械设备经营小程序商城的作用是什么

由于机械设备厂商品牌需要各地招商代理&#xff0c;因此在管理方面也需要工具进行高效管理。如今各个行业都在开展数字化转型解决行业所遇难题或通过线上销售解决传统三公里难题及品牌扩张难题、用户消费渠道少等难题&#xff0c;构建会员体系精细化管理&#xff0c;同时还需要…

MySQL 主从复制原理

文章目录 1.主从复制方式1.1 异步复制1.2 半同步复制1.3 全同步复制 2.主从复制原理3.主从复制时推还是拉&#xff1f;参考文献 主从复制是 MySQL 高可用&#xff08;备份&#xff09;和高性能&#xff08;读写分离&#xff09;的基础&#xff0c;有了这个基础&#xff0c;MySQ…

qt 开发api文档地址

直达Qt5开发api文档&#xff0c;doc.qt.io/qt-5/classes.html 打开www.qt.io后右上角的Developers点进去&#xff0c;然后Qt Documentation下面的View All C Classes 点进去就是api了。

代码随想录算法训练营Day60|单调栈01

代码随想录算法训练营Day60|单调栈01 文章目录 代码随想录算法训练营Day60|单调栈01一、739. 每日温度二、496.下一个更大元素 I 一、739. 每日温度 class Solution {public int[] dailyTemperatures(int[] temperatures) {//单调栈int lenstemperatures.length;int result[]n…

Java设计模式之状态模式

状态模式&#xff08;State Pattern&#xff09;是一种行为型设计模式&#xff0c;它允许对象在内部状态改变时改变其行为。该模式将对象的行为封装在不同的状态类中&#xff0c;使得对象在不同的状态下可以有不同的行为&#xff0c;从而实现了状态的切换和行为的变化。 在状态…

【Ascend C算子开发(入门)】——Ascend C编程模式与范式

Ascend C编程模型与范式 1.并行计算架构抽象 Ascend C编程开发的算子是运行在AI Core上的&#xff0c;所以我们需要了解一下AI Core的结构。AI Core主要包括计算单元、存储单元、搬运单元。 计算单元包括了三种计算资源&#xff1a;Scalar计算单元&#xff08;执行标量计算&…