AMDGPU backend cannot select error using thrust::sort on a 2 byte type

[markdewing]

Calling (hip)thrust::sort with uint16_t results in

fatal error: error in backend: Cannot select: t32: v8i16 = extract_subvector # D:1 t10, Constant:i32<0>
  t10: v16i16,ch = CopyFromReg # D:1 t0, Register:v16i16 %77
    t9: v16i16 = Register %77
  t28: i32 = Constant<0>
In function: _ZN7rocprim6detail16transform_kernelILj256ELj16EtPtS2_NS_8identityItEEEEvT2_mT3_T4_

This is using mainline clang (81b193f) (This commit is where I updated and built the compiler, it's not the culprit)

Source code

#include "thrust/sort.h"
#include <vector>

int main()
{
    const int NumElements = 8;
    std::vector<uint16_t> B(NumElements);
    thrust::sort(thrust::device, B.data(), B.data() + NumElements);
    return 0;
}

Compile command

clang++ \
-I/soft/compilers/rocm/rocm-5.5.0/include \
-O3 \
-x hip \
-c \
t1.cpp

[llvmbot]

@llvm/issue-subscribers-backend-amdgpu

[arsenm]

https://reviews.llvm.org/D156422