刚开始用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的资料来自学。
了解,谢版主