具有混合 CUDA 的 Profile C 应用程序

Posted 2023-03-06

【中文标题】具有混合 CUDA 的 Profile C 应用程序

【英文标题】：Profile C application with mixed CUDA

【发布时间】：2014-07-29 05:17:11

【问题描述】：

我有一个 C 程序，它的主要功能占用了总运行时间的 70%。我使用 gprof 来分析应用程序。之后，我在 CUDA 中重写了该特定函数以提高整个应用程序的运行时间。它目前正在正确给出结果，但我想知道性能。

无论如何（或工具）我可以使用新内核的运行时间来分析这个新应用程序，作为运行时间相对于整个新应用程序的百分比？我还想查看与所有其他剩余 C 函数相关的数据。我尝试使用 nvprof，但它只输出 CUDA 内核的运行时。

谢谢，

【问题讨论】：

如果您有一个程序，其功能是执行特定计算，那么该计算应该占用程序中的大部分时间。仅仅因为你优化它并不意味着它会占用更少的运行时间百分比，它实际上可能会占用更多的运行时间百分比，因为它变得更快，因此可以在相同的时间范围。

@JoachimPileborg 我不遵循你的逻辑。如果您减少在应用程序的一部分中花费的时间，t1opt

@Tom 假设您有一个应用程序每秒执行 X 次计算，并且该计算占应用程序运行时间的 50%。现在，如果计算优化为将花费的时间减半，您可以将计算的运行时间百分比减半，使其占应用程序总运行时间的 25%，但是您将如何处理另外 25% 的时间？只是闲着？为什么不让计算仍占运行时间的 50%，而是在同一时间范围内将计算次数加倍？

如果您对 gprof 感到满意，您可以在内核调用周围放置一个精简的 C 包装函数，并使用 gprof 进行分析。您可能希望在包装器末尾添加一个cudaDeviceSynchronize() 调用，具体取决于包装器中的确切代码。

@JoachimPileborg 使用您的示例，如果您将在函数中花费的时间减半，但您没有将百分比减半，您的应用程序不会因为节省的时间而空闲，它会提前完成！您所描述的是不同的（在某些方面类似于弱缩放），因为您试图在固定的时间内增加工作，而最初的问题是关于在减少的时间内完成固定数量的工作时间。

【参考方案1】：

您可以使用 NVIDIA 分析工具为您提供此信息。

运行命令行工具nvprof <app> 会给你百分比，你可以使用额外的命令行选项来进一步优化你的内核。可视化分析器 (nvvp) 将向您显示时间线以及在内核中花费的时间百分比，它还将为您提供有关如何进一步提高性能的指导（包括链接回文档以解释概念）。

请参阅the documentation 了解更多信息。

附录

在您的评论中，您说您还想查看 C 函数的配置文件。一种方法是使用 nvtx 来注释您的代码，请参阅this blog post 以了解自动化该任务的方法。或者，您可以在 nvprof 或 nvvp 中进行分析以查看整体时间线，并在 gprof 中进行分析以查看在非 GPU 代码中花费的时间。

【讨论】：

我使用了 nvprof，但它只给了我 CUDA 内核的运行时，例如，在摘要模式下，如下所示：docs.nvidia.com/cuda/profiler-users-guide/#summary-mode。我也想看看其他 C 函数。谢谢。

【参考方案2】：

嗯，你可能知道我偏爱this technique。

它会告诉您函数、代码行以及您可以识别的任何内容所花费的大致百分比。 我假设您的主程序在某些时候必须等待 CUDA 内核完成处理，因此以该等待结束的样本部分可以让您估计在 CUDA 中花费的时间。 样本不是以该等待结束，而是在做其他事情，表示做这些事情所花费的时间。

统计数据非常简单。如果一行代码或函数在堆栈上的时间为 F 的一部分，那么它负责那部分时间。因此，如果您抽取 N 个样本，则显示代码行或函数的样本数平均为 NF。标准差是 sqrt(NF(1-F))。 因此，如果 F 为 50% 或 0.5，并且您采用 20 个随机堆栈样本，您可以期望在其中 10 个样本上看到代码，标准差为 sqrt(20*0.5*0.5) = 2.24 个样本，或者介于 7 个之间和 13 个样本，最有可能在 9 到 11 之间。 换句话说，您对代码的成本有一个非常粗略的衡量，但您确切地知道什么代码具有该成本。 如果你关心速度，你主要关心成本高的东西，所以它善于指出来给你。

如果你想知道为什么 gprof 没有告诉你这些事情，有一个lot more to say about that。