Author: ronlieb Date: 2025-11-16T16:47:51-05:00 New Revision: b835c10c902a27d1423d8944534d828afbcb4f6c
URL: https://github.com/llvm/llvm-project/commit/b835c10c902a27d1423d8944534d828afbcb4f6c DIFF: https://github.com/llvm/llvm-project/commit/b835c10c902a27d1423d8944534d828afbcb4f6c.diff LOG: Revert "DAG: Allow select ptr combine for non-0 address spaces (#167909)" This reverts commit e5f499f48f2d1fddc590982da7232d08a6f8c54c. Added: Modified: llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp llvm/test/CodeGen/AMDGPU/load-select-ptr.ll llvm/test/CodeGen/AMDGPU/select-load-to-load-select-ptr-combine.ll llvm/test/CodeGen/AMDGPU/select-vectors.ll llvm/test/CodeGen/AMDGPU/select64.ll llvm/test/CodeGen/NVPTX/bf16-instructions.ll llvm/test/CodeGen/NVPTX/bf16x2-instructions.ll llvm/test/CodeGen/NVPTX/bug22246.ll llvm/test/CodeGen/NVPTX/fast-math.ll llvm/test/CodeGen/NVPTX/i1-select.ll llvm/test/CodeGen/NVPTX/i8x4-instructions.ll llvm/test/CodeGen/NVPTX/lower-byval-args.ll Removed: ################################################################################ diff --git a/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp b/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp index 6fbac0f8c8cdf..c9513611e6dcb 100644 --- a/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp +++ b/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp @@ -29033,9 +29033,9 @@ bool DAGCombiner::SimplifySelectOps(SDNode *TheSelect, SDValue LHS, // over-conservative. It would be beneficial to be able to remember // both potential memory locations. Since we are discarding // src value info, don't do the transformation if the memory - // locations are not in the same address space. - LLD->getPointerInfo().getAddrSpace() != - RLD->getPointerInfo().getAddrSpace() || + // locations are not in the default address space. + LLD->getPointerInfo().getAddrSpace() != 0 || + RLD->getPointerInfo().getAddrSpace() != 0 || // We can't produce a CMOV of a TargetFrameIndex since we won't // generate the address generation required. LLD->getBasePtr().getOpcode() == ISD::TargetFrameIndex || @@ -29117,9 +29117,6 @@ bool DAGCombiner::SimplifySelectOps(SDNode *TheSelect, SDValue LHS, // but the new load must be the minimum (most restrictive) alignment of the // inputs. Align Alignment = std::min(LLD->getAlign(), RLD->getAlign()); - unsigned AddrSpace = LLD->getAddressSpace(); - assert(AddrSpace == RLD->getAddressSpace()); - MachineMemOperand::Flags MMOFlags = LLD->getMemOperand()->getFlags(); if (!RLD->isInvariant()) MMOFlags &= ~MachineMemOperand::MOInvariant; @@ -29128,16 +29125,15 @@ bool DAGCombiner::SimplifySelectOps(SDNode *TheSelect, SDValue LHS, if (LLD->getExtensionType() == ISD::NON_EXTLOAD) { // FIXME: Discards pointer and AA info. Load = DAG.getLoad(TheSelect->getValueType(0), SDLoc(TheSelect), - LLD->getChain(), Addr, MachinePointerInfo(AddrSpace), - Alignment, MMOFlags); + LLD->getChain(), Addr, MachinePointerInfo(), Alignment, + MMOFlags); } else { // FIXME: Discards pointer and AA info. Load = DAG.getExtLoad( LLD->getExtensionType() == ISD::EXTLOAD ? RLD->getExtensionType() : LLD->getExtensionType(), SDLoc(TheSelect), TheSelect->getValueType(0), LLD->getChain(), Addr, - MachinePointerInfo(AddrSpace), LLD->getMemoryVT(), Alignment, - MMOFlags); + MachinePointerInfo(), LLD->getMemoryVT(), Alignment, MMOFlags); } // Users of the select now use the result of the load. diff --git a/llvm/test/CodeGen/AMDGPU/load-select-ptr.ll b/llvm/test/CodeGen/AMDGPU/load-select-ptr.ll index 5aabad682ad30..d9ad9590d9762 100644 --- a/llvm/test/CodeGen/AMDGPU/load-select-ptr.ll +++ b/llvm/test/CodeGen/AMDGPU/load-select-ptr.ll @@ -7,31 +7,27 @@ define amdgpu_kernel void @select_ptr_crash_i64_flat(i32 %tmp, [8 x i32], ptr %ptr0, [8 x i32], ptr %ptr1, [8 x i32], ptr addrspace(1) %ptr2) { ; GCN-LABEL: select_ptr_crash_i64_flat: ; GCN: ; %bb.0: +; GCN-NEXT: s_load_dword s6, s[8:9], 0x0 +; GCN-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x28 +; GCN-NEXT: s_load_dwordx2 s[2:3], s[8:9], 0x50 +; GCN-NEXT: s_load_dwordx2 s[4:5], s[8:9], 0x78 ; GCN-NEXT: s_add_i32 s12, s12, s17 ; GCN-NEXT: s_lshr_b32 flat_scratch_hi, s12, 8 -; GCN-NEXT: s_load_dword s2, s[8:9], 0x0 -; GCN-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x78 -; GCN-NEXT: s_add_u32 s4, s8, 40 -; GCN-NEXT: s_addc_u32 s3, s9, 0 -; GCN-NEXT: s_add_u32 s5, s8, 0x50 -; GCN-NEXT: s_addc_u32 s6, s9, 0 ; GCN-NEXT: s_waitcnt lgkmcnt(0) -; GCN-NEXT: s_cmp_eq_u32 s2, 0 -; GCN-NEXT: s_cselect_b32 s3, s3, s6 -; GCN-NEXT: s_cselect_b32 s2, s4, s5 -; GCN-NEXT: s_load_dwordx2 s[2:3], s[2:3], 0x0 +; GCN-NEXT: s_cmp_eq_u32 s6, 0 +; GCN-NEXT: s_cselect_b32 s0, s0, s2 +; GCN-NEXT: s_cselect_b32 s1, s1, s3 +; GCN-NEXT: v_mov_b32_e32 v0, s0 +; GCN-NEXT: v_mov_b32_e32 v1, s1 +; GCN-NEXT: s_add_u32 s0, s0, 4 ; GCN-NEXT: s_mov_b32 flat_scratch_lo, s13 -; GCN-NEXT: s_waitcnt lgkmcnt(0) -; GCN-NEXT: v_mov_b32_e32 v0, s2 -; GCN-NEXT: v_mov_b32_e32 v1, s3 -; GCN-NEXT: s_add_u32 s2, s2, 4 +; GCN-NEXT: s_addc_u32 s1, s1, 0 ; GCN-NEXT: flat_load_dword v0, v[0:1] -; GCN-NEXT: s_addc_u32 s3, s3, 0 -; GCN-NEXT: v_mov_b32_e32 v1, s2 -; GCN-NEXT: v_mov_b32_e32 v2, s3 +; GCN-NEXT: v_mov_b32_e32 v2, s1 +; GCN-NEXT: v_mov_b32_e32 v1, s0 ; GCN-NEXT: flat_load_dword v1, v[1:2] -; GCN-NEXT: v_mov_b32_e32 v3, s1 -; GCN-NEXT: v_mov_b32_e32 v2, s0 +; GCN-NEXT: v_mov_b32_e32 v2, s4 +; GCN-NEXT: v_mov_b32_e32 v3, s5 ; GCN-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0) ; GCN-NEXT: flat_store_dwordx2 v[2:3], v[0:1] ; GCN-NEXT: s_endpgm @@ -49,28 +45,25 @@ define amdgpu_kernel void @select_ptr_crash_i64_flat(i32 %tmp, [8 x i32], ptr %p define amdgpu_kernel void @select_ptr_crash_i64_global(i32 %tmp, [8 x i32], ptr addrspace(1) %ptr0, [8 x i32], ptr addrspace(1) %ptr1, [8 x i32], ptr addrspace(1) %ptr2) { ; GCN-LABEL: select_ptr_crash_i64_global: ; GCN: ; %bb.0: +; GCN-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x28 +; GCN-NEXT: s_load_dwordx2 s[2:3], s[8:9], 0x50 +; GCN-NEXT: s_load_dword s6, s[8:9], 0x0 +; GCN-NEXT: s_load_dwordx2 s[4:5], s[8:9], 0x78 ; GCN-NEXT: s_add_i32 s12, s12, s17 ; GCN-NEXT: s_lshr_b32 flat_scratch_hi, s12, 8 -; GCN-NEXT: s_load_dword s2, s[8:9], 0x0 -; GCN-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x78 -; GCN-NEXT: s_add_u32 s4, s8, 40 -; GCN-NEXT: s_addc_u32 s3, s9, 0 -; GCN-NEXT: s_add_u32 s5, s8, 0x50 -; GCN-NEXT: s_addc_u32 s6, s9, 0 ; GCN-NEXT: s_waitcnt lgkmcnt(0) -; GCN-NEXT: s_cmp_eq_u32 s2, 0 -; GCN-NEXT: s_cselect_b32 s3, s3, s6 -; GCN-NEXT: s_cselect_b32 s2, s4, s5 +; GCN-NEXT: s_load_dwordx2 s[0:1], s[0:1], 0x0 ; GCN-NEXT: s_load_dwordx2 s[2:3], s[2:3], 0x0 -; GCN-NEXT: v_mov_b32_e32 v0, s0 +; GCN-NEXT: s_cmp_eq_u32 s6, 0 +; GCN-NEXT: v_mov_b32_e32 v2, s4 ; GCN-NEXT: s_mov_b32 flat_scratch_lo, s13 -; GCN-NEXT: v_mov_b32_e32 v1, s1 -; GCN-NEXT: s_waitcnt lgkmcnt(0) -; GCN-NEXT: s_load_dwordx2 s[2:3], s[2:3], 0x0 +; GCN-NEXT: v_mov_b32_e32 v3, s5 ; GCN-NEXT: s_waitcnt lgkmcnt(0) -; GCN-NEXT: v_mov_b32_e32 v2, s2 -; GCN-NEXT: v_mov_b32_e32 v3, s3 -; GCN-NEXT: flat_store_dwordx2 v[0:1], v[2:3] +; GCN-NEXT: s_cselect_b32 s1, s1, s3 +; GCN-NEXT: s_cselect_b32 s0, s0, s2 +; GCN-NEXT: v_mov_b32_e32 v0, s0 +; GCN-NEXT: v_mov_b32_e32 v1, s1 +; GCN-NEXT: flat_store_dwordx2 v[2:3], v[0:1] ; GCN-NEXT: s_endpgm %tmp2 = icmp eq i32 %tmp, 0 %tmp3 = load i64, ptr addrspace(1) %ptr0, align 8 @@ -85,18 +78,22 @@ define amdgpu_kernel void @select_ptr_crash_i64_local(i32 %tmp, ptr addrspace(3) ; GCN: ; %bb.0: ; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x0 ; GCN-NEXT: s_load_dwordx2 s[4:5], s[8:9], 0x10 +; GCN-NEXT: s_mov_b32 m0, -1 ; GCN-NEXT: s_add_i32 s12, s12, s17 ; GCN-NEXT: s_lshr_b32 flat_scratch_hi, s12, 8 -; GCN-NEXT: s_mov_b32 m0, -1 ; GCN-NEXT: s_waitcnt lgkmcnt(0) -; GCN-NEXT: s_cmp_eq_u32 s0, 0 -; GCN-NEXT: s_cselect_b32 s0, s1, s2 -; GCN-NEXT: v_mov_b32_e32 v0, s0 +; GCN-NEXT: v_mov_b32_e32 v0, s1 +; GCN-NEXT: v_mov_b32_e32 v2, s2 ; GCN-NEXT: ds_read_b64 v[0:1], v0 -; GCN-NEXT: v_mov_b32_e32 v2, s4 +; GCN-NEXT: ds_read_b64 v[2:3], v2 +; GCN-NEXT: s_cmp_eq_u32 s0, 0 +; GCN-NEXT: s_cselect_b64 vcc, -1, 0 ; GCN-NEXT: s_mov_b32 flat_scratch_lo, s13 -; GCN-NEXT: v_mov_b32_e32 v3, s5 ; GCN-NEXT: s_waitcnt lgkmcnt(0) +; GCN-NEXT: v_cndmask_b32_e32 v1, v3, v1, vcc +; GCN-NEXT: v_cndmask_b32_e32 v0, v2, v0, vcc +; GCN-NEXT: v_mov_b32_e32 v2, s4 +; GCN-NEXT: v_mov_b32_e32 v3, s5 ; GCN-NEXT: flat_store_dwordx2 v[2:3], v[0:1] ; GCN-NEXT: s_endpgm %tmp2 = icmp eq i32 %tmp, 0 @@ -115,20 +112,22 @@ define amdgpu_kernel void @select_ptr_crash_i64_local_offsets(i32 %tmp, ptr addr ; GCN: ; %bb.0: ; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x0 ; GCN-NEXT: s_load_dwordx2 s[4:5], s[8:9], 0x10 +; GCN-NEXT: s_mov_b32 m0, -1 ; GCN-NEXT: s_add_i32 s12, s12, s17 ; GCN-NEXT: s_lshr_b32 flat_scratch_hi, s12, 8 -; GCN-NEXT: s_mov_b32 m0, -1 ; GCN-NEXT: s_waitcnt lgkmcnt(0) -; GCN-NEXT: s_addk_i32 s1, 0x80 -; GCN-NEXT: s_addk_i32 s2, 0x200 +; GCN-NEXT: v_mov_b32_e32 v0, s1 +; GCN-NEXT: v_mov_b32_e32 v2, s2 +; GCN-NEXT: ds_read_b64 v[0:1], v0 offset:128 +; GCN-NEXT: ds_read_b64 v[2:3], v2 offset:512 ; GCN-NEXT: s_cmp_eq_u32 s0, 0 -; GCN-NEXT: s_cselect_b32 s0, s1, s2 -; GCN-NEXT: v_mov_b32_e32 v0, s0 -; GCN-NEXT: ds_read_b64 v[0:1], v0 -; GCN-NEXT: v_mov_b32_e32 v2, s4 +; GCN-NEXT: s_cselect_b64 vcc, -1, 0 ; GCN-NEXT: s_mov_b32 flat_scratch_lo, s13 -; GCN-NEXT: v_mov_b32_e32 v3, s5 ; GCN-NEXT: s_waitcnt lgkmcnt(0) +; GCN-NEXT: v_cndmask_b32_e32 v1, v3, v1, vcc +; GCN-NEXT: v_cndmask_b32_e32 v0, v2, v0, vcc +; GCN-NEXT: v_mov_b32_e32 v2, s4 +; GCN-NEXT: v_mov_b32_e32 v3, s5 ; GCN-NEXT: flat_store_dwordx2 v[2:3], v[0:1] ; GCN-NEXT: s_endpgm %tmp2 = icmp eq i32 %tmp, 0 diff --git a/llvm/test/CodeGen/AMDGPU/select-load-to-load-select-ptr-combine.ll b/llvm/test/CodeGen/AMDGPU/select-load-to-load-select-ptr-combine.ll index cc5ae2717faf0..423fb7d52d3e3 100644 --- a/llvm/test/CodeGen/AMDGPU/select-load-to-load-select-ptr-combine.ll +++ b/llvm/test/CodeGen/AMDGPU/select-load-to-load-select-ptr-combine.ll @@ -22,12 +22,12 @@ define i32 @select_load_i32_p1(i1 %cond, ptr addrspace(1) %a, ptr addrspace(1) % ; CHECK-LABEL: select_load_i32_p1: ; CHECK: ; %bb.0: ; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) +; CHECK-NEXT: global_load_dword v5, v[1:2], off +; CHECK-NEXT: global_load_dword v6, v[3:4], off ; CHECK-NEXT: v_and_b32_e32 v0, 1, v0 ; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 1, v0 -; CHECK-NEXT: v_cndmask_b32_e32 v2, v4, v2, vcc -; CHECK-NEXT: v_cndmask_b32_e32 v1, v3, v1, vcc -; CHECK-NEXT: global_load_dword v0, v[1:2], off ; CHECK-NEXT: s_waitcnt vmcnt(0) +; CHECK-NEXT: v_cndmask_b32_e32 v0, v6, v5, vcc ; CHECK-NEXT: s_setpc_b64 s[30:31] %ld0 = load i32, ptr addrspace(1) %a %ld1 = load i32, ptr addrspace(1) %b @@ -39,11 +39,12 @@ define i32 @select_load_i32_p3(i1 %cond, ptr addrspace(3) %a, ptr addrspace(3) % ; CHECK-LABEL: select_load_i32_p3: ; CHECK: ; %bb.0: ; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) +; CHECK-NEXT: ds_read_b32 v1, v1 +; CHECK-NEXT: ds_read_b32 v2, v2 ; CHECK-NEXT: v_and_b32_e32 v0, 1, v0 ; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 1, v0 -; CHECK-NEXT: v_cndmask_b32_e32 v0, v2, v1, vcc -; CHECK-NEXT: ds_read_b32 v0, v0 ; CHECK-NEXT: s_waitcnt lgkmcnt(0) +; CHECK-NEXT: v_cndmask_b32_e32 v0, v2, v1, vcc ; CHECK-NEXT: s_setpc_b64 s[30:31] %ld0 = load i32, ptr addrspace(3) %a %ld1 = load i32, ptr addrspace(3) %b @@ -89,12 +90,12 @@ define i8 @select_load_i8_p1(i1 %cond, ptr addrspace(1) %a, ptr addrspace(1) %b) ; CHECK-LABEL: select_load_i8_p1: ; CHECK: ; %bb.0: ; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) +; CHECK-NEXT: global_load_ubyte v5, v[1:2], off +; CHECK-NEXT: global_load_ubyte v6, v[3:4], off ; CHECK-NEXT: v_and_b32_e32 v0, 1, v0 ; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 1, v0 -; CHECK-NEXT: v_cndmask_b32_e32 v2, v4, v2, vcc -; CHECK-NEXT: v_cndmask_b32_e32 v1, v3, v1, vcc -; CHECK-NEXT: global_load_ubyte v0, v[1:2], off ; CHECK-NEXT: s_waitcnt vmcnt(0) +; CHECK-NEXT: v_cndmask_b32_e32 v0, v6, v5, vcc ; CHECK-NEXT: s_setpc_b64 s[30:31] %ld0 = load i8, ptr addrspace(1) %a %ld1 = load i8, ptr addrspace(1) %b @@ -106,16 +107,12 @@ define i32 @select_load_i32_p1_offset(i1 %cond, ptr addrspace(1) %a, ptr addrspa ; CHECK-LABEL: select_load_i32_p1_offset: ; CHECK: ; %bb.0: ; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) -; CHECK-NEXT: v_add_co_u32_e32 v3, vcc, 0x100, v1 -; CHECK-NEXT: v_addc_co_u32_e32 v4, vcc, 0, v2, vcc -; CHECK-NEXT: v_add_co_u32_e32 v5, vcc, 0x200, v1 +; CHECK-NEXT: global_load_dword v3, v[1:2], off offset:256 +; CHECK-NEXT: global_load_dword v4, v[1:2], off offset:512 ; CHECK-NEXT: v_and_b32_e32 v0, 1, v0 -; CHECK-NEXT: v_addc_co_u32_e32 v1, vcc, 0, v2, vcc ; CHECK-NEXT: v_cmp_eq_u32_e32 vcc, 1, v0 -; CHECK-NEXT: v_cndmask_b32_e32 v1, v1, v4, vcc -; CHECK-NEXT: v_cndmask_b32_e32 v0, v5, v3, vcc -; CHECK-NEXT: global_load_dword v0, v[0:1], off ; CHECK-NEXT: s_waitcnt vmcnt(0) +; CHECK-NEXT: v_cndmask_b32_e32 v0, v4, v3, vcc ; CHECK-NEXT: s_setpc_b64 s[30:31] %gep.a = getelementptr i8, ptr addrspace(1) %a, i64 256 %gep.b = getelementptr i8, ptr addrspace(1) %a, i64 512 diff --git a/llvm/test/CodeGen/AMDGPU/select-vectors.ll b/llvm/test/CodeGen/AMDGPU/select-vectors.ll index e754f665c5f43..bee00f6efbd12 100644 --- a/llvm/test/CodeGen/AMDGPU/select-vectors.ll +++ b/llvm/test/CodeGen/AMDGPU/select-vectors.ll @@ -16,9 +16,9 @@ ; SelectionDAGBuilder for some reason changes the select type. ; VI: s_cselect_b64 ; VI: v_cndmask_b32 -define amdgpu_kernel void @v_select_v2i8(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(4) %b.ptr, i32 %c) #0 { +define amdgpu_kernel void @v_select_v2i8(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(1) %b.ptr, i32 %c) #0 { %a = load <2 x i8>, ptr addrspace(1) %a.ptr, align 2 - %b = load <2 x i8>, ptr addrspace(4) %b.ptr, align 2 + %b = load <2 x i8>, ptr addrspace(1) %b.ptr, align 2 %cmp = icmp eq i32 %c, 0 %select = select i1 %cmp, <2 x i8> %a, <2 x i8> %b store <2 x i8> %select, ptr addrspace(1) %out, align 2 @@ -28,9 +28,9 @@ define amdgpu_kernel void @v_select_v2i8(ptr addrspace(1) %out, ptr addrspace(1) ; GCN-LABEL: {{^}}v_select_v4i8: ; GCN: v_cndmask_b32_e32 ; GCN-NOT: cndmask -define amdgpu_kernel void @v_select_v4i8(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(4) %b.ptr, i32 %c) #0 { +define amdgpu_kernel void @v_select_v4i8(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(1) %b.ptr, i32 %c) #0 { %a = load <4 x i8>, ptr addrspace(1) %a.ptr - %b = load <4 x i8>, ptr addrspace(4) %b.ptr + %b = load <4 x i8>, ptr addrspace(1) %b.ptr %cmp = icmp eq i32 %c, 0 %select = select i1 %cmp, <4 x i8> %a, <4 x i8> %b store <4 x i8> %select, ptr addrspace(1) %out, align 4 @@ -41,9 +41,9 @@ define amdgpu_kernel void @v_select_v4i8(ptr addrspace(1) %out, ptr addrspace(1) ; GCN: v_cndmask_b32_e32 ; GCN: v_cndmask_b32_e32 ; GCN-NOT: cndmask -define amdgpu_kernel void @v_select_v8i8(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(4) %b.ptr, i32 %c) #0 { +define amdgpu_kernel void @v_select_v8i8(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(1) %b.ptr, i32 %c) #0 { %a = load <8 x i8>, ptr addrspace(1) %a.ptr - %b = load <8 x i8>, ptr addrspace(4) %b.ptr + %b = load <8 x i8>, ptr addrspace(1) %b.ptr %cmp = icmp eq i32 %c, 0 %select = select i1 %cmp, <8 x i8> %a, <8 x i8> %b store <8 x i8> %select, ptr addrspace(1) %out, align 4 @@ -56,9 +56,9 @@ define amdgpu_kernel void @v_select_v8i8(ptr addrspace(1) %out, ptr addrspace(1) ; GCN: v_cndmask_b32_e32 ; GCN: v_cndmask_b32_e32 ; GCN-NOT: cndmask -define amdgpu_kernel void @v_select_v16i8(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(4) %b.ptr, i32 %c) #0 { +define amdgpu_kernel void @v_select_v16i8(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(1) %b.ptr, i32 %c) #0 { %a = load <16 x i8>, ptr addrspace(1) %a.ptr - %b = load <16 x i8>, ptr addrspace(4) %b.ptr + %b = load <16 x i8>, ptr addrspace(1) %b.ptr %cmp = icmp eq i32 %c, 0 %select = select i1 %cmp, <16 x i8> %a, <16 x i8> %b store <16 x i8> %select, ptr addrspace(1) %out, align 4 @@ -93,16 +93,13 @@ define amdgpu_kernel void @select_v2i16(ptr addrspace(1) %out, <2 x i16> %a, <2 } ; GCN-LABEL: {{^}}v_select_v2i16: -; GCN: {{buffer|flat|global}}_load_dword v -; GCN: {{buffer|flat|global}}_load_dword v +; GCN: buffer_load_dword v +; GCN: buffer_load_dword v ; GCN: v_cndmask_b32 ; GCN-NOT: cndmask -define amdgpu_kernel void @v_select_v2i16(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(4) %b.ptr, i32 %c) #0 { - %id = call i32 @llvm.amdgcn.workitem.id.x() - %gep.a = getelementptr <2 x i16>, ptr addrspace(1) %a.ptr, i32 %id - %gep.b = getelementptr <2 x i16>, ptr addrspace(4) %b.ptr, i32 %id - %a = load <2 x i16>, ptr addrspace(1) %gep.a - %b = load <2 x i16>, ptr addrspace(4) %gep.b +define amdgpu_kernel void @v_select_v2i16(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(1) %b.ptr, i32 %c) #0 { + %a = load <2 x i16>, ptr addrspace(1) %a.ptr + %b = load <2 x i16>, ptr addrspace(1) %b.ptr %cmp = icmp eq i32 %c, 0 %select = select i1 %cmp, <2 x i16> %a, <2 x i16> %b store <2 x i16> %select, ptr addrspace(1) %out, align 4 @@ -117,9 +114,9 @@ define amdgpu_kernel void @v_select_v2i16(ptr addrspace(1) %out, ptr addrspace(1 ; VI: s_cselect_b64 ; GFX9: cndmask ; GFX9: cndmask -define amdgpu_kernel void @v_select_v3i16(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(4) %b.ptr, i32 %c) #0 { +define amdgpu_kernel void @v_select_v3i16(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(1) %b.ptr, i32 %c) #0 { %a = load <3 x i16>, ptr addrspace(1) %a.ptr - %b = load <3 x i16>, ptr addrspace(4) %b.ptr + %b = load <3 x i16>, ptr addrspace(1) %b.ptr %cmp = icmp eq i32 %c, 0 %select = select i1 %cmp, <3 x i16> %a, <3 x i16> %b store <3 x i16> %select, ptr addrspace(1) %out, align 4 @@ -130,9 +127,9 @@ define amdgpu_kernel void @v_select_v3i16(ptr addrspace(1) %out, ptr addrspace(1 ; GCN: v_cndmask_b32_e32 ; GCN: v_cndmask_b32_e32 ; GCN-NOT: cndmask -define amdgpu_kernel void @v_select_v4i16(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(4) %b.ptr, i32 %c) #0 { +define amdgpu_kernel void @v_select_v4i16(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(1) %b.ptr, i32 %c) #0 { %a = load <4 x i16>, ptr addrspace(1) %a.ptr - %b = load <4 x i16>, ptr addrspace(4) %b.ptr + %b = load <4 x i16>, ptr addrspace(1) %b.ptr %cmp = icmp eq i32 %c, 0 %select = select i1 %cmp, <4 x i16> %a, <4 x i16> %b store <4 x i16> %select, ptr addrspace(1) %out, align 4 @@ -145,9 +142,9 @@ define amdgpu_kernel void @v_select_v4i16(ptr addrspace(1) %out, ptr addrspace(1 ; GCN: v_cndmask_b32_e32 ; GCN: v_cndmask_b32_e32 ; GCN-NOT: cndmask -define amdgpu_kernel void @v_select_v8i16(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(4) %b.ptr, i32 %c) #0 { +define amdgpu_kernel void @v_select_v8i16(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(1) %b.ptr, i32 %c) #0 { %a = load <8 x i16>, ptr addrspace(1) %a.ptr - %b = load <8 x i16>, ptr addrspace(4) %b.ptr + %b = load <8 x i16>, ptr addrspace(1) %b.ptr %cmp = icmp eq i32 %c, 0 %select = select i1 %cmp, <8 x i16> %a, <8 x i16> %b store <8 x i16> %select, ptr addrspace(1) %out, align 4 @@ -164,9 +161,9 @@ define amdgpu_kernel void @v_select_v8i16(ptr addrspace(1) %out, ptr addrspace(1 ; GCN: v_cndmask_b32_e32 ; GCN: v_cndmask_b32_e32 ; GCN-NOT: cndmask -define amdgpu_kernel void @v_select_v16i16(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(4) %b.ptr, i32 %c) #0 { +define amdgpu_kernel void @v_select_v16i16(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(1) %b.ptr, i32 %c) #0 { %a = load <16 x i16>, ptr addrspace(1) %a.ptr - %b = load <16 x i16>, ptr addrspace(4) %b.ptr + %b = load <16 x i16>, ptr addrspace(1) %b.ptr %cmp = icmp eq i32 %c, 0 %select = select i1 %cmp, <16 x i16> %a, <16 x i16> %b store <16 x i16> %select, ptr addrspace(1) %out, align 4 @@ -191,9 +188,9 @@ define amdgpu_kernel void @v_select_v16i16(ptr addrspace(1) %out, ptr addrspace( ; GCN: v_cndmask_b32_e32 ; GCN: v_cndmask_b32_e32 ; GCN-NOT: cndmask -define amdgpu_kernel void @v_select_v32i16(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(4) %b.ptr, i32 %c) #0 { +define amdgpu_kernel void @v_select_v32i16(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(1) %b.ptr, i32 %c) #0 { %a = load <32 x i16>, ptr addrspace(1) %a.ptr - %b = load <32 x i16>, ptr addrspace(4) %b.ptr + %b = load <32 x i16>, ptr addrspace(1) %b.ptr %cmp = icmp eq i32 %c, 0 %select = select i1 %cmp, <32 x i16> %a, <32 x i16> %b store <32 x i16> %select, ptr addrspace(1) %out, align 4 @@ -336,7 +333,6 @@ bb: define amdgpu_kernel void @s_select_v5f32(ptr addrspace(1) %out, <5 x float> %a, <5 x float> %b, i32 %c) #0 { %cmp = icmp eq i32 %c, 0 %select = select i1 %cmp, <5 x float> %a, <5 x float> %b - call void asm "; use $0", "s"(<5 x float> %a) store <5 x float> %select, ptr addrspace(1) %out, align 16 ret void } @@ -404,7 +400,6 @@ define amdgpu_kernel void @select_v4f64(ptr addrspace(1) %out, <4 x double> %a, ; GCN: s_cselect_b32 define amdgpu_kernel void @select_v8f64(ptr addrspace(1) %out, <8 x double> %a, <8 x double> %b, i32 %c) #0 { %cmp = icmp eq i32 %c, 0 - call void asm "; use $0", "s"(<8 x double> %a) %select = select i1 %cmp, <8 x double> %a, <8 x double> %b store <8 x double> %select, ptr addrspace(1) %out, align 16 ret void @@ -413,9 +408,9 @@ define amdgpu_kernel void @select_v8f64(ptr addrspace(1) %out, <8 x double> %a, ; GCN-LABEL: {{^}}v_select_v2f16: ; GCN: v_cndmask_b32 ; GCN-NOT: cndmask -define amdgpu_kernel void @v_select_v2f16(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(4) %b.ptr, i32 %c) #0 { +define amdgpu_kernel void @v_select_v2f16(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(1) %b.ptr, i32 %c) #0 { %a = load <2 x half>, ptr addrspace(1) %a.ptr - %b = load <2 x half>, ptr addrspace(4) %b.ptr + %b = load <2 x half>, ptr addrspace(1) %b.ptr %cmp = icmp eq i32 %c, 0 %select = select i1 %cmp, <2 x half> %a, <2 x half> %b store <2 x half> %select, ptr addrspace(1) %out, align 4 @@ -426,9 +421,9 @@ define amdgpu_kernel void @v_select_v2f16(ptr addrspace(1) %out, ptr addrspace(1 ; GCN: v_cndmask_b32_e32 ; GCN: v_cndmask_b32_e32 ; GCN-NOT: cndmask -define amdgpu_kernel void @v_select_v3f16(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(4) %b.ptr, i32 %c) #0 { +define amdgpu_kernel void @v_select_v3f16(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(1) %b.ptr, i32 %c) #0 { %a = load <3 x half>, ptr addrspace(1) %a.ptr - %b = load <3 x half>, ptr addrspace(4) %b.ptr + %b = load <3 x half>, ptr addrspace(1) %b.ptr %cmp = icmp eq i32 %c, 0 %select = select i1 %cmp, <3 x half> %a, <3 x half> %b store <3 x half> %select, ptr addrspace(1) %out, align 4 @@ -439,9 +434,9 @@ define amdgpu_kernel void @v_select_v3f16(ptr addrspace(1) %out, ptr addrspace(1 ; GCN: v_cndmask_b32_e32 ; GCN: v_cndmask_b32_e32 ; GCN-NOT: cndmask -define amdgpu_kernel void @v_select_v4f16(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(4) %b.ptr, i32 %c) #0 { +define amdgpu_kernel void @v_select_v4f16(ptr addrspace(1) %out, ptr addrspace(1) %a.ptr, ptr addrspace(1) %b.ptr, i32 %c) #0 { %a = load <4 x half>, ptr addrspace(1) %a.ptr - %b = load <4 x half>, ptr addrspace(4) %b.ptr + %b = load <4 x half>, ptr addrspace(1) %b.ptr %cmp = icmp eq i32 %c, 0 %select = select i1 %cmp, <4 x half> %a, <4 x half> %b store <4 x half> %select, ptr addrspace(1) %out, align 4 diff --git a/llvm/test/CodeGen/AMDGPU/select64.ll b/llvm/test/CodeGen/AMDGPU/select64.ll index 2013db7bc7be3..de154b5ced8a8 100644 --- a/llvm/test/CodeGen/AMDGPU/select64.ll +++ b/llvm/test/CodeGen/AMDGPU/select64.ll @@ -40,10 +40,10 @@ define amdgpu_kernel void @select_trunc_i64_2(ptr addrspace(1) %out, i32 %cond, ; GCN-LABEL: {{^}}v_select_trunc_i64_2: ; GCN: s_cselect_b32 ; GCN-NOT: s_cselect_b32 -define amdgpu_kernel void @v_select_trunc_i64_2(ptr addrspace(1) %out, i32 %cond, ptr addrspace(1) %aptr, ptr addrspace(4) %bptr) nounwind { +define amdgpu_kernel void @v_select_trunc_i64_2(ptr addrspace(1) %out, i32 %cond, ptr addrspace(1) %aptr, ptr addrspace(1) %bptr) nounwind { %cmp = icmp ugt i32 %cond, 5 %a = load i64, ptr addrspace(1) %aptr, align 8 - %b = load i64, ptr addrspace(4) %bptr, align 8 + %b = load i64, ptr addrspace(1) %bptr, align 8 %sel = select i1 %cmp, i64 %a, i64 %b %trunc = trunc i64 %sel to i32 store i32 %trunc, ptr addrspace(1) %out, align 4 diff --git a/llvm/test/CodeGen/NVPTX/bf16-instructions.ll b/llvm/test/CodeGen/NVPTX/bf16-instructions.ll index 41f77b5337e6d..3626613cf8511 100644 --- a/llvm/test/CodeGen/NVPTX/bf16-instructions.ll +++ b/llvm/test/CodeGen/NVPTX/bf16-instructions.ll @@ -768,18 +768,17 @@ define bfloat @test_select_cc_bf16_f64(double %a, double %b, bfloat %c, bfloat % ; CHECK-LABEL: test_select_cc_bf16_f64( ; CHECK: { ; CHECK-NEXT: .reg .pred %p<2>; -; CHECK-NEXT: .reg .b16 %rs<2>; -; CHECK-NEXT: .reg .b64 %rd<6>; +; CHECK-NEXT: .reg .b16 %rs<4>; +; CHECK-NEXT: .reg .b64 %rd<3>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.b64 %rd1, [test_select_cc_bf16_f64_param_0]; ; CHECK-NEXT: ld.param.b64 %rd2, [test_select_cc_bf16_f64_param_1]; ; CHECK-NEXT: setp.lt.f64 %p1, %rd1, %rd2; -; CHECK-NEXT: mov.b64 %rd3, test_select_cc_bf16_f64_param_3; -; CHECK-NEXT: mov.b64 %rd4, test_select_cc_bf16_f64_param_2; -; CHECK-NEXT: selp.b64 %rd5, %rd4, %rd3, %p1; -; CHECK-NEXT: ld.param.b16 %rs1, [%rd5]; -; CHECK-NEXT: st.param.b16 [func_retval0], %rs1; +; CHECK-NEXT: ld.param.b16 %rs1, [test_select_cc_bf16_f64_param_2]; +; CHECK-NEXT: ld.param.b16 %rs2, [test_select_cc_bf16_f64_param_3]; +; CHECK-NEXT: selp.b16 %rs3, %rs1, %rs2, %p1; +; CHECK-NEXT: st.param.b16 [func_retval0], %rs3; ; CHECK-NEXT: ret; %cc = fcmp olt double %a, %b %r = select i1 %cc, bfloat %c, bfloat %d diff --git a/llvm/test/CodeGen/NVPTX/bf16x2-instructions.ll b/llvm/test/CodeGen/NVPTX/bf16x2-instructions.ll index c19e66559af86..3c6fb4b7517b8 100644 --- a/llvm/test/CodeGen/NVPTX/bf16x2-instructions.ll +++ b/llvm/test/CodeGen/NVPTX/bf16x2-instructions.ll @@ -229,18 +229,16 @@ define <2 x bfloat> @test_select(<2 x bfloat> %a, <2 x bfloat> %b, i1 zeroext %c ; CHECK: { ; CHECK-NEXT: .reg .pred %p<2>; ; CHECK-NEXT: .reg .b16 %rs<3>; -; CHECK-NEXT: .reg .b32 %r<2>; -; CHECK-NEXT: .reg .b64 %rd<4>; +; CHECK-NEXT: .reg .b32 %r<4>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.b8 %rs1, [test_select_param_2]; ; CHECK-NEXT: and.b16 %rs2, %rs1, 1; ; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0; -; CHECK-NEXT: mov.b64 %rd1, test_select_param_1; -; CHECK-NEXT: mov.b64 %rd2, test_select_param_0; -; CHECK-NEXT: selp.b64 %rd3, %rd2, %rd1, %p1; -; CHECK-NEXT: ld.param.b32 %r1, [%rd3]; -; CHECK-NEXT: st.param.b32 [func_retval0], %r1; +; CHECK-NEXT: ld.param.b32 %r1, [test_select_param_0]; +; CHECK-NEXT: ld.param.b32 %r2, [test_select_param_1]; +; CHECK-NEXT: selp.b32 %r3, %r1, %r2, %p1; +; CHECK-NEXT: st.param.b32 [func_retval0], %r3; ; CHECK-NEXT: ret; %r = select i1 %c, <2 x bfloat> %a, <2 x bfloat> %b ret <2 x bfloat> %r diff --git a/llvm/test/CodeGen/NVPTX/bug22246.ll b/llvm/test/CodeGen/NVPTX/bug22246.ll index 1d7a396f694dc..198878c1b96ff 100644 --- a/llvm/test/CodeGen/NVPTX/bug22246.ll +++ b/llvm/test/CodeGen/NVPTX/bug22246.ll @@ -9,20 +9,19 @@ define void @_Z3foobbbPb(i1 zeroext %p1, i1 zeroext %p2, i1 zeroext %p3, ptr noc ; CHECK-LABEL: _Z3foobbbPb( ; CHECK: { ; CHECK-NEXT: .reg .pred %p<2>; -; CHECK-NEXT: .reg .b16 %rs<5>; -; CHECK-NEXT: .reg .b64 %rd<5>; +; CHECK-NEXT: .reg .b16 %rs<7>; +; CHECK-NEXT: .reg .b64 %rd<2>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: // %entry ; CHECK-NEXT: ld.param.b8 %rs1, [_Z3foobbbPb_param_0]; ; CHECK-NEXT: and.b16 %rs2, %rs1, 1; ; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0; -; CHECK-NEXT: mov.b64 %rd1, _Z3foobbbPb_param_2; -; CHECK-NEXT: mov.b64 %rd2, _Z3foobbbPb_param_1; -; CHECK-NEXT: selp.b64 %rd3, %rd2, %rd1, %p1; -; CHECK-NEXT: ld.param.b8 %rs3, [%rd3]; -; CHECK-NEXT: and.b16 %rs4, %rs3, 1; -; CHECK-NEXT: ld.param.b64 %rd4, [_Z3foobbbPb_param_3]; -; CHECK-NEXT: st.b8 [%rd4], %rs4; +; CHECK-NEXT: ld.param.b8 %rs3, [_Z3foobbbPb_param_1]; +; CHECK-NEXT: ld.param.b8 %rs4, [_Z3foobbbPb_param_2]; +; CHECK-NEXT: selp.b16 %rs5, %rs3, %rs4, %p1; +; CHECK-NEXT: and.b16 %rs6, %rs5, 1; +; CHECK-NEXT: ld.param.b64 %rd1, [_Z3foobbbPb_param_3]; +; CHECK-NEXT: st.b8 [%rd1], %rs6; ; CHECK-NEXT: ret; entry: %.sink.v = select i1 %p1, i1 %p2, i1 %p3 diff --git a/llvm/test/CodeGen/NVPTX/fast-math.ll b/llvm/test/CodeGen/NVPTX/fast-math.ll index 7e778c40b8302..8561c60a46948 100644 --- a/llvm/test/CodeGen/NVPTX/fast-math.ll +++ b/llvm/test/CodeGen/NVPTX/fast-math.ll @@ -312,20 +312,18 @@ define float @repeated_div_recip_allowed_sel(i1 %pred, float %a, float %b, float ; CHECK: { ; CHECK-NEXT: .reg .pred %p<2>; ; CHECK-NEXT: .reg .b16 %rs<3>; -; CHECK-NEXT: .reg .b32 %r<4>; -; CHECK-NEXT: .reg .b64 %rd<4>; +; CHECK-NEXT: .reg .b32 %r<6>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.b8 %rs1, [repeated_div_recip_allowed_sel_param_0]; ; CHECK-NEXT: and.b16 %rs2, %rs1, 1; ; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0; -; CHECK-NEXT: mov.b64 %rd1, repeated_div_recip_allowed_sel_param_2; -; CHECK-NEXT: mov.b64 %rd2, repeated_div_recip_allowed_sel_param_1; -; CHECK-NEXT: selp.b64 %rd3, %rd2, %rd1, %p1; -; CHECK-NEXT: ld.param.b32 %r1, [%rd3]; -; CHECK-NEXT: ld.param.b32 %r2, [repeated_div_recip_allowed_sel_param_3]; -; CHECK-NEXT: div.rn.f32 %r3, %r1, %r2; -; CHECK-NEXT: st.param.b32 [func_retval0], %r3; +; CHECK-NEXT: ld.param.b32 %r1, [repeated_div_recip_allowed_sel_param_1]; +; CHECK-NEXT: ld.param.b32 %r2, [repeated_div_recip_allowed_sel_param_2]; +; CHECK-NEXT: selp.f32 %r3, %r1, %r2, %p1; +; CHECK-NEXT: ld.param.b32 %r4, [repeated_div_recip_allowed_sel_param_3]; +; CHECK-NEXT: div.rn.f32 %r5, %r3, %r4; +; CHECK-NEXT: st.param.b32 [func_retval0], %r5; ; CHECK-NEXT: ret; %x = fdiv arcp float %a, %divisor %y = fdiv arcp float %b, %divisor @@ -366,20 +364,18 @@ define float @repeated_div_recip_allowed_ftz_sel(i1 %pred, float %a, float %b, f ; CHECK: { ; CHECK-NEXT: .reg .pred %p<2>; ; CHECK-NEXT: .reg .b16 %rs<3>; -; CHECK-NEXT: .reg .b32 %r<4>; -; CHECK-NEXT: .reg .b64 %rd<4>; +; CHECK-NEXT: .reg .b32 %r<6>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.b8 %rs1, [repeated_div_recip_allowed_ftz_sel_param_0]; ; CHECK-NEXT: and.b16 %rs2, %rs1, 1; ; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0; -; CHECK-NEXT: mov.b64 %rd1, repeated_div_recip_allowed_ftz_sel_param_2; -; CHECK-NEXT: mov.b64 %rd2, repeated_div_recip_allowed_ftz_sel_param_1; -; CHECK-NEXT: selp.b64 %rd3, %rd2, %rd1, %p1; -; CHECK-NEXT: ld.param.b32 %r1, [%rd3]; -; CHECK-NEXT: ld.param.b32 %r2, [repeated_div_recip_allowed_ftz_sel_param_3]; -; CHECK-NEXT: div.rn.ftz.f32 %r3, %r1, %r2; -; CHECK-NEXT: st.param.b32 [func_retval0], %r3; +; CHECK-NEXT: ld.param.b32 %r1, [repeated_div_recip_allowed_ftz_sel_param_1]; +; CHECK-NEXT: ld.param.b32 %r2, [repeated_div_recip_allowed_ftz_sel_param_2]; +; CHECK-NEXT: selp.f32 %r3, %r1, %r2, %p1; +; CHECK-NEXT: ld.param.b32 %r4, [repeated_div_recip_allowed_ftz_sel_param_3]; +; CHECK-NEXT: div.rn.ftz.f32 %r5, %r3, %r4; +; CHECK-NEXT: st.param.b32 [func_retval0], %r5; ; CHECK-NEXT: ret; %x = fdiv arcp float %a, %divisor %y = fdiv arcp float %b, %divisor @@ -420,20 +416,18 @@ define float @repeated_div_fast_sel(i1 %pred, float %a, float %b, float %divisor ; CHECK: { ; CHECK-NEXT: .reg .pred %p<2>; ; CHECK-NEXT: .reg .b16 %rs<3>; -; CHECK-NEXT: .reg .b32 %r<4>; -; CHECK-NEXT: .reg .b64 %rd<4>; +; CHECK-NEXT: .reg .b32 %r<6>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.b8 %rs1, [repeated_div_fast_sel_param_0]; ; CHECK-NEXT: and.b16 %rs2, %rs1, 1; ; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0; -; CHECK-NEXT: mov.b64 %rd1, repeated_div_fast_sel_param_2; -; CHECK-NEXT: mov.b64 %rd2, repeated_div_fast_sel_param_1; -; CHECK-NEXT: selp.b64 %rd3, %rd2, %rd1, %p1; -; CHECK-NEXT: ld.param.b32 %r1, [%rd3]; -; CHECK-NEXT: ld.param.b32 %r2, [repeated_div_fast_sel_param_3]; -; CHECK-NEXT: div.approx.f32 %r3, %r1, %r2; -; CHECK-NEXT: st.param.b32 [func_retval0], %r3; +; CHECK-NEXT: ld.param.b32 %r1, [repeated_div_fast_sel_param_1]; +; CHECK-NEXT: ld.param.b32 %r2, [repeated_div_fast_sel_param_2]; +; CHECK-NEXT: selp.f32 %r3, %r1, %r2, %p1; +; CHECK-NEXT: ld.param.b32 %r4, [repeated_div_fast_sel_param_3]; +; CHECK-NEXT: div.approx.f32 %r5, %r3, %r4; +; CHECK-NEXT: st.param.b32 [func_retval0], %r5; ; CHECK-NEXT: ret; %x = fdiv afn float %a, %divisor %y = fdiv afn float %b, %divisor @@ -474,20 +468,18 @@ define float @repeated_div_fast_ftz_sel(i1 %pred, float %a, float %b, float %div ; CHECK: { ; CHECK-NEXT: .reg .pred %p<2>; ; CHECK-NEXT: .reg .b16 %rs<3>; -; CHECK-NEXT: .reg .b32 %r<4>; -; CHECK-NEXT: .reg .b64 %rd<4>; +; CHECK-NEXT: .reg .b32 %r<6>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.b8 %rs1, [repeated_div_fast_ftz_sel_param_0]; ; CHECK-NEXT: and.b16 %rs2, %rs1, 1; ; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0; -; CHECK-NEXT: mov.b64 %rd1, repeated_div_fast_ftz_sel_param_2; -; CHECK-NEXT: mov.b64 %rd2, repeated_div_fast_ftz_sel_param_1; -; CHECK-NEXT: selp.b64 %rd3, %rd2, %rd1, %p1; -; CHECK-NEXT: ld.param.b32 %r1, [%rd3]; -; CHECK-NEXT: ld.param.b32 %r2, [repeated_div_fast_ftz_sel_param_3]; -; CHECK-NEXT: div.approx.ftz.f32 %r3, %r1, %r2; -; CHECK-NEXT: st.param.b32 [func_retval0], %r3; +; CHECK-NEXT: ld.param.b32 %r1, [repeated_div_fast_ftz_sel_param_1]; +; CHECK-NEXT: ld.param.b32 %r2, [repeated_div_fast_ftz_sel_param_2]; +; CHECK-NEXT: selp.f32 %r3, %r1, %r2, %p1; +; CHECK-NEXT: ld.param.b32 %r4, [repeated_div_fast_ftz_sel_param_3]; +; CHECK-NEXT: div.approx.ftz.f32 %r5, %r3, %r4; +; CHECK-NEXT: st.param.b32 [func_retval0], %r5; ; CHECK-NEXT: ret; %x = fdiv afn float %a, %divisor %y = fdiv afn float %b, %divisor diff --git a/llvm/test/CodeGen/NVPTX/i1-select.ll b/llvm/test/CodeGen/NVPTX/i1-select.ll index c91a3df204d80..264f38021e1de 100644 --- a/llvm/test/CodeGen/NVPTX/i1-select.ll +++ b/llvm/test/CodeGen/NVPTX/i1-select.ll @@ -8,24 +8,21 @@ define i32 @test_select_i1_trunc(i32 %a, i32 %b, i32 %c, i32 %true, i32 %false) ; CHECK-LABEL: test_select_i1_trunc( ; CHECK: { ; CHECK-NEXT: .reg .pred %p<3>; -; CHECK-NEXT: .reg .b32 %r<6>; -; CHECK-NEXT: .reg .b64 %rd<7>; +; CHECK-NEXT: .reg .b32 %r<10>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.b32 %r1, [test_select_i1_trunc_param_0]; ; CHECK-NEXT: and.b32 %r2, %r1, 1; ; CHECK-NEXT: setp.ne.b32 %p1, %r2, 0; -; CHECK-NEXT: mov.b64 %rd1, test_select_i1_trunc_param_2; -; CHECK-NEXT: mov.b64 %rd2, test_select_i1_trunc_param_1; -; CHECK-NEXT: selp.b64 %rd3, %rd2, %rd1, %p1; -; CHECK-NEXT: ld.param.b32 %r3, [%rd3]; -; CHECK-NEXT: and.b32 %r4, %r3, 1; -; CHECK-NEXT: setp.ne.b32 %p2, %r4, 0; -; CHECK-NEXT: mov.b64 %rd4, test_select_i1_trunc_param_4; -; CHECK-NEXT: mov.b64 %rd5, test_select_i1_trunc_param_3; -; CHECK-NEXT: selp.b64 %rd6, %rd5, %rd4, %p2; -; CHECK-NEXT: ld.param.b32 %r5, [%rd6]; -; CHECK-NEXT: st.param.b32 [func_retval0], %r5; +; CHECK-NEXT: ld.param.b32 %r3, [test_select_i1_trunc_param_1]; +; CHECK-NEXT: ld.param.b32 %r4, [test_select_i1_trunc_param_2]; +; CHECK-NEXT: ld.param.b32 %r5, [test_select_i1_trunc_param_3]; +; CHECK-NEXT: selp.b32 %r6, %r3, %r4, %p1; +; CHECK-NEXT: and.b32 %r7, %r6, 1; +; CHECK-NEXT: setp.ne.b32 %p2, %r7, 0; +; CHECK-NEXT: ld.param.b32 %r8, [test_select_i1_trunc_param_4]; +; CHECK-NEXT: selp.b32 %r9, %r5, %r8, %p2; +; CHECK-NEXT: st.param.b32 [func_retval0], %r9; ; CHECK-NEXT: ret; %a_trunc = trunc i32 %a to i1 %b_trunc = trunc i32 %b to i1 @@ -39,25 +36,23 @@ define i32 @test_select_i1_trunc_2(i64 %a, i16 %b, i32 %c, i32 %true, i32 %false ; CHECK-LABEL: test_select_i1_trunc_2( ; CHECK: { ; CHECK-NEXT: .reg .pred %p<3>; -; CHECK-NEXT: .reg .b16 %rs<3>; -; CHECK-NEXT: .reg .b32 %r<2>; -; CHECK-NEXT: .reg .b64 %rd<9>; +; CHECK-NEXT: .reg .b16 %rs<5>; +; CHECK-NEXT: .reg .b32 %r<4>; +; CHECK-NEXT: .reg .b64 %rd<3>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.b64 %rd1, [test_select_i1_trunc_2_param_0]; ; CHECK-NEXT: and.b64 %rd2, %rd1, 1; ; CHECK-NEXT: setp.ne.b64 %p1, %rd2, 0; -; CHECK-NEXT: mov.b64 %rd3, test_select_i1_trunc_2_param_2; -; CHECK-NEXT: mov.b64 %rd4, test_select_i1_trunc_2_param_1; -; CHECK-NEXT: selp.b64 %rd5, %rd4, %rd3, %p1; -; CHECK-NEXT: ld.param.b16 %rs1, [%rd5]; -; CHECK-NEXT: and.b16 %rs2, %rs1, 1; -; CHECK-NEXT: setp.ne.b16 %p2, %rs2, 0; -; CHECK-NEXT: mov.b64 %rd6, test_select_i1_trunc_2_param_4; -; CHECK-NEXT: mov.b64 %rd7, test_select_i1_trunc_2_param_3; -; CHECK-NEXT: selp.b64 %rd8, %rd7, %rd6, %p2; -; CHECK-NEXT: ld.param.b32 %r1, [%rd8]; -; CHECK-NEXT: st.param.b32 [func_retval0], %r1; +; CHECK-NEXT: ld.param.b16 %rs1, [test_select_i1_trunc_2_param_1]; +; CHECK-NEXT: ld.param.b16 %rs2, [test_select_i1_trunc_2_param_2]; +; CHECK-NEXT: ld.param.b32 %r1, [test_select_i1_trunc_2_param_3]; +; CHECK-NEXT: selp.b16 %rs3, %rs1, %rs2, %p1; +; CHECK-NEXT: and.b16 %rs4, %rs3, 1; +; CHECK-NEXT: setp.ne.b16 %p2, %rs4, 0; +; CHECK-NEXT: ld.param.b32 %r2, [test_select_i1_trunc_2_param_4]; +; CHECK-NEXT: selp.b32 %r3, %r1, %r2, %p2; +; CHECK-NEXT: st.param.b32 [func_retval0], %r3; ; CHECK-NEXT: ret; %a_trunc = trunc i64 %a to i1 %b_trunc = trunc i16 %b to i1 @@ -71,8 +66,7 @@ define i32 @test_select_i1_basic(i32 %v1, i32 %v2, i32 %v3, i32 %true, i32 %fals ; CHECK-LABEL: test_select_i1_basic( ; CHECK: { ; CHECK-NEXT: .reg .pred %p<4>; -; CHECK-NEXT: .reg .b32 %r<6>; -; CHECK-NEXT: .reg .b64 %rd<6>; +; CHECK-NEXT: .reg .b32 %r<10>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.b32 %r1, [test_select_i1_basic_param_0]; @@ -81,14 +75,13 @@ define i32 @test_select_i1_basic(i32 %v1, i32 %v2, i32 %v3, i32 %true, i32 %fals ; CHECK-NEXT: setp.ne.b32 %p1, %r1, 0; ; CHECK-NEXT: ld.param.b32 %r4, [test_select_i1_basic_param_2]; ; CHECK-NEXT: setp.eq.b32 %p2, %r4, 0; +; CHECK-NEXT: ld.param.b32 %r5, [test_select_i1_basic_param_3]; ; CHECK-NEXT: setp.eq.b32 %p3, %r3, 0; -; CHECK-NEXT: mov.b64 %rd1, test_select_i1_basic_param_4; -; CHECK-NEXT: mov.b64 %rd2, test_select_i1_basic_param_3; -; CHECK-NEXT: selp.b64 %rd3, %rd2, %rd1, %p2; -; CHECK-NEXT: selp.b64 %rd4, %rd3, %rd1, %p1; -; CHECK-NEXT: selp.b64 %rd5, %rd2, %rd4, %p3; -; CHECK-NEXT: ld.param.b32 %r5, [%rd5]; -; CHECK-NEXT: st.param.b32 [func_retval0], %r5; +; CHECK-NEXT: ld.param.b32 %r6, [test_select_i1_basic_param_4]; +; CHECK-NEXT: selp.b32 %r7, %r5, %r6, %p2; +; CHECK-NEXT: selp.b32 %r8, %r7, %r6, %p1; +; CHECK-NEXT: selp.b32 %r9, %r5, %r8, %p3; +; CHECK-NEXT: st.param.b32 [func_retval0], %r9; ; CHECK-NEXT: ret; %b1 = icmp eq i32 %v1, 0 %b2 = icmp eq i32 %v2, 0 @@ -102,8 +95,7 @@ define i32 @test_select_i1_basic_folding(i32 %v1, i32 %v2, i32 %v3, i32 %true, i ; CHECK-LABEL: test_select_i1_basic_folding( ; CHECK: { ; CHECK-NEXT: .reg .pred %p<11>; -; CHECK-NEXT: .reg .b32 %r<5>; -; CHECK-NEXT: .reg .b64 %rd<4>; +; CHECK-NEXT: .reg .b32 %r<7>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.b32 %r1, [test_select_i1_basic_folding_param_0]; @@ -113,17 +105,16 @@ define i32 @test_select_i1_basic_folding(i32 %v1, i32 %v2, i32 %v3, i32 %true, i ; CHECK-NEXT: setp.eq.b32 %p3, %r2, 0; ; CHECK-NEXT: ld.param.b32 %r3, [test_select_i1_basic_folding_param_2]; ; CHECK-NEXT: setp.eq.b32 %p4, %r3, 0; +; CHECK-NEXT: ld.param.b32 %r4, [test_select_i1_basic_folding_param_3]; ; CHECK-NEXT: xor.pred %p5, %p1, %p3; +; CHECK-NEXT: ld.param.b32 %r5, [test_select_i1_basic_folding_param_4]; ; CHECK-NEXT: and.pred %p6, %p5, %p4; ; CHECK-NEXT: and.pred %p7, %p2, %p4; ; CHECK-NEXT: and.pred %p8, %p3, %p6; ; CHECK-NEXT: or.pred %p9, %p8, %p7; ; CHECK-NEXT: xor.pred %p10, %p9, %p3; -; CHECK-NEXT: mov.b64 %rd1, test_select_i1_basic_folding_param_4; -; CHECK-NEXT: mov.b64 %rd2, test_select_i1_basic_folding_param_3; -; CHECK-NEXT: selp.b64 %rd3, %rd2, %rd1, %p10; -; CHECK-NEXT: ld.param.b32 %r4, [%rd3]; -; CHECK-NEXT: st.param.b32 [func_retval0], %r4; +; CHECK-NEXT: selp.b32 %r6, %r4, %r5, %p10; +; CHECK-NEXT: st.param.b32 [func_retval0], %r6; ; CHECK-NEXT: ret; %b1 = icmp eq i32 %v1, 0 %b2 = icmp eq i32 %v2, 0 diff --git a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll index bfac2b4ffd431..40d6a07310265 100644 --- a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll +++ b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll @@ -1442,18 +1442,16 @@ define <4 x i8> @test_select(<4 x i8> %a, <4 x i8> %b, i1 zeroext %c) #0 { ; O3: { ; O3-NEXT: .reg .pred %p<2>; ; O3-NEXT: .reg .b16 %rs<3>; -; O3-NEXT: .reg .b32 %r<2>; -; O3-NEXT: .reg .b64 %rd<4>; +; O3-NEXT: .reg .b32 %r<4>; ; O3-EMPTY: ; O3-NEXT: // %bb.0: ; O3-NEXT: ld.param.b8 %rs1, [test_select_param_2]; ; O3-NEXT: and.b16 %rs2, %rs1, 1; ; O3-NEXT: setp.ne.b16 %p1, %rs2, 0; -; O3-NEXT: mov.b64 %rd1, test_select_param_1; -; O3-NEXT: mov.b64 %rd2, test_select_param_0; -; O3-NEXT: selp.b64 %rd3, %rd2, %rd1, %p1; -; O3-NEXT: ld.param.b32 %r1, [%rd3]; -; O3-NEXT: st.param.b32 [func_retval0], %r1; +; O3-NEXT: ld.param.b32 %r1, [test_select_param_0]; +; O3-NEXT: ld.param.b32 %r2, [test_select_param_1]; +; O3-NEXT: selp.b32 %r3, %r1, %r2, %p1; +; O3-NEXT: st.param.b32 [func_retval0], %r3; ; O3-NEXT: ret; %r = select i1 %c, <4 x i8> %a, <4 x i8> %b ret <4 x i8> %r diff --git a/llvm/test/CodeGen/NVPTX/lower-byval-args.ll b/llvm/test/CodeGen/NVPTX/lower-byval-args.ll index ca2914a2e8043..21257e21bea9f 100644 --- a/llvm/test/CodeGen/NVPTX/lower-byval-args.ll +++ b/llvm/test/CodeGen/NVPTX/lower-byval-args.ll @@ -584,25 +584,44 @@ define ptx_kernel void @test_select(ptr byval(i32) align 4 %input1, ptr byval(i3 ; COPY-NEXT: store i32 [[VALLOADED]], ptr [[OUT]], align 4 ; COPY-NEXT: ret void ; -; PTX-LABEL: test_select( -; PTX: { -; PTX-NEXT: .reg .pred %p<2>; -; PTX-NEXT: .reg .b16 %rs<3>; -; PTX-NEXT: .reg .b32 %r<2>; -; PTX-NEXT: .reg .b64 %rd<6>; -; PTX-EMPTY: -; PTX-NEXT: // %bb.0: // %bb -; PTX-NEXT: ld.param.b8 %rs1, [test_select_param_3]; -; PTX-NEXT: and.b16 %rs2, %rs1, 1; -; PTX-NEXT: setp.ne.b16 %p1, %rs2, 0; -; PTX-NEXT: mov.b64 %rd1, test_select_param_0; -; PTX-NEXT: ld.param.b64 %rd2, [test_select_param_2]; -; PTX-NEXT: cvta.to.global.u64 %rd3, %rd2; -; PTX-NEXT: mov.b64 %rd4, test_select_param_1; -; PTX-NEXT: selp.b64 %rd5, %rd1, %rd4, %p1; -; PTX-NEXT: ld.param.b32 %r1, [%rd5]; -; PTX-NEXT: st.global.b32 [%rd3], %r1; -; PTX-NEXT: ret; +; PTX_60-LABEL: test_select( +; PTX_60: { +; PTX_60-NEXT: .reg .pred %p<2>; +; PTX_60-NEXT: .reg .b16 %rs<3>; +; PTX_60-NEXT: .reg .b32 %r<4>; +; PTX_60-NEXT: .reg .b64 %rd<3>; +; PTX_60-EMPTY: +; PTX_60-NEXT: // %bb.0: // %bb +; PTX_60-NEXT: ld.param.b8 %rs1, [test_select_param_3]; +; PTX_60-NEXT: and.b16 %rs2, %rs1, 1; +; PTX_60-NEXT: setp.ne.b16 %p1, %rs2, 0; +; PTX_60-NEXT: ld.param.b64 %rd1, [test_select_param_2]; +; PTX_60-NEXT: cvta.to.global.u64 %rd2, %rd1; +; PTX_60-NEXT: ld.param.b32 %r1, [test_select_param_1]; +; PTX_60-NEXT: ld.param.b32 %r2, [test_select_param_0]; +; PTX_60-NEXT: selp.b32 %r3, %r2, %r1, %p1; +; PTX_60-NEXT: st.global.b32 [%rd2], %r3; +; PTX_60-NEXT: ret; +; +; PTX_70-LABEL: test_select( +; PTX_70: { +; PTX_70-NEXT: .reg .pred %p<2>; +; PTX_70-NEXT: .reg .b16 %rs<3>; +; PTX_70-NEXT: .reg .b32 %r<2>; +; PTX_70-NEXT: .reg .b64 %rd<6>; +; PTX_70-EMPTY: +; PTX_70-NEXT: // %bb.0: // %bb +; PTX_70-NEXT: ld.param.b8 %rs1, [test_select_param_3]; +; PTX_70-NEXT: and.b16 %rs2, %rs1, 1; +; PTX_70-NEXT: setp.ne.b16 %p1, %rs2, 0; +; PTX_70-NEXT: mov.b64 %rd1, test_select_param_0; +; PTX_70-NEXT: ld.param.b64 %rd2, [test_select_param_2]; +; PTX_70-NEXT: cvta.to.global.u64 %rd3, %rd2; +; PTX_70-NEXT: mov.b64 %rd4, test_select_param_1; +; PTX_70-NEXT: selp.b64 %rd5, %rd1, %rd4, %p1; +; PTX_70-NEXT: ld.param.b32 %r1, [%rd5]; +; PTX_70-NEXT: st.global.b32 [%rd3], %r1; +; PTX_70-NEXT: ret; bb: %ptrnew = select i1 %cond, ptr %input1, ptr %input2 %valloaded = load i32, ptr %ptrnew, align 4 _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
