Ped*_*dro 0 profiling cuda multiprocessor
一个简单的问题,确实:我有一个内核,它可以运行每个Streaming Multiprocessor(SM)的最大块数,并且想知道我理论上可以从中提取多少性能.理想情况下,我想知道空闲的SM周期的百分比,即所有warp在内存访问时被阻止.
我真的只是想找到这个号码.我不想要的是
Nsight Visual Studio Edition 2.1和2.2问题效率实验提供您要求的信息.应在CUDA 5.0之后的版本中将这些计数器/度量标准添加到Visual Profiler中.
从Nsight Visual Studio Edition 2.2用户指南| 分析工具| 其他分析报告| Profiler CUDA设置| 问题效率科
问题效率提供有关设备发出内核指令的能力的信息.报告的数据包括执行依赖性,符合条件的warp和SM停顿原因.
对于计算能力为2.x的设备,多处理器具有两个warp调度程序.每个warp调度程序最多可管理24个warp,每个多处理器总共48个warp.内核执行配置可能会减少运行时限制.有关入住率的信息,请参阅"实现入住率"实验.第一个调度程序负责具有奇数ID的warp,第二个调度程序负责具有偶数ID的warp.
KEPLER UPDATE:对于计算能力3.x,多处理器有四个warp调度程序.每个warp调度程序最多可管理16个warp,每个多处理器总共64个warp.
在每个指令发布时,每个调度程序将从其活动warp列表中选择一个符合条件的warp并发出指令.如果已获取指令,指令所需的执行单元可用,并且指令没有未满足的依赖性,则warp是合格的.
调度程序在多处理器中报告有关warp的以下统计信息:
活动变形 - 从多处理器上计划到完成最后一条指令为止,变换处于活动状态.每个循环的活动变形计数器增量为0-48.每个周期的最大增量由理论占用率定义.
KEPLER UPDATE范围为每循环0-64.
符合条件的Warp - 如果能够发出下一条指令,则有效的warp是合格的.不符合条件的Warp将报告问题停止原因.此计数器将在每个周期增加0-ActiveWarps.
更新在Fermi上,问题停止原因计数器仅在warp调度程序没有符合条件的warp周期更新.在Kepler上,即使warp调度程序发出指令,问题停止原因计数器也会在每个周期更新.
零符合条件的扭曲 - 如果两个调度程序都没有可以发出的扭曲,则此计数器将每个循环递增1.
一个符合条件的Warp - 如果两个调度程序中只有一个具有可以发出的warp,则此计数器会将每个周期递增1.
KEPLER更新:在Kepler上,计数器是每个调度程序,因此一个符合条件的Warp意味着warp调度程序可以发出指令.在Fermi上,两个调度器都有一个计数器,所以在Fermi上你想要一个合格的Warp计数器尽可能小.
扭曲发出孔 - 此计数器将每个循环增加不符合条件的活动扭曲数.这与Active Warps减去Eligible Warps相同.
长翘曲发出孔 - 此计数器将每个周期递增的活动warp数量不足以发出超过32个时钟周期的指令.长孔表示在长时间延迟的原因(例如障碍和内存操作)下,warp会停止.
问题失速原因 - 每个循环,每个不合格的扭曲将增加一个问题失速原因计数器.所有问题停顿原因计数器的总和等于扭曲问题漏洞. 如果尚未获取下一个汇编指令,则不合格的warp将增加指令获取停止原因.如果输入依赖性尚不可用,则执行依赖性停止原因.这可以通过增加独立指令的数量来减少.如果由于所需资源不可用或未充分利用,或者此类型的太多操作已经未完成,则当前无法进行请求时,数据请求会停止.如果数据请求构成了停顿原因的很大一部分,您还应该运行内存实验来确定是否可以针对每个请求优化现有事务,或者是否需要重新访问算法.如果纹理子系统已经被充分利用并且当前不能接受进一步的操作,则纹理失速原因.如果在__syncthreads()处阻止warp,则同步停止原因.如果此原因很大且内核执行配置仅限于少量块,则考虑将内核网格划分为更多线程块.
Visual Profiler没有可以解决您问题的计数器.在添加计数器之前,您可以使用以下计数器:
计算功能的目标和最大IPC是:
Compute Target Max
Capability IPC IPC
2.0 1.7 2.0
2.x 2.3 4.0
3.x 4.4 7.0
目标IPC用于ALU限制计算.内存绑定内核的目标IPC将更少.对于计算能力2.1设备和更高的设备,使用IPC更难,因为每个warp调度程序都可以双重发布.
| 归档时间： |
|
| 查看次数： |
1411 次 |
| 最近记录： |