[llvm] [NVPTX] Lower i1 select with logical ops in the general case (PR #135868)

[Artem Belevich via llvm-commits]

================
@@ -1,13 +1,29 @@
+; 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 datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
 target triple = "nvptx64-nvidia-cuda"
 
-; CHECK-LABEL: _Z3foobbbPb
 define void @_Z3foobbbPb(i1 zeroext %p1, i1 zeroext %p2, i1 zeroext %p3, ptr nocapture %output) {
+; CHECK-LABEL: _Z3foobbbPb(
+; CHECK:       {
+; CHECK-NEXT:    .reg .pred %p<2>;
+; CHECK-NEXT:    .reg .b16 %rs<7>;
+; CHECK-NEXT:    .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0: // %entry
+; CHECK-NEXT:    ld.param.u8 %rs1, [_Z3foobbbPb_param_0];
+; CHECK-NEXT:    and.b16 %rs2, %rs1, 1;
+; CHECK-NEXT:    setp.eq.b16 %p1, %rs2, 1;
+; CHECK-NEXT:    ld.param.u8 %rs3, [_Z3foobbbPb_param_1];
+; CHECK-NEXT:    ld.param.u8 %rs4, [_Z3foobbbPb_param_2];
+; CHECK-NEXT:    selp.b16 %rs5, %rs3, %rs4, %p1;
+; CHECK-NEXT:    and.b16 %rs6, %rs5, 1;
----------------
Artem-B wrote:

Yeah. Optimizing for NVIDIA GPUs requires reasoning about spooky interactions at a distance... Oh, well. 
It indeed looks like something ptxas can improve in the future.

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