[PyTorch][Winograd] Winograd kernel been selected has caused issue with test_Conv2d_naive_groups_cuda_float16

[junliume]

[Summary]

Winograd kernels are by design aiming performance by sacrificing numerical accuracy.
However, in this case for very small and non-practical case, selecting winograd kernels have caused test_Conv2d_naive_groups_cuda_float16 to fail.

Question:

• test_Conv2d_naive_groups_cuda_float16 has keyword naive in it, does it expect naive implementations to begin with?
• @Kirpich30000 should winograd kernels have issues with such cases? i.e. -H 6 -W 6 -k 2

MIOpenDriver convfp16 -n 2 -c 2 -H 6 -W 6 -k 2 -y 3 -x 3 -p 0 -q 0 -u 1 -v 1 -l 1 -j 1 -m conv -g 1 -F 1 -t 1 -S 0
Forward Conv solutions available: 2
- id: 84 algo: 3, time: 10 ms, ws: 0, name: ConvBinWinogradRxSf2x3g1
- id: 107 algo: 5, time: 20 ms, ws: 1280, name: ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC
MIOpen Forward Conv. Algorithm: 3, Solution: 84/ConvBinWinogradRxSf2x3g1
GPU Kernel Time Forward Conv. Elapsed: 0.015378 ms (average)
stats: name, n, c, ho, wo, x, y, k, flopCnt, bytesRead, bytesWritten, GFLOPs, GB/s, timeMs
stats: fwd-conv3x3u1, 2, 2, 4, 4, 3, 3, 2,  2304, 360, 128, 0, 0, 0.015378
Forward Convolution Verifies OK on CPU reference (0.000340009)

[Observation and Steps to reproduce]:

To Reproduce:

PYTORCH_TEST_WITH_ROCM=1 python3 nn/test_convolution.py --use-pytest --verbose -k test_Conv2d_naive_groups_cuda_float16
Docker Images:

ROCM 5.6: rocm/pytorch:rocm5.6_ubuntu20.04_py3.8_pytorch_2.0.1
PyTorch Installed at /var/lib/jenkins/pytorch/test
ROCM 5.7:  compute-artifactory.amd.com:5000/rocm-plus-docker/framework/compute-rocm-rel-5.7:86_ubuntu20.04_py3.9_pytorch_rocm5.7_internal_testing_55fbbdf
Original image: rocm/pytorch-private:86_ubuntu20.04_py3.9_pytorch_rocm5.7_internal_testing_55fbbdf
PyTorch Installed at /var/lib/jenkins/pytorch/test

NOTE: tolerance has already been raised to 1e-1, You need to run git revert e9b273df57b240f14ead07b5fda97bdf2be6673a to see the error
Expected Output:

nn/test_convolution.py::TestConvolutionNNDeviceTypeCUDA::test_Conv2d_naive_groups_cuda_float16 PASSED

Actual Output:

Mismatched elements: 47 / 128 (36.7%)
Greatest absolute difference: 0.0009765625 at index (0, 2, 2, 1) (up to 1e-05 allowed)
Greatest relative difference: 0.0999755859375 at index (0, 0, 2, 0) (up to 0.001 allowed)

[junliume]

@JehandadKhan @atamazov F.Y.I.

[JehandadKhan]

I dont think the practicality of the convolution operation is relevant here, however, if the client wishes to enforce certain numerical behavior, then they can filter out numerically inferior algorithms when MIOpen returns the result, instead of defaulting to the first element in the results of the miopenFindConvolution* call.

[atamazov]

@JehandadKhan Absolutely. However we do not know what is the actual intent of the test. The one may assume that the client wants to validate the precision of the fastest convolution available in MIOpen (i.e. the one which is normally used for computations). If this is so, then choosing some other (not the fastest, but more accurate) solution in the test does not look like a proper solution.

From the other hand, the problem config used in the test seems very small and quite far from any convolution used in the real networks. This means that we can make Winograd inapplicable for this config and the test will pass without errors, while the performance of the real networks won't be affected. However, I am far from suggesting a hack just to make some test passing ;)

In fact, we won't be able to figure out what the problem is until we see the actual values that lead to large absolute/relative differences. Let me explain with an example: Let’s say we consider a convolution that performs calculations with an error of 1ULP to be sufficient.

Let's assume that some test expects value near max FP16 (65504); max-1ULP = 65472, which gives absolute diff 32 which is far greater than 1e-05. We can easily run into similar situation with small values: the smallest positive denorm 5.97e-8, the adding +1ULP gives ~1.2e-7 and relative difference (computed as (1.2e-7 - 5.97e-8)/(1.2e-7 + 5.97e-8)/2) is ~0.168 which is again far greater than 1e-3.

Proposal:

• Ask @Kirpich30000 if precision of ConvBinWinogradRxSf2x3g1 may degrade for small configs like -n 2 -c 2 -H 6 -W 6 -k 2 -y 3 -x 3 -p 0 -q 0 -u 1 -v 1
• If the kernel is Ok, then let's investigate the test in detail.

---

💡

• The test_Conv2d_naive_groups test can be found here:
  https://github.com/ROCmSoftwarePlatform/pytorch/blob/39ff80125f5c6d240a4b59010253cb4adab5090f/test/nn/test_convolution.py#L2056
