BFloat16 support for torch.sort

[zasdfgbnm]

No description provided.

[facebook-github-bot]

💊 CI failures summary and remediations

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

---

• 2/3 failures introduced in this PR
• 1/3 broken upstream at merge base 60af6e9 on May 25 from 12:31am to 8:40pm

---

2 failures not recognized by patterns:

Job	Step	Action
Linux CI (pytorch-linux-xenial-py3.6-gcc5.4) / test	Clean up docker images	🔁 rerun
docker-pytorch-linux-bionic-cuda10.2-cudnn7-py3.6-clang9	Check if image should be built	🔁 rerun

---

🚧 1 fixed upstream failure:

These were probably caused by upstream breakages that were already fixed.

Please rebase on the viable/strict branch (expand for instructions)

If your commit is older than viable/strict, run these commands:

git fetch https://github.com/pytorch/pytorch viable/strict
git rebase FETCH_HEAD

• pytorch_linux_backward_compatibility_check_test on May 25 from 12:31am to 8:40pm (b4f3a98 - 083d3bb)
  • 🔁 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 to the (internal) Dr. CI Users group.

Click here to manually regenerate this comment.

[zasdfgbnm]

@jeffdaily Is there any way to make this hack work on ROCm? Looks like in https://github.com/ROCmSoftwarePlatform/cub-hip/blob/hip_port_1.7.4/cub/util_type.cuh#L1102-L1194, there is no specialization for __half, so I think specializing for c10::BFloat16 won't work either.

[zasdfgbnm]

OK, I was looking at an obsolete repo. The new repo does have __half specialization. But still, ROCm CI is failing:

May 13 02:15:01 In file included from /var/lib/jenkins/workspace/aten/src/ATen/native/hip/Randperm.hip:6:
May 13 02:15:01 /var/lib/jenkins/workspace/aten/src/ATen/hip/cub.cuh:89:18: error: explicit specialization of non-template struct 'FpLimits'
May 13 02:15:01 struct ::hipcub::FpLimits<c10::BFloat16>
May 13 02:15:01                  ^       ~~~~~~~~~~~~~~~
May 13 02:15:01 /var/lib/jenkins/workspace/aten/src/ATen/hip/cub.cuh:89:18: error: no struct named 'FpLimits' in namespace 'hipcub'
May 13 02:15:01 struct ::hipcub::FpLimits<c10::BFloat16>
May 13 02:15:01        ~~~~~~~~~~^
May 13 02:15:01 /var/lib/jenkins/workspace/aten/src/ATen/hip/cub.cuh:89:18: error: cannot define or redeclare 'FpLimits' here because namespace 'detail' does not enclose namespace 'hipcub'
May 13 02:15:01 struct ::hipcub::FpLimits<c10::BFloat16>
May 13 02:15:01        ~~~~~~~~~~^
May 13 02:15:01 /var/lib/jenkins/workspace/aten/src/ATen/hip/cub.cuh:102:30: error: explicit specialization of non-template struct 'NumericTraits'
May 13 02:15:01 template <> struct ::hipcub::NumericTraits<c10::BFloat16>: ::hipcub::BaseTraits<::hipcub::FLOATING_POINT, true, false, unsigned short, c10::BFloat16> {};
May 13 02:15:01                              ^            ~~~~~~~~~~~~~~~
May 13 02:15:01 /var/lib/jenkins/workspace/aten/src/ATen/hip/cub.cuh:102:30: error: no struct named 'NumericTraits' in namespace 'hipcub'
May 13 02:15:01 template <> struct ::hipcub::NumericTraits<c10::BFloat16>: ::hipcub::BaseTraits<::hipcub::FLOATING_POINT, true, false, unsigned short, c10::BFloat16> {};
May 13 02:15:01                    ~~~~~~~~~~^
May 13 02:15:01 /var/lib/jenkins/workspace/aten/src/ATen/hip/cub.cuh:102:30: error: cannot define or redeclare 'NumericTraits' here because namespace 'detail' does not enclose namespace 'hipcub'
May 13 02:15:01 template <> struct ::hipcub::NumericTraits<c10::BFloat16>: ::hipcub::BaseTraits<::hipcub::FLOATING_POINT, true, false, unsigned short, c10::BFloat16> {};
May 13 02:15:01                    ~~~~~~~~~~^
May 13 02:15:01 /var/lib/jenkins/workspace/aten/src/ATen/hip/cub.cuh:102:70: error: unknown template name 'BaseTraits'
May 13 02:15:01 template <> struct ::hipcub::NumericTraits<c10::BFloat16>: ::hipcub::BaseTraits<::hipcub::FLOATING_POINT, true, false, unsigned short, c10::BFloat16> {};
May 13 02:15:01                                                                      ^
May 13 02:15:01 /var/lib/jenkins/workspace/aten/src/ATen/hip/cub.cuh:102:91: error: no member named 'FLOATING_POINT' in namespace 'hipcub'
May 13 02:15:01 template <> struct ::hipcub::NumericTraits<c10::BFloat16>: ::hipcub::BaseTraits<::hipcub::FLOATING_POINT, true, false, unsigned short, c10::BFloat16> {};
May 13 02:15:01                                                                                 ~~~~~~~~~~^
May 13 02:15:01 8 errors generated when compiling for gfx900.
May 13 02:15:01 CMake Error at torch_hip_generated_Randperm.hip.o.cmake:192 (message):
May 13 02:15:01   Error generating file
May 13 02:15:01   /var/lib/jenkins/workspace/build/caffe2/CMakeFiles/torch_hip.dir/__/aten/src/ATen/native/hip/./torch_hip_generated_Randperm.hip.o
May 13 02:15:01 

[pruthvistony]

@zasdfgbnm ,
Can you please let us know how the compilation problem was fixed for ROCm build.

[ngimel]

It's not solved, bfloat16 is not compiled for rocm.

[jithunnair-amd]

Okay, that makes sense. It's okay to skip bfloat16 radix sort for ROCm. We'll look into adding bf16 support for radix sort into hipcub/rocprim.

[jithunnair-amd]

Internal issue to track adding bf16 support for radix sort into hipcub/rocprim: SWDEV-288399

[ngimel]

Can you also please remove fillSliceWithIndex_kernel from SortingCommon.cuh? I don't think it's used anywhere.