CUDA实践指南(二十六)

简介:

数学库:
中等优先级:每当速度超过精度时使用快速数学库。
支持两种类型的运行时数学运算。 它们可以用它们的名字加以区分:一些名称带有前置下划线,而另一些则没有(例如,__functionName()functionName())。 遵循__functionName()命名约定的函数直接映射到硬件级别。 它们速度更快,但准确度稍低(如__sinf(x)__expf(x))。 functionName()命名约定的函数较慢但具有较高的准确性(例如,sinf(x)和expf(x))。 __sinf(x)__cosf(x)__expf(x)的吞吐量远远大于sinf(x),cosf(x)和expf(x)的吞吐量。 如果参数x的大小需要减小,后者变得更加昂贵(大约慢一个数量级)。 此外,在这种情况下,自变量缩减代码使用本地内存,由于本地内存的高延迟,这会更多地影响性能。 更多详细信息可在CUDA C编程指南中找到.
还要注意,无论何时计算相同参数的正弦和余弦,sincos系列指令都应该用于优化性能:

  • __sincosf()用于单精度快速数学(参见下一段)
  • sincosf()用于常规的单精度
  • sincos()用于双精度

nvcc的-use_fast_math编译器选项强制每个functionName()调用等效的__functionName()调用。 只要准确度优先于性能,就应该使用此开关。 超验函数通常是这种情况。 请注意,此开关仅在单精度浮点上有效。
中等优先:如果可能,优先选择速度更快,更专业化的数学函数,而不是更慢,更一般的数学函数。
对于小整数幂(例如$x^2$或$x^3$),显式乘法几乎肯定比使用pow()等一般指数运算例程更快。 虽然编译器优化改进不断寻求缩小这种差距,但明确的乘法(或使用等效的专用内联函数或宏)可以具有显着的优势。 当需要几个相同基数的幂(例如,其中$x^2$和$x^5$都靠近计算)时,这个优点增加,因为这有助于编译器进行其常见的子表达式消除(CSE)优化。
对于使用基数2或10的指数,使用函数exp2()或expf2()和exp10()或expf10()而不是函数pow()或powf()。 pow()和powf()在寄存器压力和指令计数方面都是重量级的函数,这是由于在大多数幂运算中产生了许多特殊情况,并且难以在整个基数和指数的整个范围内实现良好的精度。 另一方面,函数exp2(),exp2f(),exp10()和exp10f()在性能方面与exp()和expf()类似,并且可以比它们快十倍 pow()/ powf()等价物。
对于指数为1/3的指数,使用cbrt()或cbrtf()函数而不是泛指数函数pow()或powf(),因为前者明显快于后者。 同样,对于指数为-1/3的指数,使用rcbrt()或rcbrtf()。
用sincospi()替换sinpi(),cos(π )与cospi()和sincos(π )的sin(π )。 就精度和性能而言这是有利的。 作为一个特殊的例子,为了评估正弦函数而不是弧度,使用sinpi(x / 180.0)。 类似地,当函数参数的形式为π 时,单精度函数sinpif(),cospif()和sincospif()应该替换对sinf(),cosf()和sincosf()的调用。 (性能优势sinpi()超过sin()是由于简化了参数约简;精确度优势是因为sinpi()仅隐式乘以π,有效地使用无限精确的数学π而不是单精度或双精度逼近。)
精度相关的编译器标志:
默认情况下,nvcc编译器为计算能力2.x的设备生成符合IEEE标准的代码,但它也提供了生成代码的选项,这些代码的精确度稍低但速度更快,并且更接近为早期设备生成的代码:

  • -ftz = true(非规格化数字被刷新为零)
  • -prec-div = false(不太精确的划分)
  • -prec-sqrt = false(不精确的平方根)

