[NVPTX] Set boolean contents to zero-or-one

Note that this field is a property of SelectionDAG rather than one of
the ISA. The previous value wasn't incorrect, per se, but since NVPTX
uses 1-bit predicate registers, either choice is possible.

The problem with using zero-or-negative-one manifested in issues during
instruction selection where LLVM would fold a sign-extension of a setcc
into the setcc itself, which isn't a legal operation.

We could alternatively let this happen and then custom expand it back
into a separate setcc + sext later. Or we could provide explicit
patterns for the setp/selp sequence. However, preventing it from forming
in the first place feels more appropriate to the ISA.

It does not appear to have a significant knock-on effect on codegen
tests, aside from one minor change which looks harmless:

	selp.s64 	%rd6, -1, 0, %p1;
	add.s64 	%rd7, %rd5, %rd6;

to

	selp.u64 	%rd6, 1, 0, %p1;
	sub.s64 	%rd7, %rd5, %rd6;

Author: frasercrmck

llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp

@@ -416,8 +416,8 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM,
   MaxStoresPerMemcpy = MaxStoresPerMemcpyOptSize = (unsigned) 0xFFFFFFFF;
   MaxStoresPerMemmove = MaxStoresPerMemmoveOptSize = (unsigned) 0xFFFFFFFF;
 
-  setBooleanContents(ZeroOrNegativeOneBooleanContent);
-  setBooleanVectorContents(ZeroOrNegativeOneBooleanContent);
+  setBooleanContents(ZeroOrOneBooleanContent);
+  setBooleanVectorContents(ZeroOrOneBooleanContent);
 
   // Jump is Expensive. Don't create extra control flow for 'and', 'or'
   // condition branches.

llvm/test/CodeGen/NVPTX/add-sub-128bit.ll

@@ -23,8 +23,8 @@ define i128 @test_add(i128 %a, i128 %b) {
 define i128 @test_sub(i128 %a, i128 %b) {
 ; NOCARRY:        sub.s64
 ; NOCARRY-NEXT:   setp.lt.u64
-; NOCARRY-NEXT:   selp.s64
-; NOCARRY-NEXT:   add.s64
+; NOCARRY-NEXT:   selp.u64
+; NOCARRY-NEXT:   sub.s64
 ; NOCARRY-NEXT:   sub.s64
 
 ; CARRY:          sub.cc.s64

llvm/test/CodeGen/NVPTX/sext-setcc.ll

@@ -0,0 +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_80 | FileCheck %s
+; RUN: %if ptxas-11.0 %{ llc < %s -march=nvptx64 -mcpu=sm_80 -mattr=+ptx70 | %ptxas-verify -arch=sm_80 %}
+
+define <2 x i16> @sext_setcc(ptr %p) {
+; CHECK-LABEL: sext_setcc(
+; CHECK:       {
+; CHECK-NEXT:    .reg .pred %p<3>;
+; CHECK-NEXT:    .reg .b16 %rs<5>;
+; CHECK-NEXT:    .reg .b32 %r<3>;
+; CHECK-NEXT:    .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0: // %entry
+; CHECK-NEXT:    ld.param.u64 %rd1, [sext_setcc_param_0];
+; CHECK-NEXT:    ld.u32 %r1, [%rd1];
+; CHECK-NEXT:    mov.b32 {%rs1, %rs2}, %r1;
+; CHECK-NEXT:    setp.eq.s16 %p1, %rs1, 0;
+; CHECK-NEXT:    setp.eq.s16 %p2, %rs2, 0;
+; CHECK-NEXT:    selp.s16 %rs3, -1, 0, %p2;
+; CHECK-NEXT:    selp.s16 %rs4, -1, 0, %p1;
+; CHECK-NEXT:    mov.b32 %r2, {%rs4, %rs3};
+; CHECK-NEXT:    st.param.b32 [func_retval0+0], %r2;
+; CHECK-NEXT:    ret;
+entry:
+  %v = load <2 x i16>, ptr %p, align 4
+  %cmp = icmp eq <2 x i16> %v, zeroinitializer
+  %sext = sext <2 x i1> %cmp to <2 x i16>
+  ret <2 x i16> %sext
+}