1

培训信息 / GTC 征文活动——GTC讲座认领清单

« 于: 三月 08, 2023, 11:13:17 am »

 

BrandonH	Edge AI: From Model Development to Deployment [S52125]
晓得薛定谔的喵	From Tortoise to Hare: How AI Can Turn Any Driver into a Race Car Driver [S51328]
魏武卒	Accelerating Generative AI in Biology and Healthcare [S51257]
生生	Advances in Accelerated Computing for Scientific Computing [S52137]
geo-sig	How to Write a CUDA Program [S51210]
The Galaxy Rabbit	Accelerating Transformer-Based Encoder-Decoder Language Models for Sequence-to-Sequence Tasks [S51158]
丙乙	Portable Acceleration of HPC Applications using ISO C++ — Part 1: Fundamentals* [DLIT51169]
白碧哲	How to Build a Real-time Path Tracer [S51871]
Brain Compiler	Watch Party: 如何设计优化 CUDA 程序 [WP51210]
咖啡逗	Accelerate AI Innovation with Unmatched Cloud Scale and Performance (Presented by Microsoft) [S52469]
00101010	Creating and Executing an Effective Cyberdefense Strategy in an AI-Driven Business [S51723]
PixelPassion	Deep Reinforcement Learning with Real-World Data [S51826]
Christy	Watch Party: CUDA 新特性和发展 [WP51225]
16	Accelerated AI Logistics and Route Optimization 101* [DLIT51886]
Ryan	Developing Robust Multi-Task Models for AV Perception [SE50006]
RP Cai	Addressing the AI Skills Gap [SE52129]
蔡欣	Speech-To-Speech Translation System for a Real-World Unwritten Language [S51780]
亚克西	Scaling Deep Learning Training: Fast Inter-GPU Communication with NCCL [S51111]
Ranoe:)	Jetson Edge AI Developer Days: Accelerate Edge AI With NVIDIA Jetson Software [SE52433]
d2real	Connect with the Experts: GPU-Accelerated Quantum Chemistry and Molecular Dynamics [CWES52130]
slam	Connect with the Experts: A Deep-Dive Q&A with Jetson Embedded Platform Engineers [CWES52132]
挠你个痒痒	AI 初创企业在中国市场的发展和机会 ——探索中国 AI 初创力量​ [SE52131]
。	Building Trust in AI for Autonomous Vehicles [S51934]
位位位	Introduction to Autonomous Vehicles [S51168]
Tian_DY	A Comprehensive Low-Code Solution for Large-Scale 3D Medical Image Segmentation with MONAI Core and Auto3DSeg* [DLIT51974]
Permanent Maniac	Watch Party: 基于 TensorRT 的端到端子图优化框架 [WP51416]
王威	3D by AI: Using Generative AI and NeRFs for Building Virtual Worlds [S52163]

2

深度学习讨论 / orin 的unified memory啥情况下要attach global和host吗

« 于: 十二月 03, 2022, 09:14:28 am »

问题：有了解orin 的unified memory啥情况下要attach global和host吗。发现和trt放一个线程后就需要做这个，否则不需要

3

CUDA / 想问一下cudaMemcpyDeviceToDevice，从设备端复制数据到设备端是可以的哈?

« 于: 十一月 10, 2022, 10:40:02 am »

答复：可以的。适合你从显存->显存的复制。
不过这一般用不到，因为你往往不需要直接这样做，在下一个kernel之类的地方，读取/处理后，顺便写入到另外一个地方即可。因为你要考虑啥时候需要D->D的复制，往往是为了保存原缓冲区，防止被破坏，这样新kernel可以直接在另外一个副本上操作了，但是这实际上往往直接修改kernel，写入另外一个地方即可。从而节省了这个流程。
这个和我们CUDA夏令营时候的"矩阵转置", 虽然课堂上讲了，但是实际上从来几乎不用用到一个道理，因为后者最常见的出现往往是被集成到另外一个kernel的内部。而不需要单独的这个过程。此外，D->D的传输，之前的卡（我不确定20/30系列+的情况），会自动生成一个隐形的kernel，通过该隐形的kernel，进行显存里头的复制。（而不是通过硬件的Copy Engine (DMA Engine))，所以这样的话，你自己衡量代价。因为这只是一个纯复制的kernel，而你往往可以写成读取+处理+写入的，后者更加高效。

如同可以再给前任打电话的，但是一般情况下无意义。

4

NV边缘计算 / Jetson AGX Xavier平台的部署工具链有么有提供效率评估工具？

« 于: 十月 25, 2021, 03:27:16 pm »

问：我想向您请教几个关于Jetson AGX Xavier平台的问题，非常感谢！ 1、Jetson AGX Xavier平台的部署工具链有么有提供效率评估工具？例如，我有一套算法网络，可以知道网络的结构、运算量、参数量，要求能在硬件平台上100ms内推理完成。我怎么衡量Jetson AGX Xavier平台的算力能否满足需求？Jetson AGX Xavier平台有提供相应的工具吗？ 2、Jetson AGX Xavier平台的软件栈或者TensorRT有没有运行时调度工具？ TensorRT可以调用API指定网络在哪个处理核上跑，那是否可以运行时自动调度呢（针对多个网络场景）？

