fbarrier count指标数据来源涉及内部指令依赖分析,由于一些历史原因,该指标已经失去了参考价值;分析性能可以不关注该指标;后续我们准备移除该指标,以避免对开发者有所误导
在使用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这部分工作是怎么进行迁移的。

暂无数据