1
2
数学库
pmc fbarrier count 数据
2025-09-03 15:25:42
106

在使用hipprof pmc进行我自己的算子内核性能分析时得到以下数据

kernel-name:"DoubleTensorConv3d_64x64x16"
 kernel dispatch index                     0
 gpu id                                    0
 dispatch queue id                         0
 dispatch queue index                      6
 thread id                                 5496
 grid size                                 1605632
 work group size                           256
 shared memory size                        16384
 scratch size                              0
 vgpr count                                36
 sgpr count                                48
 fbarrier count                            0
 signal handle                             0x0
 kernel time                               0.000662(s)
 processed ALU instructions                44.043566(%)
 performance                               7703.152918(Gflops)
 shared memory operation                   1756160
 shared memory bank conflict               10.910456(%)
 L1 cache unit is active                   81.986248(%)
 L1 cache unit is stalled                  41.938926(%)
 L2 cache write unit is stalled            0.000000(%)
 L2 cache hit rate                         48.340548(%)

在同样的数据输入下分析miopen的内核时得到的数据如下

kernel-name:"kme_implicitgemm_3d_fp32_fwd_64x64x16_common_ldsc"
 kernel dispatch index                     0
 gpu id                                    0
 dispatch queue id                         0
 dispatch queue index                      12
 thread id                                 5238
 grid size                                 1605632
 work group size                           256
 shared memory size                        16384
 scratch size                              0
 vgpr count                                36
 sgpr count                                80
 fbarrier count                            7416512
 signal handle                             0x0
 kernel time                               0.000504(s)
 processed ALU instructions                40.045932(%)
 performance                               5963.776000(Gflops)
 shared memory operation                   1279488
 shared memory bank conflict               5.328357(%)
 L1 cache unit is active                   87.914402(%)
 L1 cache unit is stalled                  12.125474(%)
 L2 cache write unit is stalled            0.000000(%)
 L2 cache hit rate                         49.206263(%)

我的问题是:
fbarrier count这个数据项是什么意思?是指线程同步的次数吗?如果是指线程同步的次数的话我自己的内核代码中使用了syncthreads()函数为什么fbarrier count为0?
另外请问海光团队对于miopen的适配工作是怎样进行的?据我所知AMD的matrix core指令与海光的tensor core指令的数据排布要求不同,不止tensorcore这部分工作是怎么进行迁移的。

全部评论(3)
只看楼主
网
网友bIEWUC
3楼
fbarrier count指标数据来源涉及内部指令依赖分析,由于一些历史原因,该指标已经失去了参考价值;分析性能可以不关注该指标;后续我们准备移除该指标,以避免对开发者有所误导
2025/9/5
网
网友ehVriI
2楼
海光研发的dcu tensorcore,miopen自研kernel会根据指令使用说明(具体排布),进行设计算子kernel。dtk 编译器支持miopen自研kernel的编译及适配海光新卡。
2025/9/4
A
ANYA
1楼
老师您好,非常感谢您提出的问题,相关问题已经反馈给技术老师,后续会给您答复
2025/9/4
没有更多内容了
--
楼主
Level.1
0 帖子
楼主热帖
暂无数据