如何确定内存密集型 CUDA 应用程序的最大可能线程和块？答案

【问题标题】：How to determine the maximum possible Threads and Blocks for a memory-heavy CUDA application?如何确定内存密集型 CUDA 应用程序的最大可能线程和块？
【发布时间】：2014-12-09 08:36:38
【问题描述】：

我正在尝试为我的应用程序找到线程和块的最佳值。因此我写了一个小套装来运行线程数、块大小和网格大小的可能组合。我正在处理的任务是不可并行化的，因此每个线程都在计算其独特的问题，并且需要对其独特的全局内存块进行读写访问。我还必须增加 cudaLimitStackSize 才能运行内核。当我尝试计算一次可以运行的最大线程数时遇到了问题。我的改进方法（感谢 Robert Crovella）是

threads = (freememory*0.9)/memoryperthread

其中freememory 是从cudaMemGetInfo 获得的，memoryperthread 是一个线程的全局内存要求。即使我减少常数因子，我仍然会遇到“未指定的启动失败”，我无法调试，因为调试器失败并显示Error: Internal error reported by CUDA debugger API (error=1). The application cannot be further debugged.。根据设置，此错误当我尝试不同的块大小时，我也遇到了问题。任何大于 512 个线程的块大小都会产生“请求启动的资源过多”。正如 Robert Crovella 所指出的，这可能是我的内核占用了许多寄存器（-Xptxas="-v" 报告的 63 个）的问题。由于块可以分布在多个 multiProcessorCount 中，因此我找不到任何会突然遇到块大小为 1024 的限制。

我的代码对于线程和块的小值运行良好，但我似乎无法计算我可以同时运行的最大数量。有什么方法可以正确计算这些，还是我需要根据经验进行计算？

我知道内存繁重的任务对于 CUDA 来说并不是最佳选择。我的设备是具有 Compute Capability 2.0 的 GTX480。现在我坚持使用CUDA Driver Version = 6.5, CUDA Runtime Version = 5.0。我确实使用-gencode arch=compute_20,code=sm_20 进行编译以实现计算能力。

更新：将运行时更新到 6.5 后，上述大部分问题都消失了。我将按原样保留这篇文章，因为我提到了我遇到的错误，人们在搜索他们的错误时可能会偶然发现它。为了解决大块大小的问题，我不得不减少每个线程的寄存器（-maxrregcount）。

【问题讨论】：

您知道 CUDA 6.5 附带的 cudaOccupancyMaxPotentialBlockSize() 函数吗？附带说明一下，您是否注意到您的驱动程序和运行时版本不匹配？
感谢您的评论。是的，我已经看到了，但我不确定它是否适用于我的情况，因为它没有全局内存使用选项。我知道不匹配，但我不是我正在使用的系统的主要管理员。
我不确定你的问题。我的理解是您必须处理大量数据并且您遇到的问题是您没有足够的线程来实现线程和数据之间的一对一对应。当然，您的整体thread 数量不能超过2^31-1。如果totalmemory/memoryperthread > 2^31-1 则必须增加每个线程处理的内存量。此外，您说“从此计算块大小”，但没有说如何。
通常，每个块的线程数是一个独立变量，您可以随意修复，对于 GTX480，提供的线程数少于 1024。通常会调整此变量以优化您的代码的性能。
感谢您的意见，我会相应地更新我的问题。

标签： c linux cuda

【解决方案1】：

线程 = 总内存/每线程内存

如果您对memoryperthread 的计算是准确的，这将不起作用，因为totalmemory 通常并非全部可用。由于 CUDA 运行时开销、分配粒度和其他因素，您实际可以分配的数量少于此数量。所以这会以某种方式失败，但由于你没有提供代码，所以不可能确切地说出如何。如果您从主机进行所有这些分配，例如通过cudaMalloc，那么我预计会出现错误，而不是内核未指定的启动失败。但是，如果您在内核中执行malloc 或new，那么您可能正在尝试使用返回的空指针（表示分配失败 - 即内存不足），这可能会导致未指定的启动失败。

块大小大于 512 个线程会导致“请求启动的资源过多”。

这可能是因为您没有为 cc2.0 设备进行编译，或者您的内核在每个线程中使用的寄存器比支持的多。无论如何，这肯定是一个可以解决的问题。

那么如何正确计算内核的最大可能线程和块？

通常，全局内存需求是问题的函数，而不是内核大小。如果您的全局内存需求随着内核大小而增加，那么可能有一些比率可以根据cudaMemGetInfo 报告的“可用内存”来确定（例如 90%），这应该可以提供相当安全的操作。但总的来说，如果一个程序能够容忍分配失败，那么它就是设计良好的，并且您至少应该在主机代码和设备代码上明确检查这些，而不是依靠“未指定的启动失败”来告诉您某些事情发生了错误的。这可能是由内存使用触发的任何类型的副作用错误，并且可能不是直接由于分配失败。

我建议追踪这些问题。调试问题，找到问题的根源。我认为正确的解决方案将会出现。

【讨论】：

嗨。感谢您的输入。我不知道寄存器的限制以及总内存和可用内存之间的差异。我用-gencode arch=compute_20,code=sm_20 编译。我知道我使用了很多寄存器，但除了blocksize < multiProcessorCount 之外，我看不到其他方法可以自己腾出空间。我已经在检查我在代码中执行的所有 cuda 操作时出错，但这似乎没有帮助。 cudaMalloc 永远不会发生该错误。我尝试从 freememory 计算最大总线程数。但它仍然失败。
未指定的启动失败类似于主机代码中的段错误。您的内核正在做一些非法的事情。你可以调试它。试试cuda-memcheck，或者使用 cuda 调试器，或者使用 printf，或者分而治之，或者你有什么。在 SO 问题中很难给出关于调试 cuda 代码的完整教程。对于每个线程的寄存器问题，您可以限制编译器对寄存器的使用。尝试谷歌搜索“每个线程的 cuda 寄存器”。当我这样做时，第一个打击可能对你有启发性。在 SO 上，如果您需要有关非工作代码的帮助，那么您应该显示 MCVE。
我的代码大部分是从经过彻底测试的 C 代码中复制而来的，并且在线程数和块数较低的情况下运行良好。进一步限制寄存器并没有改变任何东西。调试器仍然中断，并且 cuda-memcheck 在包含有效代码的位置抱怨 Invalid __global__ read of size 1。从我的角度来看，问题似乎仅在于我启动的线程/块的数量。我不知道是什么限制了这个数量。我会尝试做一个 MCVE。