[NVPTX] Set v2i16 SETCC to Expand #108969

frasercrmck · 2024-09-17T12:38:11Z

Note that this refers to the return type of SETCC. This operation is not
legal in PTX but was assumed as such because v2i16 is declared a legal
type. We were already expanding v4i8 SETCC.

The DAGCombiner would in certain circumstances try to fold an extension
of an illegal v2i1 SETCC (because v2i1 is illegal) into a "legal" v2i16
SETCC, which we wouldn't have patterns for.

llvmbot · 2024-09-17T12:38:50Z

@llvm/pr-subscribers-backend-nvptx

Author: Fraser Cormack (frasercrmck)

Changes

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 choices are 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;

Full diff: https://github.com/llvm/llvm-project/pull/108969.diff

3 Files Affected:

(modified) llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp (+2-2)
(modified) llvm/test/CodeGen/NVPTX/add-sub-128bit.ll (+2-2)
(added) llvm/test/CodeGen/NVPTX/sext-setcc.ll (+29)

diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index c5a40e4308860c..4495fdb45e74c9 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/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.
diff --git a/llvm/test/CodeGen/NVPTX/add-sub-128bit.ll b/llvm/test/CodeGen/NVPTX/add-sub-128bit.ll
index 9d451e90650df1..26eb5fb223a6ca 100644
--- a/llvm/test/CodeGen/NVPTX/add-sub-128bit.ll
+++ b/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
diff --git a/llvm/test/CodeGen/NVPTX/sext-setcc.ll b/llvm/test/CodeGen/NVPTX/sext-setcc.ll
new file mode 100644
index 00000000000000..4b56100f347f49
--- /dev/null
+++ b/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
+}

justinfargnoli · 2024-09-17T18:01:34Z

Notes:

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.

This is correct (source: Section 4.5 "Constants" of PTX ISA):

For predicate-type data and instructions, integer constants are allowed and are interpreted as in C, i.e., zero values are False and non-zero values are True.

Compiler Explorer reproducer.

Potentially incorrect transformation occurs during DAG optimization:

Initial selection DAG: %bb.0 'sext_setcc:entry'
SelectionDAG has 14 nodes:
  t0: ch,glue = EntryToken
              t4: v1i32,ch = load<(dereferenceable invariant load (s32) from `ptr addrspace(101) null`, addrspace 101)> t0, TargetExternalSymbol:i32'sext_setcc_param_0', undef:i32
            t5: i32 = extract_vector_elt t4, Constant:i32<0>
          t6: v2i16,ch = load<(load (s32) from %ir.p)> t0, t5, undef:i32
          t8: v2i16 = BUILD_VECTOR Constant:i16<0>, Constant:i16<0>
        t10: v2i1 = setcc t6, t8, seteq:ch
      t11: v2i16 = sign_extend t10
    t12: ch = NVPTXISD::StoreRetval<(store (s32), align 1)> t0, Constant:i32<0>, t11
  t13: ch = NVPTXISD::RET_GLUE t12



Optimized lowered selection DAG: %bb.0 'sext_setcc:entry'
SelectionDAG has 13 nodes:
  t0: ch,glue = EntryToken
            t4: v1i32,ch = load<(dereferenceable invariant load (s32) from `ptr addrspace(101) null`, addrspace 101)> t0, TargetExternalSymbol:i32'sext_setcc_param_0', undef:i32
          t5: i32 = extract_vector_elt t4, Constant:i32<0>
        t6: v2i16,ch = load<(load (s32) from %ir.p)> t0, t5, undef:i32
        t8: v2i16 = BUILD_VECTOR Constant:i16<0>, Constant:i16<0>
      t14: v2i16 = setcc t6, t8, seteq:ch
    t12: ch = NVPTXISD::StoreRetval<(store (s32), align 1)> t0, Constant:i32<0>, t14
  t13: ch = NVPTXISD::RET_GLUE t12

justinfargnoli · 2024-09-17T18:05:48Z

Where in the codebase is the sext(setcc) -> setcc DAG combine performed?

I'm curious to see how this change will disable the transformation.

justinfargnoli

Although I'm not opposed to this change for its own sake, I'm not confident it's the correct fix for this bug.

Based on my understanding, the root cause of this issue is that v2i16 setcc isn't being recognized as an illegal operation and then legalized to v2i16 sext(v2i1 setcc).

If that's correct, then I think we'd want to ensure that v2i16 setcc is correctly legalized instead of trying to ensure that the DAG combiner never generates a v2i16 setcc.

topperc · 2024-09-17T20:00:12Z

I'll take a look at this. AVX512 on X86 is similar so let me see how we handled this.

topperc · 2024-09-17T20:11:22Z

So it looks like DAGCombiner sees that getSetCCResultType returns v2i1, but then sees that v2i1 type isn't legal.

