brintos / llvm-project-archived public Read only

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