[Tuning] rel-4.5 fdb update for develop

[cderb]

No description provided.

[atamazov]

Is this from 4.5 tuning?

[junliume]

Is this from 4.5 tuning?

I think so, and the source branch name also double confirms cderb/fdb_tuning_dev-4.5 so.

[atamazov]

In performance testing right now...

[atamazov]

fdb duplicate IDs for gfx906/60

[atamazov]

Duplicate IDs for gfx906/60/HIP

It seems like the root reason of #1133 (comment) is not yet fixed.

Error messages

MIOpen(HIP): Error [ParseContents] Duplicate ID (ignored): miopenConvolutionBwdDataAlgoDirect; key: 256-28-28-1x1-64-56-56-8-0x0-2x2-1x1-0-NCHW-BF16-B
MIOpen(HIP): Error [ParseContents] Duplicate ID (ignored): miopenConvolutionBwdDataAlgoDirect; key: 64-56-56-3x3-64-56-56-8-1x1-1x1-1x1-0-NCHW-BF16-B
MIOpen(HIP): Error [ParseContents] Duplicate ID (ignored): miopenConvolutionBwdDataAlgoWinograd; key: 128-6-60-3x3-64-6-60-16-1x1-1x1-1x1-0-NCHW-FP16-B
MIOpen(HIP): Error [ParseContents] Duplicate ID (ignored): miopenConvolutionBwdDataAlgoWinograd; key: 256-7-7-1x1-832-7-7-16-0x0-1x1-1x1-0-NCHW-FP32-B
MIOpen(HIP): Error [ParseContents] Duplicate ID (ignored): miopenConvolutionFwdAlgoDirect; key: 256-14-14-1x1-1024-7-7-8-0x0-2x2-1x1-0-NCHW-BF16-F
MIOpen(HIP): Error [ParseContents] Duplicate ID (ignored): miopenConvolutionFwdAlgoDirect; key: 256-14-14-1x1-256-14-14-8-0x0-1x1-1x1-0-NCHW-FP16-F
MIOpen(HIP): Error [ParseContents] Duplicate ID (ignored): miopenConvolutionFwdAlgoImplicitGEMM; key: 128-28-28-1x1-512-28-28-8-0x0-1x1-1x1-0-NCHW-FP32-F
MIOpen(HIP): Error [ParseContents] Duplicate ID (ignored): miopenConvolutionFwdAlgoImplicitGEMM; key: 512-28-28-1x1-1024-14-14-16-0x0-2x2-1x1-0-NCHW-FP32-F
MIOpen(HIP): Error [ParseContents] Duplicate ID (ignored): miopenConvolutionFwdAlgoWinograd; key: 512-7-7-3x3-512-7-7-8-1x1-1x1-1x1-0-NCHW-FP16-F

Examples of duplicates

256-28-28-1x1-64-56-56-8-0x0-2x2-1x1-0-NCHW-BF16-B=

miopenConvolutionBwdDataAlgoDirect:
    ConvOclDirectFwd1x1,0.10016,0,miopenConvolutionBwdDataAlgoDirect,256x28x28x1x1x64x56x56x8xNCHWxBF16x0x0x2x2x1x1x1xB;
miopenConvolutionBwdDataAlgoDirect:
    ConvOclDirectFwd1x1,0.10016,0,miopenConvolutionBwdDataAlgoDirect,256x28x28x1x1x64x56x56x8xNCHWxBF16x0x0x2x2x1x1x1xB;
miopenConvolutionBwdDataAlgoGEMM:
    GemmBwd1x1_stride2,1.48128,4014080,miopenConvolutionBwdDataAlgoGEMM,256x28x28x1x1x64x56x56x8xNCHWxBF16x0x0x2x2x1x1x1xB

128-28-28-1x1-512-28-28-8-0x0-1x1-1x1-0-NCHW-FP32-F=

miopenConvolutionFwdAlgoImplicitGEMM
    ConvHipImplicitGemmV4R4Fwd,0.09984,0,miopenConvolutionFwdAlgoImplicitGEMM,128x28x28x1x1x512x28x28x8xNCHWxFP32x0x0x1x1x1x1x1xF
miopenConvolutionFwdAlgoDirect
    ConvAsm1x1U,0.11744,0,miopenConvolutionFwdAlgoDirect,128x28x28x1x1x512x28x28x8xNCHWxFP32x0x0x1x1x1x1x1xF
miopenConvolutionFwdAlgoImplicitGEMM
    ConvHipImplicitGemmV4R4Fwd,0.13424,0,miopenConvolutionFwdAlgoImplicitGEMM,128x28x28x1x1x512x28x28x8xNCHWxFP32x0x0x1x1x1x1x1xF
miopenConvolutionFwdAlgoGEMM
    GemmFwd1x1_0_1,0.13424,0,miopenConvolutionFwdAlgoGEMM,128x28x28x1x1x512x28x28x8xNCHWxFP32x0x0x1x1x1x1x1xF
miopenConvolutionFwdAlgoWinograd
    ConvBinWinogradRxSf3x2,0.231841,0,miopenConvolutionFwdAlgoWinograd,128x28x28x1x1x512x28x28x8xNCHWxFP32x0x0x1x1x1x1x1xF

[atamazov]

Missing fdb records for gfx906/60/HIP

gfx906_60.HIP.missing.fp32.txt
gfx906_60.HIP.missing.fp16.txt
gfx906_60.HIP.missing.bf16.txt

I suspect that FP32 and FP16 configs are required (not sure about BF16, though)

[atamazov]

Performance testing results

🟢 These are good!

ROCm 4.0, gfx906/60, Performance gain, based on sum of all times/directions:

		GPU time	Wall time	Aux Wall time
FP32/HIP	Default find mode	1.13	1.11	1.45
	Immediate mode	1.13	1.10	1.49
FP16/HIP	Default find mode	1.00	0.93	0.74
	Immediate mode	1.00	0.98	1.17
BF16/HIP	Default find mode	0.99	0.75	1.16
	Immediate mode	0.99	0.99	1.12
FP32/OpenCL	Default find mode	1.00	1.01	0.99
	Immediate mode	1.00	1.00	1.00

[atamazov]

Verdict

I tested this only on gfx906/60. Assuming other GPUs perform similarly, I think this can be merged provided that "Duplicate ID" issues are resolved.

[cderb]

Thanks Artem, I'll start looking into the source of the duplication.

[cderb]

fdb entries no longer contain duplicate algorithms.

Please re-review. fdb entries no longer contain duplicate algorithms.

[atamazov]

Testing on MI100 is ongoing.

[junliume]

@atamazov @JehandadKhan
Would need #1196 and #1176 for next MIOpen staging promotion
#1192 is also nice to have

[atamazov]

MI100 testing results (ROCm 4.3.1)

🟡 Looks acceptable. Develop vs cderb/fdb_tuning_dev-4.5, tested with MI100 & ROCm 4.3.1. Performance gain, based on sum of all times/directions:

		GPU time	Wall time	Aux Wall time
FP32/HIP	Default find mode	1.00	0.81	0.94
	Immediate mode	1.00	0.96	1.14
FP16/HIP	Default find mode	1.04	0.93	0.78
	Immediate mode	1.03	0.90	0.97

[atamazov]

LGTM!

[junliume]

@atamazov @JehandadKhan Would need #1196 and #1176 for next MIOpen staging promotion #1192 is also nice to have

Merged #1196 and initiated another round of develop staging, #1176 is functional update and can wait till next week's round.