16
NV边缘计算 / Re: Jetson Nano 2GB 运行命令 sudo pip install Jetson.GPIO
« 于: 二月 01, 2021, 08:51:26 pm »
1. 那个信息是”警告“不是错误
2. 完整的指令请修改为 sudo -H pip3 install Jetson.GPIO
2. 完整的指令请修改为 sudo -H pip3 install Jetson.GPIO
列出帖子
该操作将允许你查看该会员所有的帖子，注意你只能看到你有权限看到的板块的帖子。
17
CUDA / Re: 调用双GPU如何做到GPU0与GPU1之间是并行计算的？
« 于: 一月 26, 2021, 02:58:39 pm »
使用openMP做并行调度,和正常cpu程序并行一样，来做双卡并行
18
CUDA / Re: 我是做cuda实时信号处理...
« 于: 一月 08, 2021, 02:16:48 pm »
答复：这个问题比较有意思, 涉及到了GPU的吞吐率和延迟的问题. 楼主说的40Mhz采样, 得到40MB/s, 等于说这个是8-bit采样, 每1ms内采样40K次. 同时, 楼主说了需要1ms就完成一个反馈或者控制, 等于说, GPU需要每批次只处理40KB数据, 还需要在1ms内处理完. 这是问题的前提.
那么我们就需要从吞吐率和延迟两个角度, 来分析楼主的"GPU很难做到实时"的问题. 如果楼主的GPU实现, 每次kernel调用需要较大的延迟, 例如25ms, 但是吞吐率足够(例如等效每秒能完成几个GB的数据的运算), 那么移动"CPU上的控制逻辑到GPU"是没用的. 因为你的延迟在那里. 此时不如考虑损失一定的GPU的吞吐率来换取处理的延迟了, 这个需要具体问题具体分析.
而如果楼主本身的GPU一次kernel调用的延迟足够, 则他并不需要将"CPU上的逻辑移动到GPU". 哪怕一次处理KB级别的数据降低了GPU的有效利用率, 但是满足任务要求就好.最坑爹的是kernel的延迟满足不了要求, 吞吐率性能也不能满足要求, 那样我估计楼主得好好重新设计, 整体改写贵kernel了....
那么我们就需要从吞吐率和延迟两个角度, 来分析楼主的"GPU很难做到实时"的问题. 如果楼主的GPU实现, 每次kernel调用需要较大的延迟, 例如25ms, 但是吞吐率足够(例如等效每秒能完成几个GB的数据的运算), 那么移动"CPU上的控制逻辑到GPU"是没用的. 因为你的延迟在那里. 此时不如考虑损失一定的GPU的吞吐率来换取处理的延迟了, 这个需要具体问题具体分析.
而如果楼主本身的GPU一次kernel调用的延迟足够, 则他并不需要将"CPU上的逻辑移动到GPU". 哪怕一次处理KB级别的数据降低了GPU的有效利用率, 但是满足任务要求就好.最坑爹的是kernel的延迟满足不了要求, 吞吐率性能也不能满足要求, 那样我估计楼主得好好重新设计, 整体改写贵kernel了....
19
CUDA / 我是做cuda实时信号处理...
« 于: 一月 07, 2021, 09:19:14 pm »
我是做cuda实时信号处理。信号采样率40 mhz，对应40m byte/s，送入GPU，但是控制逻辑是1ms一个更新。所以每用GPU处理1ms的数据，都要传回CPU做控制。回传数据只有几k字节。这样目前很难实时跑起来。您觉得是有什么优化方向，是需要把控制逻辑也移植到cuda下么？
20
CUDA / Re: 请问一个问题。warp shuffle操作中mask参数该怎么理解呢？
« 于: 一月 06, 2021, 08:05:20 pm »
答：Mask是参与shuffle的线程(lane)标志. Mask是一个无符号的32位整数, 每一位对于了warp中的1个线程, 正好32个.其中值为1的位, 表示参与。其中值为0的位，表示不参与。请注意一定要让设定成1的位真正的参与，否则该函数的结果将是未定义的。（例如你要求线程7参与shuffle，但是对应的mask中的第7位却为0， 则结果未定义）结果的未定义是对参与者本身（线程7）和其他线程（可能使用来自线程7的源值）来说的。
21
CUDA / 请问一个问题。warp shuffle操作中mask参数该怎么理解呢？
« 于: 一月 06, 2021, 08:04:09 pm »
请问一个问题。warp shuffle操作中mask参数该怎么理解呢？我看了doc没有看明白。也没有通过profile信息看出来。可以指导一下吗？谢谢
22
CUDA / 做cuda流同步时，cudastreamquery(0)是不是有特殊含义？
« 于: 十二月 31, 2020, 03:55:55 pm »
答复：在Windows下, 使用WDDM驱动的时候, 具有立刻推送命令到显卡上的效果.等效于不存在的函数cudaStreamFlush().(该函数并不存在. Linux下不需要. TCC下也不需要. 但WDDM下需要同时并不提供，而cudaStreamQuery提供了类似的效果)
23
CUDA / DAY6：优化CUDA应用程序
« 于: 十二月 21, 2020, 01:58:05 pm »
7.优化CUDA应用程序
在完成每一轮应用程序并行化之后，开发人员可以着手优化实现以提高性能。由于可以考虑许多可能的优化，因此充分了解应用程序的需求可以帮助使过程尽可能地流畅。但是，与整个APOD一样，程序优化是一个反复的过程（确定优化机会，应用和测试优化，验证所实现的加速并重复），这意味着程序员不必花费大量资金。在看到良好的加速效果之前，先记住所有可能的优化策略。相反，可以在学习策略时逐步应用它们。
从重叠的数据传输一直到计算一直到微调浮点运算序列，优化可以应用于各种级别。可用的概要分析工具对于指导该过程非常有用，因为它们可以帮助建议开发人员进行优化工作的最佳方案，并为该指南的优化部分的相关部分提供参考。
在完成每一轮应用程序并行化之后，开发人员可以着手优化实现以提高性能。由于可以考虑许多可能的优化，因此充分了解应用程序的需求可以帮助使过程尽可能地流畅。但是，与整个APOD一样，程序优化是一个反复的过程（确定优化机会，应用和测试优化，验证所实现的加速并重复），这意味着程序员不必花费大量资金。在看到良好的加速效果之前，先记住所有可能的优化策略。相反，可以在学习策略时逐步应用它们。
从重叠的数据传输一直到计算一直到微调浮点运算序列，优化可以应用于各种级别。可用的概要分析工具对于指导该过程非常有用，因为它们可以帮助建议开发人员进行优化工作的最佳方案，并为该指南的优化部分的相关部分提供参考。
24
CUDA / Day 5：得到正确的答案
« 于: 十二月 21, 2020, 01:46:16 pm »
6.得到正确的答案
获得正确答案显然是所有计算的主要目标。在并行系统上，可能会遇到传统的面向串行编程中通常没有的困难。其中包括线程问题，由于计算浮点值而产生的意外值，以及由于CPU和GPU处理器的操作方式不同而引起的挑战。本章研究了可能影响返回数据正确性的问题，并指出了适当的解决方案。
6.1 验证
6.1.1 参考比较
对任何现有程序进行修改的正确性验证的一个关键方面是建立某种机制，从而可以将来自代表性输入的先前已知良好的参考输出与新结果进行比较。进行每次更改后，请使用适用于特定算法的任何条件来确保结果匹配。有些人会期望按位相同的结果，这并非总是可能的，尤其是在涉及浮点运算的情况下；有关数值精度，请参见数值精度和精度。对于其他算法，如果实现与某个小的epsilon中的参考相匹配，则可以认为实现是正确的。
请注意，用于验证数值结果的过程也可以轻松扩展以验证性能结果。我们希望确保我们所做的每个更改都是正确的，并且可以提高性能（以及提高多少）。作为我们的周期性APOD流程不可或缺的一部分，经常检查这些内容将有助于确保我们尽快获得理想的结果。
6.1.2 单元测试
与上述参考比较有用的对应方法是，以易于在单位级别上验证的方式构造代码本身。例如，我们可以将CUDA内核编写为许多短__device__函数的集合，而不是一个大的整体__global__函数。将每个设备功能连接在一起之前，可以对其进行单独测试。
例如，除了内核的实际计算之外，许多内核还具有用于访问内存的复杂寻址逻辑。如果我们在引入大量计算之前分别验证寻址逻辑，那么这将简化以后的调试工作。 （请注意，CUDA编译器将不有助于写入全局内存的任何设备代码视为需要消除的死代码，因此，由于寻址逻辑，我们至少必须将某些内容写到全局内存中，以便成功应用这种策略。）
更进一步，如果将大多数功能定义为__host____device__而不是__device__函数，则可以在CPU和GPU上测试这些功能，从而增强我们对功能正确性的信心，并且不会有任何意外的差异在结果中。如果存在差异，那么这些差异将尽早发现并且可以在简单功能的上下文中理解。
作为有用的副作用，如果我们希望在我们的应用程序中同时包含CPU和GPU执行路径，则此策略将使我们能够减少代码重复：如果CUDA内核的大部分工作是在__host____device__函数中完成的，我们可以轻松地从主机代码和设备代码中调用这些函数，而无需重复。
6.2 调试
CUDA-GDB是在Linux和Mac上运行的GNU调试器的端口；请参阅：https：//developer.nvidia.com/cuda-gdb。
适用于Microsoft Windows 7，Windows HPC Server 2008，Windows 8.1和Windows 10的NVIDIA Nsight Visual Studio Edition作为Microsoft Visual Studio的免费插件提供；请参阅：https：//developer.nvidia.com/nsight-visual-studio-edition。
一些第三方调试器也支持CUDA调试。有关更多详细信息，请参见：https：//developer.nvidia.com/debugging-solutions。
6.3 数值精度和精确度
错误或意外结果主要是由于浮点值的计算和存储方式导致的浮点精度问题所致。以下各节介绍了感兴趣的主要项目。在《 CUDA C ++编程指南的功能和技术规范》以及白皮书和随附的有关浮点精度和性能的网络研讨会中提供了浮点算术的其他特性，可从http://developer.nvidia.com/content/获得。精度性能浮点数和ieee 754合规NVIDIA gpus。
6.3.1 单精度与双精度
计算能力为1.3及更高版本的设备为双精度浮点值（即64位宽的值）提供了本机支持。由于前者的精度更高，并且由于舍入问题，使用双精度算术获得的结果通常会与通过单精度算术执行的相同操作不同。因此，重要的是要确保比较相似精度的值，并在一定时间内表达结果。
6.3.2 浮点数学不具有关联性
每个浮点算术运算都涉及一定数量的舍入。因此，执行算术运算的顺序很重要。如果A，B和C是浮点值，则不能保证（A + B）+ C等于符号数学中的A +（B + C）。并行化计算时，可能会更改操作顺序，因此并行结果可能与顺序结果不匹配。此限制并非特定于CUDA，而是浮点值并行计算的固有部分。
6.3.3 符合IEEE 754
所有CUDA计算设备均遵循IEEE 754标准的二进制浮点表示形式，但有一些小例外。这些异常在CUDA C ++编程指南的功能和技术规范中进行了详细说明，其结果可能导致与主机系统上计算的IEEE 754值不同。
关键区别之一是融合乘法加法（FMA）指令，该指令将乘法加法运算组合为单个指令执行。其结果通常与分别执行两个操作所获得的结果略有不同。
6.3.4 x86 80位计算
x86处理器在执行浮点计算时可以使用80位双精度扩展数学运算。这些计算的结果可能经常不同于在CUDA设备上执行的纯64位操作。要使值之间更接近匹配，请将x86主机处理器设置为使用常规双精度或单精度（分别为64位和32位）。这可以通过FLDCW x86汇编指令或等效的操作系统API来完成。
获得正确答案显然是所有计算的主要目标。在并行系统上，可能会遇到传统的面向串行编程中通常没有的困难。其中包括线程问题，由于计算浮点值而产生的意外值，以及由于CPU和GPU处理器的操作方式不同而引起的挑战。本章研究了可能影响返回数据正确性的问题，并指出了适当的解决方案。
6.1 验证
6.1.1 参考比较
对任何现有程序进行修改的正确性验证的一个关键方面是建立某种机制，从而可以将来自代表性输入的先前已知良好的参考输出与新结果进行比较。进行每次更改后，请使用适用于特定算法的任何条件来确保结果匹配。有些人会期望按位相同的结果，这并非总是可能的，尤其是在涉及浮点运算的情况下；有关数值精度，请参见数值精度和精度。对于其他算法，如果实现与某个小的epsilon中的参考相匹配，则可以认为实现是正确的。
请注意，用于验证数值结果的过程也可以轻松扩展以验证性能结果。我们希望确保我们所做的每个更改都是正确的，并且可以提高性能（以及提高多少）。作为我们的周期性APOD流程不可或缺的一部分，经常检查这些内容将有助于确保我们尽快获得理想的结果。
6.1.2 单元测试
与上述参考比较有用的对应方法是，以易于在单位级别上验证的方式构造代码本身。例如，我们可以将CUDA内核编写为许多短__device__函数的集合，而不是一个大的整体__global__函数。将每个设备功能连接在一起之前，可以对其进行单独测试。
例如，除了内核的实际计算之外，许多内核还具有用于访问内存的复杂寻址逻辑。如果我们在引入大量计算之前分别验证寻址逻辑，那么这将简化以后的调试工作。 （请注意，CUDA编译器将不有助于写入全局内存的任何设备代码视为需要消除的死代码，因此，由于寻址逻辑，我们至少必须将某些内容写到全局内存中，以便成功应用这种策略。）
更进一步，如果将大多数功能定义为__host____device__而不是__device__函数，则可以在CPU和GPU上测试这些功能，从而增强我们对功能正确性的信心，并且不会有任何意外的差异在结果中。如果存在差异，那么这些差异将尽早发现并且可以在简单功能的上下文中理解。
作为有用的副作用，如果我们希望在我们的应用程序中同时包含CPU和GPU执行路径，则此策略将使我们能够减少代码重复：如果CUDA内核的大部分工作是在__host____device__函数中完成的，我们可以轻松地从主机代码和设备代码中调用这些函数，而无需重复。
6.2 调试
CUDA-GDB是在Linux和Mac上运行的GNU调试器的端口；请参阅：https：//developer.nvidia.com/cuda-gdb。
适用于Microsoft Windows 7，Windows HPC Server 2008，Windows 8.1和Windows 10的NVIDIA Nsight Visual Studio Edition作为Microsoft Visual Studio的免费插件提供；请参阅：https：//developer.nvidia.com/nsight-visual-studio-edition。
一些第三方调试器也支持CUDA调试。有关更多详细信息，请参见：https：//developer.nvidia.com/debugging-solutions。
6.3 数值精度和精确度
错误或意外结果主要是由于浮点值的计算和存储方式导致的浮点精度问题所致。以下各节介绍了感兴趣的主要项目。在《 CUDA C ++编程指南的功能和技术规范》以及白皮书和随附的有关浮点精度和性能的网络研讨会中提供了浮点算术的其他特性，可从http://developer.nvidia.com/content/获得。精度性能浮点数和ieee 754合规NVIDIA gpus。
6.3.1 单精度与双精度
计算能力为1.3及更高版本的设备为双精度浮点值（即64位宽的值）提供了本机支持。由于前者的精度更高，并且由于舍入问题，使用双精度算术获得的结果通常会与通过单精度算术执行的相同操作不同。因此，重要的是要确保比较相似精度的值，并在一定时间内表达结果。
6.3.2 浮点数学不具有关联性
每个浮点算术运算都涉及一定数量的舍入。因此，执行算术运算的顺序很重要。如果A，B和C是浮点值，则不能保证（A + B）+ C等于符号数学中的A +（B + C）。并行化计算时，可能会更改操作顺序，因此并行结果可能与顺序结果不匹配。此限制并非特定于CUDA，而是浮点值并行计算的固有部分。
6.3.3 符合IEEE 754
所有CUDA计算设备均遵循IEEE 754标准的二进制浮点表示形式，但有一些小例外。这些异常在CUDA C ++编程指南的功能和技术规范中进行了详细说明，其结果可能导致与主机系统上计算的IEEE 754值不同。
关键区别之一是融合乘法加法（FMA）指令，该指令将乘法加法运算组合为单个指令执行。其结果通常与分别执行两个操作所获得的结果略有不同。
6.3.4 x86 80位计算
x86处理器在执行浮点计算时可以使用80位双精度扩展数学运算。这些计算的结果可能经常不同于在CUDA设备上执行的纯64位操作。要使值之间更接近匹配，请将x86主机处理器设置为使用常规双精度或单精度（分别为64位和32位）。这可以通过FLDCW x86汇编指令或等效的操作系统API来完成。
25
CUDA / Day 4:开始并行化应用程序
« 于: 十二月 21, 2020, 01:29:19 pm »
4并行化应用程序
在确定了热点并完成了设置目标和期望的基本练习之后，开发人员需要并行化代码。根据原始代码的不同，这可以像调用现有的GPU优化库（如cuBLAS、cuFFT或thrush）一样简单，也可以像向并行化编译器添加一些预处理器指令一样简单。
另一方面，一些应用程序的设计将需要一些重构来暴露它们固有的并行性。由于即使CPU架构需要公开这种并行性，以改善或简单地维持顺序应用程序的性能，CUDA并行编程语言家族（CUDA C++、CUDA FORTRAN等）的目的是使这种并行性的表达尽可能简单，同时在支持CUDA的GPU上启用操作，以实现最大并行吞吐量。
5入门
并行化顺序代码有几个关键策略。虽然如何将这些策略应用到特定应用程序的细节是一个复杂且特定于问题的主题，但是不管我们是在多核cpu上运行代码还是在cudagpu上使用，这里列出的一般主题都适用。
5.1 并行库
应用程序并行化最直接的方法是利用现有的库，这些库代表我们利用并行体系结构。CUDA工具包包括许多这样的库，它们已经针对NVIDIA CUDA gpu进行了微调，例如cuBLAS、cuFFT等。
这里的关键在于，当库与应用程序的需求很好地匹配时，它们是最有用的。例如，在其他应用程序中，使用线性代数来切换的应用程序很少。其他CUDAToolkit库也是如此：cuFFT有一个类似于FFTW的接口，等等。
还有一点是推力库，它是一个类似C++ C++模板库的并行C++模板库。Thrust提供了丰富的数据并行原语集合，如scan、sort和reduce，它们可以组合在一起，用简洁易读的源代码实现复杂的算法。通过用这些高级抽象来描述您的计算，您可以让push自由地选择最有效的实现。因此，推力可以用于CUDA应用程序的快速原型设计（程序员的生产力最为重要），以及在健壮性和绝对性能至关重要的生产中。
5.2 并行编译器
序列代码并行化的另一种常见方法是利用并行化编译器。这通常意味着使用基于指令的方法，程序员使用pragma或其他类似的符号来向编译器提供关于在哪里可以找到并行性的提示，而不需要修改或调整底层代码本身。通过向编译器公开并行性，指令允许编译器执行将计算映射到并行体系结构的详细工作。
OpenACC标准提供了一组编译器指令，以指定标准C、C++和FORTRAN中的循环和代码区域，这些代码和代码区域应该从主机CPU卸载到附加的加速器，例如CUDA GPU。ACC的详细信息由编译器的运行时管理隐式处理。
看到了吗http://www.openacc.org/了解详情。
5.3 暴露并行编码
对于需要现有并行库或并行化编译器无法提供的功能或性能的应用程序，与现有顺序代码无缝集成的并行编程语言（例如CUDA C ++）至关重要。
一旦我们在应用程序的配置文件评估中找到了热点，并确定自定义代码是最佳方法，便可以使用CUDA C ++将这部分代码的并行性作为CUDA内核公开。然后，我们可以将该内核启动到GPU上并检索结果，而无需对应用程序的其余部分进行重大重写。
当我们的应用程序的总运行时间的大部分都花在代码的几个相对隔离的部分中时，这种方法最简单。更难以并行化的是具有非常平坦的配置文件的应用程序，即，所花费的时间在整个代码库中相对平均地分布的应用程序。对于后一种应用程序，可能需要某种程度的代码重构以暴露应用程序中固有的并行性，但是请记住，这种重构工作将倾向于使所有将来的体系结构（CPU和GPU都受益），因此值得这样做付出必要的努力。
在确定了热点并完成了设置目标和期望的基本练习之后，开发人员需要并行化代码。根据原始代码的不同，这可以像调用现有的GPU优化库（如cuBLAS、cuFFT或thrush）一样简单，也可以像向并行化编译器添加一些预处理器指令一样简单。
另一方面，一些应用程序的设计将需要一些重构来暴露它们固有的并行性。由于即使CPU架构需要公开这种并行性，以改善或简单地维持顺序应用程序的性能，CUDA并行编程语言家族（CUDA C++、CUDA FORTRAN等）的目的是使这种并行性的表达尽可能简单，同时在支持CUDA的GPU上启用操作，以实现最大并行吞吐量。
5入门
并行化顺序代码有几个关键策略。虽然如何将这些策略应用到特定应用程序的细节是一个复杂且特定于问题的主题，但是不管我们是在多核cpu上运行代码还是在cudagpu上使用，这里列出的一般主题都适用。
5.1 并行库
应用程序并行化最直接的方法是利用现有的库，这些库代表我们利用并行体系结构。CUDA工具包包括许多这样的库，它们已经针对NVIDIA CUDA gpu进行了微调，例如cuBLAS、cuFFT等。
这里的关键在于，当库与应用程序的需求很好地匹配时，它们是最有用的。例如，在其他应用程序中，使用线性代数来切换的应用程序很少。其他CUDAToolkit库也是如此：cuFFT有一个类似于FFTW的接口，等等。
还有一点是推力库，它是一个类似C++ C++模板库的并行C++模板库。Thrust提供了丰富的数据并行原语集合，如scan、sort和reduce，它们可以组合在一起，用简洁易读的源代码实现复杂的算法。通过用这些高级抽象来描述您的计算，您可以让push自由地选择最有效的实现。因此，推力可以用于CUDA应用程序的快速原型设计（程序员的生产力最为重要），以及在健壮性和绝对性能至关重要的生产中。
5.2 并行编译器
序列代码并行化的另一种常见方法是利用并行化编译器。这通常意味着使用基于指令的方法，程序员使用pragma或其他类似的符号来向编译器提供关于在哪里可以找到并行性的提示，而不需要修改或调整底层代码本身。通过向编译器公开并行性，指令允许编译器执行将计算映射到并行体系结构的详细工作。
OpenACC标准提供了一组编译器指令，以指定标准C、C++和FORTRAN中的循环和代码区域，这些代码和代码区域应该从主机CPU卸载到附加的加速器，例如CUDA GPU。ACC的详细信息由编译器的运行时管理隐式处理。
看到了吗http://www.openacc.org/了解详情。
5.3 暴露并行编码
对于需要现有并行库或并行化编译器无法提供的功能或性能的应用程序，与现有顺序代码无缝集成的并行编程语言（例如CUDA C ++）至关重要。
一旦我们在应用程序的配置文件评估中找到了热点，并确定自定义代码是最佳方法，便可以使用CUDA C ++将这部分代码的并行性作为CUDA内核公开。然后，我们可以将该内核启动到GPU上并检索结果，而无需对应用程序的其余部分进行重大重写。
当我们的应用程序的总运行时间的大部分都花在代码的几个相对隔离的部分中时，这种方法最简单。更难以并行化的是具有非常平坦的配置文件的应用程序，即，所花费的时间在整个代码库中相对平均地分布的应用程序。对于后一种应用程序，可能需要某种程度的代码重构以暴露应用程序中固有的并行性，但是请记住，这种重构工作将倾向于使所有将来的体系结构（CPU和GPU都受益），因此值得这样做付出必要的努力。
26
CUDA / DAY 3:应用程序分析
« 于: 十二月 14, 2020, 09:12:01 pm »
3.1。应用Profiling
许多代码用相对较少的代码即可完成大部分工作。使用 profiler,，开发人员可以识别此类热点并开始编译候选列表以进行并行化。
3.1.1。创建配置文件
有很多可能的方法来分析代码，但是在所有情况下，目标都是相同的：识别应用程序将花费大部分执行时间的一个或多个函数。
注意：高优先级：为了最大程度地提高开发人员的工作效率，请对应用程序进行概要分析，以确定热点和瓶颈。
任何概要分析活动的最重要考虑因素是确保工作负荷是现实的-即，从测试和基于该信息的决策中获得的信息与真实数据相关。使用不切实际的工作负载会导致开发人员针对不切实际的问题大小进行优化，或者使开发人员将精力集中在错误的功能上，从而导致结果欠佳和工作浪费。
有许多工具可用于生成配置文件。以下示例基于gprof，这是来自GNU Binutils集合的Linux平台的开源探查器。
3.1.2。识别热点
在上面的示例中，我们可以清楚地看到该函数 genTimeStep（）占用应用程序总运行时间的三分之一。这应该是我们第一个并行化的候选函数。《了解扩展》讨论了我们从这种并行化中可能期望获得的潜在收益。
值得注意的是，以上示例中的其他几个功能也占据了整个运行时间的很大一部分，例如 calcStats（） 和 calcSummaryData（）。将这些功能并行化也应该增加我们的加速潜力。但是，由于APOD是一个循环过程，我们可能会选择在随后的APOD传递中并行处理这些功能，从而将我们在任何给定传递中的工作范围限制为较小的一组增量更改。
3.1.3。了解扩展
通过在CUDA上运行，应用程序将获得的性能收益的数量完全取决于它可以并行化的程度。无法充分并行化的代码应在主机上运行，除非这样做会导致主机与设备之间的传输过多。
注意：高优先级：要从CUDA中获得最大收益，请首先关注寻找并行化顺序代码的方法。
通过了解应用程序如何扩展，可以设定期望并计划增量并行化策略。强扩展和阿姆达尔定律描述了强扩展，这使我们可以为问题固定大小的加速设置上限。弱缩放和古斯塔夫森定律 描述了弱缩放，即通过增大问题大小来实现加速。在许多应用中，需要强缩放和弱缩放的组合。
3.1.3.1。强尺度和阿姆达尔定律
强扩展是衡量在固定的总体问题大小下，随着将更多的处理器添加到系统中，解决时间缩短的方法。表现出线性强扩展的应用程序的加速比等于所使用的处理器数量。
强缩放通常等同于阿姆达尔定律，该定律指定了通过并行化部分串行程序可以预期的最大加速。本质上，它指出程序的最大加速比 S为：
S = 1 ( 1 − P ) + P N
在此，P是可并行化的那部分代码所花费的总串行执行时间的一部分，N是可运行代码的并行部分的处理器数。
N越大（即处理器数量越多），则P / N分数越小。将N视为非常大的数字会更简单，这实际上将等式转换为 S = 1 / ( 1 − P ) 。现在，如果并行执行顺序程序的运行时间的3/4，则串行代码的最大加速比为1 /（1-3/4）= 4。
实际上，即使大多数应用程序确实表现出一定程度的强缩放，也不会表现出完美的线性强缩放。对于大多数目的，关键是可并行化部分P越大，潜在加速越大。相反，如果 P为小数（意味着该应用程序基本上不可并行化），则增加处理器数 N几乎没有改善性能。因此，为了在固定的问题大小下获得最大的加速，值得花更多的精力来增加P，从而最大化可并行化的代码量。
3.1.3.2。弱缩放和古斯塔夫森定律
弱扩展是一种衡量方法的解决方案，该方法用于解决方案的时间如何随着将更多处理器添加到系统中而每个处理器具有固定的问题大小而改变 ；也就是说，随着处理器数量的增加，总体问题的规模也随之增加。
弱缩放通常等同于古斯塔夫森定律，该定律指出，实际上，问题的大小与处理器数量成正比。因此，程序的最大加速比S为：
S = N + ( 1 − P ) ( 1 − N )
在此，P是可并行化的那部分代码所花费的总串行执行时间的一部分，N是可运行代码的并行部分的处理器数。
查看古斯塔夫森定律的另一种方法是，随着我们扩大系统规模，问题的大小不会保持恒定，而是执行时间。请注意，古斯塔夫森定律假设串行执行与并行执行的比率保持恒定，这反映了设置和处理较大问题的额外成本。
3.1.3.3。应用强弱缩放
了解哪种扩展类型最适合应用程序是估计加速的重要部分。对于某些应用，问题的大小将保持不变，因此仅适用于强缩放。一个例子是模拟两个分子如何相互作用，其中分子大小是固定的。
对于其他应用程序，问题的大小将增大，以填补可用的处理器。示例包括将流体或结构建模为网格或网格，以及一些蒙特卡洛模拟，其中增加问题的大小可提高准确性。
了解了应用程序配置文件后，开发人员应该了解如果计算性能发生变化，问题大小将如何变化，然后应用阿姆达尔定律或古斯塔夫森定律确定加速的上限。
许多代码用相对较少的代码即可完成大部分工作。使用 profiler,，开发人员可以识别此类热点并开始编译候选列表以进行并行化。
3.1.1。创建配置文件
有很多可能的方法来分析代码，但是在所有情况下，目标都是相同的：识别应用程序将花费大部分执行时间的一个或多个函数。
注意：高优先级：为了最大程度地提高开发人员的工作效率，请对应用程序进行概要分析，以确定热点和瓶颈。
任何概要分析活动的最重要考虑因素是确保工作负荷是现实的-即，从测试和基于该信息的决策中获得的信息与真实数据相关。使用不切实际的工作负载会导致开发人员针对不切实际的问题大小进行优化，或者使开发人员将精力集中在错误的功能上，从而导致结果欠佳和工作浪费。
有许多工具可用于生成配置文件。以下示例基于gprof，这是来自GNU Binutils集合的Linux平台的开源探查器。
程序代码: [选择]
$ gcc -O2 -g -pg myprog.c
$ gprof ./a.out > profile.txt
Each sample counts as 0.01 seconds.
% cumulative self self total
time seconds seconds calls ms/call ms/call name
33.34 0.02 0.02 7208 0.00 0.00 genTimeStep
16.67 0.03 0.01 240 0.04 0.12 calcStats
16.67 0.04 0.01 8 1.25 1.25 calcSummaryData
16.67 0.05 0.01 7 1.43 1.43 write
16.67 0.06 0.01 mcount
0.00 0.06 0.00 236 0.00 0.00 tzset
0.00 0.06 0.00 192 0.00 0.00 tolower
0.00 0.06 0.00 47 0.00 0.00 strlen
0.00 0.06 0.00 45 0.00 0.00 strchr
0.00 0.06 0.00 1 0.00 50.00 main
0.00 0.06 0.00 1 0.00 0.00 memcpy
0.00 0.06 0.00 1 0.00 10.11 print
0.00 0.06 0.00 1 0.00 0.00 profil
0.00 0.06 0.00 1 0.00 50.00 report3.1.2。识别热点
在上面的示例中，我们可以清楚地看到该函数 genTimeStep（）占用应用程序总运行时间的三分之一。这应该是我们第一个并行化的候选函数。《了解扩展》讨论了我们从这种并行化中可能期望获得的潜在收益。
值得注意的是，以上示例中的其他几个功能也占据了整个运行时间的很大一部分，例如 calcStats（） 和 calcSummaryData（）。将这些功能并行化也应该增加我们的加速潜力。但是，由于APOD是一个循环过程，我们可能会选择在随后的APOD传递中并行处理这些功能，从而将我们在任何给定传递中的工作范围限制为较小的一组增量更改。
3.1.3。了解扩展
通过在CUDA上运行，应用程序将获得的性能收益的数量完全取决于它可以并行化的程度。无法充分并行化的代码应在主机上运行，除非这样做会导致主机与设备之间的传输过多。
注意：高优先级：要从CUDA中获得最大收益，请首先关注寻找并行化顺序代码的方法。
通过了解应用程序如何扩展，可以设定期望并计划增量并行化策略。强扩展和阿姆达尔定律描述了强扩展，这使我们可以为问题固定大小的加速设置上限。弱缩放和古斯塔夫森定律 描述了弱缩放，即通过增大问题大小来实现加速。在许多应用中，需要强缩放和弱缩放的组合。
3.1.3.1。强尺度和阿姆达尔定律
强扩展是衡量在固定的总体问题大小下，随着将更多的处理器添加到系统中，解决时间缩短的方法。表现出线性强扩展的应用程序的加速比等于所使用的处理器数量。
强缩放通常等同于阿姆达尔定律，该定律指定了通过并行化部分串行程序可以预期的最大加速。本质上，它指出程序的最大加速比 S为：
S = 1 ( 1 − P ) + P N
在此，P是可并行化的那部分代码所花费的总串行执行时间的一部分，N是可运行代码的并行部分的处理器数。
N越大（即处理器数量越多），则P / N分数越小。将N视为非常大的数字会更简单，这实际上将等式转换为 S = 1 / ( 1 − P ) 。现在，如果并行执行顺序程序的运行时间的3/4，则串行代码的最大加速比为1 /（1-3/4）= 4。
实际上，即使大多数应用程序确实表现出一定程度的强缩放，也不会表现出完美的线性强缩放。对于大多数目的，关键是可并行化部分P越大，潜在加速越大。相反，如果 P为小数（意味着该应用程序基本上不可并行化），则增加处理器数 N几乎没有改善性能。因此，为了在固定的问题大小下获得最大的加速，值得花更多的精力来增加P，从而最大化可并行化的代码量。
3.1.3.2。弱缩放和古斯塔夫森定律
弱扩展是一种衡量方法的解决方案，该方法用于解决方案的时间如何随着将更多处理器添加到系统中而每个处理器具有固定的问题大小而改变 ；也就是说，随着处理器数量的增加，总体问题的规模也随之增加。
弱缩放通常等同于古斯塔夫森定律，该定律指出，实际上，问题的大小与处理器数量成正比。因此，程序的最大加速比S为：
S = N + ( 1 − P ) ( 1 − N )
在此，P是可并行化的那部分代码所花费的总串行执行时间的一部分，N是可运行代码的并行部分的处理器数。
查看古斯塔夫森定律的另一种方法是，随着我们扩大系统规模，问题的大小不会保持恒定，而是执行时间。请注意，古斯塔夫森定律假设串行执行与并行执行的比率保持恒定，这反映了设置和处理较大问题的额外成本。
3.1.3.3。应用强弱缩放
了解哪种扩展类型最适合应用程序是估计加速的重要部分。对于某些应用，问题的大小将保持不变，因此仅适用于强缩放。一个例子是模拟两个分子如何相互作用，其中分子大小是固定的。
对于其他应用程序，问题的大小将增大，以填补可用的处理器。示例包括将流体或结构建模为网格或网格，以及一些蒙特卡洛模拟，其中增加问题的大小可提高准确性。
了解了应用程序配置文件后，开发人员应该了解如果计算性能发生变化，问题大小将如何变化，然后应用阿姆达尔定律或古斯塔夫森定律确定加速的上限。
27
CUDA / DAY2：评估你的应用和异构计算
« 于: 十二月 14, 2020, 03:11:14 pm »
1.评估您的应用
从超级计算机到移动电话，现代处理器越来越依赖于并行性来提供性能。核心计算单元（包括控制，算术，寄存器和通常的一些缓存）被复制了若干次，并通过网络连接到内存。结果，所有现代处理器都需要并行代码，以实现其计算能力的良好利用。
在处理器不断发展以向程序员提供更细粒度的并行性的同时，许多现有应用程序已演变为串行代码或粗粒度并行代码（例如，数据分解为并行处理的区域，包括子区域）使用MPI共享）。为了从包括GPU在内的任何现代处理器架构中获利，第一步是评估应用程序以识别热点，确定它们是否可以并行化，并了解现在和将来的相关工作负载。
2.异构计算
CUDA编程涉及在两个不同平台上同时运行代码：具有一个或多个CPU和一个或多个启用CUDA的NVIDIA GPU设备的主机系统。
尽管NVIDIA GPU通常与图形相关联，但它们还是功能强大的算术引擎，能够并行运行数千个轻量级线程。此功能使其非常适合可以利用并行执行的计算。
但是，该设备基于与主机系统截然不同的设计，因此重要的是要了解这些差异以及它们如何确定CUDA应用程序的性能，以便有效地使用CUDA。
2.1。主机和设备之间的差异
主要区别在于线程模型和单独的物理内存：
线程资源
主机系统上的执行管道可以支持有限数量的并发线程。例如，具有两个32个核心处理器的服务器只能同时运行64个线程（如果CPU支持同时多线程，则只能运行64个线程）。相比之下， CUDA设备上最小的并行并行可执行单元包括32个线程（称为线程扭曲）。现代的NVIDIA GPU可支持高达每多并发2048个有效的线程（见特性和规格的CUDA C ++编程指南）在具有80个多处理器的GPU上，这导致超过160,000个并发活动线程。
线程数
CPU上的线程通常是重量级的实体。操作系统必须在CPU执行通道上和下交换线程，以提供多线程功能。因此，上下文切换（交换两个线程时）很慢且很昂贵。相比之下，GPU上的线程非常轻巧。在典型的系统中，成千上万的线程排队等待工作（每个线程32个线程束）。如果GPU必须在一个线程扭曲上等待，它只是开始在另一个线程扭曲上执行工作。由于将单独的寄存器分配给所有活动线程，在GPU线程之间切换时，无需交换寄存器或其他状态。资源会一直分配给每个线程，直到完成执行为止。简而言之，CPU内核旨在 每个线程一次最小化少量线程的等待时间，而GPU旨在处理大量并发的轻量级线程，以最大化吞吐量。
内存
主机系统和设备各自具有各自不同的附加物理存储器1。由于主机和设备内存是分开的，因此，主机内存中的项目有时必须在设备内存和主机内存之间进行通信，如“在 启用CUDA的设备上运行什么？”中所述。。
这些是CPU主机和GPU设备在并行编程方面的主要硬件差异。本文中其他地方也会讨论其他差异。考虑到这些差异的应用程序可以将主机和设备视为一个内聚的异构系统，其中，每个处理单元都可以用来完成其最擅长的工作：在主机上顺序工作，在设备上并行工作。
2.2。在启用CUDA的设备上运行什么？
确定在设备上运行应用程序的哪些部分时，应考虑以下问题：
该设备非常适合可以同时在多个数据元素上并行运行的计算。这通常涉及对大型数据集（例如矩阵）进行算术运算，其中可以对数千个（即使不是数百万个）元素同时执行相同的操作。这是CUDA上良好性能的要求：该软件必须使用大量（通常成千上万个）并发线程。对并行运行多个线程的支持源自CUDA对上述轻量级线程模型的使用。
要使用CUDA，必须将数据值从主机传输到设备。这些传输的性能成本很高，应将其最小化。（请参阅主机与设备之间的数据传输。）此成本有以下几个方面：
操作的复杂性应证明在设备之间来回移动数据的成本是合理的。传输数据以供少量线程短暂使用的代码几乎看不到性能收益，甚至没有收益。理想的情况是许多线程执行大量工作。
例如，将两个矩阵传输到设备以执行矩阵加法，然后将结果传输回主机将不会带来很大的性能优势。这里的问题是每个传输的数据元素执行的操作数。对于前面的过程，假设矩阵的大小为NxN，则有N 2个运算（加法）和3N 2元素传输，因此操作与元素传输的比率为1：3或O（1）。当该比率较高时，可以更轻松地实现性能优势。例如，相同矩阵的矩阵乘法需要N 3运算（乘加），因此运算与传输的元素之比为O（N），在这种情况下，矩阵越大，性能收益就越大。操作的类型是一个附加因素，因为相加的复杂性与例如三角函数不同。在确定是在主机上还是在设备上执行操作时，必须包括与设备之间进行数据传输的开销。
数据应尽可能长时间保存在设备上。由于应尽量减少传输，因此对同一数据运行多个内核的程序应优先在内核调用之间将数据保留在设备上，而不是将中间结果传输给主机，然后将其返回给设备进行后续计算。因此，在前面的示例中，由于之前的一些计算，要添加的两个矩阵已经在设备上了，或者，如果将相加的结果用于后续的计算中，则应在设备上本地执行矩阵相加。即使可以在主机上更快地执行一系列计算中的一个步骤，也应使用此方法。如果它避免主机和设备内存之间的一次或多次传输，那么即使是相对较慢的内核也可能是有利的。即使可以在主机上更快地执行一系列计算中的一个步骤，也应使用此方法。如果它避免主机和设备内存之间的一次或多次传输，那么即使是相对较慢的内核也可能是有利的。即使可以在主机上更快地执行一系列计算中的一个步骤，也应使用此方法。如果它避免主机和设备内存之间的一次或多次传输，那么即使是相对较慢的内核也可能是有利的。 主机与设备之间的数据传输 提供了更多详细信息，包括主机与设备之间以及设备内部的带宽测量。
为了获得最佳性能，设备上运行的相邻线程在内存访问中应保持一定的连贯性。某些内存访问模式使硬件能够将多个数据项的读取或写入组合并为一个操作。无法进行布局以实现 合并的数据，或者没有足够的局部性来有效地使用L1或纹理缓存，当在GPU上进行计算时，往往会看到较慢的加速。值得注意的例外是完全随机的内存访问模式。通常，应避免使用它们，因为与峰值功能相比，任何体系结构都以低效率处理这些内存访问模式。但是，与基于缓存的架构（如CPU）相比，延迟隐藏架构（如GPU）倾向于更好地应对完全随机的内存访问模式。
从超级计算机到移动电话，现代处理器越来越依赖于并行性来提供性能。核心计算单元（包括控制，算术，寄存器和通常的一些缓存）被复制了若干次，并通过网络连接到内存。结果，所有现代处理器都需要并行代码，以实现其计算能力的良好利用。
在处理器不断发展以向程序员提供更细粒度的并行性的同时，许多现有应用程序已演变为串行代码或粗粒度并行代码（例如，数据分解为并行处理的区域，包括子区域）使用MPI共享）。为了从包括GPU在内的任何现代处理器架构中获利，第一步是评估应用程序以识别热点，确定它们是否可以并行化，并了解现在和将来的相关工作负载。
2.异构计算
CUDA编程涉及在两个不同平台上同时运行代码：具有一个或多个CPU和一个或多个启用CUDA的NVIDIA GPU设备的主机系统。
尽管NVIDIA GPU通常与图形相关联，但它们还是功能强大的算术引擎，能够并行运行数千个轻量级线程。此功能使其非常适合可以利用并行执行的计算。
但是，该设备基于与主机系统截然不同的设计，因此重要的是要了解这些差异以及它们如何确定CUDA应用程序的性能，以便有效地使用CUDA。
2.1。主机和设备之间的差异
主要区别在于线程模型和单独的物理内存：
线程资源
主机系统上的执行管道可以支持有限数量的并发线程。例如，具有两个32个核心处理器的服务器只能同时运行64个线程（如果CPU支持同时多线程，则只能运行64个线程）。相比之下， CUDA设备上最小的并行并行可执行单元包括32个线程（称为线程扭曲）。现代的NVIDIA GPU可支持高达每多并发2048个有效的线程（见特性和规格的CUDA C ++编程指南）在具有80个多处理器的GPU上，这导致超过160,000个并发活动线程。
线程数
CPU上的线程通常是重量级的实体。操作系统必须在CPU执行通道上和下交换线程，以提供多线程功能。因此，上下文切换（交换两个线程时）很慢且很昂贵。相比之下，GPU上的线程非常轻巧。在典型的系统中，成千上万的线程排队等待工作（每个线程32个线程束）。如果GPU必须在一个线程扭曲上等待，它只是开始在另一个线程扭曲上执行工作。由于将单独的寄存器分配给所有活动线程，在GPU线程之间切换时，无需交换寄存器或其他状态。资源会一直分配给每个线程，直到完成执行为止。简而言之，CPU内核旨在 每个线程一次最小化少量线程的等待时间，而GPU旨在处理大量并发的轻量级线程，以最大化吞吐量。
内存
主机系统和设备各自具有各自不同的附加物理存储器1。由于主机和设备内存是分开的，因此，主机内存中的项目有时必须在设备内存和主机内存之间进行通信，如“在 启用CUDA的设备上运行什么？”中所述。。
这些是CPU主机和GPU设备在并行编程方面的主要硬件差异。本文中其他地方也会讨论其他差异。考虑到这些差异的应用程序可以将主机和设备视为一个内聚的异构系统，其中，每个处理单元都可以用来完成其最擅长的工作：在主机上顺序工作，在设备上并行工作。
2.2。在启用CUDA的设备上运行什么？
确定在设备上运行应用程序的哪些部分时，应考虑以下问题：
该设备非常适合可以同时在多个数据元素上并行运行的计算。这通常涉及对大型数据集（例如矩阵）进行算术运算，其中可以对数千个（即使不是数百万个）元素同时执行相同的操作。这是CUDA上良好性能的要求：该软件必须使用大量（通常成千上万个）并发线程。对并行运行多个线程的支持源自CUDA对上述轻量级线程模型的使用。
要使用CUDA，必须将数据值从主机传输到设备。这些传输的性能成本很高，应将其最小化。（请参阅主机与设备之间的数据传输。）此成本有以下几个方面：
操作的复杂性应证明在设备之间来回移动数据的成本是合理的。传输数据以供少量线程短暂使用的代码几乎看不到性能收益，甚至没有收益。理想的情况是许多线程执行大量工作。
例如，将两个矩阵传输到设备以执行矩阵加法，然后将结果传输回主机将不会带来很大的性能优势。这里的问题是每个传输的数据元素执行的操作数。对于前面的过程，假设矩阵的大小为NxN，则有N 2个运算（加法）和3N 2元素传输，因此操作与元素传输的比率为1：3或O（1）。当该比率较高时，可以更轻松地实现性能优势。例如，相同矩阵的矩阵乘法需要N 3运算（乘加），因此运算与传输的元素之比为O（N），在这种情况下，矩阵越大，性能收益就越大。操作的类型是一个附加因素，因为相加的复杂性与例如三角函数不同。在确定是在主机上还是在设备上执行操作时，必须包括与设备之间进行数据传输的开销。
数据应尽可能长时间保存在设备上。由于应尽量减少传输，因此对同一数据运行多个内核的程序应优先在内核调用之间将数据保留在设备上，而不是将中间结果传输给主机，然后将其返回给设备进行后续计算。因此，在前面的示例中，由于之前的一些计算，要添加的两个矩阵已经在设备上了，或者，如果将相加的结果用于后续的计算中，则应在设备上本地执行矩阵相加。即使可以在主机上更快地执行一系列计算中的一个步骤，也应使用此方法。如果它避免主机和设备内存之间的一次或多次传输，那么即使是相对较慢的内核也可能是有利的。即使可以在主机上更快地执行一系列计算中的一个步骤，也应使用此方法。如果它避免主机和设备内存之间的一次或多次传输，那么即使是相对较慢的内核也可能是有利的。即使可以在主机上更快地执行一系列计算中的一个步骤，也应使用此方法。如果它避免主机和设备内存之间的一次或多次传输，那么即使是相对较慢的内核也可能是有利的。 主机与设备之间的数据传输 提供了更多详细信息，包括主机与设备之间以及设备内部的带宽测量。
为了获得最佳性能，设备上运行的相邻线程在内存访问中应保持一定的连贯性。某些内存访问模式使硬件能够将多个数据项的读取或写入组合并为一个操作。无法进行布局以实现 合并的数据，或者没有足够的局部性来有效地使用L1或纹理缓存，当在GPU上进行计算时，往往会看到较慢的加速。值得注意的例外是完全随机的内存访问模式。通常，应避免使用它们，因为与峰值功能相比，任何体系结构都以低效率处理这些内存访问模式。但是，与基于缓存的架构（如CPU）相比，延迟隐藏架构（如GPU）倾向于更好地应对完全随机的内存访问模式。
28
这是什么文件？
这种最佳实践指南是一个手册，以帮助开发人员获得NVIDIA的最佳性能® CUDA ®图形处理器。它介绍了已建立的并行化和优化技术，并说明了可以大大简化具有CUDA功能的GPU架构编程的编码隐喻和惯用法。
尽管这些内容可以用作参考手册，但您应该了解，在探索各种编程和配置主题时，会在不同的上下文中重新讨论一些主题。因此，建议初学者阅读该指南。这种方法将极大地增进您对有效编程实践的理解，并使您能够更好地使用该指南以供以后参考。
谁应该阅读本指南？
本指南中的讨论都使用C ++编程语言，因此您应该很容易阅读C ++代码。
本指南参考并依赖于其他一些可供您参考的文档，可从CUDA网站https://docs.nvidia.com/cn/cuda/免费获得所有这些文档。以下文档是特别重要的资源：
CUDA安装指南
CUDA C ++编程指南
CUDA工具包参考手册
特别是，本指南的“优化”部分假设您已经成功下载并安装了CUDA Toolkit（如果没有，请参考所用平台的相关CUDA安装指南），并且您对CUDA C ++编程语言有基本的了解。和环境（如果没有，请参阅《CUDA C ++编程指南》）。
评估，并行化，优化，部署
本指南介绍了针对应用程序的评估，并行化，优化，部署（APOD）设计周期，旨在帮助应用程序开发人员快速识别其代码中最容易从GPU加速中受益的部分，迅速实现这一优势，并开始利用从而尽早加快生产速度。
APOD是一个循环过程：只需很少的初始投资就可以实现，测试和部署初始加速，此时，可以通过确定进一步的优化机会，查看更多的加速，然后部署甚至更快的版本来再次开始该周期。将应用程序投入生产。
评估
对于现有项目，第一步是评估应用程序，以找到负责大部分执行时间的代码部分。有了这些知识，开发人员就可以评估这些瓶颈以进行并行化，并开始研究GPU加速。
通过了解最终用户的要求和约束并应用阿姆达尔（Amdahl）和古斯塔夫森（Gustafson）的定律，开发人员可以通过加速应用程序已识别部分来确定性能改进的上限。
并行化
在确定了热点并完成了设置目标和期望的基本练习之后，开发人员需要并行处理代码。根据原始代码，这可以像调用现有的GPU优化库一样简单，例如立方玻璃， 傅立叶变换， 要么 推力，也可以像添加一些预处理器指令一样简单，作为对并行化编译器的提示。
另一方面，某些应用程序的设计将需要进行一些重构以暴露其固有的并行性。由于甚至CPU体系结构都需要公开并行性以改善或简单地维护顺序应用程序的性能，因此CUDA并行编程语言家族（CUDA C ++，CUDA Fortran等）旨在使这种并行性的表达尽可能地简单。 ，同时在支持CUDA的GPU上运行，该GPU专为实现最大并行吞吐量而设计。
优化
在完成每一轮应用程序并行化之后，开发人员可以着手优化实现以提高性能。由于可以考虑许多可能的优化，因此充分了解应用程序的需求可以帮助使过程尽可能地流畅。但是，与整个APOD一样，程序优化是一个反复的过程（确定优化机会，应用和测试优化，验证所实现的加速并重复），这意味着程序员在看到良好的加速效果之前不必花费大量的时间来记住所有可能的优化策略。相反，可以在学习策略时逐步应用它们。
从重叠的数据传输一直到计算一直到微调浮点操作序列，优化可应用于各种级别。可用的概要分析工具对于指导该过程非常有用，因为它们可以帮助建议开发人员进行优化工作的最佳方案，并为该指南的优化部分的相关部分提供参考。
部署
完成应用程序一个或多个组件的GPU加速后，可以将结果与原始预期进行比较。回想一下，初始评估步骤允许开发人员确定通过加速给定热点可获得的潜在提速上限。
在解决其他热点以提高总体速度之前，开发人员应考虑采用部分并行的实现，并将其付诸实践。这很重要，原因有很多；例如，它允许用户尽早从他们的投资中获利（提速可能是部分的，但仍然是有价值的），并且通过为开发者提供渐进而不是[行为6]性的变更，将开发人员和用户的风险降至最低。应用。
建议和最佳做法
在本指南中，针对CUDA C ++代码的设计和实现提出了具体的建议。这些建议按优先级进行分类，这是建议的效果及其范围的综合。对于大多数CUDA应用程序而言，具有实质性改进的操作具有最高的优先级，而仅影响非常特殊情况的小型优化则具有较低的优先级。
在实施低优先级建议之前，优良作法是确保已应用所有相关的高优先级建议。这种方法将为所投入的时间提供最佳结果，并避免过早优化的陷阱。
受益标准和确定优先级的范围将根据计划的性质而有所不同。在本指南中，它们代表典型情况。您的代码可能反映了不同的优先级因素。不管这种可能性如何，优良作法都是在进行低优先级项目之前验证没有忽略高优先级建议。
注意：整个指南中的代码示例都省略了错误检查，以确保简洁。但是，生产代码应该系统地检查每个API调用返回的错误代码，并通过调用检查内核启动中的失败cudaGetLastError（）。
这种最佳实践指南是一个手册，以帮助开发人员获得NVIDIA的最佳性能® CUDA ®图形处理器。它介绍了已建立的并行化和优化技术，并说明了可以大大简化具有CUDA功能的GPU架构编程的编码隐喻和惯用法。
尽管这些内容可以用作参考手册，但您应该了解，在探索各种编程和配置主题时，会在不同的上下文中重新讨论一些主题。因此，建议初学者阅读该指南。这种方法将极大地增进您对有效编程实践的理解，并使您能够更好地使用该指南以供以后参考。
谁应该阅读本指南？
本指南中的讨论都使用C ++编程语言，因此您应该很容易阅读C ++代码。
本指南参考并依赖于其他一些可供您参考的文档，可从CUDA网站https://docs.nvidia.com/cn/cuda/免费获得所有这些文档。以下文档是特别重要的资源：
CUDA安装指南
CUDA C ++编程指南
CUDA工具包参考手册
特别是，本指南的“优化”部分假设您已经成功下载并安装了CUDA Toolkit（如果没有，请参考所用平台的相关CUDA安装指南），并且您对CUDA C ++编程语言有基本的了解。和环境（如果没有，请参阅《CUDA C ++编程指南》）。
评估，并行化，优化，部署
本指南介绍了针对应用程序的评估，并行化，优化，部署（APOD）设计周期，旨在帮助应用程序开发人员快速识别其代码中最容易从GPU加速中受益的部分，迅速实现这一优势，并开始利用从而尽早加快生产速度。
APOD是一个循环过程：只需很少的初始投资就可以实现，测试和部署初始加速，此时，可以通过确定进一步的优化机会，查看更多的加速，然后部署甚至更快的版本来再次开始该周期。将应用程序投入生产。
评估
对于现有项目，第一步是评估应用程序，以找到负责大部分执行时间的代码部分。有了这些知识，开发人员就可以评估这些瓶颈以进行并行化，并开始研究GPU加速。
通过了解最终用户的要求和约束并应用阿姆达尔（Amdahl）和古斯塔夫森（Gustafson）的定律，开发人员可以通过加速应用程序已识别部分来确定性能改进的上限。
并行化
在确定了热点并完成了设置目标和期望的基本练习之后，开发人员需要并行处理代码。根据原始代码，这可以像调用现有的GPU优化库一样简单，例如立方玻璃， 傅立叶变换， 要么 推力，也可以像添加一些预处理器指令一样简单，作为对并行化编译器的提示。
另一方面，某些应用程序的设计将需要进行一些重构以暴露其固有的并行性。由于甚至CPU体系结构都需要公开并行性以改善或简单地维护顺序应用程序的性能，因此CUDA并行编程语言家族（CUDA C ++，CUDA Fortran等）旨在使这种并行性的表达尽可能地简单。 ，同时在支持CUDA的GPU上运行，该GPU专为实现最大并行吞吐量而设计。
优化
在完成每一轮应用程序并行化之后，开发人员可以着手优化实现以提高性能。由于可以考虑许多可能的优化，因此充分了解应用程序的需求可以帮助使过程尽可能地流畅。但是，与整个APOD一样，程序优化是一个反复的过程（确定优化机会，应用和测试优化，验证所实现的加速并重复），这意味着程序员在看到良好的加速效果之前不必花费大量的时间来记住所有可能的优化策略。相反，可以在学习策略时逐步应用它们。
从重叠的数据传输一直到计算一直到微调浮点操作序列，优化可应用于各种级别。可用的概要分析工具对于指导该过程非常有用，因为它们可以帮助建议开发人员进行优化工作的最佳方案，并为该指南的优化部分的相关部分提供参考。
部署
完成应用程序一个或多个组件的GPU加速后，可以将结果与原始预期进行比较。回想一下，初始评估步骤允许开发人员确定通过加速给定热点可获得的潜在提速上限。
在解决其他热点以提高总体速度之前，开发人员应考虑采用部分并行的实现，并将其付诸实践。这很重要，原因有很多；例如，它允许用户尽早从他们的投资中获利（提速可能是部分的，但仍然是有价值的），并且通过为开发者提供渐进而不是[行为6]性的变更，将开发人员和用户的风险降至最低。应用。
建议和最佳做法
在本指南中，针对CUDA C ++代码的设计和实现提出了具体的建议。这些建议按优先级进行分类，这是建议的效果及其范围的综合。对于大多数CUDA应用程序而言，具有实质性改进的操作具有最高的优先级，而仅影响非常特殊情况的小型优化则具有较低的优先级。
在实施低优先级建议之前，优良作法是确保已应用所有相关的高优先级建议。这种方法将为所投入的时间提供最佳结果，并避免过早优化的陷阱。
受益标准和确定优先级的范围将根据计划的性质而有所不同。在本指南中，它们代表典型情况。您的代码可能反映了不同的优先级因素。不管这种可能性如何，优良作法都是在进行低优先级项目之前验证没有忽略高优先级建议。
注意：整个指南中的代码示例都省略了错误检查，以确保简洁。但是，生产代码应该系统地检查每个API调用返回的错误代码，并通过调用检查内核启动中的失败cudaGetLastError（）。
29
CUDA / Cooperative Group 特殊启动核函数时的参数问题
« 于: 十一月 09, 2020, 02:45:55 pm »
我在使用Cooperative Group 特殊启动时，参数好像没有传进核函数。我的用法如下：
void *kernelArgs[] = {Type, Dist, Temp};
cudaLaunchCooperativeKernel((void*)Evolve, dimGrid, dimBlock, kernelArgs);
其中Type为char *类型，Dist和Temp都为double *类型，申请内存的方式为：cudaMallocManaged。
void *kernelArgs[] = {Type, Dist, Temp};
cudaLaunchCooperativeKernel((void*)Evolve, dimGrid, dimBlock, kernelArgs);
其中Type为char *类型，Dist和Temp都为double *类型，申请内存的方式为：cudaMallocManaged。
30
CUDA / Re: CUDA主程序中有许多核函数执行，需要在每个核函数执行完毕都使用同步指令吗？
« 于: 十月 27, 2020, 11:19:25 pm »
答：没有必要每次都让CPU对GPU上的kernel同步的.
除非需要结果, 你才需要进行一次等异步传输的同步, 以等待结果取回.
(或者干脆使用同步传输)
CPU不需要结果的话(例如后一个kernel直接使用前一个kernel的结果), 你完全可以连续发布kernel的.
除非需要结果, 你才需要进行一次等异步传输的同步, 以等待结果取回.
(或者干脆使用同步传输)
CPU不需要结果的话(例如后一个kernel直接使用前一个kernel的结果), 你完全可以连续发布kernel的.
