如何计算整个CUDA内核执行的执行线程数？

Question

我想逐渐计算整个内核执行的线程执行次数.是否有本地计数器或有没有其他方法可以做到这一点？我知道保留一个全局变量并且每个线程的增量都不会起作用,因为全局内存中的变量不保证线程的同步访问.

Answer 1

有许多方法可以衡量线程级执行效率.这个答案提供了不同收集机制的列表.Robert Crovella的答案提供了一种手动仪器方法,可以准确地收集信息.类似的技术可用于收集内核中的分歧信息.

为执行启动的线程数(静态)

gridDim.x*gridDim.y*gridDim.z*blockDim.x*blockDim.y*blockDim.z

发布的线程数

gridDim.x*gridDim.y*gridDim.z*ROUNDUP((blockDim.x*blockDim.y*blockDim.z),WARP_SIZE)

此数字包括在warp的生命周期内处于非活动状态的线程.

这可以使用PM计数器thread_seunched收集.

已执行变形指令

计数器inst_executed计算执行/退出的扭曲指令的数量.

发布变形指令

计数器inst_issued计算发出的指令数.inst_issued> = inst_executed.为了处理对窄执行单元的调度或者为了处理共享存储器和L1操作中的地址差异,将对每个执行的指令多次发出一些指令.

已执行线程指令

计数器thread_inst_executed计算执行的线程指令的数量.可以使用thread_inst_executed/inst_executed导出指标avg_threads_executed_per_instruction.此计数器的最大值为WARP_SIZE.

未预测脱线程执行的指令

计算能力2.0及更高版本的设备使用指令预测来禁用warp中线程的回写,作为发散指令的短序列的性能优化.

计数器not_predicated_off_thread_inst_executed计算所有线程执行的指令数.此计数器仅适用于计算能力3.0及更高版本的设备.

not_predicated_off_thread_inst_executed <= thread_inst_executed <= WARP_SIZE*inst_executed

由于thread_inst_executed和not_predicated_off_thread_inst_executed计数器中存在小错误,因此某些芯片上的关系会稍微偏离.

剖析

Nsight Visual Studio Edition 2.x支持收集上述计数器.

Nsight VSE 3.0支持新的指令计数实验,该实验可以收集每个SASS指令统计信息,并以表格形式或高级源,PTX或SASS代码旁边显示数据.信息从SASS汇总到高级别来源.汇总的质量取决于编译器输出高质量符号信息的能力.建议您始终同时查看源和SASS.该实验可以收集以下每条指令统计信息:

一个.inst_executed b.thread_inst_executed(或活动掩码)c.not_predicated_off_thread_inst_executed(活动谓词掩码)d.active_mask的直方图e.predicate_mask的直方图

Visual Profiler 5.0可以准确地收集上述SM计数器.nvprof可以收集并显示每个SM的详细信息.Visual Profiler 5.x不支持Nsight VSE 3.0中可用的每指令统计信息的收集.较旧版本的Visual Profiler和CUDA命令行分析器可以收集许多上述计数器,但结果可能不如5.0及以上版本的工具那样准确.