cuda - 循环展开问题答案

【问题标题】：cuda - loop unrolling issuecuda - 循环展开问题
【发布时间】：2013-03-30 14:18:16
【问题描述】：

我有一个带有#pragma unroll 80 的内核，我使用 NVIDIA GT 285 运行它，计算能力 1.3，使用网格架构：dim3 thread_block( 16, 16 ) 和 dim3 grid( 40 , 30 )，它工作正常。

当我尝试使用 NVIDIA GT 580、计算能力 2.0 和上述网格架构运行它时，它运行良好。

当我将 GT 580 上的网格架构更改为

dim3 thread_block( 32 , 32 ) 和 dim3 grid( 20 , 15 )，因此产生与上述相同数量的线程，我得到不正确的结果。

如果我在 GT 580 中删除 #pragma unroll 80 或用 #pragma unroll 1 替换它，它工作正常。如果我不这样做，那么内核就会崩溃。

有人知道为什么会这样吗？提前谢谢你

编辑：检查两个设备上的内核错误，我得到“无效参数”。当我搜索此错误的原因时，我发现当网格和块的尺寸超出其限制时会发生这种情况。但对我来说情况并非如此，因为我每个块使用 16x16=256 个线程和 40x30=1200 个总块。据我所知，这些值位于 GPU 网格的边界内，计算能力为 1.3。我想知道这是否与我遇到的循环展开问题有关。

【问题讨论】：

内核执行后你检查CUDA错误了吗？
我在主机代码中对内存副本使用错误检查，但没有收到任何错误
你能在启动内核的时候贴出代码并检查错误吗？您应该在启动内核并运行 cudaDeviceSynchronize() 后检查错误。因为我在访问 dev_feature1 时没有看到任何代码来检查边界情况。我怀疑你的代码有内存损坏。
如果它有帮助 dev_feature1 是一个_____device_____ 数组，大小为 2150400 个无符号字符，我使用 MemcpyToSymbol 从主机复制了它的值。您指的是哪种错误检查？我的印象是你不是在谈论 memcopy 错误检查
如果您在内核启动时需要关于无效参数的帮助，您需要显示内核调用行以及所有参数的定义和分配以及传递给内核的启动配置。

标签： cuda nvidia loop-unrolling

【解决方案1】：

我发现了问题所在。

在修复了一些错误后，我收到了“启动请求的资源过多”错误。对于循环展开，每个线程都需要额外的寄存器，并且我的寄存器用完了，因此错误和内核失败。每个线程需要 22 个寄存器，每个块有 1024 个线程。

通过将我的数据插入到 CUDA_Occupancy_calculator 中，它显示每个 SM 安排了 1 个块，在计算能力 2.0 设备上为整个块留下了 32678 个寄存器。

22 个寄存器*1024 个线程 = 22528 个寄存器

我使用 nvcc -arch sm_20 编译，利用了 32678 个寄存器，对于所需的 22528 来说绰绰有余，现在它工作正常。感谢大家，我了解了内核错误。

【讨论】：