在内核调用中使用断言答案

【问题标题】：Using assert within kernel invocation在内核调用中使用断言
【发布时间】：2010-01-17 08:03:06
【问题描述】：

在设备模式下的内核调用中使用断言是否有方便的方法？

【问题讨论】：

【解决方案1】：

CUDA 现在有一个原生的断言函数。使用assert(...)。如果它的参数为零，它将停止内核执行并返回一个错误。（或在 CUDA 调试中触发断点。）

确保包含“assert.h”。此外，这需要计算能力 2.x 或更高版本，并且在 MacOS 上不受支持。有关详细信息，请参阅 CUDA C 编程指南，第 B.16 节。

编程指南也包括这个例子：

#include <assert.h>
__global__ void testAssert(void)
{
   int is_one = 1;
   int should_be_one = 0;
   // This will have no effect
   assert(is_one);
   // This will halt kernel execution
   assert(should_be_one);
}
int main(int argc, char* argv[])
{
   testAssert<<<1,1>>>();
   cudaDeviceSynchronize();
   return 0;
}

【讨论】：

【解决方案2】：

#define MYASSERT(condition) \
  if (!(condition)) { return; }

MYASSERT(condition);

如果您需要更高级的东西，您可以使用cuPrintf()，注册开发人员可以从 CUDA 网站获得。

【讨论】：

宏定义的末尾不应该有分号——一般使用宏都会有那个分号。另外，考虑如下实现它，以避免它贪婪地附加到任何可能恰好位于其前面的 else 关键字：if (condition) /* do nothing */; else return
如果你在这之后的任何时候都有一个 __syncthreads()，你应该确保所有线程都达成相同的决定，否则你可能会死锁。此外，您可以在全局内存中设置一个布尔标志（例如，bool success 由主机初始化为 true）来指示事件。多个线程将false 写入标志并不重要，因为它们总是写入相同的值，因此竞争无关紧要。
@Tom 这不是真的。全局内存不是为这种用途而设计的，因此多个线程同时写入同一全局内存位置的结果是意外行为。
@Auron 可以正确地说，如果多个线程写入同一个全局内存位置，则未定义哪个线程获胜。但是，如果所有线程都写入相同的值，那么哪个获胜并不重要，因为最终结果是相同的。就像我说的，这是一场比赛，但比赛无关紧要。
@Auron 链接的答案不完整，我会在那里添加评论。如果 warp 中的多个线程写入同一位置（非原子），则未定义哪个线程执行最终写入，但该位置将被更新。这与任何其他架构上的数据竞争完全相同，获胜者是未定义的，是否所有值都将按顺序写入未定义，但位置将以一个值结束。有关详细信息，请参阅 CUDA 编程指南第 4.1 节。