Masked LSTM and GRU using cuDNN 8.1+ give randomly corrupted results and crashes, even during inference.
@gbaned Could I humbly ask you to assign some CUDA GPU developer to review this bug? It has been causing silent corruption in masked LSTM and GRU code on GPU since TensorFlow 2.5 (the first TF to use cuDNN 8.1 or newer), and it can be replicated. @sachinprasadhs,
I was able to reproduce the issue on tensorflow v2.11, [v2.12](https://colab.research.google.com/gist/tilakrayal/eaed5692e296bf592bc20505e7d4a031/lstm_gru_cudnn_bug.ipynb) and nightly. Kindly find the gist of it [here](https://colab.research.google.com/gist/tilakrayal/2fd174e9d58ca52f04b969f29f0014c3/lstm_gru_cudnn_bug_nightly.ipynb). Hi, Thanks for your detailed input and pointers, that will certainly help us to investigate further.
Before moving forward, could you please confirm if you are observing the same behavior with the below configurations using TensorFlow 2.12. 
We had observed the similar behavior with CUDA 11.2 and it was not appearing with the CUDA 11.8 version. Reference issue https://github.com/tensorflow/tensorflow/issues/59671#issuecomment-1471169393


<h4 id="gpu" data-text="GPU" role="presentation" style="box-sizing: inherit; margin: 32px 0px 16px; font-style: normal; font-variant-ligatures: normal; font-variant-caps: normal; font-variant-numeric: ; font-variant-east-asian: ; font-variant-alternates: ; font-weight: ; font-stretch: ; font-size: 16px; line-height: ; font-family: Roboto, &quot;Noto Sans&quot;, &quot;Noto Sans JP&quot;, &quot;Noto Sans KR&quot;, &quot;Noto Naskh Arabic&quot;, &quot;Noto Sans Thai&quot;, &quot;Noto Sans Hebrew&quot;, &quot;Noto Sans Bengali&quot;, sans-serif; font-optical-sizing: ; font-kerning: ; font-feature-settings: ; font-variation-settings: ; letter-spacing: normal; overflow: hidden; text-overflow: ellipsis; margin-inline-end: -40px; padding-inline-end: 40px; color: rgb(32, 33, 36); orphans: 2; text-align: start; text-indent: 0px; text-transform: none; white-space: normal; widows: 2; word-spacing: 0px; -webkit-text-stroke-width: 0px; background-color: rgb(255, 255, 255); text-decoration-thickness: initial; text-decoration-style: initial; text-decoration-color: initial;"><span class="devsite-heading" role="heading" aria-level="4" style="box-sizing: inherit;">GPU</span><button type="button" class="devsite-heading-link button-flat material-icons" aria-label="Copy link to this section: GPU" data-title="Copy link to this section: GPU" data-id="gpu" style="box-sizing: border-box; appearance: none; background: 0px center; border: 0px; border-radius: var(--devsite-button-border-radius,2px); box-shadow: none; color: var(--devsite-icon-color,var(--devsite-secondary-text-color)); cursor: pointer; display: inline-block; font-style: normal; font-variant-ligatures: ; font-variant-caps: ; font-variant-numeric: ; font-variant-east-asian: ; font-variant-alternates: ; font-weight: normal; font-stretch: ; font-size: 24px; font-family: &quot;Material Icons&quot;; font-optical-sizing: ; font-kerning: ; font-feature-settings: &quot;liga&quot;; font-variation-settings: ; height: 24px; letter-spacing: normal; line-height: 1; margin: var(--devsite-button-margin,0); margin-inline-end: var(--devsite-button-margin-x-end); max-width: var(--devsite-button-max-width,none); min-width: 36px; outline: 0px; overflow: hidden; padding: 0px 8px; text-align: center; text-decoration: none; text-overflow: ellipsis; text-transform: none; transition: background-color 0.2s ease 0s, border 0.2s ease 0s, box-shadow 0.2s ease 0s; vertical-align: bottom; white-space: nowrap; width: var(--devsite-button-width,auto); overflow-wrap: normal; direction: ltr; -webkit-font-smoothing: antialiased; opacity: 0;"></button></h4><div class="devsite-table-wrapper" style="box-sizing: inherit; margin: var(--devsite-table-margin,16px 0); padding: 0px; overflow: auto; color: rgb(32, 33, 36); font-family: Roboto, &quot;Noto Sans&quot;, &quot;Noto Sans JP&quot;, &quot;Noto Sans KR&quot;, &quot;Noto Naskh Arabic&quot;, &quot;Noto Sans Thai&quot;, &quot;Noto Sans Hebrew&quot;, &quot;Noto Sans Bengali&quot;, sans-serif; font-size: 16px; font-style: normal; font-variant-ligatures: normal; font-variant-caps: normal; font-weight: 400; letter-spacing: normal; orphans: 2; text-align: start; text-indent: 0px; text-transform: none; white-space: normal; widows: 2; word-spacing: 0px; -webkit-text-stroke-width: 0px; background-color: rgb(255, 255, 255); text-decoration-thickness: initial; text-decoration-style: initial; text-decoration-color: initial;">

Version | Python version | Compiler | Build tools | cuDNN | CUDA
-- | -- | -- | -- | -- | --
tensorflow-2.12.0 | 3.8-3.11 | GCC 9.3.1 | Bazel 5.3.0 | 8.6 | 11.8

</div> Hi @sachinprasadhs,

I confirm that I am observing the same behavior (that is, corrupted results, i.e. nonzero batch differences, plus occasional crashes with `CUDA_ERROR_ILLEGAL_ADDRESS`) when using every following configuration:
- TF nightly (2.13.0.dev20230404) from PyPI, cuDNN 8.6, CUDA 11.8
- TF 2.12.0 from PyPI, cuDNN 8.6, CUDA 11.8 (listed on https://www.tensorflow.org/install/source)
- TF 2.11.0 from PyPI, cuDNN 8.1, CUDA 11.2 (listed on https://www.tensorflow.org/install/source)
- TF 2.5.3 from PyPI, cuDNN 8.1, CUDA 11.2 (listed https://www.tensorflow.org/install/source)

On the other hand, the following configurations do not exhibit the problem:
- TF 2.4.4 from PyPI, cuDNN 8.0, CUDA 11.0
- custom build of TF 2.11.1, GCC 9, cuDNN 8.0, CUDA 11.1
- custom build of TF 2.8.0, GCC 7, cuDNN 8.0, CUDA 11.1

Note that there is quite a different codepath for RNNs in cuDNN 8.1 (new set of APIs are used), see https://github.com/tensorflow/tensorflow/blob/4f1dd6d5123f4eb6afc85fac36df09b4a8b49c83/tensorflow/compiler/xla/stream_executor/cuda/cuda_dnn.cc#L1971-L1975 so my guess is that the problem is somewhere there (while it is not in the codepath for cuDNN < 8.1).

I believe the main problem are actually the corrupted results -- even if there is no crash, the outputs are wrong for some batches. To me, the occasional crash is just an indication that the implementation is actually wrong (and this sometimes manifests in a crash, showing that there is some memory indexing error, which would be consistent with randomly corrupted results).

In the issue you mention, no RNNs are being used, so I believe these are two independent issues (because I believe the crash here is caused by the code calling the cuDNN RNN). @kaixih can you look into this? @kaixih A gentle remainder after two weeks -- I would be great if you could find the time to look into the issue (as the original code author) :pray:  Just ran the script on A6000 with the TF2.12 + cuda 12.1 + cudnn 8.9. But I couldn't see the crash even after I use 1000 steps instead of 100 steps. Need to find to try other containers with different settings. @kaixih Thanks a lot for looking in to it!

- Did you observe the changing results (i.e., non-zero max difference)? I believe it is even more serious than the crashes, because for identical batches the results can vary a lot (causing incorrect outputs of CuDNN LSTM/GRU on all PyPI versions of TF 2.5-TF 2.12).
  - My guess is that once we find out the reason for the differing results, it will solve the crashing too. My guess is that the differing results are caused by some memory overwriting, because I have observed the different results only when there are at least _two_ RNNs computed in parallel (for example a Bidirectional layer, which runs two RNNs in the same inputs).
- When I was creating the demonstration example, I found out that the larger the dimensionalities of the LSTM/GRU, the larger chance for difference and the larger change for a crash -- so maybe increasing 2048 to 8192 could make the crash happen.
  - On GeForce 1080, dimensionalities like 256 were enough to trigger different results; on A40, it had to be increased to thousands (I kept 2048 as it was enough to trigger the problem on Colab).

Cheers! Yes, I also observed the changing results. If I use the `CUDA_LAUNCH_BLOCKING=1`, the results will be all the same. I think it is a bug. Let me dig into it to confirm if this is a cuDNN issue or TF issue. I just want to add that I can observe similar behavior when using TCN code running under Tensorflow 2.12, CuDNN 8.8, CUDA 11.8 on a GV100 GPU. Our code also uses padding to adjust the sequence length for variable length time series data. Padded sequences are masked out with a masking layer. Our pipeline runs Bayesian optimization with keras tuner to find the optimal hyperparameters. We get the CUDA_ERROR_ILLEGAL_ADDRESS error somewhat randomly, it completes a few trials of BO and then crashes in the middle of an epoch after training for some batches.

the same code used to run on the same server with older software, Tensorflow 2.3 and CUDA 10.2 without these crashes. A gentle reminder after two weeks -- thanks for looking into this :pray:  :+1:  I hope this issue will be resolved soon.
The main reason is that Python 3.8 EOL is coming next year, and when Python 3.8 goes EOL, I will need to move to Python 3.9+, which will also require moving to TF2.5+, which will be affected by this issue.

Due to the EOL of Python 3.7, I considered moving to Python 3.10+, but this issue forced me to move to Python 3.8. However, only a year later I need to move to Python 3.9+. Therefore, I would appreciate it if you could resolve this issue as soon as possible. Maybe there are others in the same situation as me.

Thank you. A gentle reminder after two weeks -- thanks for looking into this :bow:  Can you try the latest cuDNN like 8.9.2: https://docs.nvidia.com/deeplearning/cudnn/release-notes/index.html?

Basically, we guess the root cause is because the TF uses the async memcopy to prepare the `devSeqLengths` and when the copy actually happens, the host side memory might already be freed, causing the corrupted content in `devSeqLengths`. So, we have some workaround from the cuDNN side to not rely on this behavior. So, can you give a shot of this cudnn version and see if the issue is gone?

 Hi, thanks a lot for all your work! It is great that just an update of cuDNN will be able to fix this.

I verified that the example from this issue (forward pass only) no longer produces any differences (and have never crashed so far) when using:
- TF 2.11, CUDA 11.2, cuDNN 8.9.1 and cuDNN 8.9.2
- TF 2.12, CUDA 11.8, cuDNN 8.9.1 and cuDNN 8.9.2

I used both RTX A4000 with cc 8.6 and Quadro P5000 with cc 6.1 for the tests.

In a few days I will test also our initial workload of training a RNN seq2seq model, which consists of a complex computation graph with quite a few RNN layers, and will report here. Leaving open until then.

BTW, in the release notes for 8.9.1, it is said that
> _Starting in cuDNN 8.9.1, the const int32_t devSeqLengths[] argument in cudnnRNNForward(), cudnnRNNBackwardData_v8(), and cudnnRNNBackwardWeights_v8() APIs will be ignored. All three functions will source variable sequence length arrays from RNN data descriptors, configured through the seqLengthArray parameter of cudnnSetRNNDataDescriptor(). The user does not need to transfer this array to device memory; the operation will be performed automatically by RNN APIs. This refinement simplifies the usage of cuDNN RNN APIs. It is also a workaround for random crashes in multi-GPU RNN training on TensorFlow._

Note that the problem was not only in multi-GPU training, but also in single-GPU training and also single-GPU inference, and the problem was not just the crashes, but also "just" silently worse results (in a few tasks we trained models for ~10 hours and performed inference on million of tokens, and we did not experience any crash, just bad results -- during inference, some of the batches produced corrupted results, while other worked fine).

:pray: Lastly, out of curiosity, is the TF implementation incorrect and you just decided to fix it on the cuDNN side, or is the problem not as simple as saying "there is a bug on this line in the TF implementation"? Was this affecting also for example PyTorch, or TF only? Yes, the quoted lines are mainly for resolving another separate issue reported to our cudnn team, which was only reproducible in the multi-GPU settings. I happened to notice it and thought the root cause might be similar to this one. From what we have seen so far, the root cause should be the corrupted `devSeqLengths` on the device side, so depending on its content, it might cause wrong results or even a crash.

At this moment, we are not sure if it is TF implementation error or not. I would argue that we need to be very careful to correctly use the cudnn rnn APIs (of course before 8.9.1). For example, currently our guess is that in TF, under some circumstances when we copy the host data to prepare the device array `devSeqLengths`, the host memory might be already freed since it is an async copy. To mitigate such usage overhead, cudnn decides to help users to prepare the `devSeqLengths` so that no explicit `devSeqLengths` is needed anymore. Not very familiar with how the pytorch calls the cudnn. But if it can guarantee the correct devSeqLengths, it should be fine. Hope this explains the backstory. Thanks a lot for the detailed response! I initially thought that the fix in 8.9.1 was motivated by this issue, but it seems it was completely independent ;-) I will report back when I try the full training, but I assume it will work fine given that the forward pass seem to be fixed with 8.9.1. This issue is stale because it has been open for 7 days with no activity. It will be closed if no further activity occurs. Thank you. Sorry for the delay, something else showed up, I will send the experiment results in a few days. I have trained a lemmatizer architecture with a three-layer character LSTM encoder, two-layer sentence-level word LSTM contextualizer, and three-layer character LSTM decoder, and with CUDA 11.8 and cuDNN <8.9.1, it either crashes or gives suboptimal results, and with cuDNN 8.9.2 the results are correct,i.e., they are virtually identical to using cuDNN 8.0.5 (which uses a different code path to cuDNN LSTM).

So I think we can close the issue :tada: Are you satisfied with the resolution of your issue?
<a href="https://docs.google.com/forms/d/e/1FAIpQLSfaP12TRhd9xSxjXZjcZFNXPGk4kc1-qMdv3gc6bEP90vY1ew/viewform?entry.85265664=Yes&entry.2137816233=https://github.com/tensorflow/tensorflow/issues/60192">Yes</a>
<a href="https://docs.google.com/forms/d/e/1FAIpQLSfaP12TRhd9xSxjXZjcZFNXPGk4kc1-qMdv3gc6bEP90vY1ew/viewform?entry.85265664=No&entry.2137816233=https://github.com/tensorflow/tensorflow/issues/60192">No</a>
 Thanks @kaixih for the help :muscle:  > On the other hand, the following configurations do not exhibit the problem:
> 
>     * TF 2.4.4 from PyPI, cuDNN 8.0, CUDA 11.0

I have used exactly this combination for my gtx1660ti but there werent deterministic results for LSTM training, even used this code:
```
os.environ['TF_DETERMINISTIC_OPS'] = '1'
os.environ['TF_CUDNN_DETERMINISM'] = '1'
```

Now I put this debugging method:
`os.environ['CUDA_LAUNCH_BLOCKING'] = '1'`
and the results begin deterministic. The only problem is that my gpu usage fall to 30%.... any ideas? ### Problem solved: Determinism LSTM bidirectional with Tensorflow-GPU:

**My config:**

> cuda_11.0.3_451.82_win10
> cudnn-11.0-windows-x64-v8.0.5.39
> tensorflow_gpu-2.4.4-cp38-cp38-win_amd64.whl
> python-3.8.7-amd64
> nvidia driver 536.23
> gpu 1660ti
> i5 3570k (base clock)

**My code:**

```
import os
import numpy as np
import tensorflow as tf
import random

os.environ['TF_DETERMINISTIC_OPS'] = '1'
os.environ['TF_CUDNN_DETERMINISM'] = '1'
os.environ["CUBLAS_WORKSPACE_CONFIG"] = ":4096:8"
tf.random.set_global_generator(tf.random.Generator.from_seed(1))

np.random.seed(1)
tf.random.set_seed(1)
random.seed(1)`
```

**Cublas vs Cuda blocking compared to Non-deterministism standard model:**

Standard LSTM bidirectional non-deterministic model

> CPU usage: 15%
> GPU usage: 100%
> Time: 100% (fastest)

`os.environ['CUDA_LAUNCH_BLOCKING'] = '1'`

> CPU usage: 50%
> GPU usage: 30%
> Time: 25%

`os.environ["CUBLAS_WORKSPACE_CONFIG"] = ":4096:8"`

> CPU usage: 50%
> GPU usage: 70%
> Time: 80%

_Interesting: The results using Cuda block or Cublas cfg are the same_
_Important: If you are using a model with Cublas then starting any cpu intensive task will slow down the training and will not return to the original speed._
<details><summary>Click to expand!</summary> 
 
 ### Issue Type

Bug

### Have you reproduced the bug with TF nightly?

Yes

### Source

binary

### Tensorflow Version

TF 2.12, TF 2.11, TF nightly, TF 2.8

### Custom Code

Yes

### OS Platform and Distribution

Linux Ubuntu 20.04, Colab

### Mobile device

_No response_

### Python version

3.9, 3.10

### Bazel version

_No response_

### GCC/Compiler version

_No response_

### CUDA/cuDNN version

cuDNN 8.1 or newer

### GPU model and memory

I have observed the issue on GeForce 1080 Ti, GeForce 3090 RTX, A40, T4 from Collab.

### Current Behaviour?

When LSTM or GRU with mask use cuDNN 8.1+ implementation, it randomly give corrupted results (even during inference), and sometimes even crash with error `CUDA_ERROR_ILLEGAL_ADDRESS`.

It has taken me quite some time to be able to reproduce the problem, but I have found it (it manifests reliably on Colab).

Important comments:

- Of course, the problem manifests only when a GPU is available (the CPU implementation is fine).

- The bug manifests only with cuDNN 8.1+, because older cuDNN use different RNN methods. I even recompiled TF 2.11 and TF 2.8 with CUDA 11.1 and cuDNN 8.0, and they work fine.

- The bug manifests only when a `mask` is passed to the GRU and LSTM (the masked RNN calls also use a specific code path).

- The problem is not data-dependent, the same batch sometimes does and sometimes does not trigger the bug (I use the same batch in the example below).

- Sometimes the RNN call end with a `CUDA_ERROR_ILLEGAL_ADDRESS` and crash the program.

- Different GPU models differ in how frequently the corruption happens -- cards with CC 6.1 seem to trigger it more often; cards with CC 8.6 seem to trigger it less, but they still do

- When the dimensionalities of the cells are larger, the problem manifests more often.

- Note that before the recent Colab update, it used cuDNN 8.0.[56] for a very long time (I assume specific Colab packages were being build), so the bug was not manifesting there; but now it does.

My initial guess is that the different RNN calls somehow share memory and sometimes overwrite it (and maybe sometimes is the memory freed from one place and being accessed from the other place, causing the crash).

### Standalone code to reproduce the issue

The Colab notebook showing the problem with TF 2.12.0: https://colab.research.google.com/drive/17a4AcbGf9CyCl4de_vPEbB3QlTwxlV1b?usp=sharing

The Colab notebook showing the problem with TF nightly https://colab.research.google.com/drive/1ONQ7EBF9iLkSmmJE3yb04nSNbhW9fnlV?usp=sharing

The code triggering the problem:

```python
import numpy as np
import tensorflow as tf
print(tf.__version__)

def create_model(use_mask: bool) -> tf.keras.Model:
    inputs = tf.keras.layers.Input(shape=[None], dtype=tf.int32)
    if use_mask:
        mask = inputs >= 0
    else:
        mask = None
    h = tf.keras.layers.Embedding(64, 2048)(tf.math.maximum(inputs, 0))
    h1 = tf.keras.layers.Bidirectional(tf.keras.layers.GRU(2048, return_sequences=True), merge_mode="sum")(h, mask=mask)
    h2 = tf.keras.layers.Bidirectional(tf.keras.layers.LSTM(2048, return_sequences=True), merge_mode="sum")(h, mask=mask)
    h = tf.keras.layers.Dense(10)(h1 + h2)
    return tf.keras.Model(inputs, h)

# Data
data = tf.data.Dataset.from_tensor_slices([[j if j <= i else -1 for j in range(64)] for i in range(64)]).batch(64)

# However, when masking is used, even prediction on GPU gives different result.
# It also sometimes crases with the error `CUDA_ERROR_ILLEGAL_ADDRESS`.
# The full error log is copied below.
# If `use_mask=False` is passed, no problem happens.

# Models
tf.keras.utils.set_random_seed(42)
model = create_model(use_mask=True)

# Run prediction
gold = None
for i in range(100):
    result = model.predict(data, verbose=0)
    if gold is None:
        gold = result
    print("Batch {}, max difference {}, mean difference {}".format(i, np.max(np.abs(gold - result)), np.mean(np.abs(gold - result))))
```


### Relevant log output

```shell
Batch 0, max difference 0.0, mean difference 0.0
Batch 1, max difference 0.0, mean difference 0.0
Batch 2, max difference 0.0, mean difference 0.0
Batch 3, max difference 0.0, mean difference 0.0
Batch 4, max difference 0.0, mean difference 0.0
Batch 5, max difference 0.0, mean difference 0.0
Batch 6, max difference 1.6253925561904907, mean difference 0.032524894922971725
Batch 7, max difference 1.5961412191390991, mean difference 0.034154243767261505
Batch 8, max difference 0.0, mean difference 0.0
Batch 9, max difference 0.13439376652240753, mean difference 0.009293164126574993
Batch 10, max difference 0.0, mean difference 0.0
Batch 11, max difference 0.0, mean difference 0.0
Batch 12, max difference 0.0, mean difference 0.0
Batch 13, max difference 0.0, mean difference 0.0
Batch 14, max difference 0.0, mean difference 0.0
Batch 15, max difference 0.0, mean difference 0.0
Batch 16, max difference 0.0, mean difference 0.0
Batch 17, max difference 0.0, mean difference 0.0
Batch 18, max difference 0.0, mean difference 0.0
Batch 19, max difference 0.0, mean difference 0.0
Batch 20, max difference 0.0900668054819107, mean difference 0.006251291837543249
Batch 21, max difference 1.516040325164795, mean difference 0.03404051065444946
Batch 22, max difference 0.0, mean difference 0.0
Batch 23, max difference 0.09510315954685211, mean difference 0.004909028299152851
Batch 24, max difference 0.0, mean difference 0.0
Batch 25, max difference 0.0, mean difference 0.0
Batch 26, max difference 0.0, mean difference 0.0
Batch 27, max difference 0.10816850513219833, mean difference 0.006311381701380014
Batch 28, max difference 0.06991208344697952, mean difference 0.004763560835272074
Batch 29, max difference 0.0, mean difference 0.0
Batch 30, max difference 0.0, mean difference 0.0
Batch 31, max difference 0.12369412928819656, mean difference 0.0054976968094706535
Batch 32, max difference 1.114426612854004, mean difference 0.01574256643652916
Batch 33, max difference 1.1078299283981323, mean difference 0.01561223715543747
Batch 34, max difference 0.0, mean difference 0.0
Batch 35, max difference 0.07631748914718628, mean difference 0.004826296120882034
Batch 36, max difference 0.10820074379444122, mean difference 0.006503588054329157
Batch 37, max difference 0.09048224240541458, mean difference 0.006238402798771858
Batch 38, max difference 0.0, mean difference 0.0
Batch 39, max difference 0.0, mean difference 0.0
Batch 40, max difference 0.1100485697388649, mean difference 0.014095092192292213
Batch 41, max difference 0.007873136550188065, mean difference 0.0007940558716654778
Batch 42, max difference 0.007873136550188065, mean difference 0.0007940558716654778
Batch 43, max difference 0.007873136550188065, mean difference 0.0007940558716654778
Batch 44, max difference 0.06849975883960724, mean difference 0.004877315368503332
Batch 45, max difference 0.007873136550188065, mean difference 0.0007940558716654778
Batch 46, max difference 0.007873136550188065, mean difference 0.0007940558716654778
Batch 47, max difference 0.09048224240541458, mean difference 0.006395397242158651
Batch 48, max difference 0.12369412928819656, mean difference 0.005698652006685734
Batch 49, max difference 0.007873136550188065, mean difference 0.0007940558716654778
Batch 50, max difference 0.007873136550188065, mean difference 0.0007940558716654778
Batch 51, max difference 1.1203747987747192, mean difference 0.016059130430221558
Batch 52, max difference 0.007873136550188065, mean difference 0.0007940558716654778
Batch 53, max difference 0.007873136550188065, mean difference 0.0007940558716654778
Batch 54, max difference 0.007873136550188065, mean difference 0.0007940558716654778
Batch 55, max difference 0.007873136550188065, mean difference 0.0007940558716654778
Batch 56, max difference 0.07624031603336334, mean difference 0.0050272345542907715
Batch 57, max difference 1.6275964975357056, mean difference 0.032646626234054565
Batch 58, max difference 0.043441060930490494, mean difference 0.0046346113085746765
Batch 59, max difference 0.007873136550188065, mean difference 0.0007940558716654778
Batch 60, max difference 0.08132395148277283, mean difference 0.0063618021085858345
```

Then the following error appeared and crashed the program:

```shell
Mar 31, 2023, 10:30:50 PM	WARNING	2023-03-31 20:30:50.700469: E tensorflow/compiler/xla/stream_executor/cuda/cuda_event.cc:29] Error polling for event status: failed to query event: CUDA_ERROR_ILLEGAL_ADDRESS: an illegal memory access was encountered
Mar 31, 2023, 10:30:50 PM	WARNING	2023-03-31 20:30:50.700546: F tensorflow/core/common_runtime/device/device_event_mgr.cc:223] Unexpected Event status: 1
```
</details>