Had it not been extended it looks like v2i1 setccs get scalarized by type legalization. Maybe you just need to mark v2i16 SETCC has Expand so LegalizeVectorOps will scalarize it?

topperc · 2024-09-17T20:15:57Z

So it looks like DAGCombiner sees that getSetCCResultType returns v2i1, but then sees that v2i1 type isn't legal.

Had it not been extended it looks like v2i1 setccs get scalarized by type legalization. Maybe you just need to mark v2i16 SETCC has Expand so LegalizeVectorOps will scalarize it?

Actually just marking it Expand stops the DAGCombiner which checks

    if (N0.hasOneUse() && TLI.isOperationLegalOrCustom(ISD::SETCC, VT) &&        
        !TLI.isOperationLegalOrCustom(ISD::SETCC, SVT)) {

where VT is v2i16 and SVT is v2i1.

Note that this refers to the return type of SETCC. This operation is not legal in PTX but was assumed as such because v2i16 is declared a legal type. We were already expanding v4i8 SETCC. The DAGCombiner would in certain circumstances try to fold an extension of an illegal v2i1 SETCC (because v2i1 is illegal) into a "legal" v2i16 SETCC, which we wouldn't have patterns for.

frasercrmck · 2024-09-18T09:53:55Z

Thanks for the help, @justinfargnoli and @topperc. I agree that that's a better approach. I've updated the PR.

I suspect ZeroOrOne might still make sense for PTX but I don't think we have enough testing to justify switching that up.

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

justinfargnoli

LGTM

Note that this refers to the return type of SETCC. This operation is not legal in PTX but was assumed as such because v2i16 is declared a legal type. We were already expanding v4i8 SETCC. The DAGCombiner would in certain circumstances try to fold an extension of an illegal v2i1 SETCC (because v2i1 is illegal) into a "legal" v2i16 SETCC, which we wouldn't have patterns for.

llvmbot added the backend:NVPTX label Sep 17, 2024

frasercrmck force-pushed the fix-nvptx-sext-setcc branch from f8eecfa to 983eeb0 Compare September 17, 2024 12:39

justinfargnoli self-requested a review September 17, 2024 18:01

justinfargnoli assigned frasercrmck Sep 17, 2024

justinfargnoli reviewed Sep 17, 2024

View reviewed changes

frasercrmck force-pushed the fix-nvptx-sext-setcc branch from 983eeb0 to 0970633 Compare September 18, 2024 09:51

frasercrmck changed the title ~~[NVPTX] Set boolean contents to zero-or-one~~ [NVPTX] Set v2i16 SETCC to Expand Sep 18, 2024

justinfargnoli reviewed Sep 18, 2024

View reviewed changes

llvm/test/CodeGen/NVPTX/sext-setcc.ll Show resolved Hide resolved

justinfargnoli approved these changes Sep 18, 2024

View reviewed changes

frasercrmck merged commit 90330e9 into llvm:main Sep 19, 2024
8 checks passed

frasercrmck deleted the fix-nvptx-sext-setcc branch September 19, 2024 06:12

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

[NVPTX] Set v2i16 SETCC to Expand #108969

[NVPTX] Set v2i16 SETCC to Expand #108969

Uh oh!

frasercrmck commented Sep 17, 2024 •

edited

Loading

Uh oh!

llvmbot commented Sep 17, 2024

Uh oh!

justinfargnoli commented Sep 17, 2024 •

edited

Loading

Uh oh!

justinfargnoli commented Sep 17, 2024 •

edited

Loading

Uh oh!

justinfargnoli left a comment •

edited

Loading

Uh oh!

topperc commented Sep 17, 2024

Uh oh!

topperc commented Sep 17, 2024

Uh oh!

topperc commented Sep 17, 2024

Uh oh!

frasercrmck commented Sep 18, 2024

Uh oh!

Uh oh!

justinfargnoli left a comment

Uh oh!

Uh oh!

Uh oh!

[NVPTX] Set v2i16 SETCC to Expand #108969

[NVPTX] Set v2i16 SETCC to Expand #108969

Uh oh!

Conversation

frasercrmck commented Sep 17, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

llvmbot commented Sep 17, 2024

Uh oh!

justinfargnoli commented Sep 17, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

justinfargnoli commented Sep 17, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

justinfargnoli left a comment • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

topperc commented Sep 17, 2024

Uh oh!

topperc commented Sep 17, 2024

Uh oh!

topperc commented Sep 17, 2024

Uh oh!

frasercrmck commented Sep 18, 2024

Uh oh!

Uh oh!

justinfargnoli left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

frasercrmck commented Sep 17, 2024 •

edited

Loading

justinfargnoli commented Sep 17, 2024 •

edited

Loading

justinfargnoli commented Sep 17, 2024 •

edited

Loading

justinfargnoli left a comment •

edited

Loading