我在DCU上优化卷积算子,按照我的性能计算方式,我的算子比miopen的算子要慢,性能要低,于是我用hipprof --pmc来分析我的算子性能。但是在分析的过程中我发现hipprof会返回一个performance的数据,而在hipprof的数据下,我的算子的performance比miopen要好。
我认为是hipprof出了问题,它返回的性能数据(miopen的3D卷积)如下:
kernel-name:"kme_implicitgemm_3d_fp32_fwd_64x64x16_common_opt_out"
kernel dispatch index 0
gpu id 0
dispatch queue id 0
dispatch queue index 12
thread id 17534
grid size 1722368
work group size 256
shared memory size 8192
scratch size 0
vgpr count 36
sgpr count 64
fbarrier count 7359360
signal handle 0x0
kernel time 0.002234(s)
processed ALU instructions 90.032578(%)
performance 10637.973887(Gflops)
shared memory operation 11787456
shared memory bank conflict 11.830472(%)
L1 cache unit is active 48.690313(%)
L1 cache unit is stalled 0.324260(%)
L2 cache write unit is stalled 0.000000(%)
L2 cache hit rate 48.102905(%)
请留意这里performane的数据,而用我的测量方法远达不到这个数据。我的测量方法是在kernel运行前后打桩,测量运行时间,与理论的浮点运算次数相除,关键代码如下:
double flopsPerConv = K * N * OH * OD * OW * 2.0 * KD * KH * KW * 1e-9f;
hipEventRecord(start, 0);
int iternum = 100;
for (int i = 0; i < iternum; i++)
{// 运行卷积
miopenConvolutionForward(
handle, &alpha, xDesc, x, wDesc, w, convDesc,
algo, &beta, yDesc, y, workspace, wsSize);
}
hipEventRecord(stop, 0);
hipEventSynchronize(stop);
std::cout << "Conv3D forward executed.\n";
float milliseconds = 0;
hipEventElapsedTime(&milliseconds, start, stop);
float timePerconv = milliseconds / iternum;
double gflops = flopsPerConv / (timePerconv / 1000.0f);
std::cout << "Conv3D forward executed in " << milliseconds / iternum << " ms." << std::endl;
std::cout << "Performance: " << gflops << " GFlops" << std::endl;
输出的性能为:
Conv3D forward executed in 2.81233 ms.
Performance: 132.266 GFlops
这跟hipprof的数据对不上,请问是哪里错了?

暂无数据
老师,您好。详细分析见下图。