diff options
Diffstat (limited to 'llvm/test/CodeGen/NVPTX/cmpxchg.ll')
| -rw-r--r-- | llvm/test/CodeGen/NVPTX/cmpxchg.ll | 180 |
1 files changed, 90 insertions, 90 deletions
diff --git a/llvm/test/CodeGen/NVPTX/cmpxchg.ll b/llvm/test/CodeGen/NVPTX/cmpxchg.ll index 25b4c74086dc..237e42394ba2 100644 --- a/llvm/test/CodeGen/NVPTX/cmpxchg.ll +++ b/llvm/test/CodeGen/NVPTX/cmpxchg.ll @@ -79,7 +79,7 @@ define i8 @relaxed_sys_i8(ptr %addr, i8 %cmp, i8 %new) { ; SM70-NEXT: // =>This Inner Loop Header: Depth=1 ; SM70-NEXT: or.b32 %r17, %r20, %r3; ; SM70-NEXT: or.b32 %r18, %r20, %r4; -; SM70-NEXT: atom.relaxed.cas.b32 %r7, [%rd1], %r18, %r17; +; SM70-NEXT: atom.relaxed.sys.cas.b32 %r7, [%rd1], %r18, %r17; ; SM70-NEXT: setp.eq.b32 %p1, %r7, %r18; ; SM70-NEXT: @%p1 bra $L__BB0_3; ; SM70-NEXT: // %bb.2: // %partword.cmpxchg.failure @@ -99,8 +99,8 @@ define i8 @relaxed_sys_i8(ptr %addr, i8 %cmp, i8 %new) { ; SM90-NEXT: .reg .b64 %rd<3>; ; SM90-EMPTY: ; SM90-NEXT: // %bb.0: -; SM90-NEXT: ld.param.b8 %rs1, [relaxed_sys_i8_param_2]; -; SM90-NEXT: ld.param.b64 %rd2, [relaxed_sys_i8_param_0]; +; SM90-NEXT: ld.param.u8 %rs1, [relaxed_sys_i8_param_2]; +; SM90-NEXT: ld.param.u64 %rd2, [relaxed_sys_i8_param_0]; ; SM90-NEXT: and.b64 %rd1, %rd2, -4; ; SM90-NEXT: cvt.u32.u64 %r9, %rd2; ; SM90-NEXT: and.b32 %r10, %r9, 3; @@ -111,9 +111,9 @@ define i8 @relaxed_sys_i8(ptr %addr, i8 %cmp, i8 %new) { ; SM90-NEXT: cvt.u32.u16 %r13, %rs1; ; SM90-NEXT: and.b32 %r14, %r13, 255; ; SM90-NEXT: shl.b32 %r3, %r14, %r1; -; SM90-NEXT: ld.param.b8 %r15, [relaxed_sys_i8_param_1]; +; SM90-NEXT: ld.param.u8 %r15, [relaxed_sys_i8_param_1]; ; SM90-NEXT: shl.b32 %r4, %r15, %r1; -; SM90-NEXT: ld.b32 %r16, [%rd1]; +; SM90-NEXT: ld.u32 %r16, [%rd1]; ; SM90-NEXT: and.b32 %r20, %r16, %r2; ; SM90-NEXT: $L__BB0_1: // %partword.cmpxchg.loop ; SM90-NEXT: // =>This Inner Loop Header: Depth=1 @@ -206,7 +206,7 @@ define i8 @acquire_sys_i8(ptr %addr, i8 %cmp, i8 %new) { ; SM70-NEXT: // =>This Inner Loop Header: Depth=1 ; SM70-NEXT: or.b32 %r17, %r20, %r3; ; SM70-NEXT: or.b32 %r18, %r20, %r4; -; SM70-NEXT: atom.relaxed.cas.b32 %r7, [%rd1], %r18, %r17; +; SM70-NEXT: atom.relaxed.sys.cas.b32 %r7, [%rd1], %r18, %r17; ; SM70-NEXT: setp.eq.b32 %p1, %r7, %r18; ; SM70-NEXT: @%p1 bra $L__BB1_3; ; SM70-NEXT: // %bb.2: // %partword.cmpxchg.failure @@ -227,8 +227,8 @@ define i8 @acquire_sys_i8(ptr %addr, i8 %cmp, i8 %new) { ; SM90-NEXT: .reg .b64 %rd<3>; ; SM90-EMPTY: ; SM90-NEXT: // %bb.0: -; SM90-NEXT: ld.param.b8 %rs1, [acquire_sys_i8_param_2]; -; SM90-NEXT: ld.param.b64 %rd2, [acquire_sys_i8_param_0]; +; SM90-NEXT: ld.param.u8 %rs1, [acquire_sys_i8_param_2]; +; SM90-NEXT: ld.param.u64 %rd2, [acquire_sys_i8_param_0]; ; SM90-NEXT: and.b64 %rd1, %rd2, -4; ; SM90-NEXT: cvt.u32.u64 %r9, %rd2; ; SM90-NEXT: and.b32 %r10, %r9, 3; @@ -239,9 +239,9 @@ define i8 @acquire_sys_i8(ptr %addr, i8 %cmp, i8 %new) { ; SM90-NEXT: cvt.u32.u16 %r13, %rs1; ; SM90-NEXT: and.b32 %r14, %r13, 255; ; SM90-NEXT: shl.b32 %r3, %r14, %r1; -; SM90-NEXT: ld.param.b8 %r15, [acquire_sys_i8_param_1]; +; SM90-NEXT: ld.param.u8 %r15, [acquire_sys_i8_param_1]; ; SM90-NEXT: shl.b32 %r4, %r15, %r1; -; SM90-NEXT: ld.b32 %r16, [%rd1]; +; SM90-NEXT: ld.u32 %r16, [%rd1]; ; SM90-NEXT: and.b32 %r20, %r16, %r2; ; SM90-NEXT: $L__BB1_1: // %partword.cmpxchg.loop ; SM90-NEXT: // =>This Inner Loop Header: Depth=1 @@ -336,7 +336,7 @@ define i8 @release_sys_i8(ptr %addr, i8 %cmp, i8 %new) { ; SM70-NEXT: // =>This Inner Loop Header: Depth=1 ; SM70-NEXT: or.b32 %r17, %r20, %r3; ; SM70-NEXT: or.b32 %r18, %r20, %r4; -; SM70-NEXT: atom.relaxed.cas.b32 %r7, [%rd1], %r18, %r17; +; SM70-NEXT: atom.relaxed.sys.cas.b32 %r7, [%rd1], %r18, %r17; ; SM70-NEXT: setp.eq.b32 %p1, %r7, %r18; ; SM70-NEXT: @%p1 bra $L__BB2_3; ; SM70-NEXT: // %bb.2: // %partword.cmpxchg.failure @@ -356,8 +356,8 @@ define i8 @release_sys_i8(ptr %addr, i8 %cmp, i8 %new) { ; SM90-NEXT: .reg .b64 %rd<3>; ; SM90-EMPTY: ; SM90-NEXT: // %bb.0: -; SM90-NEXT: ld.param.b8 %rs1, [release_sys_i8_param_2]; -; SM90-NEXT: ld.param.b64 %rd2, [release_sys_i8_param_0]; +; SM90-NEXT: ld.param.u8 %rs1, [release_sys_i8_param_2]; +; SM90-NEXT: ld.param.u64 %rd2, [release_sys_i8_param_0]; ; SM90-NEXT: fence.release.sys; ; SM90-NEXT: and.b64 %rd1, %rd2, -4; ; SM90-NEXT: cvt.u32.u64 %r9, %rd2; @@ -369,9 +369,9 @@ define i8 @release_sys_i8(ptr %addr, i8 %cmp, i8 %new) { ; SM90-NEXT: cvt.u32.u16 %r13, %rs1; ; SM90-NEXT: and.b32 %r14, %r13, 255; ; SM90-NEXT: shl.b32 %r3, %r14, %r1; -; SM90-NEXT: ld.param.b8 %r15, [release_sys_i8_param_1]; +; SM90-NEXT: ld.param.u8 %r15, [release_sys_i8_param_1]; ; SM90-NEXT: shl.b32 %r4, %r15, %r1; -; SM90-NEXT: ld.b32 %r16, [%rd1]; +; SM90-NEXT: ld.u32 %r16, [%rd1]; ; SM90-NEXT: and.b32 %r20, %r16, %r2; ; SM90-NEXT: $L__BB2_1: // %partword.cmpxchg.loop ; SM90-NEXT: // =>This Inner Loop Header: Depth=1 @@ -466,7 +466,7 @@ define i8 @acq_rel_sys_i8(ptr %addr, i8 %cmp, i8 %new) { ; SM70-NEXT: // =>This Inner Loop Header: Depth=1 ; SM70-NEXT: or.b32 %r17, %r20, %r3; ; SM70-NEXT: or.b32 %r18, %r20, %r4; -; SM70-NEXT: atom.relaxed.cas.b32 %r7, [%rd1], %r18, %r17; +; SM70-NEXT: atom.relaxed.sys.cas.b32 %r7, [%rd1], %r18, %r17; ; SM70-NEXT: setp.eq.b32 %p1, %r7, %r18; ; SM70-NEXT: @%p1 bra $L__BB3_3; ; SM70-NEXT: // %bb.2: // %partword.cmpxchg.failure @@ -487,8 +487,8 @@ define i8 @acq_rel_sys_i8(ptr %addr, i8 %cmp, i8 %new) { ; SM90-NEXT: .reg .b64 %rd<3>; ; SM90-EMPTY: ; SM90-NEXT: // %bb.0: -; SM90-NEXT: ld.param.b8 %rs1, [acq_rel_sys_i8_param_2]; -; SM90-NEXT: ld.param.b64 %rd2, [acq_rel_sys_i8_param_0]; +; SM90-NEXT: ld.param.u8 %rs1, [acq_rel_sys_i8_param_2]; +; SM90-NEXT: ld.param.u64 %rd2, [acq_rel_sys_i8_param_0]; ; SM90-NEXT: fence.release.sys; ; SM90-NEXT: and.b64 %rd1, %rd2, -4; ; SM90-NEXT: cvt.u32.u64 %r9, %rd2; @@ -500,9 +500,9 @@ define i8 @acq_rel_sys_i8(ptr %addr, i8 %cmp, i8 %new) { ; SM90-NEXT: cvt.u32.u16 %r13, %rs1; ; SM90-NEXT: and.b32 %r14, %r13, 255; ; SM90-NEXT: shl.b32 %r3, %r14, %r1; -; SM90-NEXT: ld.param.b8 %r15, [acq_rel_sys_i8_param_1]; +; SM90-NEXT: ld.param.u8 %r15, [acq_rel_sys_i8_param_1]; ; SM90-NEXT: shl.b32 %r4, %r15, %r1; -; SM90-NEXT: ld.b32 %r16, [%rd1]; +; SM90-NEXT: ld.u32 %r16, [%rd1]; ; SM90-NEXT: and.b32 %r20, %r16, %r2; ; SM90-NEXT: $L__BB3_1: // %partword.cmpxchg.loop ; SM90-NEXT: // =>This Inner Loop Header: Depth=1 @@ -598,7 +598,7 @@ define i8 @seq_cst_sys_i8(ptr %addr, i8 %cmp, i8 %new) { ; SM70-NEXT: // =>This Inner Loop Header: Depth=1 ; SM70-NEXT: or.b32 %r17, %r20, %r3; ; SM70-NEXT: or.b32 %r18, %r20, %r4; -; SM70-NEXT: atom.relaxed.cas.b32 %r7, [%rd1], %r18, %r17; +; SM70-NEXT: atom.relaxed.sys.cas.b32 %r7, [%rd1], %r18, %r17; ; SM70-NEXT: setp.eq.b32 %p1, %r7, %r18; ; SM70-NEXT: @%p1 bra $L__BB4_3; ; SM70-NEXT: // %bb.2: // %partword.cmpxchg.failure @@ -619,8 +619,8 @@ define i8 @seq_cst_sys_i8(ptr %addr, i8 %cmp, i8 %new) { ; SM90-NEXT: .reg .b64 %rd<3>; ; SM90-EMPTY: ; SM90-NEXT: // %bb.0: -; SM90-NEXT: ld.param.b8 %rs1, [seq_cst_sys_i8_param_2]; -; SM90-NEXT: ld.param.b64 %rd2, [seq_cst_sys_i8_param_0]; +; SM90-NEXT: ld.param.u8 %rs1, [seq_cst_sys_i8_param_2]; +; SM90-NEXT: ld.param.u64 %rd2, [seq_cst_sys_i8_param_0]; ; SM90-NEXT: fence.sc.sys; ; SM90-NEXT: and.b64 %rd1, %rd2, -4; ; SM90-NEXT: cvt.u32.u64 %r9, %rd2; @@ -632,9 +632,9 @@ define i8 @seq_cst_sys_i8(ptr %addr, i8 %cmp, i8 %new) { ; SM90-NEXT: cvt.u32.u16 %r13, %rs1; ; SM90-NEXT: and.b32 %r14, %r13, 255; ; SM90-NEXT: shl.b32 %r3, %r14, %r1; -; SM90-NEXT: ld.param.b8 %r15, [seq_cst_sys_i8_param_1]; +; SM90-NEXT: ld.param.u8 %r15, [seq_cst_sys_i8_param_1]; ; SM90-NEXT: shl.b32 %r4, %r15, %r1; -; SM90-NEXT: ld.b32 %r16, [%rd1]; +; SM90-NEXT: ld.u32 %r16, [%rd1]; ; SM90-NEXT: and.b32 %r20, %r16, %r2; ; SM90-NEXT: $L__BB4_1: // %partword.cmpxchg.loop ; SM90-NEXT: // =>This Inner Loop Header: Depth=1 @@ -726,7 +726,7 @@ define i16 @relaxed_sys_i16(ptr %addr, i16 %cmp, i16 %new) { ; SM70-NEXT: // =>This Inner Loop Header: Depth=1 ; SM70-NEXT: or.b32 %r16, %r19, %r3; ; SM70-NEXT: or.b32 %r17, %r19, %r4; -; SM70-NEXT: atom.relaxed.cas.b32 %r7, [%rd1], %r17, %r16; +; SM70-NEXT: atom.relaxed.sys.cas.b32 %r7, [%rd1], %r17, %r16; ; SM70-NEXT: setp.eq.b32 %p1, %r7, %r17; ; SM70-NEXT: @%p1 bra $L__BB5_3; ; SM70-NEXT: // %bb.2: // %partword.cmpxchg.failure @@ -746,10 +746,10 @@ define i16 @relaxed_sys_i16(ptr %addr, i16 %cmp, i16 %new) { ; SM90-NEXT: .reg .b64 %rd<3>; ; SM90-EMPTY: ; SM90-NEXT: // %bb.0: -; SM90-NEXT: ld.param.b16 %rs1, [relaxed_sys_i16_param_2]; -; SM90-NEXT: ld.param.b64 %rd2, [relaxed_sys_i16_param_0]; +; SM90-NEXT: ld.param.u16 %rs1, [relaxed_sys_i16_param_2]; +; SM90-NEXT: ld.param.u64 %rd2, [relaxed_sys_i16_param_0]; ; SM90-NEXT: and.b64 %rd1, %rd2, -4; -; SM90-NEXT: ld.param.b16 %r9, [relaxed_sys_i16_param_1]; +; SM90-NEXT: ld.param.u16 %r9, [relaxed_sys_i16_param_1]; ; SM90-NEXT: cvt.u32.u64 %r10, %rd2; ; SM90-NEXT: and.b32 %r11, %r10, 3; ; SM90-NEXT: shl.b32 %r1, %r11, 3; @@ -759,7 +759,7 @@ define i16 @relaxed_sys_i16(ptr %addr, i16 %cmp, i16 %new) { ; SM90-NEXT: cvt.u32.u16 %r14, %rs1; ; SM90-NEXT: shl.b32 %r3, %r14, %r1; ; SM90-NEXT: shl.b32 %r4, %r9, %r1; -; SM90-NEXT: ld.b32 %r15, [%rd1]; +; SM90-NEXT: ld.u32 %r15, [%rd1]; ; SM90-NEXT: and.b32 %r19, %r15, %r2; ; SM90-NEXT: $L__BB5_1: // %partword.cmpxchg.loop ; SM90-NEXT: // =>This Inner Loop Header: Depth=1 @@ -850,7 +850,7 @@ define i16 @acquire_sys_i16(ptr %addr, i16 %cmp, i16 %new) { ; SM70-NEXT: // =>This Inner Loop Header: Depth=1 ; SM70-NEXT: or.b32 %r16, %r19, %r3; ; SM70-NEXT: or.b32 %r17, %r19, %r4; -; SM70-NEXT: atom.relaxed.cas.b32 %r7, [%rd1], %r17, %r16; +; SM70-NEXT: atom.relaxed.sys.cas.b32 %r7, [%rd1], %r17, %r16; ; SM70-NEXT: setp.eq.b32 %p1, %r7, %r17; ; SM70-NEXT: @%p1 bra $L__BB6_3; ; SM70-NEXT: // %bb.2: // %partword.cmpxchg.failure @@ -871,10 +871,10 @@ define i16 @acquire_sys_i16(ptr %addr, i16 %cmp, i16 %new) { ; SM90-NEXT: .reg .b64 %rd<3>; ; SM90-EMPTY: ; SM90-NEXT: // %bb.0: -; SM90-NEXT: ld.param.b16 %rs1, [acquire_sys_i16_param_2]; -; SM90-NEXT: ld.param.b64 %rd2, [acquire_sys_i16_param_0]; +; SM90-NEXT: ld.param.u16 %rs1, [acquire_sys_i16_param_2]; +; SM90-NEXT: ld.param.u64 %rd2, [acquire_sys_i16_param_0]; ; SM90-NEXT: and.b64 %rd1, %rd2, -4; -; SM90-NEXT: ld.param.b16 %r9, [acquire_sys_i16_param_1]; +; SM90-NEXT: ld.param.u16 %r9, [acquire_sys_i16_param_1]; ; SM90-NEXT: cvt.u32.u64 %r10, %rd2; ; SM90-NEXT: and.b32 %r11, %r10, 3; ; SM90-NEXT: shl.b32 %r1, %r11, 3; @@ -884,7 +884,7 @@ define i16 @acquire_sys_i16(ptr %addr, i16 %cmp, i16 %new) { ; SM90-NEXT: cvt.u32.u16 %r14, %rs1; ; SM90-NEXT: shl.b32 %r3, %r14, %r1; ; SM90-NEXT: shl.b32 %r4, %r9, %r1; -; SM90-NEXT: ld.b32 %r15, [%rd1]; +; SM90-NEXT: ld.u32 %r15, [%rd1]; ; SM90-NEXT: and.b32 %r19, %r15, %r2; ; SM90-NEXT: $L__BB6_1: // %partword.cmpxchg.loop ; SM90-NEXT: // =>This Inner Loop Header: Depth=1 @@ -977,7 +977,7 @@ define i16 @release_sys_i16(ptr %addr, i16 %cmp, i16 %new) { ; SM70-NEXT: // =>This Inner Loop Header: Depth=1 ; SM70-NEXT: or.b32 %r16, %r19, %r3; ; SM70-NEXT: or.b32 %r17, %r19, %r4; -; SM70-NEXT: atom.relaxed.cas.b32 %r7, [%rd1], %r17, %r16; +; SM70-NEXT: atom.relaxed.sys.cas.b32 %r7, [%rd1], %r17, %r16; ; SM70-NEXT: setp.eq.b32 %p1, %r7, %r17; ; SM70-NEXT: @%p1 bra $L__BB7_3; ; SM70-NEXT: // %bb.2: // %partword.cmpxchg.failure @@ -997,10 +997,10 @@ define i16 @release_sys_i16(ptr %addr, i16 %cmp, i16 %new) { ; SM90-NEXT: .reg .b64 %rd<3>; ; SM90-EMPTY: ; SM90-NEXT: // %bb.0: -; SM90-NEXT: ld.param.b16 %rs1, [release_sys_i16_param_2]; -; SM90-NEXT: ld.param.b64 %rd2, [release_sys_i16_param_0]; +; SM90-NEXT: ld.param.u16 %rs1, [release_sys_i16_param_2]; +; SM90-NEXT: ld.param.u64 %rd2, [release_sys_i16_param_0]; ; SM90-NEXT: fence.release.sys; -; SM90-NEXT: ld.param.b16 %r9, [release_sys_i16_param_1]; +; SM90-NEXT: ld.param.u16 %r9, [release_sys_i16_param_1]; ; SM90-NEXT: and.b64 %rd1, %rd2, -4; ; SM90-NEXT: cvt.u32.u64 %r10, %rd2; ; SM90-NEXT: and.b32 %r11, %r10, 3; @@ -1011,7 +1011,7 @@ define i16 @release_sys_i16(ptr %addr, i16 %cmp, i16 %new) { ; SM90-NEXT: cvt.u32.u16 %r14, %rs1; ; SM90-NEXT: shl.b32 %r3, %r14, %r1; ; SM90-NEXT: shl.b32 %r4, %r9, %r1; -; SM90-NEXT: ld.b32 %r15, [%rd1]; +; SM90-NEXT: ld.u32 %r15, [%rd1]; ; SM90-NEXT: and.b32 %r19, %r15, %r2; ; SM90-NEXT: $L__BB7_1: // %partword.cmpxchg.loop ; SM90-NEXT: // =>This Inner Loop Header: Depth=1 @@ -1104,7 +1104,7 @@ define i16 @acq_rel_sys_i16(ptr %addr, i16 %cmp, i16 %new) { ; SM70-NEXT: // =>This Inner Loop Header: Depth=1 ; SM70-NEXT: or.b32 %r16, %r19, %r3; ; SM70-NEXT: or.b32 %r17, %r19, %r4; -; SM70-NEXT: atom.relaxed.cas.b32 %r7, [%rd1], %r17, %r16; +; SM70-NEXT: atom.relaxed.sys.cas.b32 %r7, [%rd1], %r17, %r16; ; SM70-NEXT: setp.eq.b32 %p1, %r7, %r17; ; SM70-NEXT: @%p1 bra $L__BB8_3; ; SM70-NEXT: // %bb.2: // %partword.cmpxchg.failure @@ -1125,10 +1125,10 @@ define i16 @acq_rel_sys_i16(ptr %addr, i16 %cmp, i16 %new) { ; SM90-NEXT: .reg .b64 %rd<3>; ; SM90-EMPTY: ; SM90-NEXT: // %bb.0: -; SM90-NEXT: ld.param.b16 %rs1, [acq_rel_sys_i16_param_2]; -; SM90-NEXT: ld.param.b64 %rd2, [acq_rel_sys_i16_param_0]; +; SM90-NEXT: ld.param.u16 %rs1, [acq_rel_sys_i16_param_2]; +; SM90-NEXT: ld.param.u64 %rd2, [acq_rel_sys_i16_param_0]; ; SM90-NEXT: fence.release.sys; -; SM90-NEXT: ld.param.b16 %r9, [acq_rel_sys_i16_param_1]; +; SM90-NEXT: ld.param.u16 %r9, [acq_rel_sys_i16_param_1]; ; SM90-NEXT: and.b64 %rd1, %rd2, -4; ; SM90-NEXT: cvt.u32.u64 %r10, %rd2; ; SM90-NEXT: and.b32 %r11, %r10, 3; @@ -1139,7 +1139,7 @@ define i16 @acq_rel_sys_i16(ptr %addr, i16 %cmp, i16 %new) { ; SM90-NEXT: cvt.u32.u16 %r14, %rs1; ; SM90-NEXT: shl.b32 %r3, %r14, %r1; ; SM90-NEXT: shl.b32 %r4, %r9, %r1; -; SM90-NEXT: ld.b32 %r15, [%rd1]; +; SM90-NEXT: ld.u32 %r15, [%rd1]; ; SM90-NEXT: and.b32 %r19, %r15, %r2; ; SM90-NEXT: $L__BB8_1: // %partword.cmpxchg.loop ; SM90-NEXT: // =>This Inner Loop Header: Depth=1 @@ -1234,7 +1234,7 @@ define i16 @seq_cst_sys_i16(ptr %addr, i16 %cmp, i16 %new) { ; SM70-NEXT: // =>This Inner Loop Header: Depth=1 ; SM70-NEXT: or.b32 %r16, %r19, %r3; ; SM70-NEXT: or.b32 %r17, %r19, %r4; -; SM70-NEXT: atom.relaxed.cas.b32 %r7, [%rd1], %r17, %r16; +; SM70-NEXT: atom.relaxed.sys.cas.b32 %r7, [%rd1], %r17, %r16; ; SM70-NEXT: setp.eq.b32 %p1, %r7, %r17; ; SM70-NEXT: @%p1 bra $L__BB9_3; ; SM70-NEXT: // %bb.2: // %partword.cmpxchg.failure @@ -1255,10 +1255,10 @@ define i16 @seq_cst_sys_i16(ptr %addr, i16 %cmp, i16 %new) { ; SM90-NEXT: .reg .b64 %rd<3>; ; SM90-EMPTY: ; SM90-NEXT: // %bb.0: -; SM90-NEXT: ld.param.b16 %rs1, [seq_cst_sys_i16_param_2]; -; SM90-NEXT: ld.param.b64 %rd2, [seq_cst_sys_i16_param_0]; +; SM90-NEXT: ld.param.u16 %rs1, [seq_cst_sys_i16_param_2]; +; SM90-NEXT: ld.param.u64 %rd2, [seq_cst_sys_i16_param_0]; ; SM90-NEXT: fence.sc.sys; -; SM90-NEXT: ld.param.b16 %r9, [seq_cst_sys_i16_param_1]; +; SM90-NEXT: ld.param.u16 %r9, [seq_cst_sys_i16_param_1]; ; SM90-NEXT: and.b64 %rd1, %rd2, -4; ; SM90-NEXT: cvt.u32.u64 %r10, %rd2; ; SM90-NEXT: and.b32 %r11, %r10, 3; @@ -1269,7 +1269,7 @@ define i16 @seq_cst_sys_i16(ptr %addr, i16 %cmp, i16 %new) { ; SM90-NEXT: cvt.u32.u16 %r14, %rs1; ; SM90-NEXT: shl.b32 %r3, %r14, %r1; ; SM90-NEXT: shl.b32 %r4, %r9, %r1; -; SM90-NEXT: ld.b32 %r15, [%rd1]; +; SM90-NEXT: ld.u32 %r15, [%rd1]; ; SM90-NEXT: and.b32 %r19, %r15, %r2; ; SM90-NEXT: $L__BB9_1: // %partword.cmpxchg.loop ; SM90-NEXT: // =>This Inner Loop Header: Depth=1 @@ -1316,7 +1316,7 @@ define i32 @relaxed_sys_i32(ptr %addr, i32 %cmp, i32 %new) { ; SM70-NEXT: ld.param.b64 %rd1, [relaxed_sys_i32_param_0]; ; SM70-NEXT: ld.param.b32 %r1, [relaxed_sys_i32_param_1]; ; SM70-NEXT: ld.param.b32 %r2, [relaxed_sys_i32_param_2]; -; SM70-NEXT: atom.relaxed.cas.b32 %r3, [%rd1], %r1, %r2; +; SM70-NEXT: atom.relaxed.sys.cas.b32 %r3, [%rd1], %r1, %r2; ; SM70-NEXT: st.param.b32 [func_retval0], %r2; ; SM70-NEXT: ret; ; SM90-LABEL: relaxed_sys_i32( @@ -1325,9 +1325,9 @@ define i32 @relaxed_sys_i32(ptr %addr, i32 %cmp, i32 %new) { ; SM90-NEXT: .reg .b64 %rd<2>; ; SM90-EMPTY: ; SM90-NEXT: // %bb.0: -; SM90-NEXT: ld.param.b64 %rd1, [relaxed_sys_i32_param_0]; -; SM90-NEXT: ld.param.b32 %r1, [relaxed_sys_i32_param_1]; -; SM90-NEXT: ld.param.b32 %r2, [relaxed_sys_i32_param_2]; +; SM90-NEXT: ld.param.u64 %rd1, [relaxed_sys_i32_param_0]; +; SM90-NEXT: ld.param.u32 %r1, [relaxed_sys_i32_param_1]; +; SM90-NEXT: ld.param.u32 %r2, [relaxed_sys_i32_param_2]; ; SM90-NEXT: atom.cas.b32 %r3, [%rd1], %r1, %r2; ; SM90-NEXT: st.param.b32 [func_retval0], %r2; ; SM90-NEXT: ret; @@ -1358,7 +1358,7 @@ define i32 @acq_rel_sys_i32(ptr %addr, i32 %cmp, i32 %new) { ; SM70-NEXT: ld.param.b64 %rd1, [acq_rel_sys_i32_param_0]; ; SM70-NEXT: ld.param.b32 %r1, [acq_rel_sys_i32_param_1]; ; SM70-NEXT: ld.param.b32 %r2, [acq_rel_sys_i32_param_2]; -; SM70-NEXT: atom.acq_rel.cas.b32 %r3, [%rd1], %r1, %r2; +; SM70-NEXT: atom.acq_rel.sys.cas.b32 %r3, [%rd1], %r1, %r2; ; SM70-NEXT: st.param.b32 [func_retval0], %r2; ; SM70-NEXT: ret; ; SM90-LABEL: acq_rel_sys_i32( @@ -1367,9 +1367,9 @@ define i32 @acq_rel_sys_i32(ptr %addr, i32 %cmp, i32 %new) { ; SM90-NEXT: .reg .b64 %rd<2>; ; SM90-EMPTY: ; SM90-NEXT: // %bb.0: -; SM90-NEXT: ld.param.b64 %rd1, [acq_rel_sys_i32_param_0]; -; SM90-NEXT: ld.param.b32 %r1, [acq_rel_sys_i32_param_1]; -; SM90-NEXT: ld.param.b32 %r2, [acq_rel_sys_i32_param_2]; +; SM90-NEXT: ld.param.u64 %rd1, [acq_rel_sys_i32_param_0]; +; SM90-NEXT: ld.param.u32 %r1, [acq_rel_sys_i32_param_1]; +; SM90-NEXT: ld.param.u32 %r2, [acq_rel_sys_i32_param_2]; ; SM90-NEXT: atom.acq_rel.cas.b32 %r3, [%rd1], %r1, %r2; ; SM90-NEXT: st.param.b32 [func_retval0], %r2; ; SM90-NEXT: ret; @@ -1400,7 +1400,7 @@ define i32 @acquire_sys_i32(ptr %addr, i32 %cmp, i32 %new) { ; SM70-NEXT: ld.param.b64 %rd1, [acquire_sys_i32_param_0]; ; SM70-NEXT: ld.param.b32 %r1, [acquire_sys_i32_param_1]; ; SM70-NEXT: ld.param.b32 %r2, [acquire_sys_i32_param_2]; -; SM70-NEXT: atom.acquire.cas.b32 %r3, [%rd1], %r1, %r2; +; SM70-NEXT: atom.acquire.sys.cas.b32 %r3, [%rd1], %r1, %r2; ; SM70-NEXT: st.param.b32 [func_retval0], %r2; ; SM70-NEXT: ret; ; SM90-LABEL: acquire_sys_i32( @@ -1409,9 +1409,9 @@ define i32 @acquire_sys_i32(ptr %addr, i32 %cmp, i32 %new) { ; SM90-NEXT: .reg .b64 %rd<2>; ; SM90-EMPTY: ; SM90-NEXT: // %bb.0: -; SM90-NEXT: ld.param.b64 %rd1, [acquire_sys_i32_param_0]; -; SM90-NEXT: ld.param.b32 %r1, [acquire_sys_i32_param_1]; -; SM90-NEXT: ld.param.b32 %r2, [acquire_sys_i32_param_2]; +; SM90-NEXT: ld.param.u64 %rd1, [acquire_sys_i32_param_0]; +; SM90-NEXT: ld.param.u32 %r1, [acquire_sys_i32_param_1]; +; SM90-NEXT: ld.param.u32 %r2, [acquire_sys_i32_param_2]; ; SM90-NEXT: atom.acquire.cas.b32 %r3, [%rd1], %r1, %r2; ; SM90-NEXT: st.param.b32 [func_retval0], %r2; ; SM90-NEXT: ret; @@ -1442,7 +1442,7 @@ define i32 @release_sys_i32(ptr %addr, i32 %cmp, i32 %new) { ; SM70-NEXT: ld.param.b64 %rd1, [release_sys_i32_param_0]; ; SM70-NEXT: ld.param.b32 %r1, [release_sys_i32_param_1]; ; SM70-NEXT: ld.param.b32 %r2, [release_sys_i32_param_2]; -; SM70-NEXT: atom.release.cas.b32 %r3, [%rd1], %r1, %r2; +; SM70-NEXT: atom.release.sys.cas.b32 %r3, [%rd1], %r1, %r2; ; SM70-NEXT: st.param.b32 [func_retval0], %r2; ; SM70-NEXT: ret; ; SM90-LABEL: release_sys_i32( @@ -1451,9 +1451,9 @@ define i32 @release_sys_i32(ptr %addr, i32 %cmp, i32 %new) { ; SM90-NEXT: .reg .b64 %rd<2>; ; SM90-EMPTY: ; SM90-NEXT: // %bb.0: -; SM90-NEXT: ld.param.b64 %rd1, [release_sys_i32_param_0]; -; SM90-NEXT: ld.param.b32 %r1, [release_sys_i32_param_1]; -; SM90-NEXT: ld.param.b32 %r2, [release_sys_i32_param_2]; +; SM90-NEXT: ld.param.u64 %rd1, [release_sys_i32_param_0]; +; SM90-NEXT: ld.param.u32 %r1, [release_sys_i32_param_1]; +; SM90-NEXT: ld.param.u32 %r2, [release_sys_i32_param_2]; ; SM90-NEXT: atom.release.cas.b32 %r3, [%rd1], %r1, %r2; ; SM90-NEXT: st.param.b32 [func_retval0], %r2; ; SM90-NEXT: ret; @@ -1486,7 +1486,7 @@ define i32 @seq_cst_sys_i32(ptr %addr, i32 %cmp, i32 %new) { ; SM70-NEXT: fence.sc.sys; ; SM70-NEXT: ld.param.b32 %r1, [seq_cst_sys_i32_param_1]; ; SM70-NEXT: ld.param.b32 %r2, [seq_cst_sys_i32_param_2]; -; SM70-NEXT: atom.acquire.cas.b32 %r3, [%rd1], %r1, %r2; +; SM70-NEXT: atom.acquire.sys.cas.b32 %r3, [%rd1], %r1, %r2; ; SM70-NEXT: st.param.b32 [func_retval0], %r2; ; SM70-NEXT: ret; ; SM90-LABEL: seq_cst_sys_i32( @@ -1495,10 +1495,10 @@ define i32 @seq_cst_sys_i32(ptr %addr, i32 %cmp, i32 %new) { ; SM90-NEXT: .reg .b64 %rd<2>; ; SM90-EMPTY: ; SM90-NEXT: // %bb.0: -; SM90-NEXT: ld.param.b64 %rd1, [seq_cst_sys_i32_param_0]; +; SM90-NEXT: ld.param.u64 %rd1, [seq_cst_sys_i32_param_0]; ; SM90-NEXT: fence.sc.sys; -; SM90-NEXT: ld.param.b32 %r1, [seq_cst_sys_i32_param_1]; -; SM90-NEXT: ld.param.b32 %r2, [seq_cst_sys_i32_param_2]; +; SM90-NEXT: ld.param.u32 %r1, [seq_cst_sys_i32_param_1]; +; SM90-NEXT: ld.param.u32 %r2, [seq_cst_sys_i32_param_2]; ; SM90-NEXT: atom.cas.b32 %r3, [%rd1], %r1, %r2; ; SM90-NEXT: fence.acquire.sys; ; SM90-NEXT: st.param.b32 [func_retval0], %r2; @@ -1529,7 +1529,7 @@ define i64 @relaxed_sys_i64(ptr %addr, i64 %cmp, i64 %new) { ; SM70-NEXT: ld.param.b64 %rd1, [relaxed_sys_i64_param_0]; ; SM70-NEXT: ld.param.b64 %rd2, [relaxed_sys_i64_param_1]; ; SM70-NEXT: ld.param.b64 %rd3, [relaxed_sys_i64_param_2]; -; SM70-NEXT: atom.relaxed.cas.b64 %rd4, [%rd1], %rd2, %rd3; +; SM70-NEXT: atom.relaxed.sys.cas.b64 %rd4, [%rd1], %rd2, %rd3; ; SM70-NEXT: st.param.b64 [func_retval0], %rd3; ; SM70-NEXT: ret; ; SM90-LABEL: relaxed_sys_i64( @@ -1537,9 +1537,9 @@ define i64 @relaxed_sys_i64(ptr %addr, i64 %cmp, i64 %new) { ; SM90-NEXT: .reg .b64 %rd<5>; ; SM90-EMPTY: ; SM90-NEXT: // %bb.0: -; SM90-NEXT: ld.param.b64 %rd1, [relaxed_sys_i64_param_0]; -; SM90-NEXT: ld.param.b64 %rd2, [relaxed_sys_i64_param_1]; -; SM90-NEXT: ld.param.b64 %rd3, [relaxed_sys_i64_param_2]; +; SM90-NEXT: ld.param.u64 %rd1, [relaxed_sys_i64_param_0]; +; SM90-NEXT: ld.param.u64 %rd2, [relaxed_sys_i64_param_1]; +; SM90-NEXT: ld.param.u64 %rd3, [relaxed_sys_i64_param_2]; ; SM90-NEXT: atom.cas.b64 %rd4, [%rd1], %rd2, %rd3; ; SM90-NEXT: st.param.b64 [func_retval0], %rd3; ; SM90-NEXT: ret; @@ -1568,7 +1568,7 @@ define i64 @acquire_sys_i64(ptr %addr, i64 %cmp, i64 %new) { ; SM70-NEXT: ld.param.b64 %rd1, [acquire_sys_i64_param_0]; ; SM70-NEXT: ld.param.b64 %rd2, [acquire_sys_i64_param_1]; ; SM70-NEXT: ld.param.b64 %rd3, [acquire_sys_i64_param_2]; -; SM70-NEXT: atom.acquire.cas.b64 %rd4, [%rd1], %rd2, %rd3; +; SM70-NEXT: atom.acquire.sys.cas.b64 %rd4, [%rd1], %rd2, %rd3; ; SM70-NEXT: st.param.b64 [func_retval0], %rd3; ; SM70-NEXT: ret; ; SM90-LABEL: acquire_sys_i64( @@ -1576,9 +1576,9 @@ define i64 @acquire_sys_i64(ptr %addr, i64 %cmp, i64 %new) { ; SM90-NEXT: .reg .b64 %rd<5>; ; SM90-EMPTY: ; SM90-NEXT: // %bb.0: -; SM90-NEXT: ld.param.b64 %rd1, [acquire_sys_i64_param_0]; -; SM90-NEXT: ld.param.b64 %rd2, [acquire_sys_i64_param_1]; -; SM90-NEXT: ld.param.b64 %rd3, [acquire_sys_i64_param_2]; +; SM90-NEXT: ld.param.u64 %rd1, [acquire_sys_i64_param_0]; +; SM90-NEXT: ld.param.u64 %rd2, [acquire_sys_i64_param_1]; +; SM90-NEXT: ld.param.u64 %rd3, [acquire_sys_i64_param_2]; ; SM90-NEXT: atom.acquire.cas.b64 %rd4, [%rd1], %rd2, %rd3; ; SM90-NEXT: st.param.b64 [func_retval0], %rd3; ; SM90-NEXT: ret; @@ -1607,7 +1607,7 @@ define i64 @acq_rel_sys_i64(ptr %addr, i64 %cmp, i64 %new) { ; SM70-NEXT: ld.param.b64 %rd1, [acq_rel_sys_i64_param_0]; ; SM70-NEXT: ld.param.b64 %rd2, [acq_rel_sys_i64_param_1]; ; SM70-NEXT: ld.param.b64 %rd3, [acq_rel_sys_i64_param_2]; -; SM70-NEXT: atom.acq_rel.cas.b64 %rd4, [%rd1], %rd2, %rd3; +; SM70-NEXT: atom.acq_rel.sys.cas.b64 %rd4, [%rd1], %rd2, %rd3; ; SM70-NEXT: st.param.b64 [func_retval0], %rd3; ; SM70-NEXT: ret; ; SM90-LABEL: acq_rel_sys_i64( @@ -1615,9 +1615,9 @@ define i64 @acq_rel_sys_i64(ptr %addr, i64 %cmp, i64 %new) { ; SM90-NEXT: .reg .b64 %rd<5>; ; SM90-EMPTY: ; SM90-NEXT: // %bb.0: -; SM90-NEXT: ld.param.b64 %rd1, [acq_rel_sys_i64_param_0]; -; SM90-NEXT: ld.param.b64 %rd2, [acq_rel_sys_i64_param_1]; -; SM90-NEXT: ld.param.b64 %rd3, [acq_rel_sys_i64_param_2]; +; SM90-NEXT: ld.param.u64 %rd1, [acq_rel_sys_i64_param_0]; +; SM90-NEXT: ld.param.u64 %rd2, [acq_rel_sys_i64_param_1]; +; SM90-NEXT: ld.param.u64 %rd3, [acq_rel_sys_i64_param_2]; ; SM90-NEXT: atom.acq_rel.cas.b64 %rd4, [%rd1], %rd2, %rd3; ; SM90-NEXT: st.param.b64 [func_retval0], %rd3; ; SM90-NEXT: ret; @@ -1646,7 +1646,7 @@ define i64 @release_sys_i64(ptr %addr, i64 %cmp, i64 %new) { ; SM70-NEXT: ld.param.b64 %rd1, [release_sys_i64_param_0]; ; SM70-NEXT: ld.param.b64 %rd2, [release_sys_i64_param_1]; ; SM70-NEXT: ld.param.b64 %rd3, [release_sys_i64_param_2]; -; SM70-NEXT: atom.release.cas.b64 %rd4, [%rd1], %rd2, %rd3; +; SM70-NEXT: atom.release.sys.cas.b64 %rd4, [%rd1], %rd2, %rd3; ; SM70-NEXT: st.param.b64 [func_retval0], %rd3; ; SM70-NEXT: ret; ; SM90-LABEL: release_sys_i64( @@ -1654,9 +1654,9 @@ define i64 @release_sys_i64(ptr %addr, i64 %cmp, i64 %new) { ; SM90-NEXT: .reg .b64 %rd<5>; ; SM90-EMPTY: ; SM90-NEXT: // %bb.0: -; SM90-NEXT: ld.param.b64 %rd1, [release_sys_i64_param_0]; -; SM90-NEXT: ld.param.b64 %rd2, [release_sys_i64_param_1]; -; SM90-NEXT: ld.param.b64 %rd3, [release_sys_i64_param_2]; +; SM90-NEXT: ld.param.u64 %rd1, [release_sys_i64_param_0]; +; SM90-NEXT: ld.param.u64 %rd2, [release_sys_i64_param_1]; +; SM90-NEXT: ld.param.u64 %rd3, [release_sys_i64_param_2]; ; SM90-NEXT: atom.release.cas.b64 %rd4, [%rd1], %rd2, %rd3; ; SM90-NEXT: st.param.b64 [func_retval0], %rd3; ; SM90-NEXT: ret; @@ -1687,7 +1687,7 @@ define i64 @seq_cst_sys_i64(ptr %addr, i64 %cmp, i64 %new) { ; SM70-NEXT: fence.sc.sys; ; SM70-NEXT: ld.param.b64 %rd2, [seq_cst_sys_i64_param_1]; ; SM70-NEXT: ld.param.b64 %rd3, [seq_cst_sys_i64_param_2]; -; SM70-NEXT: atom.acquire.cas.b64 %rd4, [%rd1], %rd2, %rd3; +; SM70-NEXT: atom.acquire.sys.cas.b64 %rd4, [%rd1], %rd2, %rd3; ; SM70-NEXT: st.param.b64 [func_retval0], %rd3; ; SM70-NEXT: ret; ; SM90-LABEL: seq_cst_sys_i64( @@ -1695,10 +1695,10 @@ define i64 @seq_cst_sys_i64(ptr %addr, i64 %cmp, i64 %new) { ; SM90-NEXT: .reg .b64 %rd<5>; ; SM90-EMPTY: ; SM90-NEXT: // %bb.0: -; SM90-NEXT: ld.param.b64 %rd1, [seq_cst_sys_i64_param_0]; +; SM90-NEXT: ld.param.u64 %rd1, [seq_cst_sys_i64_param_0]; ; SM90-NEXT: fence.sc.sys; -; SM90-NEXT: ld.param.b64 %rd2, [seq_cst_sys_i64_param_1]; -; SM90-NEXT: ld.param.b64 %rd3, [seq_cst_sys_i64_param_2]; +; SM90-NEXT: ld.param.u64 %rd2, [seq_cst_sys_i64_param_1]; +; SM90-NEXT: ld.param.u64 %rd3, [seq_cst_sys_i64_param_2]; ; SM90-NEXT: atom.cas.b64 %rd4, [%rd1], %rd2, %rd3; ; SM90-NEXT: fence.acquire.sys; ; SM90-NEXT: st.param.b64 [func_retval0], %rd3; |