Cuda 内核启动失败答案

【问题标题】：Cuda kernel launch failureCuda 内核启动失败
【发布时间】：2013-07-15 11:10:11
【问题描述】：

我正在尝试调用两个内核，如下所示

for (t=0; t<=time_total; t++)
{    
      //kernel calls
      kernel1<<<noOfBlocks,noOfThreadsPerBlock>>>(** SOME PARAMETERS **);
      checkCudaError(cudaThreadSynchronize());

      kernel2<<<noOfBlocks,noOfThreadsPerBlock>>>(** SOME PARAMETERS **);
      checkCudaError(cudaThreadSynchronize());

}

而第二个内核的结构是

var[index+0]=**SOME CALCULATION**
var[index+1]=**SOME CALCULATION**
var[index+2]=**SOME CALCULATION**

现在，当我执行此代码时，checkCudaError 不会报告任何内容，并且执行代码会给出一些输出，但 Visual Studio 会给出以下异常

First-chance exception at 0x7640c41f in **.exe: Microsoft C++ exception: cudaError_enum at memory location 0x0039f9c4..
First-chance exception at 0x7640c41f in **.exe: Microsoft C++ exception: cudaError_enum at memory location 0x0039f9c4..

当我检查 Nsight 时，它说内核 2 出现以下错误

CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES

现在的问题是内核 2 中的 var 数组给出了一些正确的行，一些是其他行值的副本，一些是垃圾。

当我这样做时

var[index+0]=3
var[index+1]=3
var[index+2]=3

var 的所有值都设置为 3

【问题讨论】：

最好给出整个代码..这样我们就可以编译、运行和检查问题..
对不起，我不能在这里生成代码。我只是想知道导致异常的原因，如果内核启动失败，这究竟意味着什么？
我已经投票决定关闭它。如果您不在这里为其他人提供必要的信息来帮助您，我看不出您如何期望任何人为您提供答案。
您可以忽略第一次机会异常。 “第一次机会异常”表示在 CUDA 运行时库内部生成的异常。异常在 CUDA 库中被捕获和处理，所以没有什么可担心的，这只是正常的行为（它已在 CUDA 5.5 中修复）。这与您的启动失败无关。
cudaThreadSynchronize() 已弃用，取而代之的是 cudaDeviceSynchronize()。 nsight 在第二次内核启动时报告错误，但您的错误检查代码没有，这让我相信您的错误检查代码已损坏。资源不足通常是由于代码请求的寄存器过多（每个线程的寄存器数乘以请求的线程数）。尝试重新编译指定 -Xptxas -v 的代码以获得详细输出，然后使用 -maxrregcount 20 重新编译（或类似的东西）尝试解决此问题以进行测试。

标签： visual-studio-2010 cuda nsight

【解决方案1】：

一些旁注：

cudaThreadSynchronize() 被 cudaDeviceSynchronize() 弃用。
nsight 在第二次内核启动时报告错误，但您的错误检查代码没有，这让我相信您的错误检查代码已损坏。

现在，关于您的问题，资源不足通常是由于代码请求的寄存器过多（每个线程的寄存器数乘以每个线程块请求的线程数。）尝试重新编译您的代码specifying-Xptxas -v获得详细的输出，然后使用-maxrregcount 20（或something like that）再次重新编译以尝试解决此问题以进行测试。

如果这“解决”了您的问题，那么您可能需要考虑以下事项：

看看是否有办法重新排序或重组代码以减少注册压力
如果不是，则将您的maxrregcount 值向上调整到大约允许您的代码根据您关心的启动配置（每个块的线程数）编译和运行的最高值。您可能还希望在此设置的不同级别对代码进行基准测试，因为它会影响占用率。通常，如果您将其设置为将编译和运行的最高值，那么您将在执行时将自己限制为每个 SM 一个线程块。这可能没问题，或者可能有一个更好的较低设置，每个 SM 驻留允许两个线程块，并且可能具有更高的性能。只有对代码进行基准测试才能说明问题。

【讨论】：