• Similar defect in PyTorch repo, dated October 2021: test_Conv2d_naive_groups_cuda_float16 (__main__.TestNNDeviceTypeCUDA) is flaky on ROCm pytorch/pytorch#66378. The proposed fix implements filtering ([ROCm] support benchmark flag for MIOpen pytorch/pytorch#77438, it attempts to use GEMM if available), but it wasn't merged.

[Kirpich30000]

Hi, this geometry should be fine with Winograd. Since the test is rather small we could analyze the actual data (reference input, filter, output and Winograd output) if the dumps are available.

[wenchenvincent]

@JehandadKhan Absolutely. However we do not know what is the actual intent of the test. The one may assume that the client wants to validate the precision of the fastest convolution available in MIOpen (i.e. the one which is normally used for computations). If this is so, then choosing some other (not the fastest, but more accurate) solution in the test does not look like a proper solution.

From the other hand, the problem config used in the test seems very small and quite far from any convolution used in the real networks. This means that we can make Winograd inapplicable for this config and the test will pass without errors, while the performance of the real networks won't be affected. However, I am far from suggesting a hack just to make some test passing ;)

In fact, we won't be able to figure out what the problem is until we see the actual values that lead to large absolute/relative differences. Let me explain with an example: Let’s say we consider a convolution that performs calculations with an error of 1ULP to be sufficient.

Let's assume that some test expects value near max FP16 (65504); max-1ULP = 65472, which gives absolute diff 32 which is far greater than 1e-05. We can easily run into similar situation with small values: the smallest positive denorm 5.97e-8, the adding +1ULP gives ~1.2e-7 and relative difference (computed as (1.2e-7 - 5.97e-8)/(1.2e-7 + 5.97e-8)/2) is ~0.168 which is again far greater than 1e-3.

Proposal:

• Ask @Kirpich30000 if precision of ConvBinWinogradRxSf2x3g1 may degrade for small configs like -n 2 -c 2 -H 6 -W 6 -k 2 -y 3 -x 3 -p 0 -q 0 -u 1 -v 1
• If the kernel is Ok, then let's investigate the test in detail.

💡

• The test_Conv2d_naive_groups test can be found here:
  https://github.com/ROCmSoftwarePlatform/pytorch/blob/39ff80125f5c6d240a4b59010253cb4adab5090f/test/nn/test_convolution.py#L2056
• Similar defect in PyTorch repo, dated October 2021: test_Conv2d_naive_groups_cuda_float16 (main.TestNNDeviceTypeCUDA) is flaky on ROCm pytorch/pytorch#66378. The proposed fix implements filtering ([ROCm] support benchmark flag for MIOpen pytorch/pytorch#77438, it attempts to use GEMM if available), but it wasn't merged.

I think this is the way to go. We need to first conduct error analysis. If this is a config that is not proper for winograd, we should just disable winograd when running the algorithm finding heuristic in MIOpen.

[CAHEK7]

I can barely imagine a real case for precision based heuristic.
There is no external api to point precision limitations, precision itself can be calculated theoretically but the real error heavily depends on the input data (I've already faced this while touched random generator for MIOpen tests).

[JehandadKhan]

@atamazov Can you please create a PR to restrict the applicability of this solver in coordination with @Kirpich30000 ?

[atamazov]

@wenchenvincent

If this is a config that is not proper for winograd, we should just disable winograd when running the algorithm finding heuristic in MIOpen.

It is proper. However, as per @JehandadKhan's suggestion, maybe we'll temporarily restrict applicability of the Winograd solver, as a workaround, until the investigation of test failure is completed and a full-blown fix is done.

[Kirpich30000]

@wenchenvincent @JehandadKhan @junliume

I've dumped output values and analyzed the data.

1. Winograd kernel doesn't have any unexpected precision loss on that test
2. The test performs convolution with G=2 using different methods (single conv operation and 2 separate nongrouped convolutions) and check if outputs of both methods are equal with some tolerance. It implicitly assumes that both methods would have the same numerical properties which is not true if different implementations would be used.
3. The test has few weak points. For example:

• It doesn't check if output is correct. The test would pass if miopen would always return zeros as output.
• If we replace Winograd with 64-bit precise implementation, then the test would still fail, because difference with implicit gemm wouldn't fit in allowed range (10e-5 abs and 10e-3 rel)

One possible solution is to change data initialization to patterns which guarantee to have bit exact results (that is possible for direct convolution and f2x3, f3x2 Winograd convolutions). If such change is not considered, then some workaround on MIOpen side is needed to peek the same algorithm for the tested convolutions. It doesn't matter if it be Winograd, igemm, FFT or any other possible approach, but it should be the same one for both methods used in the test. Some other workarounds might be possible.

[atamazov]

@wenchenvincent @JehandadKhan @junliume Right now the test has FP16 input data distribution that, in fact, requires output values to be bit exact in many cases. This is quite questionable (see #2492 (comment))

It is highly recommended to increase max abs diff form 1e-5 to 1e-3 and relative diff to 2e-2 in order to allow at least ~1ULP output deviation for FP16. That should reduce probability of false failures in this test provided that input data distribution remains the same.