| ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 |
| ; RUN: llc -mtriple=nvptx64 -verify-machineinstrs < %s | FileCheck %s |
| ; RUN: %if ptxas %{ llc -mtriple=nvptx64 -verify-machineinstrs < %s | %ptxas-verify %} |
| |
| ; Tests the following pattern: |
| ; (X & 8) != 0 --> (X & 8) >> 3 |
| |
| ; This produces incorrect code in general when boolean false is |
| ; represented as a negative one. There is however a special |
| ; case when the type has a bitsize of 1, for which the false |
| ; value will be identical regardless of the boolean representation. |
| ; Check that the optimization triggers in this case. |
| |
| define i32 @pow2_mask_cmp(i32 %x) { |
| ; CHECK-LABEL: pow2_mask_cmp( |
| ; CHECK: { |
| ; CHECK-NEXT: .reg .b16 %rs<3>; |
| ; CHECK-NEXT: .reg .b32 %r<3>; |
| ; CHECK-EMPTY: |
| ; CHECK-NEXT: // %bb.0: |
| ; CHECK-NEXT: ld.param.b8 %rs1, [pow2_mask_cmp_param_0]; |
| ; CHECK-NEXT: shr.u16 %rs2, %rs1, 3; |
| ; CHECK-NEXT: cvt.u32.u16 %r1, %rs2; |
| ; CHECK-NEXT: and.b32 %r2, %r1, 1; |
| ; CHECK-NEXT: st.param.b32 [func_retval0], %r2; |
| ; CHECK-NEXT: ret; |
| %a = and i32 %x, 8 |
| %cmp = icmp ne i32 %a, 0 |
| %r = zext i1 %cmp to i32 |
| ret i32 %r |
| } |