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

[Alex MacLean 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;
----------------
AlexMaclean wrote:

Yeah, this change is not going to be an unmitigated improvement in all cases and is highly dependent on `ptxas` behavior. I think in general this will be better when the `i1` select is using or used by other logical operations and a bit worse in some truncation cases. I think the former is much more representative of real programs and hence should be prioritized. 

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