[llvm] [NVPTX] Support slct for sign-based selects (PR #214860)

[Justin Fargnoli via llvm-commits]

================
@@ -0,0 +1,165 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s
+; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
+
+target triple = "nvptx64-nvidia-cuda"
+
+define i32 @slct_b32_s32(i32 %a, i32 %b, i32 %cond) {
+; CHECK-LABEL: slct_b32_s32(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b32 %r<7>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b32 %r1, [slct_b32_s32_param_0];
+; CHECK-NEXT:    add.s32 %r2, %r1, 1;
+; CHECK-NEXT:    ld.param.b32 %r3, [slct_b32_s32_param_1];
+; CHECK-NEXT:    not.b32 %r4, %r3;
+; CHECK-NEXT:    ld.param.b32 %r5, [slct_b32_s32_param_2];
+; CHECK-NEXT:    slct.b32.s32 %r6, %r2, %r4, %r5;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r6;
+; CHECK-NEXT:    ret;
+  %a1 = add i32 %a, 1
+  %b1 = xor i32 %b, -1
----------------
justinfargnoli wrote:

Why do we need these instructions in the test case? 

https://github.com/llvm/llvm-project/pull/214860