Repository: pytorch
Issue: conv3d has numerical issue where same input produces output that are not bit-wise identical
Labels: module: convolution,triaged,module: determinism

### 🐛 Describe the bug

A simple repro
```
import torch                                                                          
x = torch.randn(1, 4, 5, 5, 5, device="cuda")                                         
w = torch.randn(4, 8, 3, 3, 3, device="cuda")                                         
o0 = torch.ops.aten._convolution(                                                     
    x, w, None, (1, 1, 1), (0, 0, 0), (1, 1, 1), True, (0, 0, 0), 1, 0, 0, 1, 1       
)                                                                                     
o1 = torch.ops.aten._convolution(                                                     
    x, w, None, (1, 1, 1), (0, 0, 0), (1, 1, 1), True, (0, 0, 0), 1, 0, 0, 1, 1       
)                                                                                     
print((o0 - o1).abs().max())                                                          
assert(o0.equal(o1))                                                                  
torch.cuda.synchronize()
```
Here's the output.
```
root@353826729ef6:/raid/playground# python conv_repro.py 
tensor(3.8147e-06, device='cuda:0')
Traceback (most recent call last):
  File "conv_repro.py", line 11, in <module>
    assert(o0.equal(o1))
AssertionError
```

My brief investigation seems to indicate that we are running native kernels (at least I don't see anything cudnn in the log)
I'm guessing there's some reduction in arbitrary order that's causing this. Opening the issue so we have something to track the issue.
```
 Time (%)  Total Time (ns)  Instances  Avg (ns)  Med (ns)  Min (ns)  Max (ns)  StdDev (ns)                                                  Name
 --------  ---------------  ---------  --------  --------  --------  --------  -----------  ----------------------------------------------------------------------------------------------------
      8.4            84386         29    2909.9    2880.0      2816      3360        112.4  void at::native::vectorized_elementwise_kernel<4, at::native::MulFunctor<float>, at::detail::Array<…
      6.7            67522         23    2935.7    2880.0      2816      3905        218.8  void at::native::vectorized_elementwise_kernel<4, at::native::AddFunctor<float, float>, at::detail:…
      6.7            67170         18    3731.7    3680.0      3552      4576        225.5  void at::native::unrolled_elementwise_kernel<at::native::MulFunctor<float>, at::detail::Array<char*…
      5.3            53281         18    2960.1    2880.0      2848      3488        187.3  void at::native::vectorized_elementwise_kernel<4, at::native::AbsFunctor<float>, at::detail::Array<…
      5.2            52578         18    2921.0    2896.0      2816      3456        146.3  void at::native::vectorized_elementwise_kernel<4, at::native::BUnaryFunctor<at::native::CompareNEFu…
      5.1            51010         17    3000.6    2944.0      2848      3904        239.6  void at::native::vectorized_elementwise_kernel<4, at::native::BUnaryFunctor<at::native::CompareGTFu…
      4.8            47871         10    4787.1    4496.0      3968      7712       1077.8  void at::native::reduce_kernel<512, 1, at::native::ReduceOp<float, at::native::func_wrapper_t<float…
      4.4            43970          9    4885.6    4353.0      4351      8480       1355.7  void at::native::reduce_kernel<512, 1, at::native::ReduceOp<double, at::native::func_wrapper_t<doub…
      4.1            41056          9    4561.8    4416.0      4320      5663        423.9  void at::native::reduce_kernel<512, 1, at::native::ReduceOp<double, at::native::func_wrapper_t<doub…
      3.7            36928          9    4103.1    3936.0      3904      5248        432.5  void at::native::index_elementwise_kernel<128, 4, at::native::gpu_index_kernel<at::native::index_ke…
      3.2            32159          9    3573.2    3360.0      3360      4512        378.6  void at::cuda::detail::cub::DeviceSelectSweepKernel<at::cuda::detail::cub::DispatchSelectIf<at::cud…
      3.2            32094          9    3566.0    3359.0      3232      5087        589.1  void at::cuda::detail::cub::DeviceReduceSingleTileKernel<at::cuda::detail::cub::DeviceReducePolicy<…
      3.2            31778          9    3530.9    3328.0      3296      4928        529.6  void at::native::unrolled_elementwise_kernel<at::native::copy_device_to_device(at::TensorIterator&,…
      3.0            29825          9    3313.9    3168.0      3136      4576        473.8  void at::native::vectorized_elementwise_kernel<4, at::native::DivFunctor<double>, at::detail::Array…
      2.9            28705          9    3189.4    2977.0      2944      4256        439.3  void at::native::vectorized_elementwise_kernel<4, at::native::MulFunctor<bool>, at::detail::Array<c…
      2.9            28703          9    3189.2    3040.0      2976      4192        388.9  void at::native::vectorized_elementwise_kernel<4, at::native::BitwiseAndFunctor<bool>, at::detail::…
      2.8            28065          9    3118.3    2944.0      2912      3968        352.7  void at::native::vectorized_elementwise_kernel<4, at::native::CompareEqFunctor<float>, at::detail::…
      2.8            27809          5    5561.8    4705.0      4384      8640       1765.0  void at::native::reduce_kernel<256, 2, at::native::ReduceOp<float, at::native::func_wrapper_t<float…
      2.7            26976          9    2997.3    2944.0      2912      3456        174.5  void at::native::vectorized_elementwise_kernel<4, at::native::CompareNEFunctor<float>, at::detail::…
      2.7            26784          9    2976.0    2881.0      2879      3648        252.5  void at::native::vectorized_elementwise_kernel<4, at::native::ceil_kernel_cuda(at::TensorIteratorBa…
      2.6            26337          9    2926.3    2880.0      2816      3296        146.7  void at::native::vectorized_elementwise_kernel<4, at::native::MulScalarFunctor<float, float>, at::d…
      2.5            24767          9    2751.9    2688.0      2656      2976        102.4  void at::cuda::detail::cub::DeviceCompactInitKernel<at::cuda::detail::cub::ScanTileState<int, true>…
      2.5            24735          8    3091.9    2976.0      2944      3712        259.4  void at::native::vectorized_elementwise_kernel<4, at::native::BUnaryFunctor<at::native::CompareLTFu…
      1.6            15968          4    3992.0    3776.0      3616      4800        544.5  void at::native::unrolled_elementwise_kernel<at::native::AddFunctor<float, float>, at::detail::Arra…
      1.5            15230          5    3046.0    2848.0      2815      3840        445.6  void at::native::vectorized_elementwise_kernel<4, at::native::neg_kernel_cuda(at::TensorIteratorBas…
      0.9             8961          3    2987.0    2848.0      2689      3424        386.7  void at::native::vectorized_elementwise_kernel<4, at::native::FillFunctor<float>, at::detail::Array…
      0.8             8256          1    8256.0    8256.0      8256      8256          0.0  void at::native::(anonymous namespace)::RowwiseMomentsCUDAKernel<float>(long, float, float const*, …
      0.7             7360          1    7360.0    7360.0      7360      7360          0.0  void at::native::reduce_kernel<512, 1, at::native::ReduceOp<float, at::native::MeanOps<float, float…
      0.7             7168          1    7168.0    7168.0      7168      7168          0.0  _ZN2at6native77_GLOBAL__N__53_tmpxft_0000138f_00000000_7_DistributionNormal_cpp1_ii_7d80e40543distr…
      0.6             5600          1    5600.0    5600.0      5600      5600          0.0  _ZN2at6native78_GLOBAL__N__54_tmpxft_000019b5_00000000_7_DistributionUniform_cpp1_ii_f2fea07d43dist…
      0.5             4576          1    4576.0    4576.0      4576      4576          0.0  void at::native::unrolled_elementwise_kernel<at::native::MulScalarFunctor<float, float>, at::detail…
      0.4             4288          1    4288.0    4288.0      4288      4288          0.0  void at::native::vectorized_elementwise_kernel<4, at::native::reciprocal_kernel_cuda(at::TensorIter…
      0.3             3360          1    3360.0    3360.0      3360      3360          0.0  void at::native::vectorized_elementwise_kernel<4, at::native::(anonymous namespace)::clamp_min_scal…
      0.3             3328          1    3328.0    3328.0      3328      3328          0.0  void at::native::(anonymous namespace)::LayerNormForwardCUDAKernel<float>(long, float const*, float…
      0.3             3296          1    3296.0    3296.0      3296      3296          0.0  void at::native::vectorized_elementwise_kernel<4, at::native::AddScalarFunctor<float, float, 2>, at…
```

### Versions

```
Collecting environment information...
PyTorch version: 1.11.0a0+17540c5
Is debug build: True
CUDA used to build PyTorch: 11.6
ROCM used to build PyTorch: N/A

OS: Ubuntu 20.04.3 LTS (x86_64)
GCC version: (Ubuntu 9.3.0-17ubuntu1~20.04) 9.3.0
Clang version: Could not collect
CMake version: version 3.21.3
Libc version: glibc-2.31

Python version: 3.8.12 | packaged by conda-forge | (default, Oct 12 2021, 21:59:51)  [GCC 9.4.0] (64-bit runtime)
Python platform: Linux-4.15.0-122-generic-x86_64-with-glibc2.10
Is CUDA available: True
CUDA runtime version: 11.6.55
GPU models and configuration:
GPU 0: Tesla V100-DGXS-16GB
GPU 1: Tesla V100-DGXS-16GB
GPU 2: Tesla V100-DGXS-16GB
GPU 3: Tesla V100-DGXS-16GB

Nvidia driver version: 450.80.02
cuDNN version: Probably one of the following:
/usr/lib/x86_64-linux-gnu/libcudnn.so.8.3.2
/usr/lib/x86_64-linux-gnu/libcudnn_adv_infer.so.8.3.2
/usr/lib/x86_64-linux-gnu/libcudnn_adv_train.so.8.3.2
/usr/lib/x86_64-linux-gnu/libcudnn_cnn_infer.so.8.3.2
/usr/lib/x86_64-linux-gnu/libcudnn_cnn_train.so.8.3.2
/usr/lib/x86_64-linux-gnu/libcudnn_ops_infer.so.8.3.2
/usr/lib/x86_64-linux-gnu/libcudnn_ops_train.so.8.3.2
HIP runtime version: N/A
MIOpen runtime version: N/A
Is XNNPACK available: True

Versions of relevant libraries:
[pip3] mypy-extensions==0.4.3
[pip3] numpy==1.22.0
[pip3] pytorch-quantization==2.1.2
[pip3] torch==1.11.0a0+17540c5
[pip3] torchtext==0.12.0a0
[pip3] torchvision==0.12.0a0
[conda] magma-cuda110             2.5.2                         5    local
[conda] mkl                       2019.5                      281    conda-forge
[conda] mkl-include               2019.5                      281    conda-forge
[conda] mypy_extensions           0.4.3            py38h578d9bd_4    conda-forge
[conda] numpy                     1.22.0           py38h6ae9a64_0    conda-forge
[conda] pytorch-quantization      2.1.2                    pypi_0    pypi
[conda] torch                     1.11.0a0+17540c5           dev_0    <develop>
[conda] torchtext                 0.12.0a0                 pypi_0    pypi
[conda] torchvision               0.12.0a0                 pypi_0    pypi
```

cc @mruberry @kurtamohler