分析器计数器功能：
每个多处理器都有一组16个硬件计数器，通过调用__prof_trigger（）函数，应用程序可以通过单个指令来递增。
void __prof_trigger(int counter);索引计数器的每个多处理器硬件计数器每次增加一个。 计数器8至15被保留，不应被应用程序使用。
计数器0,1，...，7的值可以通过nvprof - 事件prof_trigger_0x来获得，其中x是0,1，...，7.在每次内核启动之前，所有的计数器都被重置（注意，当收集 计数器，内核启动与主机和设备之间的并行执行中提到的是同步的）
断言：
声明仅受计算能力2.x及更高版本的设备支持。 不管设备如何，它都不支持MacOS，并且加载引用Mac OS上断言功能的模块将失败。
void assert(int 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（）来重新初始化设备之前，不能再有命令发送到此设备。
如果表达式与零不同，内核执行不受影响。
例如，源文件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.断言是为了调试目的。 它们会影响性能，因此建议在生产代码中禁用它们。 通过在包含assert.h之前定义NDEBUG预处理器宏，可以在编译时禁用它们。 请注意，表达式不应该是带有副作用的表达式（例如（++ i> 0）），否则禁用断言将影响代码的功能。
格式化输出：
格式化输出仅受计算能力2.x及更高版本的设备支持。
int printf(const char *format[, arg, ...]);从内核打印格式化输出到主机端输出流。
内核中的printf（）函数的行为与标准C函数库printf（）函数类似，用户可以参考主机系统的手册页获取printf（）行为的完整描述。 本质上，以格式传入的字符串被输出到主机上的一个流中，在遇到格式说明符的地方使用参数列表进行替换。 下面列出了支持的格式说明符。
printf（）命令与任何其他设备端函数一样执行：per-thread和调用线程的上下文中。 从多线程内核来说，这意味着对每个线程执行printf（）的直接调用，并使用指定的线程数据执行。 输出字符串的多个版本将出现在主机流上，每遇到printf（）的每个线程都会出现一次。
如果只需要一个输出字符串，则由程序员决定将输出限制为单个线程（请参阅示例以了解示例）。
与C-standard printf（）不同，它返回打印的字符数，CUDA的printf（）返回解析的参数个数。 如果格式字符串后面没有参数，则返回0。 如果格式字符串为NULL，则返回-1。 如果发生内部错误，则返回-2。
