[HOTFIX][tests] WORKAROUND_ISSUE_2038: Disable validation of FP16 and BF16 in the smoke test of ConvHipImplicitGemmV4R1Fwd

[atamazov]

This is expected to unblock CI. Introduces workaround for issue #2038. Related to issue #936.

This is alternative for #2041 (see #2038 (comment)).

---

[Attribution] @junliume @johnny-keker

• https://github.com/ROCmSoftwarePlatform/MIOpen/labels/urgency_blocker
• https://github.com/ROCmSoftwarePlatform/MIOpen/labels/workaround
• https://github.com/ROCmSoftwarePlatform/MIOpen/labels/correctness
• Proposed reviewers:
  • @JehandadKhan
  • @junliume
  • @carlushuang
  • @asroy

[junliume]

@atamazov @JehandadKhan @carlushuang @asroy I prefer #2041 over this because not tuning during smoke test will mask the numerical issues we eventually wish to fix. What do you think?

[junliume]

and there is another reason to prefer #2041 over this one, since smoke_solver_ConvHipImplicitGemmV4R1Fwd is failing pretty consistently.

[2023-03-23T23:44:55.949Z]  70/324 Test #128: smoke_solver_ConvHipImplicitGemmV4R1Fwd .................................................................................................................***Failed  Error regular expression found in output. Regex=[FAILED]  6.21 sec

[2023-03-23T23:44:55.949Z] make[4]: warning: jobserver unavailable: using -j1.  Add '+' to parent make rule.

[2023-03-23T23:44:55.949Z] [  2%] Built target sqlite_memvfs

[2023-03-23T23:44:55.949Z] [  4%] Built target addkernels

[2023-03-23T23:44:55.949Z] [100%] Built target MIOpen

[2023-03-23T23:44:55.949Z] [100%] Built target test_conv2d

[2023-03-23T23:44:55.949Z] Scanning dependencies of target smoke_solver_ConvHipImplicitGemmV4R1Fwd

[2023-03-23T23:44:55.949Z] /home/jenkins/workspace/MLLibs_MIOpen_PR-2043/build/bin/test_conv2d --half --cmode conv --pmode default --group-count 1 --disable-backward-data --disable-backward-weights --input 256 32 27 27 --weights 128 32 1 1 --batch_size 256 --input_channels 32 --output_channels 128 --spatial_dim_elements 27 27 --filter_dims 1 1 --pads_strides_dilations 0 0 1 1 1 1 --trans_output_pads 0 0 --in_layout NCHW --fil_layout NCHW --out_layout NCHW --deterministic 0 --tensor_vect 0 --vector_length 1 --output_type int32 --int8_vectorize 0 

[2023-03-23T23:44:55.949Z] FAILED: 0.56234

[2023-03-23T23:44:55.949Z] Max diff: 263

[2023-03-23T23:44:55.949Z] Mismatch at 0: 30 != 124

[2023-03-23T23:44:55.949Z] Forward convolution: ConvHipImplicitGemmV4R1Fwd

[2023-03-23T23:44:55.949Z] Input tensor: 256, 32, 27, 27

[2023-03-23T23:44:55.949Z] Weights tensor: 128, 32, 1, 1

[2023-03-23T23:44:55.949Z] Output tensor: 256, 128, 27, 27

[2023-03-23T23:44:55.949Z] Filter: conv2d, miopenConvolution, miopenPaddingDefault, {0, 0}, {1, 1}, {1, 1}, 

[carlushuang]

LGTM

[atamazov]

@atamazov @JehandadKhan @carlushuang @asroy I prefer #2041 over this because not tuning during smoke test will mask the numerical issues we eventually wish to fix. What do you think?

We still do not know the reason of failures.

• It could happen due to different order of computations (the verification algorithm is sensitive to it, as we already know). See (B) at Mismatch in ConvHipImplicitGemmV4R1Fwd  #2038 (comment).
• Or there is a real problem in the kernel. We need investigation from our convolution kernels' specialists, preferably authors of the kernel. See (A) at Mismatch in ConvHipImplicitGemmV4R1Fwd  #2038 (comment).

I will update both this PR and #2041. Then you'll be able to select the one which better suits our needs.

• [HOTFIX] WORKAROUND_ISSUE_2038: Disable FP16 and BF16 in ConvHipImplicitGemmV4R1Fwd for MI100/200 and all new targets. #2041 will disable both BF16 and FP16, so there is some risk of performance regression.
• This PR will disable validation in the BF16 and FP16 smoke tests (which is camouflaging the numerical issues indeed, IF THEY REALLY EXIST)

[atamazov]

@junliume @JehandadKhan @carlushuang @asroy PR is ready for the next round of reviewing and CI testing.

[junliume]

@atamazov @JehandadKhan let's merge this PR (test only) and making #2041 standing by in case of further actions needed to disable this solver more completely.