[MI100][Fp16] ConvAsmImplicitGemmGTCDynamicWrwXdlops fails validation

[atamazov]

gfx908, any ROCm:

Two cases from test_conv_igemm_dynamic_xdlops_wrw:

• -n 400 -c 256 -H 7 -W 7 -k 1024 -y 7 -x 7 -p 0 -q 0

$ MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicWrwXdlops \
./bin/test_conv2d --half --cmode conv --pmode default --group-count 1 --disable-forward --disable-backward-data \
--input 400, 256, 7, 7 --weights 1024, 256, 7, 7 --pads_strides_dilations 0 0 1 1 1 1 --trans_output_pads 0 0
...
FAILED: 0.14228
...
# Same config
MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicWrwXdlops \
./bin/MIOpenDriver convfp16 -w 1 -t 1 -i 1 -V 1 \
-n 400 -c 256 -H 7 -W 7 -k 1024 -y 7 -x 7 -p 0 -q 0 -F 4
...
Backward Convolution Weights Failed: 0.588887 > 0.082

• -n 400 -c 256 -H 1 -W 1 -k 1024 -y 1 -x 1 -p 0 -q 0

$ MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicWrwXdlops \
./bin/test_conv2d --half --cmode conv --pmode default --group-count 1 --disable-forward --disable-backward-data \
--input 400, 256, 1, 1 --weights 1024, 256, 1, 1 --pads_strides_dilations 0 0 1 1 1 1 --trans_output_pads 0 0
...
FAILED: 0.18679

What needs to be done:

• (1) W/A for test to unblock CI @atamazov
• (2) Check kernels.
  • If kernels are correct, then this may be a verification issue, similar to [iGemmfwd][test_conv2d][gfx906][half] Verification failed #936
  • Otherwise ConvAsmImplicitGemmGTCDynamicWrwXdlops must be fixed ASAP

This is release critical issue!

[shurale-nkn]

ConvAsmImplicitGemmGTCDynamicFwdXdlops may also fail under test_conv2d

FAILED: 0.109962
Iteration: 0
Forward convolution: ConvAsmImplicitGemmGTCDynamicFwdXdlops
Input tensor: 8, 32, 14, 14
Weights tensor: 32, 32, 1, 1
Output tensor: 8, 32, 16, 16
Filter: conv2d, miopenConvolution, miopenPaddingDefault, {1, 1}, {1, 1}, {1, 1}, 
Max diff: 87
Mismatch at 16: 0 != -24

[atamazov]

@shurale-nkn This ticket is about WrW.

[shaojiewang]

Root cause is found.

[atamazov]

Good!