另一个更积极的选项是-use_fast_math,它强制每个functionName()调用等效的__functionName()调用。 这使得代码运行速度更快,代价是精度和准确性降低。
内存指令:
高优先级:尽量减少全局内存的使用。 尽可能优先使用共享内存访问。
内存指令包括从共享,本地或全局内存读取或写入的任何指令。 访问未缓存的本地或全局内存时,存在400至600个时钟周期的内存延迟。 例如,以下示例代码中的赋值运算符具有较高的吞吐量,但关键的是,从全局内存中读取数据的延迟时间为400至600个时钟周期:

__shared__ float shared[32];
__device__ float device[32];
shared[threadIdx.x] = device[threadIdx.x];

如果在等待全局内存访问完成时可以发出足够的独立算术指令,那么线程调度程序可以隐藏大部分全局内存延迟。 但是,最好尽可能避免访问全局内存。

目录
相关文章
|
缓存 Shell Linux
【Shell 命令集合 链接器(linker)工具】Linux ld命令 将目标文件与库链接为可执行文件或库文件
【Shell 命令集合 链接器(linker)工具】Linux ld命令 将目标文件与库链接为可执行文件或库文件
460 0
|
机器学习/深度学习
【从零开始学习深度学习】23. CNN中的多通道输入及多通道输出计算方式及1X1卷积层介绍
【从零开始学习深度学习】23. CNN中的多通道输入及多通道输出计算方式及1X1卷积层介绍
【从零开始学习深度学习】23. CNN中的多通道输入及多通道输出计算方式及1X1卷积层介绍
|
9月前
|
人工智能 并行计算 程序员
【AI系统】SIMD & SIMT 与芯片架构
本文深入解析了SIMD(单指令多数据)与SIMT(单指令多线程)的计算本质及其在AI芯片中的应用,特别是NVIDIA CUDA如何实现这两种计算模式。SIMD通过单指令对多个数据进行操作,提高数据并行处理能力;而SIMT则在GPU上实现了多线程并行,每个线程独立执行相同指令,增强了灵活性和性能。文章详细探讨了两者的硬件结构、编程模型及硬件执行模型的区别与联系,为理解现代AI计算架构提供了理论基础。
1030 12
Element UI 自定义/修改下拉弹窗的样式(如级联选择器的下拉弹窗样式)
Element UI 自定义/修改下拉弹窗的样式(如级联选择器的下拉弹窗样式)
1046 0
|
人工智能 弹性计算 并行计算
技术改变AI发展:CUDA Graph优化的底层原理分析(GPU底层技术系列一)
随着人工智能(AI)的迅速发展,越来越多的应用需要巨大的GPU计算资源。CUDA是一种并行计算平台和编程模型,由Nvidia推出,可利用GPU的强大处理能力进行加速计算。
106191 1
|
Web App开发 人工智能 自然语言处理
2023 年最好的36款 AI 生产力工具(七)
本文主要展示了36 款 AI 应用,可以帮助读者更快、更好地工作。每个人都在与ChatGPT交流,从完整的博客文章到特定代码行的功能都在询问。其结果令人惊叹。虽然我们仍在探索如何将这项技术纳入我们的工作流程中,但明显的是,人工智能工具正在改变游戏规则。尽管ChatGPT是目前最受欢迎的,但它远不是首款进入市场的人工智能应用程序。经过Zapier团队的大量研究和测试,总结出了以下36款能够改变工作方式的人工智能生产力工具。
259 1
|
并行计算 C++ Windows
windows10下visual studio 2019安装以及cuda11配置
windows10下visual studio 2019安装以及cuda11配置
1652 0
|
Unix Linux Apache
(转)五种开源协议的对比分析BSD, Apache 2.0, GPL,LGPL,MIT
现今存在的开源协议很多,而经过Open Source Initiative组织通过批准的开源协议目前有58种(http://www.opensource.org/licenses /alphabetical)。
5677 0
|
关系型数据库 MySQL 数据库
MySQL 和 PostgreSQL,我到底选择哪个?
MySQL 和 PostgreSQL,我到底选择哪个?
14594 0
|
弹性计算 虚拟化 异构计算
阿里云GPU服务器价格表(Nvidia M40/P100/P4/V100)
阿里云GPU服务器价格表(Nvidia M40/P100/P4/V100)
1341 0