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