CUDA学习笔记(七)Kernel性能调节

这篇具有很好参考价值的文章主要介绍了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中寻求一个平衡。

总结文章来源地址https://www.toymoban.com/news/detail-722195.html

  • 在大多数情形下,并不存在唯一的metric可以精确的优化性能。
  • 哪个metric或者event对性能的影响大是由kernel具体的代码决定的。
  • 在众多相关的metric和event中寻求一个平衡。
  • Grid/blcok heuristics(启发) 为调节性能提供了不错的切入点。

到了这里,关于CUDA学习笔记(七)Kernel性能调节的文章就介绍完了。如果您还想了解更多内容,请在右上角搜索TOY模板网以前的文章或继续浏览下面的相关文章,希望大家以后多多支持TOY模板网!

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处: 如若内容造成侵权/违法违规/事实不符,请点击违法举报进行投诉反馈,一经查实,立即删除!

领支付宝红包 赞助服务器费用

相关文章

  • 出现错误(已解决)RuntimeError: CUDA error: no kernel image is available for execution on the device CUDA ker

    为什么把警告po出来,是因为警告可以让我们了解一些有用信息。 首先警告里的内容不可忽略,翻译过来就是NVIDIA RTX GeForce 3060Ti(我使用的服务器)支持的CUDA的算力为8.6,与当前的pytorch的版本不匹配。说白了就是CUDA和pytorch版本不一致。 当前的pytorch版本支持的CUDA的算力为

    2024年02月10日
    浏览(52)
  • 性能优化-OpenCL kernel 开发

    「发表于知乎专栏《移动端算法优化》」 本文主要介绍OpenCL的 Kernel,包括代码的实例以及使用注意的详解。 🎬个人简介:一个全栈工程师的升级之路! 📋个人专栏:高性能(HPC)开发基础教程 🎀CSDN主页 发狂的小花 🌄人生秘诀:学习的本质就是极致重复! 目录 一、概述

    2024年01月23日
    浏览(84)
  • RuntimeError: CUDA error: no kernel image is available for execution on the device

    导致的原因一般都是显卡算力和cuda或者torch版本不匹配 比如在conda中安装的pytorch=1.5.0 cuda=10.2 错误:RuntimeError: CUDA error: no kernel image is available for execution on the device 参考pytorch 报错 RuntimeError: CUDA error: no kernel image is available for execution on the device_可豌豆的博客-CSDN博客 则应该安装

    2024年02月15日
    浏览(60)
  • CUDA kernel errors might be asynchronously reported at some other API call 错误解决

    CUDA kernel errors might be asynchronously reported at some other API call 在运行基于pytorch的深度学习项目时,有时候会遇到上述错误,并且在报错时没有定位到正确的位置。 这里查阅了很多网上的相关资料,说是分类数目和模型里的实际分类数目不匹配,大家可以仔细查看一下这个。也有

    2024年02月13日
    浏览(37)
  • 小研究 - Java虚拟机垃圾收集器的性能分析与调节

    垃圾收集器是Java虚拟机(JVM)的核心组成部分之一,对Java虚拟机的性能有非常重要的影响。本文将介绍GC的工作原理以及对象回收算法,重点介绍JVM的分段回收技术;剖析JVM自带的GC性能分析工具;阐述如何通过命令行参数调节GC的运行,提

    2024年02月11日
    浏览(35)
  • 解决RuntimeError: CUDA error: no kernel image is available for execution on the deviceCUDA

    解决RuntimeError: CUDA error: no kernel image is available for execution on the deviceCUDA kernel errors might be asynchronously reported at some other API call,so the stacktrace below might be incorrect. 在服务器复现代码的时候,遇到了上述错误,解决办法如下。 .bashrc文件在服务器上初始页面的配置文件的地方 参考:

    2024年02月16日
    浏览(51)
  • [已解决]RuntimeError: CUDA error: no kernel image is available for execution on the device

    在ubuntu服务器上用python炼丹的时候遇到的两个问题,一个warning和一个runtimeErro,我的环境是用conda配置的,我就切换了一下环境,然后切回来就报这两个错误,期间啥也没干,之前重新安装opencv疯狂报错也是这种样子的。 warning warning:NVIDIA GeForce RTX 3090 with CUDA capability sm_86

    2024年02月02日
    浏览(70)
  • RuntimeError:CUDA error:no kernel image is available for execution on the device报错解决(亲测)

    调试Transformer网络,安装完timm包之后,运行程序时报错 CUDA error:no kernel image is available for execution on the device ,如图所示: 网上对于该错误说啥的都有,因为这是第一次遇到这个错误,之前训练CNN也正常,排除显卡算力低,不支持高版本CUDA问题。看来看去,这位博主说的有道

    2024年02月11日
    浏览(50)
  • 当出现RuntimeError:CUDA error:no kernel image is available for execution on the device 问题时候的pytorch安装方法

    当出现一个明显的特征就是出现: RuntimeError:CUDA error:no kernel image is av ailable for execution on the device 这就说明你的显卡太低了 可以到这个路径下C:Program FilesNVIDIA GPU Computing ToolkitCUDAv11.1extrasdemo_suite, 找到deviceQuenry.exe这个文件拖到cmd命令窗口运行可以看到自身电脑的算力  从

    2024年02月01日
    浏览(65)
  • 学习笔记:在Anaconda环境下装Pytorch,CUDA,pillow,numpy等库并寻找适配的Python版本

    目录 1.在Anaconda Prompt创建新环境 2.去pytorch官网上查找环境中Python所对应的pytorch版本并下载相关包 查询电脑独立NVIDIA显卡所适配的CUDA版本 CUDA版本查询 3.部分库版本的安装与修改  本人在Anaconda下创建的新环境为 python 3.7.0 pytorch 1.8.0 pillow 9.5.0 numpy 1.21.5 能够正常运行 如果我这

    2024年02月05日
    浏览(45)

觉得文章有用就打赏一下文章作者

支付宝扫一扫打赏

博客赞助

微信扫一扫打赏

请作者喝杯咖啡吧~博客赞助

支付宝扫一扫领取红包,优惠每天领

二维码1

领取红包

二维码2

领红包