刚开始用profiler分析程序，我用的版本是5.0
可能是不熟悉的缘故，用起来不是很顺手
比如我就是没找到active block的数据，这倒是也可以自己算出来，但也应该直接由提供的吧？
我在论坛上下载了profiler4.0用户指南中文版，里面介绍的是4.0版本，我发现比5.更清楚
如4.0图[attach]3254[/attach]
我在5.0里面就没有找到如此清楚的描述
因此我想问版主
(1)profiler5.0怎么查acitve block的数据（不是自己算）
(2)如何查register不够用导致的local memory的使用量
(3)如下图中很多参数我不知道什么意思，像local memory overhead、shared memory replay overhead等等，跪求版主指点
楼主您好，曾经有多种资料（无论是官方的还是非官方的资料）使用过active这个名词。
在解答您的问题前，我重点将profiler的历史版本的active用词给予解释：
在老版本(compute profiler)中，active threads/blocks实际上是resident threads/blocks, 并且此数值直接等于理论值。
在新版本(NVIDIA Visual Profiler)中， active是指resident的warps/threads中，处于可执行态的warps/threads。换句话说，active warps是SM上resident warps的一部分，而active threads是驻留的线程中的一部分（即驻留的不一定是active的)
那么根据这个：
（1）如果您需要老版定义的active blocks per sm，您应当直接使用理论值。我建议你此时使用resident blocks以便和大部分资料一致。
（2）如果您需要使用新版定义的active blocks per sm(即平均每周期内有几个blocks可以处于可执行的状态）, 那么您无法直接取得，但是您可以使用active warps 除以 elapsed_cycles 除于 SM个数 除以 每block的warps数目。
感谢您的来访。
楼主您好，
关于您的第二个问题：“如何查register不够用导致的local memory的使用量”
可能在编译参数中指定–ptxas-options=-v后，观察ptxas的输出，查找如下行：
XX bytes stack frame, XX bytes spill stores, XX bytes spill loads
他们应该共同构成了local memory的使用，
前者应该是对__device__的调用，以及大数组的时候用的local memory。
后者可能是针对寄存器不足被交换出来而使用的local memory.
但我不知道他们的具体含义(例如XX Bytes spill stores是一共使用了XX字节的local memory作为交换用呢，还是一共交换写入了XX字节），手册也没说明。但这行大致可以参考。
（或者您也可以参考profiler给出的local store和local read的指令数）
关于您的第三个问题，local memory replay overhead和shared memory replay overhead是指1 - local/shared memory读写指令数 / 实际issue的总数， 最后乘以100%。（因为某些原因，会导致访问local/shared memory的执行被发射多次才能执行一次的(例如cache miss或者bank conflict等)。
这统计了实际issue的总数中，多少比例是无用的。
感谢版主，信息量很大，我仔细消化消化
大致懂了。
也就是说，假如是C2050卡，有14个SM，如果一共有42个bolck，并且由于资源(如shared memory, register等)限制每个SM中只能有两个block同时在计算。
则按照新版本的说法：每个SM中的resident block数是3个，每个SM中的active block数是2个。
不知我是否理解正确？
楼主您好，
（1）如果在14 SM上同时能存在42个blocks, 那么每SM的resident blocks数目是3
（2）关于第二点，实际上我们不讨论active blocks(这个单位太粗糙，而且blocks大小是可变的），我们这里用active warps来代替（这个是所有SM上的所有周期的积分值），active warps / SM总数 / 总的时间(周期数)您可以得到一个值，这个值在贵卡上应该>0同时小于等于48(贵卡一个SM最多能驻留48个warps, 如果总是都处于就绪态，那么是48）。这衡量了平均每周期多少个warps是可以被调度执行的（但这个不代表性能），您可以用这个指标来衡量TLP的程度（但不代表性能）。
以及，回到您的问题，每个kernel的算法的平均active warps数目(每SM每周期)是多少这个真心不一定的。要看您的具体算法运行效果如何。
以及，再次强调一遍，这个不代表性能。较真的说，这个和性能完全无关。
感谢版主，从您的评论当中，我发现讨论active warps比active blocks要更合理，是因为一个warp中的不同线程是同时发射到SM上进行计算的缘故吗？
还有一个问题，如果我定义了block大小为(16,16)，那么一个block有8个warp，如果计算资源(如shared memory，register等)限制每个SM至多有12个warp同时计算，那么在实际计算时，每个SM在同一时刻是有8个warp(同属一个block)在运行，还是同一时刻有12个warp在运行？
我从版主的评述中得到的理解，并且归纳总结是：不同的warp按照随机的顺序发射到SM运行，如果没有block内部同步的话（__syncthreads()），这些warp完成计算的顺序也是随机的，可能出现全部的block，每个block都有一部分warp已完成计算和另一部分warp未完成的情况；而block内部需要同步的话，同一个block内部的不同warp之间要等待同步，因此一个block中的不同warp是同步完成的(假设同步操作在kernel函数的最后)，而不像之前那么随机。
以及(哈哈，版主风格)，对于我最初提到的register和local memory使用情况的问题。
在修改程序前，我在编译参数中加入–ptxas-options=-v后，出现
88 bytes stack frame, 188 bytes spill stores, 176 bytes spill loads
这表明肯定是有使用local memory了吧？
在修改程序后(将原先的一个kernel分割成多个kernel)，结果
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
这表明肯定是没有使用local memory了吧？
楼主您好，
如果资源允许最多12个warps/sm, 但是您的block有8个warp, 那么此SM上只能同时驻留一个block(剩下的能放下4个warps的资源将不被使用，浪费了，因为不能上半个block)
我没有说过“不同的warp按照随机的顺序发射到SM运行…”这段话，并表示无法理解这段话，因为无法为您解释说明此段落。
关于“88 bytes stack frame, 188 bytes spill stores, 176 bytes spill loads”
–这个肯定是用了。但我无法确定一共有多大local memory/thread. 是88B呢？还是其他？前者包含一些正常的非寄存器果断而导致的local memory使用呢？还是也包含寄存器交换呢？这个我无法确定。唯一能确定的是您的确使用了local memory. (您也可以从profiler里看，将产生local store/load)
关于“0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads"
–这个应该是没有。但我不确定。因为我没有关于ptxas输出格式的资料。但是您可以通过profiler看(local store和load的指令数，这个如果都是0,那么肯定没有）
感谢深夜来访。
这里加上了我的一些推测，呵呵，感谢版主，浪费了多出来的warp资源，真是可惜
嗯，我查看了local store和load，在程序修改前local store and load有很大的数，修改后都为零。十分感谢版主指点。
还想再麻烦一下版主，profiler里面有哪些常用的，比较重要的参数？能否简单的说明一下？
因为Metric和event里面的参数实在太多，抓不住重点，而且许多连什么意思的不知道。再次谢过版主。
这个真心无法回答。
（1）提供的数据都是衡量贵程序的各个方面的，很难说哪些有用，哪些没用/不重要。（
请问楼主您的眼、耳、鼻、舌、手，哪些比较重要？哪些是常用的？）
（2）逐个讲解里面的每个数据的作用超出了本论坛的工作范围，您可以直接阅读profiler的资料来自学。
了解，谢版主