2105 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -march=nvptx64 -mcpu=sm_60 -mattr=+ptx50 | FileCheck %s --check-prefix=SM603; RUN: %if ptxas-sm_60 && ptxas-isa-5.0 %{ llc < %s -march=nvptx64 -mcpu=sm_60 -mattr=+ptx50 | %ptxas-verify -arch=sm_60 %}4 5define i8 @monotonic_monotonic_i8_global_cta(ptr addrspace(1) %addr, i8 %cmp, i8 %new) {6; SM60-LABEL: monotonic_monotonic_i8_global_cta(7; SM60: {8; SM60-NEXT: .reg .pred %p<3>;9; SM60-NEXT: .reg .b16 %rs<2>;10; SM60-NEXT: .reg .b32 %r<17>;11; SM60-NEXT: .reg .b64 %rd<3>;12; SM60-EMPTY:13; SM60-NEXT: // %bb.0:14; SM60-NEXT: ld.param.b8 %rs1, [monotonic_monotonic_i8_global_cta_param_2];15; SM60-NEXT: ld.param.b64 %rd2, [monotonic_monotonic_i8_global_cta_param_0];16; SM60-NEXT: and.b64 %rd1, %rd2, -4;17; SM60-NEXT: ld.param.b8 %r7, [monotonic_monotonic_i8_global_cta_param_1];18; SM60-NEXT: cvt.u32.u64 %r8, %rd2;19; SM60-NEXT: and.b32 %r9, %r8, 3;20; SM60-NEXT: shl.b32 %r1, %r9, 3;21; SM60-NEXT: mov.b32 %r10, 255;22; SM60-NEXT: shl.b32 %r11, %r10, %r1;23; SM60-NEXT: not.b32 %r2, %r11;24; SM60-NEXT: cvt.u32.u16 %r12, %rs1;25; SM60-NEXT: shl.b32 %r3, %r12, %r1;26; SM60-NEXT: shl.b32 %r4, %r7, %r1;27; SM60-NEXT: ld.global.b32 %r13, [%rd1];28; SM60-NEXT: and.b32 %r16, %r13, %r2;29; SM60-NEXT: $L__BB0_1: // %partword.cmpxchg.loop30; SM60-NEXT: // =>This Inner Loop Header: Depth=131; SM60-NEXT: or.b32 %r14, %r16, %r3;32; SM60-NEXT: or.b32 %r15, %r16, %r4;33; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;34; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;35; SM60-NEXT: @%p1 bra $L__BB0_3;36; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure37; SM60-NEXT: // in Loop: Header=BB0_1 Depth=138; SM60-NEXT: and.b32 %r6, %r5, %r2;39; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;40; SM60-NEXT: mov.b32 %r16, %r6;41; SM60-NEXT: @%p2 bra $L__BB0_1;42; SM60-NEXT: $L__BB0_3: // %partword.cmpxchg.end43; SM60-NEXT: st.param.b32 [func_retval0], %r12;44; SM60-NEXT: ret;45 %pairold = cmpxchg ptr addrspace(1) %addr, i8 %cmp, i8 %new syncscope("block") monotonic monotonic46 ret i8 %new47}48 49define i8 @monotonic_acquire_i8_global_cta(ptr addrspace(1) %addr, i8 %cmp, i8 %new) {50; SM60-LABEL: monotonic_acquire_i8_global_cta(51; SM60: {52; SM60-NEXT: .reg .pred %p<3>;53; SM60-NEXT: .reg .b16 %rs<2>;54; SM60-NEXT: .reg .b32 %r<17>;55; SM60-NEXT: .reg .b64 %rd<3>;56; SM60-EMPTY:57; SM60-NEXT: // %bb.0:58; SM60-NEXT: ld.param.b8 %rs1, [monotonic_acquire_i8_global_cta_param_2];59; SM60-NEXT: ld.param.b64 %rd2, [monotonic_acquire_i8_global_cta_param_0];60; SM60-NEXT: and.b64 %rd1, %rd2, -4;61; SM60-NEXT: ld.param.b8 %r7, [monotonic_acquire_i8_global_cta_param_1];62; SM60-NEXT: cvt.u32.u64 %r8, %rd2;63; SM60-NEXT: and.b32 %r9, %r8, 3;64; SM60-NEXT: shl.b32 %r1, %r9, 3;65; SM60-NEXT: mov.b32 %r10, 255;66; SM60-NEXT: shl.b32 %r11, %r10, %r1;67; SM60-NEXT: not.b32 %r2, %r11;68; SM60-NEXT: cvt.u32.u16 %r12, %rs1;69; SM60-NEXT: shl.b32 %r3, %r12, %r1;70; SM60-NEXT: shl.b32 %r4, %r7, %r1;71; SM60-NEXT: ld.global.b32 %r13, [%rd1];72; SM60-NEXT: and.b32 %r16, %r13, %r2;73; SM60-NEXT: $L__BB1_1: // %partword.cmpxchg.loop74; SM60-NEXT: // =>This Inner Loop Header: Depth=175; SM60-NEXT: or.b32 %r14, %r16, %r3;76; SM60-NEXT: or.b32 %r15, %r16, %r4;77; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;78; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;79; SM60-NEXT: @%p1 bra $L__BB1_3;80; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure81; SM60-NEXT: // in Loop: Header=BB1_1 Depth=182; SM60-NEXT: and.b32 %r6, %r5, %r2;83; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;84; SM60-NEXT: mov.b32 %r16, %r6;85; SM60-NEXT: @%p2 bra $L__BB1_1;86; SM60-NEXT: $L__BB1_3: // %partword.cmpxchg.end87; SM60-NEXT: membar.cta;88; SM60-NEXT: st.param.b32 [func_retval0], %r12;89; SM60-NEXT: ret;90 %pairold = cmpxchg ptr addrspace(1) %addr, i8 %cmp, i8 %new syncscope("block") monotonic acquire91 ret i8 %new92}93 94define i8 @monotonic_seq_cst_i8_global_cta(ptr addrspace(1) %addr, i8 %cmp, i8 %new) {95; SM60-LABEL: monotonic_seq_cst_i8_global_cta(96; SM60: {97; SM60-NEXT: .reg .pred %p<3>;98; SM60-NEXT: .reg .b16 %rs<2>;99; SM60-NEXT: .reg .b32 %r<17>;100; SM60-NEXT: .reg .b64 %rd<3>;101; SM60-EMPTY:102; SM60-NEXT: // %bb.0:103; SM60-NEXT: ld.param.b8 %rs1, [monotonic_seq_cst_i8_global_cta_param_2];104; SM60-NEXT: ld.param.b64 %rd2, [monotonic_seq_cst_i8_global_cta_param_0];105; SM60-NEXT: membar.cta;106; SM60-NEXT: ld.param.b8 %r7, [monotonic_seq_cst_i8_global_cta_param_1];107; SM60-NEXT: and.b64 %rd1, %rd2, -4;108; SM60-NEXT: cvt.u32.u64 %r8, %rd2;109; SM60-NEXT: and.b32 %r9, %r8, 3;110; SM60-NEXT: shl.b32 %r1, %r9, 3;111; SM60-NEXT: mov.b32 %r10, 255;112; SM60-NEXT: shl.b32 %r11, %r10, %r1;113; SM60-NEXT: not.b32 %r2, %r11;114; SM60-NEXT: cvt.u32.u16 %r12, %rs1;115; SM60-NEXT: shl.b32 %r3, %r12, %r1;116; SM60-NEXT: shl.b32 %r4, %r7, %r1;117; SM60-NEXT: ld.global.b32 %r13, [%rd1];118; SM60-NEXT: and.b32 %r16, %r13, %r2;119; SM60-NEXT: $L__BB2_1: // %partword.cmpxchg.loop120; SM60-NEXT: // =>This Inner Loop Header: Depth=1121; SM60-NEXT: or.b32 %r14, %r16, %r3;122; SM60-NEXT: or.b32 %r15, %r16, %r4;123; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;124; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;125; SM60-NEXT: @%p1 bra $L__BB2_3;126; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure127; SM60-NEXT: // in Loop: Header=BB2_1 Depth=1128; SM60-NEXT: and.b32 %r6, %r5, %r2;129; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;130; SM60-NEXT: mov.b32 %r16, %r6;131; SM60-NEXT: @%p2 bra $L__BB2_1;132; SM60-NEXT: $L__BB2_3: // %partword.cmpxchg.end133; SM60-NEXT: membar.cta;134; SM60-NEXT: st.param.b32 [func_retval0], %r12;135; SM60-NEXT: ret;136 %pairold = cmpxchg ptr addrspace(1) %addr, i8 %cmp, i8 %new syncscope("block") monotonic seq_cst137 ret i8 %new138}139 140define i8 @acquire_monotonic_i8_global_cta(ptr addrspace(1) %addr, i8 %cmp, i8 %new) {141; SM60-LABEL: acquire_monotonic_i8_global_cta(142; SM60: {143; SM60-NEXT: .reg .pred %p<3>;144; SM60-NEXT: .reg .b16 %rs<2>;145; SM60-NEXT: .reg .b32 %r<17>;146; SM60-NEXT: .reg .b64 %rd<3>;147; SM60-EMPTY:148; SM60-NEXT: // %bb.0:149; SM60-NEXT: ld.param.b8 %rs1, [acquire_monotonic_i8_global_cta_param_2];150; SM60-NEXT: ld.param.b64 %rd2, [acquire_monotonic_i8_global_cta_param_0];151; SM60-NEXT: and.b64 %rd1, %rd2, -4;152; SM60-NEXT: ld.param.b8 %r7, [acquire_monotonic_i8_global_cta_param_1];153; SM60-NEXT: cvt.u32.u64 %r8, %rd2;154; SM60-NEXT: and.b32 %r9, %r8, 3;155; SM60-NEXT: shl.b32 %r1, %r9, 3;156; SM60-NEXT: mov.b32 %r10, 255;157; SM60-NEXT: shl.b32 %r11, %r10, %r1;158; SM60-NEXT: not.b32 %r2, %r11;159; SM60-NEXT: cvt.u32.u16 %r12, %rs1;160; SM60-NEXT: shl.b32 %r3, %r12, %r1;161; SM60-NEXT: shl.b32 %r4, %r7, %r1;162; SM60-NEXT: ld.global.b32 %r13, [%rd1];163; SM60-NEXT: and.b32 %r16, %r13, %r2;164; SM60-NEXT: $L__BB3_1: // %partword.cmpxchg.loop165; SM60-NEXT: // =>This Inner Loop Header: Depth=1166; SM60-NEXT: or.b32 %r14, %r16, %r3;167; SM60-NEXT: or.b32 %r15, %r16, %r4;168; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;169; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;170; SM60-NEXT: @%p1 bra $L__BB3_3;171; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure172; SM60-NEXT: // in Loop: Header=BB3_1 Depth=1173; SM60-NEXT: and.b32 %r6, %r5, %r2;174; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;175; SM60-NEXT: mov.b32 %r16, %r6;176; SM60-NEXT: @%p2 bra $L__BB3_1;177; SM60-NEXT: $L__BB3_3: // %partword.cmpxchg.end178; SM60-NEXT: membar.cta;179; SM60-NEXT: st.param.b32 [func_retval0], %r12;180; SM60-NEXT: ret;181 %pairold = cmpxchg ptr addrspace(1) %addr, i8 %cmp, i8 %new syncscope("block") acquire monotonic182 ret i8 %new183}184 185define i8 @acquire_acquire_i8_global_cta(ptr addrspace(1) %addr, i8 %cmp, i8 %new) {186; SM60-LABEL: acquire_acquire_i8_global_cta(187; SM60: {188; SM60-NEXT: .reg .pred %p<3>;189; SM60-NEXT: .reg .b16 %rs<2>;190; SM60-NEXT: .reg .b32 %r<17>;191; SM60-NEXT: .reg .b64 %rd<3>;192; SM60-EMPTY:193; SM60-NEXT: // %bb.0:194; SM60-NEXT: ld.param.b8 %rs1, [acquire_acquire_i8_global_cta_param_2];195; SM60-NEXT: ld.param.b64 %rd2, [acquire_acquire_i8_global_cta_param_0];196; SM60-NEXT: and.b64 %rd1, %rd2, -4;197; SM60-NEXT: ld.param.b8 %r7, [acquire_acquire_i8_global_cta_param_1];198; SM60-NEXT: cvt.u32.u64 %r8, %rd2;199; SM60-NEXT: and.b32 %r9, %r8, 3;200; SM60-NEXT: shl.b32 %r1, %r9, 3;201; SM60-NEXT: mov.b32 %r10, 255;202; SM60-NEXT: shl.b32 %r11, %r10, %r1;203; SM60-NEXT: not.b32 %r2, %r11;204; SM60-NEXT: cvt.u32.u16 %r12, %rs1;205; SM60-NEXT: shl.b32 %r3, %r12, %r1;206; SM60-NEXT: shl.b32 %r4, %r7, %r1;207; SM60-NEXT: ld.global.b32 %r13, [%rd1];208; SM60-NEXT: and.b32 %r16, %r13, %r2;209; SM60-NEXT: $L__BB4_1: // %partword.cmpxchg.loop210; SM60-NEXT: // =>This Inner Loop Header: Depth=1211; SM60-NEXT: or.b32 %r14, %r16, %r3;212; SM60-NEXT: or.b32 %r15, %r16, %r4;213; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;214; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;215; SM60-NEXT: @%p1 bra $L__BB4_3;216; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure217; SM60-NEXT: // in Loop: Header=BB4_1 Depth=1218; SM60-NEXT: and.b32 %r6, %r5, %r2;219; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;220; SM60-NEXT: mov.b32 %r16, %r6;221; SM60-NEXT: @%p2 bra $L__BB4_1;222; SM60-NEXT: $L__BB4_3: // %partword.cmpxchg.end223; SM60-NEXT: membar.cta;224; SM60-NEXT: st.param.b32 [func_retval0], %r12;225; SM60-NEXT: ret;226 %pairold = cmpxchg ptr addrspace(1) %addr, i8 %cmp, i8 %new syncscope("block") acquire acquire227 ret i8 %new228}229 230define i8 @acquire_seq_cst_i8_global_cta(ptr addrspace(1) %addr, i8 %cmp, i8 %new) {231; SM60-LABEL: acquire_seq_cst_i8_global_cta(232; SM60: {233; SM60-NEXT: .reg .pred %p<3>;234; SM60-NEXT: .reg .b16 %rs<2>;235; SM60-NEXT: .reg .b32 %r<17>;236; SM60-NEXT: .reg .b64 %rd<3>;237; SM60-EMPTY:238; SM60-NEXT: // %bb.0:239; SM60-NEXT: ld.param.b8 %rs1, [acquire_seq_cst_i8_global_cta_param_2];240; SM60-NEXT: ld.param.b64 %rd2, [acquire_seq_cst_i8_global_cta_param_0];241; SM60-NEXT: membar.cta;242; SM60-NEXT: ld.param.b8 %r7, [acquire_seq_cst_i8_global_cta_param_1];243; SM60-NEXT: and.b64 %rd1, %rd2, -4;244; SM60-NEXT: cvt.u32.u64 %r8, %rd2;245; SM60-NEXT: and.b32 %r9, %r8, 3;246; SM60-NEXT: shl.b32 %r1, %r9, 3;247; SM60-NEXT: mov.b32 %r10, 255;248; SM60-NEXT: shl.b32 %r11, %r10, %r1;249; SM60-NEXT: not.b32 %r2, %r11;250; SM60-NEXT: cvt.u32.u16 %r12, %rs1;251; SM60-NEXT: shl.b32 %r3, %r12, %r1;252; SM60-NEXT: shl.b32 %r4, %r7, %r1;253; SM60-NEXT: ld.global.b32 %r13, [%rd1];254; SM60-NEXT: and.b32 %r16, %r13, %r2;255; SM60-NEXT: $L__BB5_1: // %partword.cmpxchg.loop256; SM60-NEXT: // =>This Inner Loop Header: Depth=1257; SM60-NEXT: or.b32 %r14, %r16, %r3;258; SM60-NEXT: or.b32 %r15, %r16, %r4;259; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;260; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;261; SM60-NEXT: @%p1 bra $L__BB5_3;262; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure263; SM60-NEXT: // in Loop: Header=BB5_1 Depth=1264; SM60-NEXT: and.b32 %r6, %r5, %r2;265; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;266; SM60-NEXT: mov.b32 %r16, %r6;267; SM60-NEXT: @%p2 bra $L__BB5_1;268; SM60-NEXT: $L__BB5_3: // %partword.cmpxchg.end269; SM60-NEXT: membar.cta;270; SM60-NEXT: st.param.b32 [func_retval0], %r12;271; SM60-NEXT: ret;272 %pairold = cmpxchg ptr addrspace(1) %addr, i8 %cmp, i8 %new syncscope("block") acquire seq_cst273 ret i8 %new274}275 276define i8 @release_monotonic_i8_global_cta(ptr addrspace(1) %addr, i8 %cmp, i8 %new) {277; SM60-LABEL: release_monotonic_i8_global_cta(278; SM60: {279; SM60-NEXT: .reg .pred %p<3>;280; SM60-NEXT: .reg .b16 %rs<2>;281; SM60-NEXT: .reg .b32 %r<17>;282; SM60-NEXT: .reg .b64 %rd<3>;283; SM60-EMPTY:284; SM60-NEXT: // %bb.0:285; SM60-NEXT: ld.param.b8 %rs1, [release_monotonic_i8_global_cta_param_2];286; SM60-NEXT: ld.param.b64 %rd2, [release_monotonic_i8_global_cta_param_0];287; SM60-NEXT: membar.cta;288; SM60-NEXT: ld.param.b8 %r7, [release_monotonic_i8_global_cta_param_1];289; SM60-NEXT: and.b64 %rd1, %rd2, -4;290; SM60-NEXT: cvt.u32.u64 %r8, %rd2;291; SM60-NEXT: and.b32 %r9, %r8, 3;292; SM60-NEXT: shl.b32 %r1, %r9, 3;293; SM60-NEXT: mov.b32 %r10, 255;294; SM60-NEXT: shl.b32 %r11, %r10, %r1;295; SM60-NEXT: not.b32 %r2, %r11;296; SM60-NEXT: cvt.u32.u16 %r12, %rs1;297; SM60-NEXT: shl.b32 %r3, %r12, %r1;298; SM60-NEXT: shl.b32 %r4, %r7, %r1;299; SM60-NEXT: ld.global.b32 %r13, [%rd1];300; SM60-NEXT: and.b32 %r16, %r13, %r2;301; SM60-NEXT: $L__BB6_1: // %partword.cmpxchg.loop302; SM60-NEXT: // =>This Inner Loop Header: Depth=1303; SM60-NEXT: or.b32 %r14, %r16, %r3;304; SM60-NEXT: or.b32 %r15, %r16, %r4;305; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;306; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;307; SM60-NEXT: @%p1 bra $L__BB6_3;308; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure309; SM60-NEXT: // in Loop: Header=BB6_1 Depth=1310; SM60-NEXT: and.b32 %r6, %r5, %r2;311; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;312; SM60-NEXT: mov.b32 %r16, %r6;313; SM60-NEXT: @%p2 bra $L__BB6_1;314; SM60-NEXT: $L__BB6_3: // %partword.cmpxchg.end315; SM60-NEXT: st.param.b32 [func_retval0], %r12;316; SM60-NEXT: ret;317 %pairold = cmpxchg ptr addrspace(1) %addr, i8 %cmp, i8 %new syncscope("block") release monotonic318 ret i8 %new319}320 321define i8 @release_acquire_i8_global_cta(ptr addrspace(1) %addr, i8 %cmp, i8 %new) {322; SM60-LABEL: release_acquire_i8_global_cta(323; SM60: {324; SM60-NEXT: .reg .pred %p<3>;325; SM60-NEXT: .reg .b16 %rs<2>;326; SM60-NEXT: .reg .b32 %r<17>;327; SM60-NEXT: .reg .b64 %rd<3>;328; SM60-EMPTY:329; SM60-NEXT: // %bb.0:330; SM60-NEXT: ld.param.b8 %rs1, [release_acquire_i8_global_cta_param_2];331; SM60-NEXT: ld.param.b64 %rd2, [release_acquire_i8_global_cta_param_0];332; SM60-NEXT: membar.cta;333; SM60-NEXT: ld.param.b8 %r7, [release_acquire_i8_global_cta_param_1];334; SM60-NEXT: and.b64 %rd1, %rd2, -4;335; SM60-NEXT: cvt.u32.u64 %r8, %rd2;336; SM60-NEXT: and.b32 %r9, %r8, 3;337; SM60-NEXT: shl.b32 %r1, %r9, 3;338; SM60-NEXT: mov.b32 %r10, 255;339; SM60-NEXT: shl.b32 %r11, %r10, %r1;340; SM60-NEXT: not.b32 %r2, %r11;341; SM60-NEXT: cvt.u32.u16 %r12, %rs1;342; SM60-NEXT: shl.b32 %r3, %r12, %r1;343; SM60-NEXT: shl.b32 %r4, %r7, %r1;344; SM60-NEXT: ld.global.b32 %r13, [%rd1];345; SM60-NEXT: and.b32 %r16, %r13, %r2;346; SM60-NEXT: $L__BB7_1: // %partword.cmpxchg.loop347; SM60-NEXT: // =>This Inner Loop Header: Depth=1348; SM60-NEXT: or.b32 %r14, %r16, %r3;349; SM60-NEXT: or.b32 %r15, %r16, %r4;350; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;351; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;352; SM60-NEXT: @%p1 bra $L__BB7_3;353; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure354; SM60-NEXT: // in Loop: Header=BB7_1 Depth=1355; SM60-NEXT: and.b32 %r6, %r5, %r2;356; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;357; SM60-NEXT: mov.b32 %r16, %r6;358; SM60-NEXT: @%p2 bra $L__BB7_1;359; SM60-NEXT: $L__BB7_3: // %partword.cmpxchg.end360; SM60-NEXT: membar.cta;361; SM60-NEXT: st.param.b32 [func_retval0], %r12;362; SM60-NEXT: ret;363 %pairold = cmpxchg ptr addrspace(1) %addr, i8 %cmp, i8 %new syncscope("block") release acquire364 ret i8 %new365}366 367define i8 @release_seq_cst_i8_global_cta(ptr addrspace(1) %addr, i8 %cmp, i8 %new) {368; SM60-LABEL: release_seq_cst_i8_global_cta(369; SM60: {370; SM60-NEXT: .reg .pred %p<3>;371; SM60-NEXT: .reg .b16 %rs<2>;372; SM60-NEXT: .reg .b32 %r<17>;373; SM60-NEXT: .reg .b64 %rd<3>;374; SM60-EMPTY:375; SM60-NEXT: // %bb.0:376; SM60-NEXT: ld.param.b8 %rs1, [release_seq_cst_i8_global_cta_param_2];377; SM60-NEXT: ld.param.b64 %rd2, [release_seq_cst_i8_global_cta_param_0];378; SM60-NEXT: membar.cta;379; SM60-NEXT: ld.param.b8 %r7, [release_seq_cst_i8_global_cta_param_1];380; SM60-NEXT: and.b64 %rd1, %rd2, -4;381; SM60-NEXT: cvt.u32.u64 %r8, %rd2;382; SM60-NEXT: and.b32 %r9, %r8, 3;383; SM60-NEXT: shl.b32 %r1, %r9, 3;384; SM60-NEXT: mov.b32 %r10, 255;385; SM60-NEXT: shl.b32 %r11, %r10, %r1;386; SM60-NEXT: not.b32 %r2, %r11;387; SM60-NEXT: cvt.u32.u16 %r12, %rs1;388; SM60-NEXT: shl.b32 %r3, %r12, %r1;389; SM60-NEXT: shl.b32 %r4, %r7, %r1;390; SM60-NEXT: ld.global.b32 %r13, [%rd1];391; SM60-NEXT: and.b32 %r16, %r13, %r2;392; SM60-NEXT: $L__BB8_1: // %partword.cmpxchg.loop393; SM60-NEXT: // =>This Inner Loop Header: Depth=1394; SM60-NEXT: or.b32 %r14, %r16, %r3;395; SM60-NEXT: or.b32 %r15, %r16, %r4;396; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;397; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;398; SM60-NEXT: @%p1 bra $L__BB8_3;399; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure400; SM60-NEXT: // in Loop: Header=BB8_1 Depth=1401; SM60-NEXT: and.b32 %r6, %r5, %r2;402; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;403; SM60-NEXT: mov.b32 %r16, %r6;404; SM60-NEXT: @%p2 bra $L__BB8_1;405; SM60-NEXT: $L__BB8_3: // %partword.cmpxchg.end406; SM60-NEXT: membar.cta;407; SM60-NEXT: st.param.b32 [func_retval0], %r12;408; SM60-NEXT: ret;409 %pairold = cmpxchg ptr addrspace(1) %addr, i8 %cmp, i8 %new syncscope("block") release seq_cst410 ret i8 %new411}412 413define i8 @acq_rel_monotonic_i8_global_cta(ptr addrspace(1) %addr, i8 %cmp, i8 %new) {414; SM60-LABEL: acq_rel_monotonic_i8_global_cta(415; SM60: {416; SM60-NEXT: .reg .pred %p<3>;417; SM60-NEXT: .reg .b16 %rs<2>;418; SM60-NEXT: .reg .b32 %r<17>;419; SM60-NEXT: .reg .b64 %rd<3>;420; SM60-EMPTY:421; SM60-NEXT: // %bb.0:422; SM60-NEXT: ld.param.b8 %rs1, [acq_rel_monotonic_i8_global_cta_param_2];423; SM60-NEXT: ld.param.b64 %rd2, [acq_rel_monotonic_i8_global_cta_param_0];424; SM60-NEXT: membar.cta;425; SM60-NEXT: ld.param.b8 %r7, [acq_rel_monotonic_i8_global_cta_param_1];426; SM60-NEXT: and.b64 %rd1, %rd2, -4;427; SM60-NEXT: cvt.u32.u64 %r8, %rd2;428; SM60-NEXT: and.b32 %r9, %r8, 3;429; SM60-NEXT: shl.b32 %r1, %r9, 3;430; SM60-NEXT: mov.b32 %r10, 255;431; SM60-NEXT: shl.b32 %r11, %r10, %r1;432; SM60-NEXT: not.b32 %r2, %r11;433; SM60-NEXT: cvt.u32.u16 %r12, %rs1;434; SM60-NEXT: shl.b32 %r3, %r12, %r1;435; SM60-NEXT: shl.b32 %r4, %r7, %r1;436; SM60-NEXT: ld.global.b32 %r13, [%rd1];437; SM60-NEXT: and.b32 %r16, %r13, %r2;438; SM60-NEXT: $L__BB9_1: // %partword.cmpxchg.loop439; SM60-NEXT: // =>This Inner Loop Header: Depth=1440; SM60-NEXT: or.b32 %r14, %r16, %r3;441; SM60-NEXT: or.b32 %r15, %r16, %r4;442; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;443; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;444; SM60-NEXT: @%p1 bra $L__BB9_3;445; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure446; SM60-NEXT: // in Loop: Header=BB9_1 Depth=1447; SM60-NEXT: and.b32 %r6, %r5, %r2;448; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;449; SM60-NEXT: mov.b32 %r16, %r6;450; SM60-NEXT: @%p2 bra $L__BB9_1;451; SM60-NEXT: $L__BB9_3: // %partword.cmpxchg.end452; SM60-NEXT: membar.cta;453; SM60-NEXT: st.param.b32 [func_retval0], %r12;454; SM60-NEXT: ret;455 %pairold = cmpxchg ptr addrspace(1) %addr, i8 %cmp, i8 %new syncscope("block") acq_rel monotonic456 ret i8 %new457}458 459define i8 @acq_rel_acquire_i8_global_cta(ptr addrspace(1) %addr, i8 %cmp, i8 %new) {460; SM60-LABEL: acq_rel_acquire_i8_global_cta(461; SM60: {462; SM60-NEXT: .reg .pred %p<3>;463; SM60-NEXT: .reg .b16 %rs<2>;464; SM60-NEXT: .reg .b32 %r<17>;465; SM60-NEXT: .reg .b64 %rd<3>;466; SM60-EMPTY:467; SM60-NEXT: // %bb.0:468; SM60-NEXT: ld.param.b8 %rs1, [acq_rel_acquire_i8_global_cta_param_2];469; SM60-NEXT: ld.param.b64 %rd2, [acq_rel_acquire_i8_global_cta_param_0];470; SM60-NEXT: membar.cta;471; SM60-NEXT: ld.param.b8 %r7, [acq_rel_acquire_i8_global_cta_param_1];472; SM60-NEXT: and.b64 %rd1, %rd2, -4;473; SM60-NEXT: cvt.u32.u64 %r8, %rd2;474; SM60-NEXT: and.b32 %r9, %r8, 3;475; SM60-NEXT: shl.b32 %r1, %r9, 3;476; SM60-NEXT: mov.b32 %r10, 255;477; SM60-NEXT: shl.b32 %r11, %r10, %r1;478; SM60-NEXT: not.b32 %r2, %r11;479; SM60-NEXT: cvt.u32.u16 %r12, %rs1;480; SM60-NEXT: shl.b32 %r3, %r12, %r1;481; SM60-NEXT: shl.b32 %r4, %r7, %r1;482; SM60-NEXT: ld.global.b32 %r13, [%rd1];483; SM60-NEXT: and.b32 %r16, %r13, %r2;484; SM60-NEXT: $L__BB10_1: // %partword.cmpxchg.loop485; SM60-NEXT: // =>This Inner Loop Header: Depth=1486; SM60-NEXT: or.b32 %r14, %r16, %r3;487; SM60-NEXT: or.b32 %r15, %r16, %r4;488; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;489; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;490; SM60-NEXT: @%p1 bra $L__BB10_3;491; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure492; SM60-NEXT: // in Loop: Header=BB10_1 Depth=1493; SM60-NEXT: and.b32 %r6, %r5, %r2;494; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;495; SM60-NEXT: mov.b32 %r16, %r6;496; SM60-NEXT: @%p2 bra $L__BB10_1;497; SM60-NEXT: $L__BB10_3: // %partword.cmpxchg.end498; SM60-NEXT: membar.cta;499; SM60-NEXT: st.param.b32 [func_retval0], %r12;500; SM60-NEXT: ret;501 %pairold = cmpxchg ptr addrspace(1) %addr, i8 %cmp, i8 %new syncscope("block") acq_rel acquire502 ret i8 %new503}504 505define i8 @acq_rel_seq_cst_i8_global_cta(ptr addrspace(1) %addr, i8 %cmp, i8 %new) {506; SM60-LABEL: acq_rel_seq_cst_i8_global_cta(507; SM60: {508; SM60-NEXT: .reg .pred %p<3>;509; SM60-NEXT: .reg .b16 %rs<2>;510; SM60-NEXT: .reg .b32 %r<17>;511; SM60-NEXT: .reg .b64 %rd<3>;512; SM60-EMPTY:513; SM60-NEXT: // %bb.0:514; SM60-NEXT: ld.param.b8 %rs1, [acq_rel_seq_cst_i8_global_cta_param_2];515; SM60-NEXT: ld.param.b64 %rd2, [acq_rel_seq_cst_i8_global_cta_param_0];516; SM60-NEXT: membar.cta;517; SM60-NEXT: ld.param.b8 %r7, [acq_rel_seq_cst_i8_global_cta_param_1];518; SM60-NEXT: and.b64 %rd1, %rd2, -4;519; SM60-NEXT: cvt.u32.u64 %r8, %rd2;520; SM60-NEXT: and.b32 %r9, %r8, 3;521; SM60-NEXT: shl.b32 %r1, %r9, 3;522; SM60-NEXT: mov.b32 %r10, 255;523; SM60-NEXT: shl.b32 %r11, %r10, %r1;524; SM60-NEXT: not.b32 %r2, %r11;525; SM60-NEXT: cvt.u32.u16 %r12, %rs1;526; SM60-NEXT: shl.b32 %r3, %r12, %r1;527; SM60-NEXT: shl.b32 %r4, %r7, %r1;528; SM60-NEXT: ld.global.b32 %r13, [%rd1];529; SM60-NEXT: and.b32 %r16, %r13, %r2;530; SM60-NEXT: $L__BB11_1: // %partword.cmpxchg.loop531; SM60-NEXT: // =>This Inner Loop Header: Depth=1532; SM60-NEXT: or.b32 %r14, %r16, %r3;533; SM60-NEXT: or.b32 %r15, %r16, %r4;534; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;535; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;536; SM60-NEXT: @%p1 bra $L__BB11_3;537; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure538; SM60-NEXT: // in Loop: Header=BB11_1 Depth=1539; SM60-NEXT: and.b32 %r6, %r5, %r2;540; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;541; SM60-NEXT: mov.b32 %r16, %r6;542; SM60-NEXT: @%p2 bra $L__BB11_1;543; SM60-NEXT: $L__BB11_3: // %partword.cmpxchg.end544; SM60-NEXT: membar.cta;545; SM60-NEXT: st.param.b32 [func_retval0], %r12;546; SM60-NEXT: ret;547 %pairold = cmpxchg ptr addrspace(1) %addr, i8 %cmp, i8 %new syncscope("block") acq_rel seq_cst548 ret i8 %new549}550 551define i8 @seq_cst_monotonic_i8_global_cta(ptr addrspace(1) %addr, i8 %cmp, i8 %new) {552; SM60-LABEL: seq_cst_monotonic_i8_global_cta(553; SM60: {554; SM60-NEXT: .reg .pred %p<3>;555; SM60-NEXT: .reg .b16 %rs<2>;556; SM60-NEXT: .reg .b32 %r<17>;557; SM60-NEXT: .reg .b64 %rd<3>;558; SM60-EMPTY:559; SM60-NEXT: // %bb.0:560; SM60-NEXT: ld.param.b8 %rs1, [seq_cst_monotonic_i8_global_cta_param_2];561; SM60-NEXT: ld.param.b64 %rd2, [seq_cst_monotonic_i8_global_cta_param_0];562; SM60-NEXT: membar.cta;563; SM60-NEXT: ld.param.b8 %r7, [seq_cst_monotonic_i8_global_cta_param_1];564; SM60-NEXT: and.b64 %rd1, %rd2, -4;565; SM60-NEXT: cvt.u32.u64 %r8, %rd2;566; SM60-NEXT: and.b32 %r9, %r8, 3;567; SM60-NEXT: shl.b32 %r1, %r9, 3;568; SM60-NEXT: mov.b32 %r10, 255;569; SM60-NEXT: shl.b32 %r11, %r10, %r1;570; SM60-NEXT: not.b32 %r2, %r11;571; SM60-NEXT: cvt.u32.u16 %r12, %rs1;572; SM60-NEXT: shl.b32 %r3, %r12, %r1;573; SM60-NEXT: shl.b32 %r4, %r7, %r1;574; SM60-NEXT: ld.global.b32 %r13, [%rd1];575; SM60-NEXT: and.b32 %r16, %r13, %r2;576; SM60-NEXT: $L__BB12_1: // %partword.cmpxchg.loop577; SM60-NEXT: // =>This Inner Loop Header: Depth=1578; SM60-NEXT: or.b32 %r14, %r16, %r3;579; SM60-NEXT: or.b32 %r15, %r16, %r4;580; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;581; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;582; SM60-NEXT: @%p1 bra $L__BB12_3;583; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure584; SM60-NEXT: // in Loop: Header=BB12_1 Depth=1585; SM60-NEXT: and.b32 %r6, %r5, %r2;586; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;587; SM60-NEXT: mov.b32 %r16, %r6;588; SM60-NEXT: @%p2 bra $L__BB12_1;589; SM60-NEXT: $L__BB12_3: // %partword.cmpxchg.end590; SM60-NEXT: membar.cta;591; SM60-NEXT: st.param.b32 [func_retval0], %r12;592; SM60-NEXT: ret;593 %pairold = cmpxchg ptr addrspace(1) %addr, i8 %cmp, i8 %new syncscope("block") seq_cst monotonic594 ret i8 %new595}596 597define i8 @seq_cst_acquire_i8_global_cta(ptr addrspace(1) %addr, i8 %cmp, i8 %new) {598; SM60-LABEL: seq_cst_acquire_i8_global_cta(599; SM60: {600; SM60-NEXT: .reg .pred %p<3>;601; SM60-NEXT: .reg .b16 %rs<2>;602; SM60-NEXT: .reg .b32 %r<17>;603; SM60-NEXT: .reg .b64 %rd<3>;604; SM60-EMPTY:605; SM60-NEXT: // %bb.0:606; SM60-NEXT: ld.param.b8 %rs1, [seq_cst_acquire_i8_global_cta_param_2];607; SM60-NEXT: ld.param.b64 %rd2, [seq_cst_acquire_i8_global_cta_param_0];608; SM60-NEXT: membar.cta;609; SM60-NEXT: ld.param.b8 %r7, [seq_cst_acquire_i8_global_cta_param_1];610; SM60-NEXT: and.b64 %rd1, %rd2, -4;611; SM60-NEXT: cvt.u32.u64 %r8, %rd2;612; SM60-NEXT: and.b32 %r9, %r8, 3;613; SM60-NEXT: shl.b32 %r1, %r9, 3;614; SM60-NEXT: mov.b32 %r10, 255;615; SM60-NEXT: shl.b32 %r11, %r10, %r1;616; SM60-NEXT: not.b32 %r2, %r11;617; SM60-NEXT: cvt.u32.u16 %r12, %rs1;618; SM60-NEXT: shl.b32 %r3, %r12, %r1;619; SM60-NEXT: shl.b32 %r4, %r7, %r1;620; SM60-NEXT: ld.global.b32 %r13, [%rd1];621; SM60-NEXT: and.b32 %r16, %r13, %r2;622; SM60-NEXT: $L__BB13_1: // %partword.cmpxchg.loop623; SM60-NEXT: // =>This Inner Loop Header: Depth=1624; SM60-NEXT: or.b32 %r14, %r16, %r3;625; SM60-NEXT: or.b32 %r15, %r16, %r4;626; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;627; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;628; SM60-NEXT: @%p1 bra $L__BB13_3;629; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure630; SM60-NEXT: // in Loop: Header=BB13_1 Depth=1631; SM60-NEXT: and.b32 %r6, %r5, %r2;632; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;633; SM60-NEXT: mov.b32 %r16, %r6;634; SM60-NEXT: @%p2 bra $L__BB13_1;635; SM60-NEXT: $L__BB13_3: // %partword.cmpxchg.end636; SM60-NEXT: membar.cta;637; SM60-NEXT: st.param.b32 [func_retval0], %r12;638; SM60-NEXT: ret;639 %pairold = cmpxchg ptr addrspace(1) %addr, i8 %cmp, i8 %new syncscope("block") seq_cst acquire640 ret i8 %new641}642 643define i8 @seq_cst_seq_cst_i8_global_cta(ptr addrspace(1) %addr, i8 %cmp, i8 %new) {644; SM60-LABEL: seq_cst_seq_cst_i8_global_cta(645; SM60: {646; SM60-NEXT: .reg .pred %p<3>;647; SM60-NEXT: .reg .b16 %rs<2>;648; SM60-NEXT: .reg .b32 %r<17>;649; SM60-NEXT: .reg .b64 %rd<3>;650; SM60-EMPTY:651; SM60-NEXT: // %bb.0:652; SM60-NEXT: ld.param.b8 %rs1, [seq_cst_seq_cst_i8_global_cta_param_2];653; SM60-NEXT: ld.param.b64 %rd2, [seq_cst_seq_cst_i8_global_cta_param_0];654; SM60-NEXT: membar.cta;655; SM60-NEXT: ld.param.b8 %r7, [seq_cst_seq_cst_i8_global_cta_param_1];656; SM60-NEXT: and.b64 %rd1, %rd2, -4;657; SM60-NEXT: cvt.u32.u64 %r8, %rd2;658; SM60-NEXT: and.b32 %r9, %r8, 3;659; SM60-NEXT: shl.b32 %r1, %r9, 3;660; SM60-NEXT: mov.b32 %r10, 255;661; SM60-NEXT: shl.b32 %r11, %r10, %r1;662; SM60-NEXT: not.b32 %r2, %r11;663; SM60-NEXT: cvt.u32.u16 %r12, %rs1;664; SM60-NEXT: shl.b32 %r3, %r12, %r1;665; SM60-NEXT: shl.b32 %r4, %r7, %r1;666; SM60-NEXT: ld.global.b32 %r13, [%rd1];667; SM60-NEXT: and.b32 %r16, %r13, %r2;668; SM60-NEXT: $L__BB14_1: // %partword.cmpxchg.loop669; SM60-NEXT: // =>This Inner Loop Header: Depth=1670; SM60-NEXT: or.b32 %r14, %r16, %r3;671; SM60-NEXT: or.b32 %r15, %r16, %r4;672; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;673; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;674; SM60-NEXT: @%p1 bra $L__BB14_3;675; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure676; SM60-NEXT: // in Loop: Header=BB14_1 Depth=1677; SM60-NEXT: and.b32 %r6, %r5, %r2;678; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;679; SM60-NEXT: mov.b32 %r16, %r6;680; SM60-NEXT: @%p2 bra $L__BB14_1;681; SM60-NEXT: $L__BB14_3: // %partword.cmpxchg.end682; SM60-NEXT: membar.cta;683; SM60-NEXT: st.param.b32 [func_retval0], %r12;684; SM60-NEXT: ret;685 %pairold = cmpxchg ptr addrspace(1) %addr, i8 %cmp, i8 %new syncscope("block") seq_cst seq_cst686 ret i8 %new687}688 689define i16 @monotonic_monotonic_i16_global_cta(ptr addrspace(1) %addr, i16 %cmp, i16 %new) {690; SM60-LABEL: monotonic_monotonic_i16_global_cta(691; SM60: {692; SM60-NEXT: .reg .pred %p<3>;693; SM60-NEXT: .reg .b16 %rs<2>;694; SM60-NEXT: .reg .b32 %r<17>;695; SM60-NEXT: .reg .b64 %rd<3>;696; SM60-EMPTY:697; SM60-NEXT: // %bb.0:698; SM60-NEXT: ld.param.b16 %rs1, [monotonic_monotonic_i16_global_cta_param_2];699; SM60-NEXT: ld.param.b64 %rd2, [monotonic_monotonic_i16_global_cta_param_0];700; SM60-NEXT: and.b64 %rd1, %rd2, -4;701; SM60-NEXT: ld.param.b16 %r7, [monotonic_monotonic_i16_global_cta_param_1];702; SM60-NEXT: cvt.u32.u64 %r8, %rd2;703; SM60-NEXT: and.b32 %r9, %r8, 3;704; SM60-NEXT: shl.b32 %r1, %r9, 3;705; SM60-NEXT: mov.b32 %r10, 65535;706; SM60-NEXT: shl.b32 %r11, %r10, %r1;707; SM60-NEXT: not.b32 %r2, %r11;708; SM60-NEXT: cvt.u32.u16 %r12, %rs1;709; SM60-NEXT: shl.b32 %r3, %r12, %r1;710; SM60-NEXT: shl.b32 %r4, %r7, %r1;711; SM60-NEXT: ld.global.b32 %r13, [%rd1];712; SM60-NEXT: and.b32 %r16, %r13, %r2;713; SM60-NEXT: $L__BB15_1: // %partword.cmpxchg.loop714; SM60-NEXT: // =>This Inner Loop Header: Depth=1715; SM60-NEXT: or.b32 %r14, %r16, %r3;716; SM60-NEXT: or.b32 %r15, %r16, %r4;717; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;718; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;719; SM60-NEXT: @%p1 bra $L__BB15_3;720; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure721; SM60-NEXT: // in Loop: Header=BB15_1 Depth=1722; SM60-NEXT: and.b32 %r6, %r5, %r2;723; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;724; SM60-NEXT: mov.b32 %r16, %r6;725; SM60-NEXT: @%p2 bra $L__BB15_1;726; SM60-NEXT: $L__BB15_3: // %partword.cmpxchg.end727; SM60-NEXT: st.param.b32 [func_retval0], %r12;728; SM60-NEXT: ret;729 %pairold = cmpxchg ptr addrspace(1) %addr, i16 %cmp, i16 %new syncscope("block") monotonic monotonic730 ret i16 %new731}732 733define i16 @monotonic_acquire_i16_global_cta(ptr addrspace(1) %addr, i16 %cmp, i16 %new) {734; SM60-LABEL: monotonic_acquire_i16_global_cta(735; SM60: {736; SM60-NEXT: .reg .pred %p<3>;737; SM60-NEXT: .reg .b16 %rs<2>;738; SM60-NEXT: .reg .b32 %r<17>;739; SM60-NEXT: .reg .b64 %rd<3>;740; SM60-EMPTY:741; SM60-NEXT: // %bb.0:742; SM60-NEXT: ld.param.b16 %rs1, [monotonic_acquire_i16_global_cta_param_2];743; SM60-NEXT: ld.param.b64 %rd2, [monotonic_acquire_i16_global_cta_param_0];744; SM60-NEXT: and.b64 %rd1, %rd2, -4;745; SM60-NEXT: ld.param.b16 %r7, [monotonic_acquire_i16_global_cta_param_1];746; SM60-NEXT: cvt.u32.u64 %r8, %rd2;747; SM60-NEXT: and.b32 %r9, %r8, 3;748; SM60-NEXT: shl.b32 %r1, %r9, 3;749; SM60-NEXT: mov.b32 %r10, 65535;750; SM60-NEXT: shl.b32 %r11, %r10, %r1;751; SM60-NEXT: not.b32 %r2, %r11;752; SM60-NEXT: cvt.u32.u16 %r12, %rs1;753; SM60-NEXT: shl.b32 %r3, %r12, %r1;754; SM60-NEXT: shl.b32 %r4, %r7, %r1;755; SM60-NEXT: ld.global.b32 %r13, [%rd1];756; SM60-NEXT: and.b32 %r16, %r13, %r2;757; SM60-NEXT: $L__BB16_1: // %partword.cmpxchg.loop758; SM60-NEXT: // =>This Inner Loop Header: Depth=1759; SM60-NEXT: or.b32 %r14, %r16, %r3;760; SM60-NEXT: or.b32 %r15, %r16, %r4;761; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;762; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;763; SM60-NEXT: @%p1 bra $L__BB16_3;764; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure765; SM60-NEXT: // in Loop: Header=BB16_1 Depth=1766; SM60-NEXT: and.b32 %r6, %r5, %r2;767; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;768; SM60-NEXT: mov.b32 %r16, %r6;769; SM60-NEXT: @%p2 bra $L__BB16_1;770; SM60-NEXT: $L__BB16_3: // %partword.cmpxchg.end771; SM60-NEXT: membar.cta;772; SM60-NEXT: st.param.b32 [func_retval0], %r12;773; SM60-NEXT: ret;774 %pairold = cmpxchg ptr addrspace(1) %addr, i16 %cmp, i16 %new syncscope("block") monotonic acquire775 ret i16 %new776}777 778define i16 @monotonic_seq_cst_i16_global_cta(ptr addrspace(1) %addr, i16 %cmp, i16 %new) {779; SM60-LABEL: monotonic_seq_cst_i16_global_cta(780; SM60: {781; SM60-NEXT: .reg .pred %p<3>;782; SM60-NEXT: .reg .b16 %rs<2>;783; SM60-NEXT: .reg .b32 %r<17>;784; SM60-NEXT: .reg .b64 %rd<3>;785; SM60-EMPTY:786; SM60-NEXT: // %bb.0:787; SM60-NEXT: ld.param.b16 %rs1, [monotonic_seq_cst_i16_global_cta_param_2];788; SM60-NEXT: ld.param.b64 %rd2, [monotonic_seq_cst_i16_global_cta_param_0];789; SM60-NEXT: membar.cta;790; SM60-NEXT: ld.param.b16 %r7, [monotonic_seq_cst_i16_global_cta_param_1];791; SM60-NEXT: and.b64 %rd1, %rd2, -4;792; SM60-NEXT: cvt.u32.u64 %r8, %rd2;793; SM60-NEXT: and.b32 %r9, %r8, 3;794; SM60-NEXT: shl.b32 %r1, %r9, 3;795; SM60-NEXT: mov.b32 %r10, 65535;796; SM60-NEXT: shl.b32 %r11, %r10, %r1;797; SM60-NEXT: not.b32 %r2, %r11;798; SM60-NEXT: cvt.u32.u16 %r12, %rs1;799; SM60-NEXT: shl.b32 %r3, %r12, %r1;800; SM60-NEXT: shl.b32 %r4, %r7, %r1;801; SM60-NEXT: ld.global.b32 %r13, [%rd1];802; SM60-NEXT: and.b32 %r16, %r13, %r2;803; SM60-NEXT: $L__BB17_1: // %partword.cmpxchg.loop804; SM60-NEXT: // =>This Inner Loop Header: Depth=1805; SM60-NEXT: or.b32 %r14, %r16, %r3;806; SM60-NEXT: or.b32 %r15, %r16, %r4;807; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;808; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;809; SM60-NEXT: @%p1 bra $L__BB17_3;810; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure811; SM60-NEXT: // in Loop: Header=BB17_1 Depth=1812; SM60-NEXT: and.b32 %r6, %r5, %r2;813; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;814; SM60-NEXT: mov.b32 %r16, %r6;815; SM60-NEXT: @%p2 bra $L__BB17_1;816; SM60-NEXT: $L__BB17_3: // %partword.cmpxchg.end817; SM60-NEXT: membar.cta;818; SM60-NEXT: st.param.b32 [func_retval0], %r12;819; SM60-NEXT: ret;820 %pairold = cmpxchg ptr addrspace(1) %addr, i16 %cmp, i16 %new syncscope("block") monotonic seq_cst821 ret i16 %new822}823 824define i16 @acquire_monotonic_i16_global_cta(ptr addrspace(1) %addr, i16 %cmp, i16 %new) {825; SM60-LABEL: acquire_monotonic_i16_global_cta(826; SM60: {827; SM60-NEXT: .reg .pred %p<3>;828; SM60-NEXT: .reg .b16 %rs<2>;829; SM60-NEXT: .reg .b32 %r<17>;830; SM60-NEXT: .reg .b64 %rd<3>;831; SM60-EMPTY:832; SM60-NEXT: // %bb.0:833; SM60-NEXT: ld.param.b16 %rs1, [acquire_monotonic_i16_global_cta_param_2];834; SM60-NEXT: ld.param.b64 %rd2, [acquire_monotonic_i16_global_cta_param_0];835; SM60-NEXT: and.b64 %rd1, %rd2, -4;836; SM60-NEXT: ld.param.b16 %r7, [acquire_monotonic_i16_global_cta_param_1];837; SM60-NEXT: cvt.u32.u64 %r8, %rd2;838; SM60-NEXT: and.b32 %r9, %r8, 3;839; SM60-NEXT: shl.b32 %r1, %r9, 3;840; SM60-NEXT: mov.b32 %r10, 65535;841; SM60-NEXT: shl.b32 %r11, %r10, %r1;842; SM60-NEXT: not.b32 %r2, %r11;843; SM60-NEXT: cvt.u32.u16 %r12, %rs1;844; SM60-NEXT: shl.b32 %r3, %r12, %r1;845; SM60-NEXT: shl.b32 %r4, %r7, %r1;846; SM60-NEXT: ld.global.b32 %r13, [%rd1];847; SM60-NEXT: and.b32 %r16, %r13, %r2;848; SM60-NEXT: $L__BB18_1: // %partword.cmpxchg.loop849; SM60-NEXT: // =>This Inner Loop Header: Depth=1850; SM60-NEXT: or.b32 %r14, %r16, %r3;851; SM60-NEXT: or.b32 %r15, %r16, %r4;852; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;853; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;854; SM60-NEXT: @%p1 bra $L__BB18_3;855; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure856; SM60-NEXT: // in Loop: Header=BB18_1 Depth=1857; SM60-NEXT: and.b32 %r6, %r5, %r2;858; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;859; SM60-NEXT: mov.b32 %r16, %r6;860; SM60-NEXT: @%p2 bra $L__BB18_1;861; SM60-NEXT: $L__BB18_3: // %partword.cmpxchg.end862; SM60-NEXT: membar.cta;863; SM60-NEXT: st.param.b32 [func_retval0], %r12;864; SM60-NEXT: ret;865 %pairold = cmpxchg ptr addrspace(1) %addr, i16 %cmp, i16 %new syncscope("block") acquire monotonic866 ret i16 %new867}868 869define i16 @acquire_acquire_i16_global_cta(ptr addrspace(1) %addr, i16 %cmp, i16 %new) {870; SM60-LABEL: acquire_acquire_i16_global_cta(871; SM60: {872; SM60-NEXT: .reg .pred %p<3>;873; SM60-NEXT: .reg .b16 %rs<2>;874; SM60-NEXT: .reg .b32 %r<17>;875; SM60-NEXT: .reg .b64 %rd<3>;876; SM60-EMPTY:877; SM60-NEXT: // %bb.0:878; SM60-NEXT: ld.param.b16 %rs1, [acquire_acquire_i16_global_cta_param_2];879; SM60-NEXT: ld.param.b64 %rd2, [acquire_acquire_i16_global_cta_param_0];880; SM60-NEXT: and.b64 %rd1, %rd2, -4;881; SM60-NEXT: ld.param.b16 %r7, [acquire_acquire_i16_global_cta_param_1];882; SM60-NEXT: cvt.u32.u64 %r8, %rd2;883; SM60-NEXT: and.b32 %r9, %r8, 3;884; SM60-NEXT: shl.b32 %r1, %r9, 3;885; SM60-NEXT: mov.b32 %r10, 65535;886; SM60-NEXT: shl.b32 %r11, %r10, %r1;887; SM60-NEXT: not.b32 %r2, %r11;888; SM60-NEXT: cvt.u32.u16 %r12, %rs1;889; SM60-NEXT: shl.b32 %r3, %r12, %r1;890; SM60-NEXT: shl.b32 %r4, %r7, %r1;891; SM60-NEXT: ld.global.b32 %r13, [%rd1];892; SM60-NEXT: and.b32 %r16, %r13, %r2;893; SM60-NEXT: $L__BB19_1: // %partword.cmpxchg.loop894; SM60-NEXT: // =>This Inner Loop Header: Depth=1895; SM60-NEXT: or.b32 %r14, %r16, %r3;896; SM60-NEXT: or.b32 %r15, %r16, %r4;897; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;898; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;899; SM60-NEXT: @%p1 bra $L__BB19_3;900; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure901; SM60-NEXT: // in Loop: Header=BB19_1 Depth=1902; SM60-NEXT: and.b32 %r6, %r5, %r2;903; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;904; SM60-NEXT: mov.b32 %r16, %r6;905; SM60-NEXT: @%p2 bra $L__BB19_1;906; SM60-NEXT: $L__BB19_3: // %partword.cmpxchg.end907; SM60-NEXT: membar.cta;908; SM60-NEXT: st.param.b32 [func_retval0], %r12;909; SM60-NEXT: ret;910 %pairold = cmpxchg ptr addrspace(1) %addr, i16 %cmp, i16 %new syncscope("block") acquire acquire911 ret i16 %new912}913 914define i16 @acquire_seq_cst_i16_global_cta(ptr addrspace(1) %addr, i16 %cmp, i16 %new) {915; SM60-LABEL: acquire_seq_cst_i16_global_cta(916; SM60: {917; SM60-NEXT: .reg .pred %p<3>;918; SM60-NEXT: .reg .b16 %rs<2>;919; SM60-NEXT: .reg .b32 %r<17>;920; SM60-NEXT: .reg .b64 %rd<3>;921; SM60-EMPTY:922; SM60-NEXT: // %bb.0:923; SM60-NEXT: ld.param.b16 %rs1, [acquire_seq_cst_i16_global_cta_param_2];924; SM60-NEXT: ld.param.b64 %rd2, [acquire_seq_cst_i16_global_cta_param_0];925; SM60-NEXT: membar.cta;926; SM60-NEXT: ld.param.b16 %r7, [acquire_seq_cst_i16_global_cta_param_1];927; SM60-NEXT: and.b64 %rd1, %rd2, -4;928; SM60-NEXT: cvt.u32.u64 %r8, %rd2;929; SM60-NEXT: and.b32 %r9, %r8, 3;930; SM60-NEXT: shl.b32 %r1, %r9, 3;931; SM60-NEXT: mov.b32 %r10, 65535;932; SM60-NEXT: shl.b32 %r11, %r10, %r1;933; SM60-NEXT: not.b32 %r2, %r11;934; SM60-NEXT: cvt.u32.u16 %r12, %rs1;935; SM60-NEXT: shl.b32 %r3, %r12, %r1;936; SM60-NEXT: shl.b32 %r4, %r7, %r1;937; SM60-NEXT: ld.global.b32 %r13, [%rd1];938; SM60-NEXT: and.b32 %r16, %r13, %r2;939; SM60-NEXT: $L__BB20_1: // %partword.cmpxchg.loop940; SM60-NEXT: // =>This Inner Loop Header: Depth=1941; SM60-NEXT: or.b32 %r14, %r16, %r3;942; SM60-NEXT: or.b32 %r15, %r16, %r4;943; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;944; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;945; SM60-NEXT: @%p1 bra $L__BB20_3;946; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure947; SM60-NEXT: // in Loop: Header=BB20_1 Depth=1948; SM60-NEXT: and.b32 %r6, %r5, %r2;949; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;950; SM60-NEXT: mov.b32 %r16, %r6;951; SM60-NEXT: @%p2 bra $L__BB20_1;952; SM60-NEXT: $L__BB20_3: // %partword.cmpxchg.end953; SM60-NEXT: membar.cta;954; SM60-NEXT: st.param.b32 [func_retval0], %r12;955; SM60-NEXT: ret;956 %pairold = cmpxchg ptr addrspace(1) %addr, i16 %cmp, i16 %new syncscope("block") acquire seq_cst957 ret i16 %new958}959 960define i16 @release_monotonic_i16_global_cta(ptr addrspace(1) %addr, i16 %cmp, i16 %new) {961; SM60-LABEL: release_monotonic_i16_global_cta(962; SM60: {963; SM60-NEXT: .reg .pred %p<3>;964; SM60-NEXT: .reg .b16 %rs<2>;965; SM60-NEXT: .reg .b32 %r<17>;966; SM60-NEXT: .reg .b64 %rd<3>;967; SM60-EMPTY:968; SM60-NEXT: // %bb.0:969; SM60-NEXT: ld.param.b16 %rs1, [release_monotonic_i16_global_cta_param_2];970; SM60-NEXT: ld.param.b64 %rd2, [release_monotonic_i16_global_cta_param_0];971; SM60-NEXT: membar.cta;972; SM60-NEXT: ld.param.b16 %r7, [release_monotonic_i16_global_cta_param_1];973; SM60-NEXT: and.b64 %rd1, %rd2, -4;974; SM60-NEXT: cvt.u32.u64 %r8, %rd2;975; SM60-NEXT: and.b32 %r9, %r8, 3;976; SM60-NEXT: shl.b32 %r1, %r9, 3;977; SM60-NEXT: mov.b32 %r10, 65535;978; SM60-NEXT: shl.b32 %r11, %r10, %r1;979; SM60-NEXT: not.b32 %r2, %r11;980; SM60-NEXT: cvt.u32.u16 %r12, %rs1;981; SM60-NEXT: shl.b32 %r3, %r12, %r1;982; SM60-NEXT: shl.b32 %r4, %r7, %r1;983; SM60-NEXT: ld.global.b32 %r13, [%rd1];984; SM60-NEXT: and.b32 %r16, %r13, %r2;985; SM60-NEXT: $L__BB21_1: // %partword.cmpxchg.loop986; SM60-NEXT: // =>This Inner Loop Header: Depth=1987; SM60-NEXT: or.b32 %r14, %r16, %r3;988; SM60-NEXT: or.b32 %r15, %r16, %r4;989; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;990; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;991; SM60-NEXT: @%p1 bra $L__BB21_3;992; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure993; SM60-NEXT: // in Loop: Header=BB21_1 Depth=1994; SM60-NEXT: and.b32 %r6, %r5, %r2;995; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;996; SM60-NEXT: mov.b32 %r16, %r6;997; SM60-NEXT: @%p2 bra $L__BB21_1;998; SM60-NEXT: $L__BB21_3: // %partword.cmpxchg.end999; SM60-NEXT: st.param.b32 [func_retval0], %r12;1000; SM60-NEXT: ret;1001 %pairold = cmpxchg ptr addrspace(1) %addr, i16 %cmp, i16 %new syncscope("block") release monotonic1002 ret i16 %new1003}1004 1005define i16 @release_acquire_i16_global_cta(ptr addrspace(1) %addr, i16 %cmp, i16 %new) {1006; SM60-LABEL: release_acquire_i16_global_cta(1007; SM60: {1008; SM60-NEXT: .reg .pred %p<3>;1009; SM60-NEXT: .reg .b16 %rs<2>;1010; SM60-NEXT: .reg .b32 %r<17>;1011; SM60-NEXT: .reg .b64 %rd<3>;1012; SM60-EMPTY:1013; SM60-NEXT: // %bb.0:1014; SM60-NEXT: ld.param.b16 %rs1, [release_acquire_i16_global_cta_param_2];1015; SM60-NEXT: ld.param.b64 %rd2, [release_acquire_i16_global_cta_param_0];1016; SM60-NEXT: membar.cta;1017; SM60-NEXT: ld.param.b16 %r7, [release_acquire_i16_global_cta_param_1];1018; SM60-NEXT: and.b64 %rd1, %rd2, -4;1019; SM60-NEXT: cvt.u32.u64 %r8, %rd2;1020; SM60-NEXT: and.b32 %r9, %r8, 3;1021; SM60-NEXT: shl.b32 %r1, %r9, 3;1022; SM60-NEXT: mov.b32 %r10, 65535;1023; SM60-NEXT: shl.b32 %r11, %r10, %r1;1024; SM60-NEXT: not.b32 %r2, %r11;1025; SM60-NEXT: cvt.u32.u16 %r12, %rs1;1026; SM60-NEXT: shl.b32 %r3, %r12, %r1;1027; SM60-NEXT: shl.b32 %r4, %r7, %r1;1028; SM60-NEXT: ld.global.b32 %r13, [%rd1];1029; SM60-NEXT: and.b32 %r16, %r13, %r2;1030; SM60-NEXT: $L__BB22_1: // %partword.cmpxchg.loop1031; SM60-NEXT: // =>This Inner Loop Header: Depth=11032; SM60-NEXT: or.b32 %r14, %r16, %r3;1033; SM60-NEXT: or.b32 %r15, %r16, %r4;1034; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;1035; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;1036; SM60-NEXT: @%p1 bra $L__BB22_3;1037; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure1038; SM60-NEXT: // in Loop: Header=BB22_1 Depth=11039; SM60-NEXT: and.b32 %r6, %r5, %r2;1040; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;1041; SM60-NEXT: mov.b32 %r16, %r6;1042; SM60-NEXT: @%p2 bra $L__BB22_1;1043; SM60-NEXT: $L__BB22_3: // %partword.cmpxchg.end1044; SM60-NEXT: membar.cta;1045; SM60-NEXT: st.param.b32 [func_retval0], %r12;1046; SM60-NEXT: ret;1047 %pairold = cmpxchg ptr addrspace(1) %addr, i16 %cmp, i16 %new syncscope("block") release acquire1048 ret i16 %new1049}1050 1051define i16 @release_seq_cst_i16_global_cta(ptr addrspace(1) %addr, i16 %cmp, i16 %new) {1052; SM60-LABEL: release_seq_cst_i16_global_cta(1053; SM60: {1054; SM60-NEXT: .reg .pred %p<3>;1055; SM60-NEXT: .reg .b16 %rs<2>;1056; SM60-NEXT: .reg .b32 %r<17>;1057; SM60-NEXT: .reg .b64 %rd<3>;1058; SM60-EMPTY:1059; SM60-NEXT: // %bb.0:1060; SM60-NEXT: ld.param.b16 %rs1, [release_seq_cst_i16_global_cta_param_2];1061; SM60-NEXT: ld.param.b64 %rd2, [release_seq_cst_i16_global_cta_param_0];1062; SM60-NEXT: membar.cta;1063; SM60-NEXT: ld.param.b16 %r7, [release_seq_cst_i16_global_cta_param_1];1064; SM60-NEXT: and.b64 %rd1, %rd2, -4;1065; SM60-NEXT: cvt.u32.u64 %r8, %rd2;1066; SM60-NEXT: and.b32 %r9, %r8, 3;1067; SM60-NEXT: shl.b32 %r1, %r9, 3;1068; SM60-NEXT: mov.b32 %r10, 65535;1069; SM60-NEXT: shl.b32 %r11, %r10, %r1;1070; SM60-NEXT: not.b32 %r2, %r11;1071; SM60-NEXT: cvt.u32.u16 %r12, %rs1;1072; SM60-NEXT: shl.b32 %r3, %r12, %r1;1073; SM60-NEXT: shl.b32 %r4, %r7, %r1;1074; SM60-NEXT: ld.global.b32 %r13, [%rd1];1075; SM60-NEXT: and.b32 %r16, %r13, %r2;1076; SM60-NEXT: $L__BB23_1: // %partword.cmpxchg.loop1077; SM60-NEXT: // =>This Inner Loop Header: Depth=11078; SM60-NEXT: or.b32 %r14, %r16, %r3;1079; SM60-NEXT: or.b32 %r15, %r16, %r4;1080; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;1081; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;1082; SM60-NEXT: @%p1 bra $L__BB23_3;1083; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure1084; SM60-NEXT: // in Loop: Header=BB23_1 Depth=11085; SM60-NEXT: and.b32 %r6, %r5, %r2;1086; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;1087; SM60-NEXT: mov.b32 %r16, %r6;1088; SM60-NEXT: @%p2 bra $L__BB23_1;1089; SM60-NEXT: $L__BB23_3: // %partword.cmpxchg.end1090; SM60-NEXT: membar.cta;1091; SM60-NEXT: st.param.b32 [func_retval0], %r12;1092; SM60-NEXT: ret;1093 %pairold = cmpxchg ptr addrspace(1) %addr, i16 %cmp, i16 %new syncscope("block") release seq_cst1094 ret i16 %new1095}1096 1097define i16 @acq_rel_monotonic_i16_global_cta(ptr addrspace(1) %addr, i16 %cmp, i16 %new) {1098; SM60-LABEL: acq_rel_monotonic_i16_global_cta(1099; SM60: {1100; SM60-NEXT: .reg .pred %p<3>;1101; SM60-NEXT: .reg .b16 %rs<2>;1102; SM60-NEXT: .reg .b32 %r<17>;1103; SM60-NEXT: .reg .b64 %rd<3>;1104; SM60-EMPTY:1105; SM60-NEXT: // %bb.0:1106; SM60-NEXT: ld.param.b16 %rs1, [acq_rel_monotonic_i16_global_cta_param_2];1107; SM60-NEXT: ld.param.b64 %rd2, [acq_rel_monotonic_i16_global_cta_param_0];1108; SM60-NEXT: membar.cta;1109; SM60-NEXT: ld.param.b16 %r7, [acq_rel_monotonic_i16_global_cta_param_1];1110; SM60-NEXT: and.b64 %rd1, %rd2, -4;1111; SM60-NEXT: cvt.u32.u64 %r8, %rd2;1112; SM60-NEXT: and.b32 %r9, %r8, 3;1113; SM60-NEXT: shl.b32 %r1, %r9, 3;1114; SM60-NEXT: mov.b32 %r10, 65535;1115; SM60-NEXT: shl.b32 %r11, %r10, %r1;1116; SM60-NEXT: not.b32 %r2, %r11;1117; SM60-NEXT: cvt.u32.u16 %r12, %rs1;1118; SM60-NEXT: shl.b32 %r3, %r12, %r1;1119; SM60-NEXT: shl.b32 %r4, %r7, %r1;1120; SM60-NEXT: ld.global.b32 %r13, [%rd1];1121; SM60-NEXT: and.b32 %r16, %r13, %r2;1122; SM60-NEXT: $L__BB24_1: // %partword.cmpxchg.loop1123; SM60-NEXT: // =>This Inner Loop Header: Depth=11124; SM60-NEXT: or.b32 %r14, %r16, %r3;1125; SM60-NEXT: or.b32 %r15, %r16, %r4;1126; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;1127; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;1128; SM60-NEXT: @%p1 bra $L__BB24_3;1129; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure1130; SM60-NEXT: // in Loop: Header=BB24_1 Depth=11131; SM60-NEXT: and.b32 %r6, %r5, %r2;1132; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;1133; SM60-NEXT: mov.b32 %r16, %r6;1134; SM60-NEXT: @%p2 bra $L__BB24_1;1135; SM60-NEXT: $L__BB24_3: // %partword.cmpxchg.end1136; SM60-NEXT: membar.cta;1137; SM60-NEXT: st.param.b32 [func_retval0], %r12;1138; SM60-NEXT: ret;1139 %pairold = cmpxchg ptr addrspace(1) %addr, i16 %cmp, i16 %new syncscope("block") acq_rel monotonic1140 ret i16 %new1141}1142 1143define i16 @acq_rel_acquire_i16_global_cta(ptr addrspace(1) %addr, i16 %cmp, i16 %new) {1144; SM60-LABEL: acq_rel_acquire_i16_global_cta(1145; SM60: {1146; SM60-NEXT: .reg .pred %p<3>;1147; SM60-NEXT: .reg .b16 %rs<2>;1148; SM60-NEXT: .reg .b32 %r<17>;1149; SM60-NEXT: .reg .b64 %rd<3>;1150; SM60-EMPTY:1151; SM60-NEXT: // %bb.0:1152; SM60-NEXT: ld.param.b16 %rs1, [acq_rel_acquire_i16_global_cta_param_2];1153; SM60-NEXT: ld.param.b64 %rd2, [acq_rel_acquire_i16_global_cta_param_0];1154; SM60-NEXT: membar.cta;1155; SM60-NEXT: ld.param.b16 %r7, [acq_rel_acquire_i16_global_cta_param_1];1156; SM60-NEXT: and.b64 %rd1, %rd2, -4;1157; SM60-NEXT: cvt.u32.u64 %r8, %rd2;1158; SM60-NEXT: and.b32 %r9, %r8, 3;1159; SM60-NEXT: shl.b32 %r1, %r9, 3;1160; SM60-NEXT: mov.b32 %r10, 65535;1161; SM60-NEXT: shl.b32 %r11, %r10, %r1;1162; SM60-NEXT: not.b32 %r2, %r11;1163; SM60-NEXT: cvt.u32.u16 %r12, %rs1;1164; SM60-NEXT: shl.b32 %r3, %r12, %r1;1165; SM60-NEXT: shl.b32 %r4, %r7, %r1;1166; SM60-NEXT: ld.global.b32 %r13, [%rd1];1167; SM60-NEXT: and.b32 %r16, %r13, %r2;1168; SM60-NEXT: $L__BB25_1: // %partword.cmpxchg.loop1169; SM60-NEXT: // =>This Inner Loop Header: Depth=11170; SM60-NEXT: or.b32 %r14, %r16, %r3;1171; SM60-NEXT: or.b32 %r15, %r16, %r4;1172; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;1173; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;1174; SM60-NEXT: @%p1 bra $L__BB25_3;1175; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure1176; SM60-NEXT: // in Loop: Header=BB25_1 Depth=11177; SM60-NEXT: and.b32 %r6, %r5, %r2;1178; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;1179; SM60-NEXT: mov.b32 %r16, %r6;1180; SM60-NEXT: @%p2 bra $L__BB25_1;1181; SM60-NEXT: $L__BB25_3: // %partword.cmpxchg.end1182; SM60-NEXT: membar.cta;1183; SM60-NEXT: st.param.b32 [func_retval0], %r12;1184; SM60-NEXT: ret;1185 %pairold = cmpxchg ptr addrspace(1) %addr, i16 %cmp, i16 %new syncscope("block") acq_rel acquire1186 ret i16 %new1187}1188 1189define i16 @acq_rel_seq_cst_i16_global_cta(ptr addrspace(1) %addr, i16 %cmp, i16 %new) {1190; SM60-LABEL: acq_rel_seq_cst_i16_global_cta(1191; SM60: {1192; SM60-NEXT: .reg .pred %p<3>;1193; SM60-NEXT: .reg .b16 %rs<2>;1194; SM60-NEXT: .reg .b32 %r<17>;1195; SM60-NEXT: .reg .b64 %rd<3>;1196; SM60-EMPTY:1197; SM60-NEXT: // %bb.0:1198; SM60-NEXT: ld.param.b16 %rs1, [acq_rel_seq_cst_i16_global_cta_param_2];1199; SM60-NEXT: ld.param.b64 %rd2, [acq_rel_seq_cst_i16_global_cta_param_0];1200; SM60-NEXT: membar.cta;1201; SM60-NEXT: ld.param.b16 %r7, [acq_rel_seq_cst_i16_global_cta_param_1];1202; SM60-NEXT: and.b64 %rd1, %rd2, -4;1203; SM60-NEXT: cvt.u32.u64 %r8, %rd2;1204; SM60-NEXT: and.b32 %r9, %r8, 3;1205; SM60-NEXT: shl.b32 %r1, %r9, 3;1206; SM60-NEXT: mov.b32 %r10, 65535;1207; SM60-NEXT: shl.b32 %r11, %r10, %r1;1208; SM60-NEXT: not.b32 %r2, %r11;1209; SM60-NEXT: cvt.u32.u16 %r12, %rs1;1210; SM60-NEXT: shl.b32 %r3, %r12, %r1;1211; SM60-NEXT: shl.b32 %r4, %r7, %r1;1212; SM60-NEXT: ld.global.b32 %r13, [%rd1];1213; SM60-NEXT: and.b32 %r16, %r13, %r2;1214; SM60-NEXT: $L__BB26_1: // %partword.cmpxchg.loop1215; SM60-NEXT: // =>This Inner Loop Header: Depth=11216; SM60-NEXT: or.b32 %r14, %r16, %r3;1217; SM60-NEXT: or.b32 %r15, %r16, %r4;1218; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;1219; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;1220; SM60-NEXT: @%p1 bra $L__BB26_3;1221; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure1222; SM60-NEXT: // in Loop: Header=BB26_1 Depth=11223; SM60-NEXT: and.b32 %r6, %r5, %r2;1224; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;1225; SM60-NEXT: mov.b32 %r16, %r6;1226; SM60-NEXT: @%p2 bra $L__BB26_1;1227; SM60-NEXT: $L__BB26_3: // %partword.cmpxchg.end1228; SM60-NEXT: membar.cta;1229; SM60-NEXT: st.param.b32 [func_retval0], %r12;1230; SM60-NEXT: ret;1231 %pairold = cmpxchg ptr addrspace(1) %addr, i16 %cmp, i16 %new syncscope("block") acq_rel seq_cst1232 ret i16 %new1233}1234 1235define i16 @seq_cst_monotonic_i16_global_cta(ptr addrspace(1) %addr, i16 %cmp, i16 %new) {1236; SM60-LABEL: seq_cst_monotonic_i16_global_cta(1237; SM60: {1238; SM60-NEXT: .reg .pred %p<3>;1239; SM60-NEXT: .reg .b16 %rs<2>;1240; SM60-NEXT: .reg .b32 %r<17>;1241; SM60-NEXT: .reg .b64 %rd<3>;1242; SM60-EMPTY:1243; SM60-NEXT: // %bb.0:1244; SM60-NEXT: ld.param.b16 %rs1, [seq_cst_monotonic_i16_global_cta_param_2];1245; SM60-NEXT: ld.param.b64 %rd2, [seq_cst_monotonic_i16_global_cta_param_0];1246; SM60-NEXT: membar.cta;1247; SM60-NEXT: ld.param.b16 %r7, [seq_cst_monotonic_i16_global_cta_param_1];1248; SM60-NEXT: and.b64 %rd1, %rd2, -4;1249; SM60-NEXT: cvt.u32.u64 %r8, %rd2;1250; SM60-NEXT: and.b32 %r9, %r8, 3;1251; SM60-NEXT: shl.b32 %r1, %r9, 3;1252; SM60-NEXT: mov.b32 %r10, 65535;1253; SM60-NEXT: shl.b32 %r11, %r10, %r1;1254; SM60-NEXT: not.b32 %r2, %r11;1255; SM60-NEXT: cvt.u32.u16 %r12, %rs1;1256; SM60-NEXT: shl.b32 %r3, %r12, %r1;1257; SM60-NEXT: shl.b32 %r4, %r7, %r1;1258; SM60-NEXT: ld.global.b32 %r13, [%rd1];1259; SM60-NEXT: and.b32 %r16, %r13, %r2;1260; SM60-NEXT: $L__BB27_1: // %partword.cmpxchg.loop1261; SM60-NEXT: // =>This Inner Loop Header: Depth=11262; SM60-NEXT: or.b32 %r14, %r16, %r3;1263; SM60-NEXT: or.b32 %r15, %r16, %r4;1264; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;1265; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;1266; SM60-NEXT: @%p1 bra $L__BB27_3;1267; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure1268; SM60-NEXT: // in Loop: Header=BB27_1 Depth=11269; SM60-NEXT: and.b32 %r6, %r5, %r2;1270; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;1271; SM60-NEXT: mov.b32 %r16, %r6;1272; SM60-NEXT: @%p2 bra $L__BB27_1;1273; SM60-NEXT: $L__BB27_3: // %partword.cmpxchg.end1274; SM60-NEXT: membar.cta;1275; SM60-NEXT: st.param.b32 [func_retval0], %r12;1276; SM60-NEXT: ret;1277 %pairold = cmpxchg ptr addrspace(1) %addr, i16 %cmp, i16 %new syncscope("block") seq_cst monotonic1278 ret i16 %new1279}1280 1281define i16 @seq_cst_acquire_i16_global_cta(ptr addrspace(1) %addr, i16 %cmp, i16 %new) {1282; SM60-LABEL: seq_cst_acquire_i16_global_cta(1283; SM60: {1284; SM60-NEXT: .reg .pred %p<3>;1285; SM60-NEXT: .reg .b16 %rs<2>;1286; SM60-NEXT: .reg .b32 %r<17>;1287; SM60-NEXT: .reg .b64 %rd<3>;1288; SM60-EMPTY:1289; SM60-NEXT: // %bb.0:1290; SM60-NEXT: ld.param.b16 %rs1, [seq_cst_acquire_i16_global_cta_param_2];1291; SM60-NEXT: ld.param.b64 %rd2, [seq_cst_acquire_i16_global_cta_param_0];1292; SM60-NEXT: membar.cta;1293; SM60-NEXT: ld.param.b16 %r7, [seq_cst_acquire_i16_global_cta_param_1];1294; SM60-NEXT: and.b64 %rd1, %rd2, -4;1295; SM60-NEXT: cvt.u32.u64 %r8, %rd2;1296; SM60-NEXT: and.b32 %r9, %r8, 3;1297; SM60-NEXT: shl.b32 %r1, %r9, 3;1298; SM60-NEXT: mov.b32 %r10, 65535;1299; SM60-NEXT: shl.b32 %r11, %r10, %r1;1300; SM60-NEXT: not.b32 %r2, %r11;1301; SM60-NEXT: cvt.u32.u16 %r12, %rs1;1302; SM60-NEXT: shl.b32 %r3, %r12, %r1;1303; SM60-NEXT: shl.b32 %r4, %r7, %r1;1304; SM60-NEXT: ld.global.b32 %r13, [%rd1];1305; SM60-NEXT: and.b32 %r16, %r13, %r2;1306; SM60-NEXT: $L__BB28_1: // %partword.cmpxchg.loop1307; SM60-NEXT: // =>This Inner Loop Header: Depth=11308; SM60-NEXT: or.b32 %r14, %r16, %r3;1309; SM60-NEXT: or.b32 %r15, %r16, %r4;1310; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;1311; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;1312; SM60-NEXT: @%p1 bra $L__BB28_3;1313; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure1314; SM60-NEXT: // in Loop: Header=BB28_1 Depth=11315; SM60-NEXT: and.b32 %r6, %r5, %r2;1316; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;1317; SM60-NEXT: mov.b32 %r16, %r6;1318; SM60-NEXT: @%p2 bra $L__BB28_1;1319; SM60-NEXT: $L__BB28_3: // %partword.cmpxchg.end1320; SM60-NEXT: membar.cta;1321; SM60-NEXT: st.param.b32 [func_retval0], %r12;1322; SM60-NEXT: ret;1323 %pairold = cmpxchg ptr addrspace(1) %addr, i16 %cmp, i16 %new syncscope("block") seq_cst acquire1324 ret i16 %new1325}1326 1327define i16 @seq_cst_seq_cst_i16_global_cta(ptr addrspace(1) %addr, i16 %cmp, i16 %new) {1328; SM60-LABEL: seq_cst_seq_cst_i16_global_cta(1329; SM60: {1330; SM60-NEXT: .reg .pred %p<3>;1331; SM60-NEXT: .reg .b16 %rs<2>;1332; SM60-NEXT: .reg .b32 %r<17>;1333; SM60-NEXT: .reg .b64 %rd<3>;1334; SM60-EMPTY:1335; SM60-NEXT: // %bb.0:1336; SM60-NEXT: ld.param.b16 %rs1, [seq_cst_seq_cst_i16_global_cta_param_2];1337; SM60-NEXT: ld.param.b64 %rd2, [seq_cst_seq_cst_i16_global_cta_param_0];1338; SM60-NEXT: membar.cta;1339; SM60-NEXT: ld.param.b16 %r7, [seq_cst_seq_cst_i16_global_cta_param_1];1340; SM60-NEXT: and.b64 %rd1, %rd2, -4;1341; SM60-NEXT: cvt.u32.u64 %r8, %rd2;1342; SM60-NEXT: and.b32 %r9, %r8, 3;1343; SM60-NEXT: shl.b32 %r1, %r9, 3;1344; SM60-NEXT: mov.b32 %r10, 65535;1345; SM60-NEXT: shl.b32 %r11, %r10, %r1;1346; SM60-NEXT: not.b32 %r2, %r11;1347; SM60-NEXT: cvt.u32.u16 %r12, %rs1;1348; SM60-NEXT: shl.b32 %r3, %r12, %r1;1349; SM60-NEXT: shl.b32 %r4, %r7, %r1;1350; SM60-NEXT: ld.global.b32 %r13, [%rd1];1351; SM60-NEXT: and.b32 %r16, %r13, %r2;1352; SM60-NEXT: $L__BB29_1: // %partword.cmpxchg.loop1353; SM60-NEXT: // =>This Inner Loop Header: Depth=11354; SM60-NEXT: or.b32 %r14, %r16, %r3;1355; SM60-NEXT: or.b32 %r15, %r16, %r4;1356; SM60-NEXT: atom.cta.global.cas.b32 %r5, [%rd1], %r15, %r14;1357; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;1358; SM60-NEXT: @%p1 bra $L__BB29_3;1359; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure1360; SM60-NEXT: // in Loop: Header=BB29_1 Depth=11361; SM60-NEXT: and.b32 %r6, %r5, %r2;1362; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;1363; SM60-NEXT: mov.b32 %r16, %r6;1364; SM60-NEXT: @%p2 bra $L__BB29_1;1365; SM60-NEXT: $L__BB29_3: // %partword.cmpxchg.end1366; SM60-NEXT: membar.cta;1367; SM60-NEXT: st.param.b32 [func_retval0], %r12;1368; SM60-NEXT: ret;1369 %pairold = cmpxchg ptr addrspace(1) %addr, i16 %cmp, i16 %new syncscope("block") seq_cst seq_cst1370 ret i16 %new1371}1372 1373define i32 @monotonic_monotonic_i32_global_cta(ptr addrspace(1) %addr, i32 %cmp, i32 %new) {1374; SM60-LABEL: monotonic_monotonic_i32_global_cta(1375; SM60: {1376; SM60-NEXT: .reg .b32 %r<4>;1377; SM60-NEXT: .reg .b64 %rd<2>;1378; SM60-EMPTY:1379; SM60-NEXT: // %bb.0:1380; SM60-NEXT: ld.param.b64 %rd1, [monotonic_monotonic_i32_global_cta_param_0];1381; SM60-NEXT: ld.param.b32 %r1, [monotonic_monotonic_i32_global_cta_param_1];1382; SM60-NEXT: ld.param.b32 %r2, [monotonic_monotonic_i32_global_cta_param_2];1383; SM60-NEXT: atom.cta.global.cas.b32 %r3, [%rd1], %r1, %r2;1384; SM60-NEXT: st.param.b32 [func_retval0], %r2;1385; SM60-NEXT: ret;1386 %pairold = cmpxchg ptr addrspace(1) %addr, i32 %cmp, i32 %new syncscope("block") monotonic monotonic1387 ret i32 %new1388}1389 1390define i32 @monotonic_acquire_i32_global_cta(ptr addrspace(1) %addr, i32 %cmp, i32 %new) {1391; SM60-LABEL: monotonic_acquire_i32_global_cta(1392; SM60: {1393; SM60-NEXT: .reg .b32 %r<4>;1394; SM60-NEXT: .reg .b64 %rd<2>;1395; SM60-EMPTY:1396; SM60-NEXT: // %bb.0:1397; SM60-NEXT: ld.param.b64 %rd1, [monotonic_acquire_i32_global_cta_param_0];1398; SM60-NEXT: ld.param.b32 %r1, [monotonic_acquire_i32_global_cta_param_1];1399; SM60-NEXT: ld.param.b32 %r2, [monotonic_acquire_i32_global_cta_param_2];1400; SM60-NEXT: atom.cta.global.cas.b32 %r3, [%rd1], %r1, %r2;1401; SM60-NEXT: st.param.b32 [func_retval0], %r2;1402; SM60-NEXT: ret;1403 %pairold = cmpxchg ptr addrspace(1) %addr, i32 %cmp, i32 %new syncscope("block") monotonic acquire1404 ret i32 %new1405}1406 1407define i32 @monotonic_seq_cst_i32_global_cta(ptr addrspace(1) %addr, i32 %cmp, i32 %new) {1408; SM60-LABEL: monotonic_seq_cst_i32_global_cta(1409; SM60: {1410; SM60-NEXT: .reg .b32 %r<4>;1411; SM60-NEXT: .reg .b64 %rd<2>;1412; SM60-EMPTY:1413; SM60-NEXT: // %bb.0:1414; SM60-NEXT: ld.param.b64 %rd1, [monotonic_seq_cst_i32_global_cta_param_0];1415; SM60-NEXT: membar.cta;1416; SM60-NEXT: ld.param.b32 %r1, [monotonic_seq_cst_i32_global_cta_param_1];1417; SM60-NEXT: ld.param.b32 %r2, [monotonic_seq_cst_i32_global_cta_param_2];1418; SM60-NEXT: atom.cta.global.cas.b32 %r3, [%rd1], %r1, %r2;1419; SM60-NEXT: st.param.b32 [func_retval0], %r2;1420; SM60-NEXT: ret;1421 %pairold = cmpxchg ptr addrspace(1) %addr, i32 %cmp, i32 %new syncscope("block") monotonic seq_cst1422 ret i32 %new1423}1424 1425define i32 @acquire_monotonic_i32_global_cta(ptr addrspace(1) %addr, i32 %cmp, i32 %new) {1426; SM60-LABEL: acquire_monotonic_i32_global_cta(1427; SM60: {1428; SM60-NEXT: .reg .b32 %r<4>;1429; SM60-NEXT: .reg .b64 %rd<2>;1430; SM60-EMPTY:1431; SM60-NEXT: // %bb.0:1432; SM60-NEXT: ld.param.b64 %rd1, [acquire_monotonic_i32_global_cta_param_0];1433; SM60-NEXT: ld.param.b32 %r1, [acquire_monotonic_i32_global_cta_param_1];1434; SM60-NEXT: ld.param.b32 %r2, [acquire_monotonic_i32_global_cta_param_2];1435; SM60-NEXT: atom.cta.global.cas.b32 %r3, [%rd1], %r1, %r2;1436; SM60-NEXT: st.param.b32 [func_retval0], %r2;1437; SM60-NEXT: ret;1438 %pairold = cmpxchg ptr addrspace(1) %addr, i32 %cmp, i32 %new syncscope("block") acquire monotonic1439 ret i32 %new1440}1441 1442define i32 @acquire_acquire_i32_global_cta(ptr addrspace(1) %addr, i32 %cmp, i32 %new) {1443; SM60-LABEL: acquire_acquire_i32_global_cta(1444; SM60: {1445; SM60-NEXT: .reg .b32 %r<4>;1446; SM60-NEXT: .reg .b64 %rd<2>;1447; SM60-EMPTY:1448; SM60-NEXT: // %bb.0:1449; SM60-NEXT: ld.param.b64 %rd1, [acquire_acquire_i32_global_cta_param_0];1450; SM60-NEXT: ld.param.b32 %r1, [acquire_acquire_i32_global_cta_param_1];1451; SM60-NEXT: ld.param.b32 %r2, [acquire_acquire_i32_global_cta_param_2];1452; SM60-NEXT: atom.cta.global.cas.b32 %r3, [%rd1], %r1, %r2;1453; SM60-NEXT: st.param.b32 [func_retval0], %r2;1454; SM60-NEXT: ret;1455 %pairold = cmpxchg ptr addrspace(1) %addr, i32 %cmp, i32 %new syncscope("block") acquire acquire1456 ret i32 %new1457}1458 1459define i32 @acquire_seq_cst_i32_global_cta(ptr addrspace(1) %addr, i32 %cmp, i32 %new) {1460; SM60-LABEL: acquire_seq_cst_i32_global_cta(1461; SM60: {1462; SM60-NEXT: .reg .b32 %r<4>;1463; SM60-NEXT: .reg .b64 %rd<2>;1464; SM60-EMPTY:1465; SM60-NEXT: // %bb.0:1466; SM60-NEXT: ld.param.b64 %rd1, [acquire_seq_cst_i32_global_cta_param_0];1467; SM60-NEXT: membar.cta;1468; SM60-NEXT: ld.param.b32 %r1, [acquire_seq_cst_i32_global_cta_param_1];1469; SM60-NEXT: ld.param.b32 %r2, [acquire_seq_cst_i32_global_cta_param_2];1470; SM60-NEXT: atom.cta.global.cas.b32 %r3, [%rd1], %r1, %r2;1471; SM60-NEXT: st.param.b32 [func_retval0], %r2;1472; SM60-NEXT: ret;1473 %pairold = cmpxchg ptr addrspace(1) %addr, i32 %cmp, i32 %new syncscope("block") acquire seq_cst1474 ret i32 %new1475}1476 1477define i32 @release_monotonic_i32_global_cta(ptr addrspace(1) %addr, i32 %cmp, i32 %new) {1478; SM60-LABEL: release_monotonic_i32_global_cta(1479; SM60: {1480; SM60-NEXT: .reg .b32 %r<4>;1481; SM60-NEXT: .reg .b64 %rd<2>;1482; SM60-EMPTY:1483; SM60-NEXT: // %bb.0:1484; SM60-NEXT: ld.param.b64 %rd1, [release_monotonic_i32_global_cta_param_0];1485; SM60-NEXT: ld.param.b32 %r1, [release_monotonic_i32_global_cta_param_1];1486; SM60-NEXT: ld.param.b32 %r2, [release_monotonic_i32_global_cta_param_2];1487; SM60-NEXT: atom.cta.global.cas.b32 %r3, [%rd1], %r1, %r2;1488; SM60-NEXT: st.param.b32 [func_retval0], %r2;1489; SM60-NEXT: ret;1490 %pairold = cmpxchg ptr addrspace(1) %addr, i32 %cmp, i32 %new syncscope("block") release monotonic1491 ret i32 %new1492}1493 1494define i32 @release_acquire_i32_global_cta(ptr addrspace(1) %addr, i32 %cmp, i32 %new) {1495; SM60-LABEL: release_acquire_i32_global_cta(1496; SM60: {1497; SM60-NEXT: .reg .b32 %r<4>;1498; SM60-NEXT: .reg .b64 %rd<2>;1499; SM60-EMPTY:1500; SM60-NEXT: // %bb.0:1501; SM60-NEXT: ld.param.b64 %rd1, [release_acquire_i32_global_cta_param_0];1502; SM60-NEXT: ld.param.b32 %r1, [release_acquire_i32_global_cta_param_1];1503; SM60-NEXT: ld.param.b32 %r2, [release_acquire_i32_global_cta_param_2];1504; SM60-NEXT: atom.cta.global.cas.b32 %r3, [%rd1], %r1, %r2;1505; SM60-NEXT: st.param.b32 [func_retval0], %r2;1506; SM60-NEXT: ret;1507 %pairold = cmpxchg ptr addrspace(1) %addr, i32 %cmp, i32 %new syncscope("block") release acquire1508 ret i32 %new1509}1510 1511define i32 @release_seq_cst_i32_global_cta(ptr addrspace(1) %addr, i32 %cmp, i32 %new) {1512; SM60-LABEL: release_seq_cst_i32_global_cta(1513; SM60: {1514; SM60-NEXT: .reg .b32 %r<4>;1515; SM60-NEXT: .reg .b64 %rd<2>;1516; SM60-EMPTY:1517; SM60-NEXT: // %bb.0:1518; SM60-NEXT: ld.param.b64 %rd1, [release_seq_cst_i32_global_cta_param_0];1519; SM60-NEXT: membar.cta;1520; SM60-NEXT: ld.param.b32 %r1, [release_seq_cst_i32_global_cta_param_1];1521; SM60-NEXT: ld.param.b32 %r2, [release_seq_cst_i32_global_cta_param_2];1522; SM60-NEXT: atom.cta.global.cas.b32 %r3, [%rd1], %r1, %r2;1523; SM60-NEXT: st.param.b32 [func_retval0], %r2;1524; SM60-NEXT: ret;1525 %pairold = cmpxchg ptr addrspace(1) %addr, i32 %cmp, i32 %new syncscope("block") release seq_cst1526 ret i32 %new1527}1528 1529define i32 @acq_rel_monotonic_i32_global_cta(ptr addrspace(1) %addr, i32 %cmp, i32 %new) {1530; SM60-LABEL: acq_rel_monotonic_i32_global_cta(1531; SM60: {1532; SM60-NEXT: .reg .b32 %r<4>;1533; SM60-NEXT: .reg .b64 %rd<2>;1534; SM60-EMPTY:1535; SM60-NEXT: // %bb.0:1536; SM60-NEXT: ld.param.b64 %rd1, [acq_rel_monotonic_i32_global_cta_param_0];1537; SM60-NEXT: ld.param.b32 %r1, [acq_rel_monotonic_i32_global_cta_param_1];1538; SM60-NEXT: ld.param.b32 %r2, [acq_rel_monotonic_i32_global_cta_param_2];1539; SM60-NEXT: atom.cta.global.cas.b32 %r3, [%rd1], %r1, %r2;1540; SM60-NEXT: st.param.b32 [func_retval0], %r2;1541; SM60-NEXT: ret;1542 %pairold = cmpxchg ptr addrspace(1) %addr, i32 %cmp, i32 %new syncscope("block") acq_rel monotonic1543 ret i32 %new1544}1545 1546define i32 @acq_rel_acquire_i32_global_cta(ptr addrspace(1) %addr, i32 %cmp, i32 %new) {1547; SM60-LABEL: acq_rel_acquire_i32_global_cta(1548; SM60: {1549; SM60-NEXT: .reg .b32 %r<4>;1550; SM60-NEXT: .reg .b64 %rd<2>;1551; SM60-EMPTY:1552; SM60-NEXT: // %bb.0:1553; SM60-NEXT: ld.param.b64 %rd1, [acq_rel_acquire_i32_global_cta_param_0];1554; SM60-NEXT: ld.param.b32 %r1, [acq_rel_acquire_i32_global_cta_param_1];1555; SM60-NEXT: ld.param.b32 %r2, [acq_rel_acquire_i32_global_cta_param_2];1556; SM60-NEXT: atom.cta.global.cas.b32 %r3, [%rd1], %r1, %r2;1557; SM60-NEXT: st.param.b32 [func_retval0], %r2;1558; SM60-NEXT: ret;1559 %pairold = cmpxchg ptr addrspace(1) %addr, i32 %cmp, i32 %new syncscope("block") acq_rel acquire1560 ret i32 %new1561}1562 1563define i32 @acq_rel_seq_cst_i32_global_cta(ptr addrspace(1) %addr, i32 %cmp, i32 %new) {1564; SM60-LABEL: acq_rel_seq_cst_i32_global_cta(1565; SM60: {1566; SM60-NEXT: .reg .b32 %r<4>;1567; SM60-NEXT: .reg .b64 %rd<2>;1568; SM60-EMPTY:1569; SM60-NEXT: // %bb.0:1570; SM60-NEXT: ld.param.b64 %rd1, [acq_rel_seq_cst_i32_global_cta_param_0];1571; SM60-NEXT: membar.cta;1572; SM60-NEXT: ld.param.b32 %r1, [acq_rel_seq_cst_i32_global_cta_param_1];1573; SM60-NEXT: ld.param.b32 %r2, [acq_rel_seq_cst_i32_global_cta_param_2];1574; SM60-NEXT: atom.cta.global.cas.b32 %r3, [%rd1], %r1, %r2;1575; SM60-NEXT: st.param.b32 [func_retval0], %r2;1576; SM60-NEXT: ret;1577 %pairold = cmpxchg ptr addrspace(1) %addr, i32 %cmp, i32 %new syncscope("block") acq_rel seq_cst1578 ret i32 %new1579}1580 1581define i32 @seq_cst_monotonic_i32_global_cta(ptr addrspace(1) %addr, i32 %cmp, i32 %new) {1582; SM60-LABEL: seq_cst_monotonic_i32_global_cta(1583; SM60: {1584; SM60-NEXT: .reg .b32 %r<4>;1585; SM60-NEXT: .reg .b64 %rd<2>;1586; SM60-EMPTY:1587; SM60-NEXT: // %bb.0:1588; SM60-NEXT: ld.param.b64 %rd1, [seq_cst_monotonic_i32_global_cta_param_0];1589; SM60-NEXT: membar.cta;1590; SM60-NEXT: ld.param.b32 %r1, [seq_cst_monotonic_i32_global_cta_param_1];1591; SM60-NEXT: ld.param.b32 %r2, [seq_cst_monotonic_i32_global_cta_param_2];1592; SM60-NEXT: atom.cta.global.cas.b32 %r3, [%rd1], %r1, %r2;1593; SM60-NEXT: st.param.b32 [func_retval0], %r2;1594; SM60-NEXT: ret;1595 %pairold = cmpxchg ptr addrspace(1) %addr, i32 %cmp, i32 %new syncscope("block") seq_cst monotonic1596 ret i32 %new1597}1598 1599define i32 @seq_cst_acquire_i32_global_cta(ptr addrspace(1) %addr, i32 %cmp, i32 %new) {1600; SM60-LABEL: seq_cst_acquire_i32_global_cta(1601; SM60: {1602; SM60-NEXT: .reg .b32 %r<4>;1603; SM60-NEXT: .reg .b64 %rd<2>;1604; SM60-EMPTY:1605; SM60-NEXT: // %bb.0:1606; SM60-NEXT: ld.param.b64 %rd1, [seq_cst_acquire_i32_global_cta_param_0];1607; SM60-NEXT: membar.cta;1608; SM60-NEXT: ld.param.b32 %r1, [seq_cst_acquire_i32_global_cta_param_1];1609; SM60-NEXT: ld.param.b32 %r2, [seq_cst_acquire_i32_global_cta_param_2];1610; SM60-NEXT: atom.cta.global.cas.b32 %r3, [%rd1], %r1, %r2;1611; SM60-NEXT: st.param.b32 [func_retval0], %r2;1612; SM60-NEXT: ret;1613 %pairold = cmpxchg ptr addrspace(1) %addr, i32 %cmp, i32 %new syncscope("block") seq_cst acquire1614 ret i32 %new1615}1616 1617define i32 @seq_cst_seq_cst_i32_global_cta(ptr addrspace(1) %addr, i32 %cmp, i32 %new) {1618; SM60-LABEL: seq_cst_seq_cst_i32_global_cta(1619; SM60: {1620; SM60-NEXT: .reg .b32 %r<4>;1621; SM60-NEXT: .reg .b64 %rd<2>;1622; SM60-EMPTY:1623; SM60-NEXT: // %bb.0:1624; SM60-NEXT: ld.param.b64 %rd1, [seq_cst_seq_cst_i32_global_cta_param_0];1625; SM60-NEXT: membar.cta;1626; SM60-NEXT: ld.param.b32 %r1, [seq_cst_seq_cst_i32_global_cta_param_1];1627; SM60-NEXT: ld.param.b32 %r2, [seq_cst_seq_cst_i32_global_cta_param_2];1628; SM60-NEXT: atom.cta.global.cas.b32 %r3, [%rd1], %r1, %r2;1629; SM60-NEXT: st.param.b32 [func_retval0], %r2;1630; SM60-NEXT: ret;1631 %pairold = cmpxchg ptr addrspace(1) %addr, i32 %cmp, i32 %new syncscope("block") seq_cst seq_cst1632 ret i32 %new1633}1634 1635define i64 @monotonic_monotonic_i64_global_cta(ptr addrspace(1) %addr, i64 %cmp, i64 %new) {1636; SM60-LABEL: monotonic_monotonic_i64_global_cta(1637; SM60: {1638; SM60-NEXT: .reg .b64 %rd<5>;1639; SM60-EMPTY:1640; SM60-NEXT: // %bb.0:1641; SM60-NEXT: ld.param.b64 %rd1, [monotonic_monotonic_i64_global_cta_param_0];1642; SM60-NEXT: ld.param.b64 %rd2, [monotonic_monotonic_i64_global_cta_param_1];1643; SM60-NEXT: ld.param.b64 %rd3, [monotonic_monotonic_i64_global_cta_param_2];1644; SM60-NEXT: atom.cta.global.cas.b64 %rd4, [%rd1], %rd2, %rd3;1645; SM60-NEXT: st.param.b64 [func_retval0], %rd3;1646; SM60-NEXT: ret;1647 %pairold = cmpxchg ptr addrspace(1) %addr, i64 %cmp, i64 %new syncscope("block") monotonic monotonic1648 ret i64 %new1649}1650 1651define i64 @monotonic_acquire_i64_global_cta(ptr addrspace(1) %addr, i64 %cmp, i64 %new) {1652; SM60-LABEL: monotonic_acquire_i64_global_cta(1653; SM60: {1654; SM60-NEXT: .reg .b64 %rd<5>;1655; SM60-EMPTY:1656; SM60-NEXT: // %bb.0:1657; SM60-NEXT: ld.param.b64 %rd1, [monotonic_acquire_i64_global_cta_param_0];1658; SM60-NEXT: ld.param.b64 %rd2, [monotonic_acquire_i64_global_cta_param_1];1659; SM60-NEXT: ld.param.b64 %rd3, [monotonic_acquire_i64_global_cta_param_2];1660; SM60-NEXT: atom.cta.global.cas.b64 %rd4, [%rd1], %rd2, %rd3;1661; SM60-NEXT: st.param.b64 [func_retval0], %rd3;1662; SM60-NEXT: ret;1663 %pairold = cmpxchg ptr addrspace(1) %addr, i64 %cmp, i64 %new syncscope("block") monotonic acquire1664 ret i64 %new1665}1666 1667define i64 @monotonic_seq_cst_i64_global_cta(ptr addrspace(1) %addr, i64 %cmp, i64 %new) {1668; SM60-LABEL: monotonic_seq_cst_i64_global_cta(1669; SM60: {1670; SM60-NEXT: .reg .b64 %rd<5>;1671; SM60-EMPTY:1672; SM60-NEXT: // %bb.0:1673; SM60-NEXT: ld.param.b64 %rd1, [monotonic_seq_cst_i64_global_cta_param_0];1674; SM60-NEXT: membar.cta;1675; SM60-NEXT: ld.param.b64 %rd2, [monotonic_seq_cst_i64_global_cta_param_1];1676; SM60-NEXT: ld.param.b64 %rd3, [monotonic_seq_cst_i64_global_cta_param_2];1677; SM60-NEXT: atom.cta.global.cas.b64 %rd4, [%rd1], %rd2, %rd3;1678; SM60-NEXT: st.param.b64 [func_retval0], %rd3;1679; SM60-NEXT: ret;1680 %pairold = cmpxchg ptr addrspace(1) %addr, i64 %cmp, i64 %new syncscope("block") monotonic seq_cst1681 ret i64 %new1682}1683 1684define i64 @acquire_monotonic_i64_global_cta(ptr addrspace(1) %addr, i64 %cmp, i64 %new) {1685; SM60-LABEL: acquire_monotonic_i64_global_cta(1686; SM60: {1687; SM60-NEXT: .reg .b64 %rd<5>;1688; SM60-EMPTY:1689; SM60-NEXT: // %bb.0:1690; SM60-NEXT: ld.param.b64 %rd1, [acquire_monotonic_i64_global_cta_param_0];1691; SM60-NEXT: ld.param.b64 %rd2, [acquire_monotonic_i64_global_cta_param_1];1692; SM60-NEXT: ld.param.b64 %rd3, [acquire_monotonic_i64_global_cta_param_2];1693; SM60-NEXT: atom.cta.global.cas.b64 %rd4, [%rd1], %rd2, %rd3;1694; SM60-NEXT: st.param.b64 [func_retval0], %rd3;1695; SM60-NEXT: ret;1696 %pairold = cmpxchg ptr addrspace(1) %addr, i64 %cmp, i64 %new syncscope("block") acquire monotonic1697 ret i64 %new1698}1699 1700define i64 @acquire_acquire_i64_global_cta(ptr addrspace(1) %addr, i64 %cmp, i64 %new) {1701; SM60-LABEL: acquire_acquire_i64_global_cta(1702; SM60: {1703; SM60-NEXT: .reg .b64 %rd<5>;1704; SM60-EMPTY:1705; SM60-NEXT: // %bb.0:1706; SM60-NEXT: ld.param.b64 %rd1, [acquire_acquire_i64_global_cta_param_0];1707; SM60-NEXT: ld.param.b64 %rd2, [acquire_acquire_i64_global_cta_param_1];1708; SM60-NEXT: ld.param.b64 %rd3, [acquire_acquire_i64_global_cta_param_2];1709; SM60-NEXT: atom.cta.global.cas.b64 %rd4, [%rd1], %rd2, %rd3;1710; SM60-NEXT: st.param.b64 [func_retval0], %rd3;1711; SM60-NEXT: ret;1712 %pairold = cmpxchg ptr addrspace(1) %addr, i64 %cmp, i64 %new syncscope("block") acquire acquire1713 ret i64 %new1714}1715 1716define i64 @acquire_seq_cst_i64_global_cta(ptr addrspace(1) %addr, i64 %cmp, i64 %new) {1717; SM60-LABEL: acquire_seq_cst_i64_global_cta(1718; SM60: {1719; SM60-NEXT: .reg .b64 %rd<5>;1720; SM60-EMPTY:1721; SM60-NEXT: // %bb.0:1722; SM60-NEXT: ld.param.b64 %rd1, [acquire_seq_cst_i64_global_cta_param_0];1723; SM60-NEXT: membar.cta;1724; SM60-NEXT: ld.param.b64 %rd2, [acquire_seq_cst_i64_global_cta_param_1];1725; SM60-NEXT: ld.param.b64 %rd3, [acquire_seq_cst_i64_global_cta_param_2];1726; SM60-NEXT: atom.cta.global.cas.b64 %rd4, [%rd1], %rd2, %rd3;1727; SM60-NEXT: st.param.b64 [func_retval0], %rd3;1728; SM60-NEXT: ret;1729 %pairold = cmpxchg ptr addrspace(1) %addr, i64 %cmp, i64 %new syncscope("block") acquire seq_cst1730 ret i64 %new1731}1732 1733define i64 @release_monotonic_i64_global_cta(ptr addrspace(1) %addr, i64 %cmp, i64 %new) {1734; SM60-LABEL: release_monotonic_i64_global_cta(1735; SM60: {1736; SM60-NEXT: .reg .b64 %rd<5>;1737; SM60-EMPTY:1738; SM60-NEXT: // %bb.0:1739; SM60-NEXT: ld.param.b64 %rd1, [release_monotonic_i64_global_cta_param_0];1740; SM60-NEXT: ld.param.b64 %rd2, [release_monotonic_i64_global_cta_param_1];1741; SM60-NEXT: ld.param.b64 %rd3, [release_monotonic_i64_global_cta_param_2];1742; SM60-NEXT: atom.cta.global.cas.b64 %rd4, [%rd1], %rd2, %rd3;1743; SM60-NEXT: st.param.b64 [func_retval0], %rd3;1744; SM60-NEXT: ret;1745 %pairold = cmpxchg ptr addrspace(1) %addr, i64 %cmp, i64 %new syncscope("block") release monotonic1746 ret i64 %new1747}1748 1749define i64 @release_acquire_i64_global_cta(ptr addrspace(1) %addr, i64 %cmp, i64 %new) {1750; SM60-LABEL: release_acquire_i64_global_cta(1751; SM60: {1752; SM60-NEXT: .reg .b64 %rd<5>;1753; SM60-EMPTY:1754; SM60-NEXT: // %bb.0:1755; SM60-NEXT: ld.param.b64 %rd1, [release_acquire_i64_global_cta_param_0];1756; SM60-NEXT: ld.param.b64 %rd2, [release_acquire_i64_global_cta_param_1];1757; SM60-NEXT: ld.param.b64 %rd3, [release_acquire_i64_global_cta_param_2];1758; SM60-NEXT: atom.cta.global.cas.b64 %rd4, [%rd1], %rd2, %rd3;1759; SM60-NEXT: st.param.b64 [func_retval0], %rd3;1760; SM60-NEXT: ret;1761 %pairold = cmpxchg ptr addrspace(1) %addr, i64 %cmp, i64 %new syncscope("block") release acquire1762 ret i64 %new1763}1764 1765define i64 @release_seq_cst_i64_global_cta(ptr addrspace(1) %addr, i64 %cmp, i64 %new) {1766; SM60-LABEL: release_seq_cst_i64_global_cta(1767; SM60: {1768; SM60-NEXT: .reg .b64 %rd<5>;1769; SM60-EMPTY:1770; SM60-NEXT: // %bb.0:1771; SM60-NEXT: ld.param.b64 %rd1, [release_seq_cst_i64_global_cta_param_0];1772; SM60-NEXT: membar.cta;1773; SM60-NEXT: ld.param.b64 %rd2, [release_seq_cst_i64_global_cta_param_1];1774; SM60-NEXT: ld.param.b64 %rd3, [release_seq_cst_i64_global_cta_param_2];1775; SM60-NEXT: atom.cta.global.cas.b64 %rd4, [%rd1], %rd2, %rd3;1776; SM60-NEXT: st.param.b64 [func_retval0], %rd3;1777; SM60-NEXT: ret;1778 %pairold = cmpxchg ptr addrspace(1) %addr, i64 %cmp, i64 %new syncscope("block") release seq_cst1779 ret i64 %new1780}1781 1782define i64 @acq_rel_monotonic_i64_global_cta(ptr addrspace(1) %addr, i64 %cmp, i64 %new) {1783; SM60-LABEL: acq_rel_monotonic_i64_global_cta(1784; SM60: {1785; SM60-NEXT: .reg .b64 %rd<5>;1786; SM60-EMPTY:1787; SM60-NEXT: // %bb.0:1788; SM60-NEXT: ld.param.b64 %rd1, [acq_rel_monotonic_i64_global_cta_param_0];1789; SM60-NEXT: ld.param.b64 %rd2, [acq_rel_monotonic_i64_global_cta_param_1];1790; SM60-NEXT: ld.param.b64 %rd3, [acq_rel_monotonic_i64_global_cta_param_2];1791; SM60-NEXT: atom.cta.global.cas.b64 %rd4, [%rd1], %rd2, %rd3;1792; SM60-NEXT: st.param.b64 [func_retval0], %rd3;1793; SM60-NEXT: ret;1794 %pairold = cmpxchg ptr addrspace(1) %addr, i64 %cmp, i64 %new syncscope("block") acq_rel monotonic1795 ret i64 %new1796}1797 1798define i64 @acq_rel_acquire_i64_global_cta(ptr addrspace(1) %addr, i64 %cmp, i64 %new) {1799; SM60-LABEL: acq_rel_acquire_i64_global_cta(1800; SM60: {1801; SM60-NEXT: .reg .b64 %rd<5>;1802; SM60-EMPTY:1803; SM60-NEXT: // %bb.0:1804; SM60-NEXT: ld.param.b64 %rd1, [acq_rel_acquire_i64_global_cta_param_0];1805; SM60-NEXT: ld.param.b64 %rd2, [acq_rel_acquire_i64_global_cta_param_1];1806; SM60-NEXT: ld.param.b64 %rd3, [acq_rel_acquire_i64_global_cta_param_2];1807; SM60-NEXT: atom.cta.global.cas.b64 %rd4, [%rd1], %rd2, %rd3;1808; SM60-NEXT: st.param.b64 [func_retval0], %rd3;1809; SM60-NEXT: ret;1810 %pairold = cmpxchg ptr addrspace(1) %addr, i64 %cmp, i64 %new syncscope("block") acq_rel acquire1811 ret i64 %new1812}1813 1814define i64 @acq_rel_seq_cst_i64_global_cta(ptr addrspace(1) %addr, i64 %cmp, i64 %new) {1815; SM60-LABEL: acq_rel_seq_cst_i64_global_cta(1816; SM60: {1817; SM60-NEXT: .reg .b64 %rd<5>;1818; SM60-EMPTY:1819; SM60-NEXT: // %bb.0:1820; SM60-NEXT: ld.param.b64 %rd1, [acq_rel_seq_cst_i64_global_cta_param_0];1821; SM60-NEXT: membar.cta;1822; SM60-NEXT: ld.param.b64 %rd2, [acq_rel_seq_cst_i64_global_cta_param_1];1823; SM60-NEXT: ld.param.b64 %rd3, [acq_rel_seq_cst_i64_global_cta_param_2];1824; SM60-NEXT: atom.cta.global.cas.b64 %rd4, [%rd1], %rd2, %rd3;1825; SM60-NEXT: st.param.b64 [func_retval0], %rd3;1826; SM60-NEXT: ret;1827 %pairold = cmpxchg ptr addrspace(1) %addr, i64 %cmp, i64 %new syncscope("block") acq_rel seq_cst1828 ret i64 %new1829}1830 1831define i64 @seq_cst_monotonic_i64_global_cta(ptr addrspace(1) %addr, i64 %cmp, i64 %new) {1832; SM60-LABEL: seq_cst_monotonic_i64_global_cta(1833; SM60: {1834; SM60-NEXT: .reg .b64 %rd<5>;1835; SM60-EMPTY:1836; SM60-NEXT: // %bb.0:1837; SM60-NEXT: ld.param.b64 %rd1, [seq_cst_monotonic_i64_global_cta_param_0];1838; SM60-NEXT: membar.cta;1839; SM60-NEXT: ld.param.b64 %rd2, [seq_cst_monotonic_i64_global_cta_param_1];1840; SM60-NEXT: ld.param.b64 %rd3, [seq_cst_monotonic_i64_global_cta_param_2];1841; SM60-NEXT: atom.cta.global.cas.b64 %rd4, [%rd1], %rd2, %rd3;1842; SM60-NEXT: st.param.b64 [func_retval0], %rd3;1843; SM60-NEXT: ret;1844 %pairold = cmpxchg ptr addrspace(1) %addr, i64 %cmp, i64 %new syncscope("block") seq_cst monotonic1845 ret i64 %new1846}1847 1848define i64 @seq_cst_acquire_i64_global_cta(ptr addrspace(1) %addr, i64 %cmp, i64 %new) {1849; SM60-LABEL: seq_cst_acquire_i64_global_cta(1850; SM60: {1851; SM60-NEXT: .reg .b64 %rd<5>;1852; SM60-EMPTY:1853; SM60-NEXT: // %bb.0:1854; SM60-NEXT: ld.param.b64 %rd1, [seq_cst_acquire_i64_global_cta_param_0];1855; SM60-NEXT: membar.cta;1856; SM60-NEXT: ld.param.b64 %rd2, [seq_cst_acquire_i64_global_cta_param_1];1857; SM60-NEXT: ld.param.b64 %rd3, [seq_cst_acquire_i64_global_cta_param_2];1858; SM60-NEXT: atom.cta.global.cas.b64 %rd4, [%rd1], %rd2, %rd3;1859; SM60-NEXT: st.param.b64 [func_retval0], %rd3;1860; SM60-NEXT: ret;1861 %pairold = cmpxchg ptr addrspace(1) %addr, i64 %cmp, i64 %new syncscope("block") seq_cst acquire1862 ret i64 %new1863}1864 1865define i64 @seq_cst_seq_cst_i64_global_cta(ptr addrspace(1) %addr, i64 %cmp, i64 %new) {1866; SM60-LABEL: seq_cst_seq_cst_i64_global_cta(1867; SM60: {1868; SM60-NEXT: .reg .b64 %rd<5>;1869; SM60-EMPTY:1870; SM60-NEXT: // %bb.0:1871; SM60-NEXT: ld.param.b64 %rd1, [seq_cst_seq_cst_i64_global_cta_param_0];1872; SM60-NEXT: membar.cta;1873; SM60-NEXT: ld.param.b64 %rd2, [seq_cst_seq_cst_i64_global_cta_param_1];1874; SM60-NEXT: ld.param.b64 %rd3, [seq_cst_seq_cst_i64_global_cta_param_2];1875; SM60-NEXT: atom.cta.global.cas.b64 %rd4, [%rd1], %rd2, %rd3;1876; SM60-NEXT: st.param.b64 [func_retval0], %rd3;1877; SM60-NEXT: ret;1878 %pairold = cmpxchg ptr addrspace(1) %addr, i64 %cmp, i64 %new syncscope("block") seq_cst seq_cst1879 ret i64 %new1880}1881 1882define i8 @acq_rel_acquire_i8_global(ptr addrspace(1) %addr, i8 %cmp, i8 %new) {1883; SM60-LABEL: acq_rel_acquire_i8_global(1884; SM60: {1885; SM60-NEXT: .reg .pred %p<3>;1886; SM60-NEXT: .reg .b16 %rs<2>;1887; SM60-NEXT: .reg .b32 %r<17>;1888; SM60-NEXT: .reg .b64 %rd<3>;1889; SM60-EMPTY:1890; SM60-NEXT: // %bb.0:1891; SM60-NEXT: ld.param.b8 %rs1, [acq_rel_acquire_i8_global_param_2];1892; SM60-NEXT: ld.param.b64 %rd2, [acq_rel_acquire_i8_global_param_0];1893; SM60-NEXT: membar.sys;1894; SM60-NEXT: ld.param.b8 %r7, [acq_rel_acquire_i8_global_param_1];1895; SM60-NEXT: and.b64 %rd1, %rd2, -4;1896; SM60-NEXT: cvt.u32.u64 %r8, %rd2;1897; SM60-NEXT: and.b32 %r9, %r8, 3;1898; SM60-NEXT: shl.b32 %r1, %r9, 3;1899; SM60-NEXT: mov.b32 %r10, 255;1900; SM60-NEXT: shl.b32 %r11, %r10, %r1;1901; SM60-NEXT: not.b32 %r2, %r11;1902; SM60-NEXT: cvt.u32.u16 %r12, %rs1;1903; SM60-NEXT: shl.b32 %r3, %r12, %r1;1904; SM60-NEXT: shl.b32 %r4, %r7, %r1;1905; SM60-NEXT: ld.global.b32 %r13, [%rd1];1906; SM60-NEXT: and.b32 %r16, %r13, %r2;1907; SM60-NEXT: $L__BB60_1: // %partword.cmpxchg.loop1908; SM60-NEXT: // =>This Inner Loop Header: Depth=11909; SM60-NEXT: or.b32 %r14, %r16, %r3;1910; SM60-NEXT: or.b32 %r15, %r16, %r4;1911; SM60-NEXT: atom.sys.global.cas.b32 %r5, [%rd1], %r15, %r14;1912; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;1913; SM60-NEXT: @%p1 bra $L__BB60_3;1914; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure1915; SM60-NEXT: // in Loop: Header=BB60_1 Depth=11916; SM60-NEXT: and.b32 %r6, %r5, %r2;1917; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;1918; SM60-NEXT: mov.b32 %r16, %r6;1919; SM60-NEXT: @%p2 bra $L__BB60_1;1920; SM60-NEXT: $L__BB60_3: // %partword.cmpxchg.end1921; SM60-NEXT: membar.sys;1922; SM60-NEXT: st.param.b32 [func_retval0], %r12;1923; SM60-NEXT: ret;1924 %pairold = cmpxchg ptr addrspace(1) %addr, i8 %cmp, i8 %new acq_rel acquire1925 ret i8 %new1926}1927 1928define i32 @acq_rel_acquire_i32_global(ptr addrspace(1) %addr, i32 %cmp, i32 %new) {1929; SM60-LABEL: acq_rel_acquire_i32_global(1930; SM60: {1931; SM60-NEXT: .reg .b32 %r<4>;1932; SM60-NEXT: .reg .b64 %rd<2>;1933; SM60-EMPTY:1934; SM60-NEXT: // %bb.0:1935; SM60-NEXT: ld.param.b64 %rd1, [acq_rel_acquire_i32_global_param_0];1936; SM60-NEXT: ld.param.b32 %r1, [acq_rel_acquire_i32_global_param_1];1937; SM60-NEXT: ld.param.b32 %r2, [acq_rel_acquire_i32_global_param_2];1938; SM60-NEXT: atom.sys.global.cas.b32 %r3, [%rd1], %r1, %r2;1939; SM60-NEXT: st.param.b32 [func_retval0], %r2;1940; SM60-NEXT: ret;1941 %pairold = cmpxchg ptr addrspace(1) %addr, i32 %cmp, i32 %new acq_rel acquire1942 ret i32 %new1943}1944 1945define i32 @acq_rel_acquire_i32_global_sys(ptr addrspace(1) %addr, i32 %cmp, i32 %new) {1946; SM60-LABEL: acq_rel_acquire_i32_global_sys(1947; SM60: {1948; SM60-NEXT: .reg .b32 %r<4>;1949; SM60-NEXT: .reg .b64 %rd<2>;1950; SM60-EMPTY:1951; SM60-NEXT: // %bb.0:1952; SM60-NEXT: ld.param.b64 %rd1, [acq_rel_acquire_i32_global_sys_param_0];1953; SM60-NEXT: ld.param.b32 %r1, [acq_rel_acquire_i32_global_sys_param_1];1954; SM60-NEXT: ld.param.b32 %r2, [acq_rel_acquire_i32_global_sys_param_2];1955; SM60-NEXT: atom.sys.global.cas.b32 %r3, [%rd1], %r1, %r2;1956; SM60-NEXT: st.param.b32 [func_retval0], %r2;1957; SM60-NEXT: ret;1958 %pairold = cmpxchg ptr addrspace(1) %addr, i32 %cmp, i32 %new syncscope("") acq_rel acquire1959 ret i32 %new1960}1961 1962define i32 @acq_rel_acquire_i32_global_gpu(ptr addrspace(1) %addr, i32 %cmp, i32 %new) {1963; SM60-LABEL: acq_rel_acquire_i32_global_gpu(1964; SM60: {1965; SM60-NEXT: .reg .b32 %r<4>;1966; SM60-NEXT: .reg .b64 %rd<2>;1967; SM60-EMPTY:1968; SM60-NEXT: // %bb.0:1969; SM60-NEXT: ld.param.b64 %rd1, [acq_rel_acquire_i32_global_gpu_param_0];1970; SM60-NEXT: ld.param.b32 %r1, [acq_rel_acquire_i32_global_gpu_param_1];1971; SM60-NEXT: ld.param.b32 %r2, [acq_rel_acquire_i32_global_gpu_param_2];1972; SM60-NEXT: atom.gpu.global.cas.b32 %r3, [%rd1], %r1, %r2;1973; SM60-NEXT: st.param.b32 [func_retval0], %r2;1974; SM60-NEXT: ret;1975 %pairold = cmpxchg ptr addrspace(1) %addr, i32 %cmp, i32 %new syncscope("device") acq_rel acquire1976 ret i32 %new1977}1978 1979define i8 @acq_rel_acquire_i8_generic_cta(ptr %addr, i8 %cmp, i8 %new) {1980; SM60-LABEL: acq_rel_acquire_i8_generic_cta(1981; SM60: {1982; SM60-NEXT: .reg .pred %p<3>;1983; SM60-NEXT: .reg .b16 %rs<2>;1984; SM60-NEXT: .reg .b32 %r<17>;1985; SM60-NEXT: .reg .b64 %rd<3>;1986; SM60-EMPTY:1987; SM60-NEXT: // %bb.0:1988; SM60-NEXT: ld.param.b8 %rs1, [acq_rel_acquire_i8_generic_cta_param_2];1989; SM60-NEXT: ld.param.b64 %rd2, [acq_rel_acquire_i8_generic_cta_param_0];1990; SM60-NEXT: membar.cta;1991; SM60-NEXT: ld.param.b8 %r7, [acq_rel_acquire_i8_generic_cta_param_1];1992; SM60-NEXT: and.b64 %rd1, %rd2, -4;1993; SM60-NEXT: cvt.u32.u64 %r8, %rd2;1994; SM60-NEXT: and.b32 %r9, %r8, 3;1995; SM60-NEXT: shl.b32 %r1, %r9, 3;1996; SM60-NEXT: mov.b32 %r10, 255;1997; SM60-NEXT: shl.b32 %r11, %r10, %r1;1998; SM60-NEXT: not.b32 %r2, %r11;1999; SM60-NEXT: cvt.u32.u16 %r12, %rs1;2000; SM60-NEXT: shl.b32 %r3, %r12, %r1;2001; SM60-NEXT: shl.b32 %r4, %r7, %r1;2002; SM60-NEXT: ld.b32 %r13, [%rd1];2003; SM60-NEXT: and.b32 %r16, %r13, %r2;2004; SM60-NEXT: $L__BB64_1: // %partword.cmpxchg.loop2005; SM60-NEXT: // =>This Inner Loop Header: Depth=12006; SM60-NEXT: or.b32 %r14, %r16, %r3;2007; SM60-NEXT: or.b32 %r15, %r16, %r4;2008; SM60-NEXT: atom.cta.cas.b32 %r5, [%rd1], %r15, %r14;2009; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;2010; SM60-NEXT: @%p1 bra $L__BB64_3;2011; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure2012; SM60-NEXT: // in Loop: Header=BB64_1 Depth=12013; SM60-NEXT: and.b32 %r6, %r5, %r2;2014; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;2015; SM60-NEXT: mov.b32 %r16, %r6;2016; SM60-NEXT: @%p2 bra $L__BB64_1;2017; SM60-NEXT: $L__BB64_3: // %partword.cmpxchg.end2018; SM60-NEXT: membar.cta;2019; SM60-NEXT: st.param.b32 [func_retval0], %r12;2020; SM60-NEXT: ret;2021 %pairold = cmpxchg ptr %addr, i8 %cmp, i8 %new syncscope("block") acq_rel acquire2022 ret i8 %new2023}2024 2025define i8 @acq_rel_acquire_i8_shared_cta(ptr addrspace(3) %addr, i8 %cmp, i8 %new) {2026; SM60-LABEL: acq_rel_acquire_i8_shared_cta(2027; SM60: {2028; SM60-NEXT: .reg .pred %p<3>;2029; SM60-NEXT: .reg .b16 %rs<2>;2030; SM60-NEXT: .reg .b32 %r<17>;2031; SM60-NEXT: .reg .b64 %rd<3>;2032; SM60-EMPTY:2033; SM60-NEXT: // %bb.0:2034; SM60-NEXT: ld.param.b8 %rs1, [acq_rel_acquire_i8_shared_cta_param_2];2035; SM60-NEXT: ld.param.b64 %rd2, [acq_rel_acquire_i8_shared_cta_param_0];2036; SM60-NEXT: membar.cta;2037; SM60-NEXT: ld.param.b8 %r7, [acq_rel_acquire_i8_shared_cta_param_1];2038; SM60-NEXT: and.b64 %rd1, %rd2, -4;2039; SM60-NEXT: cvt.u32.u64 %r8, %rd2;2040; SM60-NEXT: and.b32 %r9, %r8, 3;2041; SM60-NEXT: shl.b32 %r1, %r9, 3;2042; SM60-NEXT: mov.b32 %r10, 255;2043; SM60-NEXT: shl.b32 %r11, %r10, %r1;2044; SM60-NEXT: not.b32 %r2, %r11;2045; SM60-NEXT: cvt.u32.u16 %r12, %rs1;2046; SM60-NEXT: shl.b32 %r3, %r12, %r1;2047; SM60-NEXT: shl.b32 %r4, %r7, %r1;2048; SM60-NEXT: ld.shared.b32 %r13, [%rd1];2049; SM60-NEXT: and.b32 %r16, %r13, %r2;2050; SM60-NEXT: $L__BB65_1: // %partword.cmpxchg.loop2051; SM60-NEXT: // =>This Inner Loop Header: Depth=12052; SM60-NEXT: or.b32 %r14, %r16, %r3;2053; SM60-NEXT: or.b32 %r15, %r16, %r4;2054; SM60-NEXT: atom.cta.shared.cas.b32 %r5, [%rd1], %r15, %r14;2055; SM60-NEXT: setp.eq.b32 %p1, %r5, %r15;2056; SM60-NEXT: @%p1 bra $L__BB65_3;2057; SM60-NEXT: // %bb.2: // %partword.cmpxchg.failure2058; SM60-NEXT: // in Loop: Header=BB65_1 Depth=12059; SM60-NEXT: and.b32 %r6, %r5, %r2;2060; SM60-NEXT: setp.ne.b32 %p2, %r16, %r6;2061; SM60-NEXT: mov.b32 %r16, %r6;2062; SM60-NEXT: @%p2 bra $L__BB65_1;2063; SM60-NEXT: $L__BB65_3: // %partword.cmpxchg.end2064; SM60-NEXT: membar.cta;2065; SM60-NEXT: st.param.b32 [func_retval0], %r12;2066; SM60-NEXT: ret;2067 %pairold = cmpxchg ptr addrspace(3) %addr, i8 %cmp, i8 %new syncscope("block") acq_rel acquire2068 ret i8 %new2069}2070 2071define i32 @acq_rel_acquire_i32_generic_cta(ptr %addr, i32 %cmp, i32 %new) {2072; SM60-LABEL: acq_rel_acquire_i32_generic_cta(2073; SM60: {2074; SM60-NEXT: .reg .b32 %r<4>;2075; SM60-NEXT: .reg .b64 %rd<2>;2076; SM60-EMPTY:2077; SM60-NEXT: // %bb.0:2078; SM60-NEXT: ld.param.b64 %rd1, [acq_rel_acquire_i32_generic_cta_param_0];2079; SM60-NEXT: ld.param.b32 %r1, [acq_rel_acquire_i32_generic_cta_param_1];2080; SM60-NEXT: ld.param.b32 %r2, [acq_rel_acquire_i32_generic_cta_param_2];2081; SM60-NEXT: atom.cta.cas.b32 %r3, [%rd1], %r1, %r2;2082; SM60-NEXT: st.param.b32 [func_retval0], %r2;2083; SM60-NEXT: ret;2084 %pairold = cmpxchg ptr %addr, i32 %cmp, i32 %new syncscope("block") acq_rel acquire2085 ret i32 %new2086}2087 2088define i32 @acq_rel_acquire_i32_shared_cta(ptr addrspace(3) %addr, i32 %cmp, i32 %new) {2089; SM60-LABEL: acq_rel_acquire_i32_shared_cta(2090; SM60: {2091; SM60-NEXT: .reg .b32 %r<4>;2092; SM60-NEXT: .reg .b64 %rd<2>;2093; SM60-EMPTY:2094; SM60-NEXT: // %bb.0:2095; SM60-NEXT: ld.param.b64 %rd1, [acq_rel_acquire_i32_shared_cta_param_0];2096; SM60-NEXT: ld.param.b32 %r1, [acq_rel_acquire_i32_shared_cta_param_1];2097; SM60-NEXT: ld.param.b32 %r2, [acq_rel_acquire_i32_shared_cta_param_2];2098; SM60-NEXT: atom.cta.shared.cas.b32 %r3, [%rd1], %r1, %r2;2099; SM60-NEXT: st.param.b32 [func_retval0], %r2;2100; SM60-NEXT: ret;2101 %pairold = cmpxchg ptr addrspace(3) %addr, i32 %cmp, i32 %new syncscope("block") acq_rel acquire2102 ret i32 %new2103}2104 2105