7.31. 性能分析计数器函数
每个多处理器都配备了一组16个硬件计数器，应用程序可以通过调用__prof_trigger()函数，用一条指令来递增这些计数器。
void __prof_trigger(int counter);
每个warp将索引为counter的每多处理器硬件计数器递增1。计数器8至15为保留值，应用程序不应使用。
计数器0、1、…、7的值可以通过nvprof工具使用nvprof –events prof_trigger_0x命令获取，其中x取值为0、1、…、7。所有计数器在每次内核启动前都会重置（请注意，在收集计数器数据时，内核启动是同步进行的）。
7.32. 断言
断言功能仅支持计算能力2.x及更高版本的设备。
void assert(int expression);
如果expression等于零，则停止内核执行。如果程序在调试器中运行，这将触发断点，调试器可用于检查设备的当前状态。否则，对于expression等于零的每个线程，在通过cudaDeviceSynchronize()、cudaStreamSynchronize()或cudaEventSynchronize()与主机同步后，会向stderr打印一条消息。该消息的格式如下：
<filename>:<line number>:<function>:
block: [blockId.x,blockId.x,blockIdx.z],
thread: [threadIdx.x,threadIdx.y,threadIdx.z]
Assertion `<expression>` failed.
针对同一设备的任何后续主机端同步调用都将返回cudaErrorAssert。在调用cudaDeviceReset()重新初始化设备之前，无法向该设备发送更多命令。
如果expression不为零，则内核执行不受影响。
例如，以下程序来自源文件 test.cu
#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;
}
将输出：
test.cu:19: void testAssert(): block: [0,0,0], thread: [0,0,0] Assertion `should_be_one` failed.
断言(Assertions)主要用于调试目的。它们可能会影响性能，因此建议在生产代码中禁用断言。通过在包含assert.h头文件之前定义NDEBUG预处理器宏，可以在编译时禁用断言。需要注意的是，expression不应是具有副作用的表达式(例如类似(++i > 0)的表达式)，否则禁用断言会影响代码的功能。
7.33. 陷阱函数
可以通过从任意设备线程调用__trap()函数来启动陷阱操作。
void __trap();
内核执行被中止，并在主机程序中引发中断。
7.34. 断点函数
可以通过从任意设备线程调用__brkpt()函数来暂停内核函数的执行。
void __brkpt();