5

NV边缘计算 / 你们测评过jetson系列“硬件编码推流-公网传输-解码”这个loop的时延么？

« 于: 九月 22, 2021, 06:40:24 pm »

你们测评过jetson系列“硬件编码推流-公网传输-解码”这个loop的时延么？

6

CUDA / 算力与功耗在GPU应用上的关系是？

« 于: 四月 16, 2021, 04:07:19 pm »

答复：一般来说, 随着制程的进步, 和GPU架构的发展. 往往单位算力下, 所需要的功率, 在随着年份的增加而下降，例如从40nm到7nm的制程, 代表了前者; 而从GPU的SP自身运算FP16, 到引入TensorCore进行计算, 则代表了后者，无论是前者还是后者, 都代表了随着时间和GPU代数的增加, 单位功耗下的算力性能的增加, 即能效在增加，此外, 软件方面的进步, 例如CUDNN的版本增加, 则在除了前面说过的两个硬件的方面, 提供了另外的在单位功耗下, 提升算力的渠道。

7

CUDA / 问个CUDA并行上的小白问题，既然SM只能同时处理一个WARP，那是不是有的SP处于闲置？

« 于: 二月 22, 2021, 02:52:13 pm »

问：问个CUDA并行上的小白问题，既然SM只能同时处理一个WARP，那是不是有的SP处于闲置？GPU的SM每次只能运行一个warp（32个线程），比如1080Ti有28个SM，每个SM里有128个SP，但是resident thread只有28*32 = 896个，那么1080Ti总共128*28=3584个SP（CUDA核心）不就没有同时并行起来。那SM中这么多SP的意义是提供WARP的切换吗？

8

CUDA / CUDA编程中从主机传递一个数组给设备

« 于: 二月 18, 2021, 04:55:16 pm »

CUDA编程中从主机传递一个数组给设备，当数组长度较小时结果正确，当数组长度较大时报非法内存访问。已经排除了GPU设备线程块不够和开辟数组的问题，还能是什么问题?

9

CUDA / 我是做cuda实时信号处理...

« 于: 一月 07, 2021, 09:19:14 pm »

我是做cuda实时信号处理。信号采样率40 mhz，对应40m byte/s，送入GPU，但是控制逻辑是1ms一个更新。所以每用GPU处理1ms的数据，都要传回CPU做控制。回传数据只有几k字节。这样目前很难实时跑起来。您觉得是有什么优化方向，是需要把控制逻辑也移植到cuda下么？

10

CUDA / 请问一个问题。warp shuffle操作中mask参数该怎么理解呢？

« 于: 一月 06, 2021, 08:04:09 pm »

请问一个问题。warp shuffle操作中mask参数该怎么理解呢？我看了doc没有看明白。也没有通过profile信息看出来。可以指导一下吗？谢谢

11

CUDA / 做cuda流同步时，cudastreamquery(0)是不是有特殊含义？

« 于: 十二月 31, 2020, 03:55:55 pm »

答复：在Windows下, 使用WDDM驱动的时候, 具有立刻推送命令到显卡上的效果.等效于不存在的函数cudaStreamFlush().(该函数并不存在. Linux下不需要. TCC下也不需要. 但WDDM下需要同时并不提供，而cudaStreamQuery提供了类似的效果)

12

CUDA / DAY6：优化CUDA应用程序

« 于: 十二月 21, 2020, 01:58:05 pm »

7.优化CUDA应用程序
在完成每一轮应用程序并行化之后，开发人员可以着手优化实现以提高性能。由于可以考虑许多可能的优化，因此充分了解应用程序的需求可以帮助使过程尽可能地流畅。但是，与整个APOD一样，程序优化是一个反复的过程（确定优化机会，应用和测试优化，验证所实现的加速并重复），这意味着程序员不必花费大量资金。在看到良好的加速效果之前，先记住所有可能的优化策略。相反，可以在学习策略时逐步应用它们。

从重叠的数据传输一直到计算一直到微调浮点运算序列，优化可以应用于各种级别。可用的概要分析工具对于指导该过程非常有用，因为它们可以帮助建议开发人员进行优化工作的最佳方案，并为该指南的优化部分的相关部分提供参考。

13

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来完成。

14

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都受益），因此值得这样做付出必要的努力。

15

CUDA / DAY 3:应用程序分析

« 于: 十二月 14, 2020, 09:12:01 pm »

3.1。应用Profiling
许多代码用相对较少的代码即可完成大部分工作。使用 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  report
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。应用强弱缩放
了解哪种扩展类型最适合应用程序是估计加速的重要部分。对于某些应用，问题的大小将保持不变，因此仅适用于强缩放。一个例子是模拟两个分子如何相互作用，其中分子大小是固定的。

对于其他应用程序，问题的大小将增大，以填补可用的处理器。示例包括将流体或结构建模为网格或网格，以及一些蒙特卡洛模拟，其中增加问题的大小可提高准确性。

了解了应用程序配置文件后，开发人员应该了解如果计算性能发生变化，问题大小将如何变化，然后应用阿姆达尔定律或古斯塔夫森定律确定加速的上限。