Fix incorrect CUDA torch.nn.Embedding result when max_norm is not None and indices are not sorted

[kurtamohler]

Sorting indices before calling thrust::unique fixes the issue.
Fixes #44792

[dr-ci]

💊 CI failures summary and remediations

As of commit c4edc84 (more details on the Dr. CI page):

---

Commit c4edc84 was recently pushed. Waiting for builds...

---

This comment was automatically generated by Dr. CI (expand for details).

Follow this link to opt-out of these comments for your Pull Requests.

Please report bugs/suggestions on the GitHub issue tracker or post in the (internal) Dr. CI Users group.

See how this bot performed.

This comment has been revised 41 times.

[dr-ci]

💊 CI failures summary and remediations

As of commit 18dbfb8ee4 (more details on the Dr. CI page):

---

✅ None of the CI failures appear to be your fault 💚

---

• 8/8 broken upstream at merge base 76c185d since Sep 23

---

🚧 8 ongoing upstream failures:

These were probably caused by upstream breakages that are not fixed yet:

• pytorch_linux_xenial_py3_6_gcc5_4_ge_config_legacy_test since Sep 23
  • 🔁 rerun
• pytorch_linux_xenial_py3_6_gcc5_4_ge_config_profiling_test since Sep 23
  • 🔁 rerun
• pytorch_windows_vs2019_py36_cpu_build since Sep 23
  • 🔁 rerun
• pytorch_windows_vs2019_py36_cuda11.0_build since Sep 23
  • 🔁 rerun
• binary_windows_libtorch_3_7_cpu_debug_build since Sep 23
  • 🔁 rerun
• pytorch_macos_10_13_py3_build since Sep 23
  • 🔁 rerun
• binary_windows_libtorch_3_7_cpu_release_build since Sep 23
  • 🔁 rerun
• pytorch_windows_vs2019_py36_cuda10.1_build since Sep 23
  • 🔁 rerun

---

This comment was automatically generated by Dr. CI (expand for details).

Follow this link to opt-out of these comments for your Pull Requests.

Please report bugs/suggestions on the GitHub issue tracker or post in the (internal) Dr. CI Users group.

See how this bot performed.

[kurtamohler]

I think this is ready. The CI failures didn't seem to be my fault, but I'm rerunning them.

[codecov]

Codecov Report

Merging #45248 into master will decrease coverage by 0.00%.
The diff coverage is n/a.

@@            Coverage Diff             @@
##           master   #45248      +/-   ##
==========================================
- Coverage   68.28%   68.27%   -0.01%     
==========================================
  Files         410      410              
  Lines       53306    53306              
==========================================
- Hits        36398    36397       -1     
- Misses      16908    16909       +1     

Impacted Files	Coverage Δ
torch/testing/_internal/expecttest.py	77.55% <0.00%> (-1.03%)	⬇️

---

Continue to review full report at Codecov.

Legend - Click here to learn more
Δ = absolute <relative> (impact), ø = not affected, ? = missing data
Powered by Codecov. Last update acca11b...c4edc84. Read the comment docs.

[mruberry]

So I'm curious about your thinking here, @kurtamohler,

• when duplicative entries are present on CUDA aren't the results just wrong? They may be nondeterministically wrong, but I thought the determinism flag was for algorithms which are correct-ish but whose values may vary (due to slight numerical instability, for example, or if there's a multiplicity of valid results).
• assuming this was a determinism issue for the moment, it's a determinism issue that only triggers when the input has a certain structure. Yet this PR suggests disallowing the behavior on all CUDA inputs, not just CUDA inputs with this property. If the determinism flag is set in cases like this should it do the extra work of reviewing the inputs' structure?

[ngimel]

@kurtamohler as the comment here says, thrust::unique does not work on unsorted unputs https://github.com/pytorch/pytorch/blob/master/aten/src/ATen/native/cuda/Embedding.cu#L350-L356, so you should sort inputs before calling thrust on them. It may be extra work (if there are no repeating indices), but otherwise there's no way to guarantee that incorrect results won't be produced.
After the inputs are sorted, the behavior will no longer be nondeterministic, so it won't be necessary to set the flag.

[kurtamohler]

Thanks @ngimel! I've made that change and wrote a test based on the repro from the issue description.

[facebook-github-bot]

@ngimel has imported this pull request. If you are a Facebook employee, you can view this diff on Phabricator.

[facebook-github-bot]

@ngimel merged this pull request in 66505b6.