1702 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --default-march nvptx64 --version 52; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_32 | FileCheck %s --check-prefixes=SM30,CHECK3; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_32 | %ptxas-verify %}4; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_70 -mattr=+ptx63 | FileCheck %s --check-prefixes=SM70,CHECK5; RUN: %if ptxas-sm_70 && ptxas-isa-6.3 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_70 -mattr=+ptx63 | %ptxas-verify -arch=sm_70 %}6 7; TODO: these are system scope, but are compiled to gpu scope..8; TODO: these are seq_cst, but are compiled to relaxed..9 10 11; CHECK-LABEL: relaxed_sys_i812define i8 @relaxed_sys_i8(ptr %addr, i8 %cmp, i8 %new) {13; SM30-LABEL: relaxed_sys_i8(14; SM30: {15; SM30-NEXT: .reg .pred %p<3>;16; SM30-NEXT: .reg .b16 %rs<2>;17; SM30-NEXT: .reg .b32 %r<17>;18; SM30-NEXT: .reg .b64 %rd<3>;19; SM30-EMPTY:20; SM30-NEXT: // %bb.0:21; SM30-NEXT: ld.param.b8 %rs1, [relaxed_sys_i8_param_2];22; SM30-NEXT: ld.param.b64 %rd2, [relaxed_sys_i8_param_0];23; SM30-NEXT: and.b64 %rd1, %rd2, -4;24; SM30-NEXT: ld.param.b8 %r7, [relaxed_sys_i8_param_1];25; SM30-NEXT: cvt.u32.u64 %r8, %rd2;26; SM30-NEXT: and.b32 %r9, %r8, 3;27; SM30-NEXT: shl.b32 %r1, %r9, 3;28; SM30-NEXT: mov.b32 %r10, 255;29; SM30-NEXT: shl.b32 %r11, %r10, %r1;30; SM30-NEXT: not.b32 %r2, %r11;31; SM30-NEXT: cvt.u32.u16 %r12, %rs1;32; SM30-NEXT: shl.b32 %r3, %r12, %r1;33; SM30-NEXT: shl.b32 %r4, %r7, %r1;34; SM30-NEXT: ld.b32 %r13, [%rd1];35; SM30-NEXT: and.b32 %r16, %r13, %r2;36; SM30-NEXT: $L__BB0_1: // %partword.cmpxchg.loop37; SM30-NEXT: // =>This Inner Loop Header: Depth=138; SM30-NEXT: or.b32 %r14, %r16, %r3;39; SM30-NEXT: or.b32 %r15, %r16, %r4;40; SM30-NEXT: atom.cas.b32 %r5, [%rd1], %r15, %r14;41; SM30-NEXT: setp.eq.b32 %p1, %r5, %r15;42; SM30-NEXT: @%p1 bra $L__BB0_3;43; SM30-NEXT: // %bb.2: // %partword.cmpxchg.failure44; SM30-NEXT: // in Loop: Header=BB0_1 Depth=145; SM30-NEXT: and.b32 %r6, %r5, %r2;46; SM30-NEXT: setp.ne.b32 %p2, %r16, %r6;47; SM30-NEXT: mov.b32 %r16, %r6;48; SM30-NEXT: @%p2 bra $L__BB0_1;49; SM30-NEXT: $L__BB0_3: // %partword.cmpxchg.end50; SM30-NEXT: st.param.b32 [func_retval0], %r12;51; SM30-NEXT: ret;52;53; SM70-LABEL: relaxed_sys_i8(54; SM70: {55; SM70-NEXT: .reg .pred %p<3>;56; SM70-NEXT: .reg .b16 %rs<2>;57; SM70-NEXT: .reg .b32 %r<17>;58; SM70-NEXT: .reg .b64 %rd<3>;59; SM70-EMPTY:60; SM70-NEXT: // %bb.0:61; SM70-NEXT: ld.param.b8 %rs1, [relaxed_sys_i8_param_2];62; SM70-NEXT: ld.param.b64 %rd2, [relaxed_sys_i8_param_0];63; SM70-NEXT: and.b64 %rd1, %rd2, -4;64; SM70-NEXT: ld.param.b8 %r7, [relaxed_sys_i8_param_1];65; SM70-NEXT: cvt.u32.u64 %r8, %rd2;66; SM70-NEXT: and.b32 %r9, %r8, 3;67; SM70-NEXT: shl.b32 %r1, %r9, 3;68; SM70-NEXT: mov.b32 %r10, 255;69; SM70-NEXT: shl.b32 %r11, %r10, %r1;70; SM70-NEXT: not.b32 %r2, %r11;71; SM70-NEXT: cvt.u32.u16 %r12, %rs1;72; SM70-NEXT: shl.b32 %r3, %r12, %r1;73; SM70-NEXT: shl.b32 %r4, %r7, %r1;74; SM70-NEXT: ld.b32 %r13, [%rd1];75; SM70-NEXT: and.b32 %r16, %r13, %r2;76; SM70-NEXT: $L__BB0_1: // %partword.cmpxchg.loop77; SM70-NEXT: // =>This Inner Loop Header: Depth=178; SM70-NEXT: or.b32 %r14, %r16, %r3;79; SM70-NEXT: or.b32 %r15, %r16, %r4;80; SM70-NEXT: atom.relaxed.sys.cas.b32 %r5, [%rd1], %r15, %r14;81; SM70-NEXT: setp.eq.b32 %p1, %r5, %r15;82; SM70-NEXT: @%p1 bra $L__BB0_3;83; SM70-NEXT: // %bb.2: // %partword.cmpxchg.failure84; SM70-NEXT: // in Loop: Header=BB0_1 Depth=185; SM70-NEXT: and.b32 %r6, %r5, %r2;86; SM70-NEXT: setp.ne.b32 %p2, %r16, %r6;87; SM70-NEXT: mov.b32 %r16, %r6;88; SM70-NEXT: @%p2 bra $L__BB0_1;89; SM70-NEXT: $L__BB0_3: // %partword.cmpxchg.end90; SM70-NEXT: st.param.b32 [func_retval0], %r12;91; SM70-NEXT: ret;92; SM90-LABEL: relaxed_sys_i8(93; SM90: {94; SM90-NEXT: .reg .pred %p<3>;95; SM90-NEXT: .reg .b16 %rs<2>;96; SM90-NEXT: .reg .b32 %r<21>;97; SM90-NEXT: .reg .b64 %rd<3>;98; SM90-EMPTY:99; SM90-NEXT: // %bb.0:100; SM90-NEXT: ld.param.u8 %rs1, [relaxed_sys_i8_param_2];101; SM90-NEXT: ld.param.u64 %rd2, [relaxed_sys_i8_param_0];102; SM90-NEXT: and.b64 %rd1, %rd2, -4;103; SM90-NEXT: cvt.u32.u64 %r9, %rd2;104; SM90-NEXT: and.b32 %r10, %r9, 3;105; SM90-NEXT: shl.b32 %r1, %r10, 3;106; SM90-NEXT: mov.b32 %r11, 255;107; SM90-NEXT: shl.b32 %r12, %r11, %r1;108; SM90-NEXT: not.b32 %r2, %r12;109; SM90-NEXT: cvt.u32.u16 %r13, %rs1;110; SM90-NEXT: and.b32 %r14, %r13, 255;111; SM90-NEXT: shl.b32 %r3, %r14, %r1;112; SM90-NEXT: ld.param.u8 %r15, [relaxed_sys_i8_param_1];113; SM90-NEXT: shl.b32 %r4, %r15, %r1;114; SM90-NEXT: ld.u32 %r16, [%rd1];115; SM90-NEXT: and.b32 %r20, %r16, %r2;116; SM90-NEXT: $L__BB0_1: // %partword.cmpxchg.loop117; SM90-NEXT: // =>This Inner Loop Header: Depth=1118; SM90-NEXT: or.b32 %r17, %r20, %r3;119; SM90-NEXT: or.b32 %r18, %r20, %r4;120; SM90-NEXT: atom.cas.b32 %r7, [%rd1], %r18, %r17;121; SM90-NEXT: setp.eq.s32 %p1, %r7, %r18;122; SM90-NEXT: @%p1 bra $L__BB0_3;123; SM90-NEXT: // %bb.2: // %partword.cmpxchg.failure124; SM90-NEXT: // in Loop: Header=BB0_1 Depth=1125; SM90-NEXT: and.b32 %r8, %r7, %r2;126; SM90-NEXT: setp.ne.s32 %p2, %r20, %r8;127; SM90-NEXT: mov.u32 %r20, %r8;128; SM90-NEXT: @%p2 bra $L__BB0_1;129; SM90-NEXT: $L__BB0_3: // %partword.cmpxchg.end130; SM90-NEXT: st.param.b32 [func_retval0], %r13;131; SM90-NEXT: ret;132 %pairold = cmpxchg ptr %addr, i8 %cmp, i8 %new monotonic monotonic133 ret i8 %new134}135 136define i8 @acquire_sys_i8(ptr %addr, i8 %cmp, i8 %new) {137; SM30-LABEL: acquire_sys_i8(138; SM30: {139; SM30-NEXT: .reg .pred %p<3>;140; SM30-NEXT: .reg .b16 %rs<2>;141; SM30-NEXT: .reg .b32 %r<17>;142; SM30-NEXT: .reg .b64 %rd<3>;143; SM30-EMPTY:144; SM30-NEXT: // %bb.0:145; SM30-NEXT: ld.param.b8 %rs1, [acquire_sys_i8_param_2];146; SM30-NEXT: ld.param.b64 %rd2, [acquire_sys_i8_param_0];147; SM30-NEXT: and.b64 %rd1, %rd2, -4;148; SM30-NEXT: ld.param.b8 %r7, [acquire_sys_i8_param_1];149; SM30-NEXT: cvt.u32.u64 %r8, %rd2;150; SM30-NEXT: and.b32 %r9, %r8, 3;151; SM30-NEXT: shl.b32 %r1, %r9, 3;152; SM30-NEXT: mov.b32 %r10, 255;153; SM30-NEXT: shl.b32 %r11, %r10, %r1;154; SM30-NEXT: not.b32 %r2, %r11;155; SM30-NEXT: cvt.u32.u16 %r12, %rs1;156; SM30-NEXT: shl.b32 %r3, %r12, %r1;157; SM30-NEXT: shl.b32 %r4, %r7, %r1;158; SM30-NEXT: ld.b32 %r13, [%rd1];159; SM30-NEXT: and.b32 %r16, %r13, %r2;160; SM30-NEXT: $L__BB1_1: // %partword.cmpxchg.loop161; SM30-NEXT: // =>This Inner Loop Header: Depth=1162; SM30-NEXT: or.b32 %r14, %r16, %r3;163; SM30-NEXT: or.b32 %r15, %r16, %r4;164; SM30-NEXT: atom.cas.b32 %r5, [%rd1], %r15, %r14;165; SM30-NEXT: setp.eq.b32 %p1, %r5, %r15;166; SM30-NEXT: @%p1 bra $L__BB1_3;167; SM30-NEXT: // %bb.2: // %partword.cmpxchg.failure168; SM30-NEXT: // in Loop: Header=BB1_1 Depth=1169; SM30-NEXT: and.b32 %r6, %r5, %r2;170; SM30-NEXT: setp.ne.b32 %p2, %r16, %r6;171; SM30-NEXT: mov.b32 %r16, %r6;172; SM30-NEXT: @%p2 bra $L__BB1_1;173; SM30-NEXT: $L__BB1_3: // %partword.cmpxchg.end174; SM30-NEXT: membar.sys;175; SM30-NEXT: st.param.b32 [func_retval0], %r12;176; SM30-NEXT: ret;177;178; SM70-LABEL: acquire_sys_i8(179; SM70: {180; SM70-NEXT: .reg .pred %p<3>;181; SM70-NEXT: .reg .b16 %rs<2>;182; SM70-NEXT: .reg .b32 %r<17>;183; SM70-NEXT: .reg .b64 %rd<3>;184; SM70-EMPTY:185; SM70-NEXT: // %bb.0:186; SM70-NEXT: ld.param.b8 %rs1, [acquire_sys_i8_param_2];187; SM70-NEXT: ld.param.b64 %rd2, [acquire_sys_i8_param_0];188; SM70-NEXT: and.b64 %rd1, %rd2, -4;189; SM70-NEXT: ld.param.b8 %r7, [acquire_sys_i8_param_1];190; SM70-NEXT: cvt.u32.u64 %r8, %rd2;191; SM70-NEXT: and.b32 %r9, %r8, 3;192; SM70-NEXT: shl.b32 %r1, %r9, 3;193; SM70-NEXT: mov.b32 %r10, 255;194; SM70-NEXT: shl.b32 %r11, %r10, %r1;195; SM70-NEXT: not.b32 %r2, %r11;196; SM70-NEXT: cvt.u32.u16 %r12, %rs1;197; SM70-NEXT: shl.b32 %r3, %r12, %r1;198; SM70-NEXT: shl.b32 %r4, %r7, %r1;199; SM70-NEXT: ld.b32 %r13, [%rd1];200; SM70-NEXT: and.b32 %r16, %r13, %r2;201; SM70-NEXT: $L__BB1_1: // %partword.cmpxchg.loop202; SM70-NEXT: // =>This Inner Loop Header: Depth=1203; SM70-NEXT: or.b32 %r14, %r16, %r3;204; SM70-NEXT: or.b32 %r15, %r16, %r4;205; SM70-NEXT: atom.relaxed.sys.cas.b32 %r5, [%rd1], %r15, %r14;206; SM70-NEXT: setp.eq.b32 %p1, %r5, %r15;207; SM70-NEXT: @%p1 bra $L__BB1_3;208; SM70-NEXT: // %bb.2: // %partword.cmpxchg.failure209; SM70-NEXT: // in Loop: Header=BB1_1 Depth=1210; SM70-NEXT: and.b32 %r6, %r5, %r2;211; SM70-NEXT: setp.ne.b32 %p2, %r16, %r6;212; SM70-NEXT: mov.b32 %r16, %r6;213; SM70-NEXT: @%p2 bra $L__BB1_1;214; SM70-NEXT: $L__BB1_3: // %partword.cmpxchg.end215; SM70-NEXT: fence.acq_rel.sys;216; SM70-NEXT: st.param.b32 [func_retval0], %r12;217; SM70-NEXT: ret;218; SM90-LABEL: acquire_sys_i8(219; SM90: {220; SM90-NEXT: .reg .pred %p<3>;221; SM90-NEXT: .reg .b16 %rs<2>;222; SM90-NEXT: .reg .b32 %r<21>;223; SM90-NEXT: .reg .b64 %rd<3>;224; SM90-EMPTY:225; SM90-NEXT: // %bb.0:226; SM90-NEXT: ld.param.u8 %rs1, [acquire_sys_i8_param_2];227; SM90-NEXT: ld.param.u64 %rd2, [acquire_sys_i8_param_0];228; SM90-NEXT: and.b64 %rd1, %rd2, -4;229; SM90-NEXT: cvt.u32.u64 %r9, %rd2;230; SM90-NEXT: and.b32 %r10, %r9, 3;231; SM90-NEXT: shl.b32 %r1, %r10, 3;232; SM90-NEXT: mov.b32 %r11, 255;233; SM90-NEXT: shl.b32 %r12, %r11, %r1;234; SM90-NEXT: not.b32 %r2, %r12;235; SM90-NEXT: cvt.u32.u16 %r13, %rs1;236; SM90-NEXT: and.b32 %r14, %r13, 255;237; SM90-NEXT: shl.b32 %r3, %r14, %r1;238; SM90-NEXT: ld.param.u8 %r15, [acquire_sys_i8_param_1];239; SM90-NEXT: shl.b32 %r4, %r15, %r1;240; SM90-NEXT: ld.u32 %r16, [%rd1];241; SM90-NEXT: and.b32 %r20, %r16, %r2;242; SM90-NEXT: $L__BB1_1: // %partword.cmpxchg.loop243; SM90-NEXT: // =>This Inner Loop Header: Depth=1244; SM90-NEXT: or.b32 %r17, %r20, %r3;245; SM90-NEXT: or.b32 %r18, %r20, %r4;246; SM90-NEXT: atom.cas.b32 %r7, [%rd1], %r18, %r17;247; SM90-NEXT: setp.eq.s32 %p1, %r7, %r18;248; SM90-NEXT: @%p1 bra $L__BB1_3;249; SM90-NEXT: // %bb.2: // %partword.cmpxchg.failure250; SM90-NEXT: // in Loop: Header=BB1_1 Depth=1251; SM90-NEXT: and.b32 %r8, %r7, %r2;252; SM90-NEXT: setp.ne.s32 %p2, %r20, %r8;253; SM90-NEXT: mov.u32 %r20, %r8;254; SM90-NEXT: @%p2 bra $L__BB1_1;255; SM90-NEXT: $L__BB1_3: // %partword.cmpxchg.end256; SM90-NEXT: fence.acquire.sys;257; SM90-NEXT: st.param.b32 [func_retval0], %r13;258; SM90-NEXT: ret;259 %pairold = cmpxchg ptr %addr, i8 %cmp, i8 %new acquire acquire260 ret i8 %new261}262 263define i8 @release_sys_i8(ptr %addr, i8 %cmp, i8 %new) {264; SM30-LABEL: release_sys_i8(265; SM30: {266; SM30-NEXT: .reg .pred %p<3>;267; SM30-NEXT: .reg .b16 %rs<2>;268; SM30-NEXT: .reg .b32 %r<17>;269; SM30-NEXT: .reg .b64 %rd<3>;270; SM30-EMPTY:271; SM30-NEXT: // %bb.0:272; SM30-NEXT: ld.param.b8 %rs1, [release_sys_i8_param_2];273; SM30-NEXT: ld.param.b64 %rd2, [release_sys_i8_param_0];274; SM30-NEXT: membar.sys;275; SM30-NEXT: ld.param.b8 %r7, [release_sys_i8_param_1];276; SM30-NEXT: and.b64 %rd1, %rd2, -4;277; SM30-NEXT: cvt.u32.u64 %r8, %rd2;278; SM30-NEXT: and.b32 %r9, %r8, 3;279; SM30-NEXT: shl.b32 %r1, %r9, 3;280; SM30-NEXT: mov.b32 %r10, 255;281; SM30-NEXT: shl.b32 %r11, %r10, %r1;282; SM30-NEXT: not.b32 %r2, %r11;283; SM30-NEXT: cvt.u32.u16 %r12, %rs1;284; SM30-NEXT: shl.b32 %r3, %r12, %r1;285; SM30-NEXT: shl.b32 %r4, %r7, %r1;286; SM30-NEXT: ld.b32 %r13, [%rd1];287; SM30-NEXT: and.b32 %r16, %r13, %r2;288; SM30-NEXT: $L__BB2_1: // %partword.cmpxchg.loop289; SM30-NEXT: // =>This Inner Loop Header: Depth=1290; SM30-NEXT: or.b32 %r14, %r16, %r3;291; SM30-NEXT: or.b32 %r15, %r16, %r4;292; SM30-NEXT: atom.cas.b32 %r5, [%rd1], %r15, %r14;293; SM30-NEXT: setp.eq.b32 %p1, %r5, %r15;294; SM30-NEXT: @%p1 bra $L__BB2_3;295; SM30-NEXT: // %bb.2: // %partword.cmpxchg.failure296; SM30-NEXT: // in Loop: Header=BB2_1 Depth=1297; SM30-NEXT: and.b32 %r6, %r5, %r2;298; SM30-NEXT: setp.ne.b32 %p2, %r16, %r6;299; SM30-NEXT: mov.b32 %r16, %r6;300; SM30-NEXT: @%p2 bra $L__BB2_1;301; SM30-NEXT: $L__BB2_3: // %partword.cmpxchg.end302; SM30-NEXT: st.param.b32 [func_retval0], %r12;303; SM30-NEXT: ret;304;305; SM70-LABEL: release_sys_i8(306; SM70: {307; SM70-NEXT: .reg .pred %p<3>;308; SM70-NEXT: .reg .b16 %rs<2>;309; SM70-NEXT: .reg .b32 %r<17>;310; SM70-NEXT: .reg .b64 %rd<3>;311; SM70-EMPTY:312; SM70-NEXT: // %bb.0:313; SM70-NEXT: ld.param.b8 %rs1, [release_sys_i8_param_2];314; SM70-NEXT: ld.param.b64 %rd2, [release_sys_i8_param_0];315; SM70-NEXT: fence.acq_rel.sys;316; SM70-NEXT: ld.param.b8 %r7, [release_sys_i8_param_1];317; SM70-NEXT: and.b64 %rd1, %rd2, -4;318; SM70-NEXT: cvt.u32.u64 %r8, %rd2;319; SM70-NEXT: and.b32 %r9, %r8, 3;320; SM70-NEXT: shl.b32 %r1, %r9, 3;321; SM70-NEXT: mov.b32 %r10, 255;322; SM70-NEXT: shl.b32 %r11, %r10, %r1;323; SM70-NEXT: not.b32 %r2, %r11;324; SM70-NEXT: cvt.u32.u16 %r12, %rs1;325; SM70-NEXT: shl.b32 %r3, %r12, %r1;326; SM70-NEXT: shl.b32 %r4, %r7, %r1;327; SM70-NEXT: ld.b32 %r13, [%rd1];328; SM70-NEXT: and.b32 %r16, %r13, %r2;329; SM70-NEXT: $L__BB2_1: // %partword.cmpxchg.loop330; SM70-NEXT: // =>This Inner Loop Header: Depth=1331; SM70-NEXT: or.b32 %r14, %r16, %r3;332; SM70-NEXT: or.b32 %r15, %r16, %r4;333; SM70-NEXT: atom.relaxed.sys.cas.b32 %r5, [%rd1], %r15, %r14;334; SM70-NEXT: setp.eq.b32 %p1, %r5, %r15;335; SM70-NEXT: @%p1 bra $L__BB2_3;336; SM70-NEXT: // %bb.2: // %partword.cmpxchg.failure337; SM70-NEXT: // in Loop: Header=BB2_1 Depth=1338; SM70-NEXT: and.b32 %r6, %r5, %r2;339; SM70-NEXT: setp.ne.b32 %p2, %r16, %r6;340; SM70-NEXT: mov.b32 %r16, %r6;341; SM70-NEXT: @%p2 bra $L__BB2_1;342; SM70-NEXT: $L__BB2_3: // %partword.cmpxchg.end343; SM70-NEXT: st.param.b32 [func_retval0], %r12;344; SM70-NEXT: ret;345; SM90-LABEL: release_sys_i8(346; SM90: {347; SM90-NEXT: .reg .pred %p<3>;348; SM90-NEXT: .reg .b16 %rs<2>;349; SM90-NEXT: .reg .b32 %r<21>;350; SM90-NEXT: .reg .b64 %rd<3>;351; SM90-EMPTY:352; SM90-NEXT: // %bb.0:353; SM90-NEXT: ld.param.u8 %rs1, [release_sys_i8_param_2];354; SM90-NEXT: ld.param.u64 %rd2, [release_sys_i8_param_0];355; SM90-NEXT: fence.release.sys;356; SM90-NEXT: and.b64 %rd1, %rd2, -4;357; SM90-NEXT: cvt.u32.u64 %r9, %rd2;358; SM90-NEXT: and.b32 %r10, %r9, 3;359; SM90-NEXT: shl.b32 %r1, %r10, 3;360; SM90-NEXT: mov.b32 %r11, 255;361; SM90-NEXT: shl.b32 %r12, %r11, %r1;362; SM90-NEXT: not.b32 %r2, %r12;363; SM90-NEXT: cvt.u32.u16 %r13, %rs1;364; SM90-NEXT: and.b32 %r14, %r13, 255;365; SM90-NEXT: shl.b32 %r3, %r14, %r1;366; SM90-NEXT: ld.param.u8 %r15, [release_sys_i8_param_1];367; SM90-NEXT: shl.b32 %r4, %r15, %r1;368; SM90-NEXT: ld.u32 %r16, [%rd1];369; SM90-NEXT: and.b32 %r20, %r16, %r2;370; SM90-NEXT: $L__BB2_1: // %partword.cmpxchg.loop371; SM90-NEXT: // =>This Inner Loop Header: Depth=1372; SM90-NEXT: or.b32 %r17, %r20, %r3;373; SM90-NEXT: or.b32 %r18, %r20, %r4;374; SM90-NEXT: atom.cas.b32 %r7, [%rd1], %r18, %r17;375; SM90-NEXT: setp.eq.s32 %p1, %r7, %r18;376; SM90-NEXT: @%p1 bra $L__BB2_3;377; SM90-NEXT: // %bb.2: // %partword.cmpxchg.failure378; SM90-NEXT: // in Loop: Header=BB2_1 Depth=1379; SM90-NEXT: and.b32 %r8, %r7, %r2;380; SM90-NEXT: setp.ne.s32 %p2, %r20, %r8;381; SM90-NEXT: mov.u32 %r20, %r8;382; SM90-NEXT: @%p2 bra $L__BB2_1;383; SM90-NEXT: $L__BB2_3: // %partword.cmpxchg.end384; SM90-NEXT: st.param.b32 [func_retval0], %r13;385; SM90-NEXT: ret;386 %pairold = cmpxchg ptr %addr, i8 %cmp, i8 %new release monotonic387 ret i8 %new388}389 390define i8 @acq_rel_sys_i8(ptr %addr, i8 %cmp, i8 %new) {391; SM30-LABEL: acq_rel_sys_i8(392; SM30: {393; SM30-NEXT: .reg .pred %p<3>;394; SM30-NEXT: .reg .b16 %rs<2>;395; SM30-NEXT: .reg .b32 %r<17>;396; SM30-NEXT: .reg .b64 %rd<3>;397; SM30-EMPTY:398; SM30-NEXT: // %bb.0:399; SM30-NEXT: ld.param.b8 %rs1, [acq_rel_sys_i8_param_2];400; SM30-NEXT: ld.param.b64 %rd2, [acq_rel_sys_i8_param_0];401; SM30-NEXT: membar.sys;402; SM30-NEXT: ld.param.b8 %r7, [acq_rel_sys_i8_param_1];403; SM30-NEXT: and.b64 %rd1, %rd2, -4;404; SM30-NEXT: cvt.u32.u64 %r8, %rd2;405; SM30-NEXT: and.b32 %r9, %r8, 3;406; SM30-NEXT: shl.b32 %r1, %r9, 3;407; SM30-NEXT: mov.b32 %r10, 255;408; SM30-NEXT: shl.b32 %r11, %r10, %r1;409; SM30-NEXT: not.b32 %r2, %r11;410; SM30-NEXT: cvt.u32.u16 %r12, %rs1;411; SM30-NEXT: shl.b32 %r3, %r12, %r1;412; SM30-NEXT: shl.b32 %r4, %r7, %r1;413; SM30-NEXT: ld.b32 %r13, [%rd1];414; SM30-NEXT: and.b32 %r16, %r13, %r2;415; SM30-NEXT: $L__BB3_1: // %partword.cmpxchg.loop416; SM30-NEXT: // =>This Inner Loop Header: Depth=1417; SM30-NEXT: or.b32 %r14, %r16, %r3;418; SM30-NEXT: or.b32 %r15, %r16, %r4;419; SM30-NEXT: atom.cas.b32 %r5, [%rd1], %r15, %r14;420; SM30-NEXT: setp.eq.b32 %p1, %r5, %r15;421; SM30-NEXT: @%p1 bra $L__BB3_3;422; SM30-NEXT: // %bb.2: // %partword.cmpxchg.failure423; SM30-NEXT: // in Loop: Header=BB3_1 Depth=1424; SM30-NEXT: and.b32 %r6, %r5, %r2;425; SM30-NEXT: setp.ne.b32 %p2, %r16, %r6;426; SM30-NEXT: mov.b32 %r16, %r6;427; SM30-NEXT: @%p2 bra $L__BB3_1;428; SM30-NEXT: $L__BB3_3: // %partword.cmpxchg.end429; SM30-NEXT: membar.sys;430; SM30-NEXT: st.param.b32 [func_retval0], %r12;431; SM30-NEXT: ret;432;433; SM70-LABEL: acq_rel_sys_i8(434; SM70: {435; SM70-NEXT: .reg .pred %p<3>;436; SM70-NEXT: .reg .b16 %rs<2>;437; SM70-NEXT: .reg .b32 %r<17>;438; SM70-NEXT: .reg .b64 %rd<3>;439; SM70-EMPTY:440; SM70-NEXT: // %bb.0:441; SM70-NEXT: ld.param.b8 %rs1, [acq_rel_sys_i8_param_2];442; SM70-NEXT: ld.param.b64 %rd2, [acq_rel_sys_i8_param_0];443; SM70-NEXT: fence.acq_rel.sys;444; SM70-NEXT: ld.param.b8 %r7, [acq_rel_sys_i8_param_1];445; SM70-NEXT: and.b64 %rd1, %rd2, -4;446; SM70-NEXT: cvt.u32.u64 %r8, %rd2;447; SM70-NEXT: and.b32 %r9, %r8, 3;448; SM70-NEXT: shl.b32 %r1, %r9, 3;449; SM70-NEXT: mov.b32 %r10, 255;450; SM70-NEXT: shl.b32 %r11, %r10, %r1;451; SM70-NEXT: not.b32 %r2, %r11;452; SM70-NEXT: cvt.u32.u16 %r12, %rs1;453; SM70-NEXT: shl.b32 %r3, %r12, %r1;454; SM70-NEXT: shl.b32 %r4, %r7, %r1;455; SM70-NEXT: ld.b32 %r13, [%rd1];456; SM70-NEXT: and.b32 %r16, %r13, %r2;457; SM70-NEXT: $L__BB3_1: // %partword.cmpxchg.loop458; SM70-NEXT: // =>This Inner Loop Header: Depth=1459; SM70-NEXT: or.b32 %r14, %r16, %r3;460; SM70-NEXT: or.b32 %r15, %r16, %r4;461; SM70-NEXT: atom.relaxed.sys.cas.b32 %r5, [%rd1], %r15, %r14;462; SM70-NEXT: setp.eq.b32 %p1, %r5, %r15;463; SM70-NEXT: @%p1 bra $L__BB3_3;464; SM70-NEXT: // %bb.2: // %partword.cmpxchg.failure465; SM70-NEXT: // in Loop: Header=BB3_1 Depth=1466; SM70-NEXT: and.b32 %r6, %r5, %r2;467; SM70-NEXT: setp.ne.b32 %p2, %r16, %r6;468; SM70-NEXT: mov.b32 %r16, %r6;469; SM70-NEXT: @%p2 bra $L__BB3_1;470; SM70-NEXT: $L__BB3_3: // %partword.cmpxchg.end471; SM70-NEXT: fence.acq_rel.sys;472; SM70-NEXT: st.param.b32 [func_retval0], %r12;473; SM70-NEXT: ret;474; SM90-LABEL: acq_rel_sys_i8(475; SM90: {476; SM90-NEXT: .reg .pred %p<3>;477; SM90-NEXT: .reg .b16 %rs<2>;478; SM90-NEXT: .reg .b32 %r<21>;479; SM90-NEXT: .reg .b64 %rd<3>;480; SM90-EMPTY:481; SM90-NEXT: // %bb.0:482; SM90-NEXT: ld.param.u8 %rs1, [acq_rel_sys_i8_param_2];483; SM90-NEXT: ld.param.u64 %rd2, [acq_rel_sys_i8_param_0];484; SM90-NEXT: fence.release.sys;485; SM90-NEXT: and.b64 %rd1, %rd2, -4;486; SM90-NEXT: cvt.u32.u64 %r9, %rd2;487; SM90-NEXT: and.b32 %r10, %r9, 3;488; SM90-NEXT: shl.b32 %r1, %r10, 3;489; SM90-NEXT: mov.b32 %r11, 255;490; SM90-NEXT: shl.b32 %r12, %r11, %r1;491; SM90-NEXT: not.b32 %r2, %r12;492; SM90-NEXT: cvt.u32.u16 %r13, %rs1;493; SM90-NEXT: and.b32 %r14, %r13, 255;494; SM90-NEXT: shl.b32 %r3, %r14, %r1;495; SM90-NEXT: ld.param.u8 %r15, [acq_rel_sys_i8_param_1];496; SM90-NEXT: shl.b32 %r4, %r15, %r1;497; SM90-NEXT: ld.u32 %r16, [%rd1];498; SM90-NEXT: and.b32 %r20, %r16, %r2;499; SM90-NEXT: $L__BB3_1: // %partword.cmpxchg.loop500; SM90-NEXT: // =>This Inner Loop Header: Depth=1501; SM90-NEXT: or.b32 %r17, %r20, %r3;502; SM90-NEXT: or.b32 %r18, %r20, %r4;503; SM90-NEXT: atom.cas.b32 %r7, [%rd1], %r18, %r17;504; SM90-NEXT: setp.eq.s32 %p1, %r7, %r18;505; SM90-NEXT: @%p1 bra $L__BB3_3;506; SM90-NEXT: // %bb.2: // %partword.cmpxchg.failure507; SM90-NEXT: // in Loop: Header=BB3_1 Depth=1508; SM90-NEXT: and.b32 %r8, %r7, %r2;509; SM90-NEXT: setp.ne.s32 %p2, %r20, %r8;510; SM90-NEXT: mov.u32 %r20, %r8;511; SM90-NEXT: @%p2 bra $L__BB3_1;512; SM90-NEXT: $L__BB3_3: // %partword.cmpxchg.end513; SM90-NEXT: fence.acquire.sys;514; SM90-NEXT: st.param.b32 [func_retval0], %r13;515; SM90-NEXT: ret;516 %pairold = cmpxchg ptr %addr, i8 %cmp, i8 %new acq_rel acquire517 ret i8 %new518}519 520define i8 @seq_cst_sys_i8(ptr %addr, i8 %cmp, i8 %new) {521; SM30-LABEL: seq_cst_sys_i8(522; SM30: {523; SM30-NEXT: .reg .pred %p<3>;524; SM30-NEXT: .reg .b16 %rs<2>;525; SM30-NEXT: .reg .b32 %r<17>;526; SM30-NEXT: .reg .b64 %rd<3>;527; SM30-EMPTY:528; SM30-NEXT: // %bb.0:529; SM30-NEXT: ld.param.b8 %rs1, [seq_cst_sys_i8_param_2];530; SM30-NEXT: ld.param.b64 %rd2, [seq_cst_sys_i8_param_0];531; SM30-NEXT: membar.sys;532; SM30-NEXT: ld.param.b8 %r7, [seq_cst_sys_i8_param_1];533; SM30-NEXT: and.b64 %rd1, %rd2, -4;534; SM30-NEXT: cvt.u32.u64 %r8, %rd2;535; SM30-NEXT: and.b32 %r9, %r8, 3;536; SM30-NEXT: shl.b32 %r1, %r9, 3;537; SM30-NEXT: mov.b32 %r10, 255;538; SM30-NEXT: shl.b32 %r11, %r10, %r1;539; SM30-NEXT: not.b32 %r2, %r11;540; SM30-NEXT: cvt.u32.u16 %r12, %rs1;541; SM30-NEXT: shl.b32 %r3, %r12, %r1;542; SM30-NEXT: shl.b32 %r4, %r7, %r1;543; SM30-NEXT: ld.b32 %r13, [%rd1];544; SM30-NEXT: and.b32 %r16, %r13, %r2;545; SM30-NEXT: $L__BB4_1: // %partword.cmpxchg.loop546; SM30-NEXT: // =>This Inner Loop Header: Depth=1547; SM30-NEXT: or.b32 %r14, %r16, %r3;548; SM30-NEXT: or.b32 %r15, %r16, %r4;549; SM30-NEXT: atom.cas.b32 %r5, [%rd1], %r15, %r14;550; SM30-NEXT: setp.eq.b32 %p1, %r5, %r15;551; SM30-NEXT: @%p1 bra $L__BB4_3;552; SM30-NEXT: // %bb.2: // %partword.cmpxchg.failure553; SM30-NEXT: // in Loop: Header=BB4_1 Depth=1554; SM30-NEXT: and.b32 %r6, %r5, %r2;555; SM30-NEXT: setp.ne.b32 %p2, %r16, %r6;556; SM30-NEXT: mov.b32 %r16, %r6;557; SM30-NEXT: @%p2 bra $L__BB4_1;558; SM30-NEXT: $L__BB4_3: // %partword.cmpxchg.end559; SM30-NEXT: membar.sys;560; SM30-NEXT: st.param.b32 [func_retval0], %r12;561; SM30-NEXT: ret;562;563; SM70-LABEL: seq_cst_sys_i8(564; SM70: {565; SM70-NEXT: .reg .pred %p<3>;566; SM70-NEXT: .reg .b16 %rs<2>;567; SM70-NEXT: .reg .b32 %r<17>;568; SM70-NEXT: .reg .b64 %rd<3>;569; SM70-EMPTY:570; SM70-NEXT: // %bb.0:571; SM70-NEXT: ld.param.b8 %rs1, [seq_cst_sys_i8_param_2];572; SM70-NEXT: ld.param.b64 %rd2, [seq_cst_sys_i8_param_0];573; SM70-NEXT: fence.sc.sys;574; SM70-NEXT: ld.param.b8 %r7, [seq_cst_sys_i8_param_1];575; SM70-NEXT: and.b64 %rd1, %rd2, -4;576; SM70-NEXT: cvt.u32.u64 %r8, %rd2;577; SM70-NEXT: and.b32 %r9, %r8, 3;578; SM70-NEXT: shl.b32 %r1, %r9, 3;579; SM70-NEXT: mov.b32 %r10, 255;580; SM70-NEXT: shl.b32 %r11, %r10, %r1;581; SM70-NEXT: not.b32 %r2, %r11;582; SM70-NEXT: cvt.u32.u16 %r12, %rs1;583; SM70-NEXT: shl.b32 %r3, %r12, %r1;584; SM70-NEXT: shl.b32 %r4, %r7, %r1;585; SM70-NEXT: ld.b32 %r13, [%rd1];586; SM70-NEXT: and.b32 %r16, %r13, %r2;587; SM70-NEXT: $L__BB4_1: // %partword.cmpxchg.loop588; SM70-NEXT: // =>This Inner Loop Header: Depth=1589; SM70-NEXT: or.b32 %r14, %r16, %r3;590; SM70-NEXT: or.b32 %r15, %r16, %r4;591; SM70-NEXT: atom.relaxed.sys.cas.b32 %r5, [%rd1], %r15, %r14;592; SM70-NEXT: setp.eq.b32 %p1, %r5, %r15;593; SM70-NEXT: @%p1 bra $L__BB4_3;594; SM70-NEXT: // %bb.2: // %partword.cmpxchg.failure595; SM70-NEXT: // in Loop: Header=BB4_1 Depth=1596; SM70-NEXT: and.b32 %r6, %r5, %r2;597; SM70-NEXT: setp.ne.b32 %p2, %r16, %r6;598; SM70-NEXT: mov.b32 %r16, %r6;599; SM70-NEXT: @%p2 bra $L__BB4_1;600; SM70-NEXT: $L__BB4_3: // %partword.cmpxchg.end601; SM70-NEXT: fence.acq_rel.sys;602; SM70-NEXT: st.param.b32 [func_retval0], %r12;603; SM70-NEXT: ret;604; SM90-LABEL: seq_cst_sys_i8(605; SM90: {606; SM90-NEXT: .reg .pred %p<3>;607; SM90-NEXT: .reg .b16 %rs<2>;608; SM90-NEXT: .reg .b32 %r<21>;609; SM90-NEXT: .reg .b64 %rd<3>;610; SM90-EMPTY:611; SM90-NEXT: // %bb.0:612; SM90-NEXT: ld.param.u8 %rs1, [seq_cst_sys_i8_param_2];613; SM90-NEXT: ld.param.u64 %rd2, [seq_cst_sys_i8_param_0];614; SM90-NEXT: fence.sc.sys;615; SM90-NEXT: and.b64 %rd1, %rd2, -4;616; SM90-NEXT: cvt.u32.u64 %r9, %rd2;617; SM90-NEXT: and.b32 %r10, %r9, 3;618; SM90-NEXT: shl.b32 %r1, %r10, 3;619; SM90-NEXT: mov.b32 %r11, 255;620; SM90-NEXT: shl.b32 %r12, %r11, %r1;621; SM90-NEXT: not.b32 %r2, %r12;622; SM90-NEXT: cvt.u32.u16 %r13, %rs1;623; SM90-NEXT: and.b32 %r14, %r13, 255;624; SM90-NEXT: shl.b32 %r3, %r14, %r1;625; SM90-NEXT: ld.param.u8 %r15, [seq_cst_sys_i8_param_1];626; SM90-NEXT: shl.b32 %r4, %r15, %r1;627; SM90-NEXT: ld.u32 %r16, [%rd1];628; SM90-NEXT: and.b32 %r20, %r16, %r2;629; SM90-NEXT: $L__BB4_1: // %partword.cmpxchg.loop630; SM90-NEXT: // =>This Inner Loop Header: Depth=1631; SM90-NEXT: or.b32 %r17, %r20, %r3;632; SM90-NEXT: or.b32 %r18, %r20, %r4;633; SM90-NEXT: atom.cas.b32 %r7, [%rd1], %r18, %r17;634; SM90-NEXT: setp.eq.s32 %p1, %r7, %r18;635; SM90-NEXT: @%p1 bra $L__BB4_3;636; SM90-NEXT: // %bb.2: // %partword.cmpxchg.failure637; SM90-NEXT: // in Loop: Header=BB4_1 Depth=1638; SM90-NEXT: and.b32 %r8, %r7, %r2;639; SM90-NEXT: setp.ne.s32 %p2, %r20, %r8;640; SM90-NEXT: mov.u32 %r20, %r8;641; SM90-NEXT: @%p2 bra $L__BB4_1;642; SM90-NEXT: $L__BB4_3: // %partword.cmpxchg.end643; SM90-NEXT: fence.acquire.sys;644; SM90-NEXT: st.param.b32 [func_retval0], %r13;645; SM90-NEXT: ret;646 %pairold = cmpxchg ptr %addr, i8 %cmp, i8 %new seq_cst seq_cst647 ret i8 %new648}649 650; CHECK-LABEL: relaxed_sys_i16651define i16 @relaxed_sys_i16(ptr %addr, i16 %cmp, i16 %new) {652; SM30-LABEL: relaxed_sys_i16(653; SM30: {654; SM30-NEXT: .reg .pred %p<3>;655; SM30-NEXT: .reg .b16 %rs<2>;656; SM30-NEXT: .reg .b32 %r<17>;657; SM30-NEXT: .reg .b64 %rd<3>;658; SM30-EMPTY:659; SM30-NEXT: // %bb.0:660; SM30-NEXT: ld.param.b16 %rs1, [relaxed_sys_i16_param_2];661; SM30-NEXT: ld.param.b64 %rd2, [relaxed_sys_i16_param_0];662; SM30-NEXT: and.b64 %rd1, %rd2, -4;663; SM30-NEXT: ld.param.b16 %r7, [relaxed_sys_i16_param_1];664; SM30-NEXT: cvt.u32.u64 %r8, %rd2;665; SM30-NEXT: and.b32 %r9, %r8, 3;666; SM30-NEXT: shl.b32 %r1, %r9, 3;667; SM30-NEXT: mov.b32 %r10, 65535;668; SM30-NEXT: shl.b32 %r11, %r10, %r1;669; SM30-NEXT: not.b32 %r2, %r11;670; SM30-NEXT: cvt.u32.u16 %r12, %rs1;671; SM30-NEXT: shl.b32 %r3, %r12, %r1;672; SM30-NEXT: shl.b32 %r4, %r7, %r1;673; SM30-NEXT: ld.b32 %r13, [%rd1];674; SM30-NEXT: and.b32 %r16, %r13, %r2;675; SM30-NEXT: $L__BB5_1: // %partword.cmpxchg.loop676; SM30-NEXT: // =>This Inner Loop Header: Depth=1677; SM30-NEXT: or.b32 %r14, %r16, %r3;678; SM30-NEXT: or.b32 %r15, %r16, %r4;679; SM30-NEXT: atom.cas.b32 %r5, [%rd1], %r15, %r14;680; SM30-NEXT: setp.eq.b32 %p1, %r5, %r15;681; SM30-NEXT: @%p1 bra $L__BB5_3;682; SM30-NEXT: // %bb.2: // %partword.cmpxchg.failure683; SM30-NEXT: // in Loop: Header=BB5_1 Depth=1684; SM30-NEXT: and.b32 %r6, %r5, %r2;685; SM30-NEXT: setp.ne.b32 %p2, %r16, %r6;686; SM30-NEXT: mov.b32 %r16, %r6;687; SM30-NEXT: @%p2 bra $L__BB5_1;688; SM30-NEXT: $L__BB5_3: // %partword.cmpxchg.end689; SM30-NEXT: st.param.b32 [func_retval0], %r12;690; SM30-NEXT: ret;691;692; SM70-LABEL: relaxed_sys_i16(693; SM70: {694; SM70-NEXT: .reg .pred %p<3>;695; SM70-NEXT: .reg .b16 %rs<2>;696; SM70-NEXT: .reg .b32 %r<17>;697; SM70-NEXT: .reg .b64 %rd<3>;698; SM70-EMPTY:699; SM70-NEXT: // %bb.0:700; SM70-NEXT: ld.param.b16 %rs1, [relaxed_sys_i16_param_2];701; SM70-NEXT: ld.param.b64 %rd2, [relaxed_sys_i16_param_0];702; SM70-NEXT: and.b64 %rd1, %rd2, -4;703; SM70-NEXT: ld.param.b16 %r7, [relaxed_sys_i16_param_1];704; SM70-NEXT: cvt.u32.u64 %r8, %rd2;705; SM70-NEXT: and.b32 %r9, %r8, 3;706; SM70-NEXT: shl.b32 %r1, %r9, 3;707; SM70-NEXT: mov.b32 %r10, 65535;708; SM70-NEXT: shl.b32 %r11, %r10, %r1;709; SM70-NEXT: not.b32 %r2, %r11;710; SM70-NEXT: cvt.u32.u16 %r12, %rs1;711; SM70-NEXT: shl.b32 %r3, %r12, %r1;712; SM70-NEXT: shl.b32 %r4, %r7, %r1;713; SM70-NEXT: ld.b32 %r13, [%rd1];714; SM70-NEXT: and.b32 %r16, %r13, %r2;715; SM70-NEXT: $L__BB5_1: // %partword.cmpxchg.loop716; SM70-NEXT: // =>This Inner Loop Header: Depth=1717; SM70-NEXT: or.b32 %r14, %r16, %r3;718; SM70-NEXT: or.b32 %r15, %r16, %r4;719; SM70-NEXT: atom.relaxed.sys.cas.b32 %r5, [%rd1], %r15, %r14;720; SM70-NEXT: setp.eq.b32 %p1, %r5, %r15;721; SM70-NEXT: @%p1 bra $L__BB5_3;722; SM70-NEXT: // %bb.2: // %partword.cmpxchg.failure723; SM70-NEXT: // in Loop: Header=BB5_1 Depth=1724; SM70-NEXT: and.b32 %r6, %r5, %r2;725; SM70-NEXT: setp.ne.b32 %p2, %r16, %r6;726; SM70-NEXT: mov.b32 %r16, %r6;727; SM70-NEXT: @%p2 bra $L__BB5_1;728; SM70-NEXT: $L__BB5_3: // %partword.cmpxchg.end729; SM70-NEXT: st.param.b32 [func_retval0], %r12;730; SM70-NEXT: ret;731; SM90-LABEL: relaxed_sys_i16(732; SM90: {733; SM90-NEXT: .reg .pred %p<3>;734; SM90-NEXT: .reg .b16 %rs<2>;735; SM90-NEXT: .reg .b32 %r<20>;736; SM90-NEXT: .reg .b64 %rd<3>;737; SM90-EMPTY:738; SM90-NEXT: // %bb.0:739; SM90-NEXT: ld.param.u16 %rs1, [relaxed_sys_i16_param_2];740; SM90-NEXT: ld.param.u64 %rd2, [relaxed_sys_i16_param_0];741; SM90-NEXT: and.b64 %rd1, %rd2, -4;742; SM90-NEXT: ld.param.u16 %r9, [relaxed_sys_i16_param_1];743; SM90-NEXT: cvt.u32.u64 %r10, %rd2;744; SM90-NEXT: and.b32 %r11, %r10, 3;745; SM90-NEXT: shl.b32 %r1, %r11, 3;746; SM90-NEXT: mov.b32 %r12, 65535;747; SM90-NEXT: shl.b32 %r13, %r12, %r1;748; SM90-NEXT: not.b32 %r2, %r13;749; SM90-NEXT: cvt.u32.u16 %r14, %rs1;750; SM90-NEXT: shl.b32 %r3, %r14, %r1;751; SM90-NEXT: shl.b32 %r4, %r9, %r1;752; SM90-NEXT: ld.u32 %r15, [%rd1];753; SM90-NEXT: and.b32 %r19, %r15, %r2;754; SM90-NEXT: $L__BB5_1: // %partword.cmpxchg.loop755; SM90-NEXT: // =>This Inner Loop Header: Depth=1756; SM90-NEXT: or.b32 %r16, %r19, %r3;757; SM90-NEXT: or.b32 %r17, %r19, %r4;758; SM90-NEXT: atom.cas.b32 %r7, [%rd1], %r17, %r16;759; SM90-NEXT: setp.eq.s32 %p1, %r7, %r17;760; SM90-NEXT: @%p1 bra $L__BB5_3;761; SM90-NEXT: // %bb.2: // %partword.cmpxchg.failure762; SM90-NEXT: // in Loop: Header=BB5_1 Depth=1763; SM90-NEXT: and.b32 %r8, %r7, %r2;764; SM90-NEXT: setp.ne.s32 %p2, %r19, %r8;765; SM90-NEXT: mov.u32 %r19, %r8;766; SM90-NEXT: @%p2 bra $L__BB5_1;767; SM90-NEXT: $L__BB5_3: // %partword.cmpxchg.end768; SM90-NEXT: st.param.b32 [func_retval0], %r14;769; SM90-NEXT: ret;770 %pairold = cmpxchg ptr %addr, i16 %cmp, i16 %new monotonic monotonic771 ret i16 %new772}773 774define i16 @acquire_sys_i16(ptr %addr, i16 %cmp, i16 %new) {775; SM30-LABEL: acquire_sys_i16(776; SM30: {777; SM30-NEXT: .reg .pred %p<3>;778; SM30-NEXT: .reg .b16 %rs<2>;779; SM30-NEXT: .reg .b32 %r<17>;780; SM30-NEXT: .reg .b64 %rd<3>;781; SM30-EMPTY:782; SM30-NEXT: // %bb.0:783; SM30-NEXT: ld.param.b16 %rs1, [acquire_sys_i16_param_2];784; SM30-NEXT: ld.param.b64 %rd2, [acquire_sys_i16_param_0];785; SM30-NEXT: and.b64 %rd1, %rd2, -4;786; SM30-NEXT: ld.param.b16 %r7, [acquire_sys_i16_param_1];787; SM30-NEXT: cvt.u32.u64 %r8, %rd2;788; SM30-NEXT: and.b32 %r9, %r8, 3;789; SM30-NEXT: shl.b32 %r1, %r9, 3;790; SM30-NEXT: mov.b32 %r10, 65535;791; SM30-NEXT: shl.b32 %r11, %r10, %r1;792; SM30-NEXT: not.b32 %r2, %r11;793; SM30-NEXT: cvt.u32.u16 %r12, %rs1;794; SM30-NEXT: shl.b32 %r3, %r12, %r1;795; SM30-NEXT: shl.b32 %r4, %r7, %r1;796; SM30-NEXT: ld.b32 %r13, [%rd1];797; SM30-NEXT: and.b32 %r16, %r13, %r2;798; SM30-NEXT: $L__BB6_1: // %partword.cmpxchg.loop799; SM30-NEXT: // =>This Inner Loop Header: Depth=1800; SM30-NEXT: or.b32 %r14, %r16, %r3;801; SM30-NEXT: or.b32 %r15, %r16, %r4;802; SM30-NEXT: atom.cas.b32 %r5, [%rd1], %r15, %r14;803; SM30-NEXT: setp.eq.b32 %p1, %r5, %r15;804; SM30-NEXT: @%p1 bra $L__BB6_3;805; SM30-NEXT: // %bb.2: // %partword.cmpxchg.failure806; SM30-NEXT: // in Loop: Header=BB6_1 Depth=1807; SM30-NEXT: and.b32 %r6, %r5, %r2;808; SM30-NEXT: setp.ne.b32 %p2, %r16, %r6;809; SM30-NEXT: mov.b32 %r16, %r6;810; SM30-NEXT: @%p2 bra $L__BB6_1;811; SM30-NEXT: $L__BB6_3: // %partword.cmpxchg.end812; SM30-NEXT: membar.sys;813; SM30-NEXT: st.param.b32 [func_retval0], %r12;814; SM30-NEXT: ret;815;816; SM70-LABEL: acquire_sys_i16(817; SM70: {818; SM70-NEXT: .reg .pred %p<3>;819; SM70-NEXT: .reg .b16 %rs<2>;820; SM70-NEXT: .reg .b32 %r<17>;821; SM70-NEXT: .reg .b64 %rd<3>;822; SM70-EMPTY:823; SM70-NEXT: // %bb.0:824; SM70-NEXT: ld.param.b16 %rs1, [acquire_sys_i16_param_2];825; SM70-NEXT: ld.param.b64 %rd2, [acquire_sys_i16_param_0];826; SM70-NEXT: and.b64 %rd1, %rd2, -4;827; SM70-NEXT: ld.param.b16 %r7, [acquire_sys_i16_param_1];828; SM70-NEXT: cvt.u32.u64 %r8, %rd2;829; SM70-NEXT: and.b32 %r9, %r8, 3;830; SM70-NEXT: shl.b32 %r1, %r9, 3;831; SM70-NEXT: mov.b32 %r10, 65535;832; SM70-NEXT: shl.b32 %r11, %r10, %r1;833; SM70-NEXT: not.b32 %r2, %r11;834; SM70-NEXT: cvt.u32.u16 %r12, %rs1;835; SM70-NEXT: shl.b32 %r3, %r12, %r1;836; SM70-NEXT: shl.b32 %r4, %r7, %r1;837; SM70-NEXT: ld.b32 %r13, [%rd1];838; SM70-NEXT: and.b32 %r16, %r13, %r2;839; SM70-NEXT: $L__BB6_1: // %partword.cmpxchg.loop840; SM70-NEXT: // =>This Inner Loop Header: Depth=1841; SM70-NEXT: or.b32 %r14, %r16, %r3;842; SM70-NEXT: or.b32 %r15, %r16, %r4;843; SM70-NEXT: atom.relaxed.sys.cas.b32 %r5, [%rd1], %r15, %r14;844; SM70-NEXT: setp.eq.b32 %p1, %r5, %r15;845; SM70-NEXT: @%p1 bra $L__BB6_3;846; SM70-NEXT: // %bb.2: // %partword.cmpxchg.failure847; SM70-NEXT: // in Loop: Header=BB6_1 Depth=1848; SM70-NEXT: and.b32 %r6, %r5, %r2;849; SM70-NEXT: setp.ne.b32 %p2, %r16, %r6;850; SM70-NEXT: mov.b32 %r16, %r6;851; SM70-NEXT: @%p2 bra $L__BB6_1;852; SM70-NEXT: $L__BB6_3: // %partword.cmpxchg.end853; SM70-NEXT: fence.acq_rel.sys;854; SM70-NEXT: st.param.b32 [func_retval0], %r12;855; SM70-NEXT: ret;856; SM90-LABEL: acquire_sys_i16(857; SM90: {858; SM90-NEXT: .reg .pred %p<3>;859; SM90-NEXT: .reg .b16 %rs<2>;860; SM90-NEXT: .reg .b32 %r<20>;861; SM90-NEXT: .reg .b64 %rd<3>;862; SM90-EMPTY:863; SM90-NEXT: // %bb.0:864; SM90-NEXT: ld.param.u16 %rs1, [acquire_sys_i16_param_2];865; SM90-NEXT: ld.param.u64 %rd2, [acquire_sys_i16_param_0];866; SM90-NEXT: and.b64 %rd1, %rd2, -4;867; SM90-NEXT: ld.param.u16 %r9, [acquire_sys_i16_param_1];868; SM90-NEXT: cvt.u32.u64 %r10, %rd2;869; SM90-NEXT: and.b32 %r11, %r10, 3;870; SM90-NEXT: shl.b32 %r1, %r11, 3;871; SM90-NEXT: mov.b32 %r12, 65535;872; SM90-NEXT: shl.b32 %r13, %r12, %r1;873; SM90-NEXT: not.b32 %r2, %r13;874; SM90-NEXT: cvt.u32.u16 %r14, %rs1;875; SM90-NEXT: shl.b32 %r3, %r14, %r1;876; SM90-NEXT: shl.b32 %r4, %r9, %r1;877; SM90-NEXT: ld.u32 %r15, [%rd1];878; SM90-NEXT: and.b32 %r19, %r15, %r2;879; SM90-NEXT: $L__BB6_1: // %partword.cmpxchg.loop880; SM90-NEXT: // =>This Inner Loop Header: Depth=1881; SM90-NEXT: or.b32 %r16, %r19, %r3;882; SM90-NEXT: or.b32 %r17, %r19, %r4;883; SM90-NEXT: atom.cas.b32 %r7, [%rd1], %r17, %r16;884; SM90-NEXT: setp.eq.s32 %p1, %r7, %r17;885; SM90-NEXT: @%p1 bra $L__BB6_3;886; SM90-NEXT: // %bb.2: // %partword.cmpxchg.failure887; SM90-NEXT: // in Loop: Header=BB6_1 Depth=1888; SM90-NEXT: and.b32 %r8, %r7, %r2;889; SM90-NEXT: setp.ne.s32 %p2, %r19, %r8;890; SM90-NEXT: mov.u32 %r19, %r8;891; SM90-NEXT: @%p2 bra $L__BB6_1;892; SM90-NEXT: $L__BB6_3: // %partword.cmpxchg.end893; SM90-NEXT: fence.acquire.sys;894; SM90-NEXT: st.param.b32 [func_retval0], %r14;895; SM90-NEXT: ret;896 %pairold = cmpxchg ptr %addr, i16 %cmp, i16 %new acquire acquire897 ret i16 %new898}899 900define i16 @release_sys_i16(ptr %addr, i16 %cmp, i16 %new) {901; SM30-LABEL: release_sys_i16(902; SM30: {903; SM30-NEXT: .reg .pred %p<3>;904; SM30-NEXT: .reg .b16 %rs<2>;905; SM30-NEXT: .reg .b32 %r<17>;906; SM30-NEXT: .reg .b64 %rd<3>;907; SM30-EMPTY:908; SM30-NEXT: // %bb.0:909; SM30-NEXT: ld.param.b16 %rs1, [release_sys_i16_param_2];910; SM30-NEXT: ld.param.b64 %rd2, [release_sys_i16_param_0];911; SM30-NEXT: membar.sys;912; SM30-NEXT: ld.param.b16 %r7, [release_sys_i16_param_1];913; SM30-NEXT: and.b64 %rd1, %rd2, -4;914; SM30-NEXT: cvt.u32.u64 %r8, %rd2;915; SM30-NEXT: and.b32 %r9, %r8, 3;916; SM30-NEXT: shl.b32 %r1, %r9, 3;917; SM30-NEXT: mov.b32 %r10, 65535;918; SM30-NEXT: shl.b32 %r11, %r10, %r1;919; SM30-NEXT: not.b32 %r2, %r11;920; SM30-NEXT: cvt.u32.u16 %r12, %rs1;921; SM30-NEXT: shl.b32 %r3, %r12, %r1;922; SM30-NEXT: shl.b32 %r4, %r7, %r1;923; SM30-NEXT: ld.b32 %r13, [%rd1];924; SM30-NEXT: and.b32 %r16, %r13, %r2;925; SM30-NEXT: $L__BB7_1: // %partword.cmpxchg.loop926; SM30-NEXT: // =>This Inner Loop Header: Depth=1927; SM30-NEXT: or.b32 %r14, %r16, %r3;928; SM30-NEXT: or.b32 %r15, %r16, %r4;929; SM30-NEXT: atom.cas.b32 %r5, [%rd1], %r15, %r14;930; SM30-NEXT: setp.eq.b32 %p1, %r5, %r15;931; SM30-NEXT: @%p1 bra $L__BB7_3;932; SM30-NEXT: // %bb.2: // %partword.cmpxchg.failure933; SM30-NEXT: // in Loop: Header=BB7_1 Depth=1934; SM30-NEXT: and.b32 %r6, %r5, %r2;935; SM30-NEXT: setp.ne.b32 %p2, %r16, %r6;936; SM30-NEXT: mov.b32 %r16, %r6;937; SM30-NEXT: @%p2 bra $L__BB7_1;938; SM30-NEXT: $L__BB7_3: // %partword.cmpxchg.end939; SM30-NEXT: st.param.b32 [func_retval0], %r12;940; SM30-NEXT: ret;941;942; SM70-LABEL: release_sys_i16(943; SM70: {944; SM70-NEXT: .reg .pred %p<3>;945; SM70-NEXT: .reg .b16 %rs<2>;946; SM70-NEXT: .reg .b32 %r<17>;947; SM70-NEXT: .reg .b64 %rd<3>;948; SM70-EMPTY:949; SM70-NEXT: // %bb.0:950; SM70-NEXT: ld.param.b16 %rs1, [release_sys_i16_param_2];951; SM70-NEXT: ld.param.b64 %rd2, [release_sys_i16_param_0];952; SM70-NEXT: fence.acq_rel.sys;953; SM70-NEXT: ld.param.b16 %r7, [release_sys_i16_param_1];954; SM70-NEXT: and.b64 %rd1, %rd2, -4;955; SM70-NEXT: cvt.u32.u64 %r8, %rd2;956; SM70-NEXT: and.b32 %r9, %r8, 3;957; SM70-NEXT: shl.b32 %r1, %r9, 3;958; SM70-NEXT: mov.b32 %r10, 65535;959; SM70-NEXT: shl.b32 %r11, %r10, %r1;960; SM70-NEXT: not.b32 %r2, %r11;961; SM70-NEXT: cvt.u32.u16 %r12, %rs1;962; SM70-NEXT: shl.b32 %r3, %r12, %r1;963; SM70-NEXT: shl.b32 %r4, %r7, %r1;964; SM70-NEXT: ld.b32 %r13, [%rd1];965; SM70-NEXT: and.b32 %r16, %r13, %r2;966; SM70-NEXT: $L__BB7_1: // %partword.cmpxchg.loop967; SM70-NEXT: // =>This Inner Loop Header: Depth=1968; SM70-NEXT: or.b32 %r14, %r16, %r3;969; SM70-NEXT: or.b32 %r15, %r16, %r4;970; SM70-NEXT: atom.relaxed.sys.cas.b32 %r5, [%rd1], %r15, %r14;971; SM70-NEXT: setp.eq.b32 %p1, %r5, %r15;972; SM70-NEXT: @%p1 bra $L__BB7_3;973; SM70-NEXT: // %bb.2: // %partword.cmpxchg.failure974; SM70-NEXT: // in Loop: Header=BB7_1 Depth=1975; SM70-NEXT: and.b32 %r6, %r5, %r2;976; SM70-NEXT: setp.ne.b32 %p2, %r16, %r6;977; SM70-NEXT: mov.b32 %r16, %r6;978; SM70-NEXT: @%p2 bra $L__BB7_1;979; SM70-NEXT: $L__BB7_3: // %partword.cmpxchg.end980; SM70-NEXT: st.param.b32 [func_retval0], %r12;981; SM70-NEXT: ret;982; SM90-LABEL: release_sys_i16(983; SM90: {984; SM90-NEXT: .reg .pred %p<3>;985; SM90-NEXT: .reg .b16 %rs<2>;986; SM90-NEXT: .reg .b32 %r<20>;987; SM90-NEXT: .reg .b64 %rd<3>;988; SM90-EMPTY:989; SM90-NEXT: // %bb.0:990; SM90-NEXT: ld.param.u16 %rs1, [release_sys_i16_param_2];991; SM90-NEXT: ld.param.u64 %rd2, [release_sys_i16_param_0];992; SM90-NEXT: fence.release.sys;993; SM90-NEXT: ld.param.u16 %r9, [release_sys_i16_param_1];994; SM90-NEXT: and.b64 %rd1, %rd2, -4;995; SM90-NEXT: cvt.u32.u64 %r10, %rd2;996; SM90-NEXT: and.b32 %r11, %r10, 3;997; SM90-NEXT: shl.b32 %r1, %r11, 3;998; SM90-NEXT: mov.b32 %r12, 65535;999; SM90-NEXT: shl.b32 %r13, %r12, %r1;1000; SM90-NEXT: not.b32 %r2, %r13;1001; SM90-NEXT: cvt.u32.u16 %r14, %rs1;1002; SM90-NEXT: shl.b32 %r3, %r14, %r1;1003; SM90-NEXT: shl.b32 %r4, %r9, %r1;1004; SM90-NEXT: ld.u32 %r15, [%rd1];1005; SM90-NEXT: and.b32 %r19, %r15, %r2;1006; SM90-NEXT: $L__BB7_1: // %partword.cmpxchg.loop1007; SM90-NEXT: // =>This Inner Loop Header: Depth=11008; SM90-NEXT: or.b32 %r16, %r19, %r3;1009; SM90-NEXT: or.b32 %r17, %r19, %r4;1010; SM90-NEXT: atom.cas.b32 %r7, [%rd1], %r17, %r16;1011; SM90-NEXT: setp.eq.s32 %p1, %r7, %r17;1012; SM90-NEXT: @%p1 bra $L__BB7_3;1013; SM90-NEXT: // %bb.2: // %partword.cmpxchg.failure1014; SM90-NEXT: // in Loop: Header=BB7_1 Depth=11015; SM90-NEXT: and.b32 %r8, %r7, %r2;1016; SM90-NEXT: setp.ne.s32 %p2, %r19, %r8;1017; SM90-NEXT: mov.u32 %r19, %r8;1018; SM90-NEXT: @%p2 bra $L__BB7_1;1019; SM90-NEXT: $L__BB7_3: // %partword.cmpxchg.end1020; SM90-NEXT: st.param.b32 [func_retval0], %r14;1021; SM90-NEXT: ret;1022 %pairold = cmpxchg ptr %addr, i16 %cmp, i16 %new release monotonic1023 ret i16 %new1024}1025 1026define i16 @acq_rel_sys_i16(ptr %addr, i16 %cmp, i16 %new) {1027; SM30-LABEL: acq_rel_sys_i16(1028; SM30: {1029; SM30-NEXT: .reg .pred %p<3>;1030; SM30-NEXT: .reg .b16 %rs<2>;1031; SM30-NEXT: .reg .b32 %r<17>;1032; SM30-NEXT: .reg .b64 %rd<3>;1033; SM30-EMPTY:1034; SM30-NEXT: // %bb.0:1035; SM30-NEXT: ld.param.b16 %rs1, [acq_rel_sys_i16_param_2];1036; SM30-NEXT: ld.param.b64 %rd2, [acq_rel_sys_i16_param_0];1037; SM30-NEXT: membar.sys;1038; SM30-NEXT: ld.param.b16 %r7, [acq_rel_sys_i16_param_1];1039; SM30-NEXT: and.b64 %rd1, %rd2, -4;1040; SM30-NEXT: cvt.u32.u64 %r8, %rd2;1041; SM30-NEXT: and.b32 %r9, %r8, 3;1042; SM30-NEXT: shl.b32 %r1, %r9, 3;1043; SM30-NEXT: mov.b32 %r10, 65535;1044; SM30-NEXT: shl.b32 %r11, %r10, %r1;1045; SM30-NEXT: not.b32 %r2, %r11;1046; SM30-NEXT: cvt.u32.u16 %r12, %rs1;1047; SM30-NEXT: shl.b32 %r3, %r12, %r1;1048; SM30-NEXT: shl.b32 %r4, %r7, %r1;1049; SM30-NEXT: ld.b32 %r13, [%rd1];1050; SM30-NEXT: and.b32 %r16, %r13, %r2;1051; SM30-NEXT: $L__BB8_1: // %partword.cmpxchg.loop1052; SM30-NEXT: // =>This Inner Loop Header: Depth=11053; SM30-NEXT: or.b32 %r14, %r16, %r3;1054; SM30-NEXT: or.b32 %r15, %r16, %r4;1055; SM30-NEXT: atom.cas.b32 %r5, [%rd1], %r15, %r14;1056; SM30-NEXT: setp.eq.b32 %p1, %r5, %r15;1057; SM30-NEXT: @%p1 bra $L__BB8_3;1058; SM30-NEXT: // %bb.2: // %partword.cmpxchg.failure1059; SM30-NEXT: // in Loop: Header=BB8_1 Depth=11060; SM30-NEXT: and.b32 %r6, %r5, %r2;1061; SM30-NEXT: setp.ne.b32 %p2, %r16, %r6;1062; SM30-NEXT: mov.b32 %r16, %r6;1063; SM30-NEXT: @%p2 bra $L__BB8_1;1064; SM30-NEXT: $L__BB8_3: // %partword.cmpxchg.end1065; SM30-NEXT: membar.sys;1066; SM30-NEXT: st.param.b32 [func_retval0], %r12;1067; SM30-NEXT: ret;1068;1069; SM70-LABEL: acq_rel_sys_i16(1070; SM70: {1071; SM70-NEXT: .reg .pred %p<3>;1072; SM70-NEXT: .reg .b16 %rs<2>;1073; SM70-NEXT: .reg .b32 %r<17>;1074; SM70-NEXT: .reg .b64 %rd<3>;1075; SM70-EMPTY:1076; SM70-NEXT: // %bb.0:1077; SM70-NEXT: ld.param.b16 %rs1, [acq_rel_sys_i16_param_2];1078; SM70-NEXT: ld.param.b64 %rd2, [acq_rel_sys_i16_param_0];1079; SM70-NEXT: fence.acq_rel.sys;1080; SM70-NEXT: ld.param.b16 %r7, [acq_rel_sys_i16_param_1];1081; SM70-NEXT: and.b64 %rd1, %rd2, -4;1082; SM70-NEXT: cvt.u32.u64 %r8, %rd2;1083; SM70-NEXT: and.b32 %r9, %r8, 3;1084; SM70-NEXT: shl.b32 %r1, %r9, 3;1085; SM70-NEXT: mov.b32 %r10, 65535;1086; SM70-NEXT: shl.b32 %r11, %r10, %r1;1087; SM70-NEXT: not.b32 %r2, %r11;1088; SM70-NEXT: cvt.u32.u16 %r12, %rs1;1089; SM70-NEXT: shl.b32 %r3, %r12, %r1;1090; SM70-NEXT: shl.b32 %r4, %r7, %r1;1091; SM70-NEXT: ld.b32 %r13, [%rd1];1092; SM70-NEXT: and.b32 %r16, %r13, %r2;1093; SM70-NEXT: $L__BB8_1: // %partword.cmpxchg.loop1094; SM70-NEXT: // =>This Inner Loop Header: Depth=11095; SM70-NEXT: or.b32 %r14, %r16, %r3;1096; SM70-NEXT: or.b32 %r15, %r16, %r4;1097; SM70-NEXT: atom.relaxed.sys.cas.b32 %r5, [%rd1], %r15, %r14;1098; SM70-NEXT: setp.eq.b32 %p1, %r5, %r15;1099; SM70-NEXT: @%p1 bra $L__BB8_3;1100; SM70-NEXT: // %bb.2: // %partword.cmpxchg.failure1101; SM70-NEXT: // in Loop: Header=BB8_1 Depth=11102; SM70-NEXT: and.b32 %r6, %r5, %r2;1103; SM70-NEXT: setp.ne.b32 %p2, %r16, %r6;1104; SM70-NEXT: mov.b32 %r16, %r6;1105; SM70-NEXT: @%p2 bra $L__BB8_1;1106; SM70-NEXT: $L__BB8_3: // %partword.cmpxchg.end1107; SM70-NEXT: fence.acq_rel.sys;1108; SM70-NEXT: st.param.b32 [func_retval0], %r12;1109; SM70-NEXT: ret;1110; SM90-LABEL: acq_rel_sys_i16(1111; SM90: {1112; SM90-NEXT: .reg .pred %p<3>;1113; SM90-NEXT: .reg .b16 %rs<2>;1114; SM90-NEXT: .reg .b32 %r<20>;1115; SM90-NEXT: .reg .b64 %rd<3>;1116; SM90-EMPTY:1117; SM90-NEXT: // %bb.0:1118; SM90-NEXT: ld.param.u16 %rs1, [acq_rel_sys_i16_param_2];1119; SM90-NEXT: ld.param.u64 %rd2, [acq_rel_sys_i16_param_0];1120; SM90-NEXT: fence.release.sys;1121; SM90-NEXT: ld.param.u16 %r9, [acq_rel_sys_i16_param_1];1122; SM90-NEXT: and.b64 %rd1, %rd2, -4;1123; SM90-NEXT: cvt.u32.u64 %r10, %rd2;1124; SM90-NEXT: and.b32 %r11, %r10, 3;1125; SM90-NEXT: shl.b32 %r1, %r11, 3;1126; SM90-NEXT: mov.b32 %r12, 65535;1127; SM90-NEXT: shl.b32 %r13, %r12, %r1;1128; SM90-NEXT: not.b32 %r2, %r13;1129; SM90-NEXT: cvt.u32.u16 %r14, %rs1;1130; SM90-NEXT: shl.b32 %r3, %r14, %r1;1131; SM90-NEXT: shl.b32 %r4, %r9, %r1;1132; SM90-NEXT: ld.u32 %r15, [%rd1];1133; SM90-NEXT: and.b32 %r19, %r15, %r2;1134; SM90-NEXT: $L__BB8_1: // %partword.cmpxchg.loop1135; SM90-NEXT: // =>This Inner Loop Header: Depth=11136; SM90-NEXT: or.b32 %r16, %r19, %r3;1137; SM90-NEXT: or.b32 %r17, %r19, %r4;1138; SM90-NEXT: atom.cas.b32 %r7, [%rd1], %r17, %r16;1139; SM90-NEXT: setp.eq.s32 %p1, %r7, %r17;1140; SM90-NEXT: @%p1 bra $L__BB8_3;1141; SM90-NEXT: // %bb.2: // %partword.cmpxchg.failure1142; SM90-NEXT: // in Loop: Header=BB8_1 Depth=11143; SM90-NEXT: and.b32 %r8, %r7, %r2;1144; SM90-NEXT: setp.ne.s32 %p2, %r19, %r8;1145; SM90-NEXT: mov.u32 %r19, %r8;1146; SM90-NEXT: @%p2 bra $L__BB8_1;1147; SM90-NEXT: $L__BB8_3: // %partword.cmpxchg.end1148; SM90-NEXT: fence.acquire.sys;1149; SM90-NEXT: st.param.b32 [func_retval0], %r14;1150; SM90-NEXT: ret;1151 %pairold = cmpxchg ptr %addr, i16 %cmp, i16 %new acq_rel acquire1152 ret i16 %new1153}1154 1155; CHECK-LABEL: seq_cst_sys_i161156define i16 @seq_cst_sys_i16(ptr %addr, i16 %cmp, i16 %new) {1157; SM30-LABEL: seq_cst_sys_i16(1158; SM30: {1159; SM30-NEXT: .reg .pred %p<3>;1160; SM30-NEXT: .reg .b16 %rs<2>;1161; SM30-NEXT: .reg .b32 %r<17>;1162; SM30-NEXT: .reg .b64 %rd<3>;1163; SM30-EMPTY:1164; SM30-NEXT: // %bb.0:1165; SM30-NEXT: ld.param.b16 %rs1, [seq_cst_sys_i16_param_2];1166; SM30-NEXT: ld.param.b64 %rd2, [seq_cst_sys_i16_param_0];1167; SM30-NEXT: membar.sys;1168; SM30-NEXT: ld.param.b16 %r7, [seq_cst_sys_i16_param_1];1169; SM30-NEXT: and.b64 %rd1, %rd2, -4;1170; SM30-NEXT: cvt.u32.u64 %r8, %rd2;1171; SM30-NEXT: and.b32 %r9, %r8, 3;1172; SM30-NEXT: shl.b32 %r1, %r9, 3;1173; SM30-NEXT: mov.b32 %r10, 65535;1174; SM30-NEXT: shl.b32 %r11, %r10, %r1;1175; SM30-NEXT: not.b32 %r2, %r11;1176; SM30-NEXT: cvt.u32.u16 %r12, %rs1;1177; SM30-NEXT: shl.b32 %r3, %r12, %r1;1178; SM30-NEXT: shl.b32 %r4, %r7, %r1;1179; SM30-NEXT: ld.b32 %r13, [%rd1];1180; SM30-NEXT: and.b32 %r16, %r13, %r2;1181; SM30-NEXT: $L__BB9_1: // %partword.cmpxchg.loop1182; SM30-NEXT: // =>This Inner Loop Header: Depth=11183; SM30-NEXT: or.b32 %r14, %r16, %r3;1184; SM30-NEXT: or.b32 %r15, %r16, %r4;1185; SM30-NEXT: atom.cas.b32 %r5, [%rd1], %r15, %r14;1186; SM30-NEXT: setp.eq.b32 %p1, %r5, %r15;1187; SM30-NEXT: @%p1 bra $L__BB9_3;1188; SM30-NEXT: // %bb.2: // %partword.cmpxchg.failure1189; SM30-NEXT: // in Loop: Header=BB9_1 Depth=11190; SM30-NEXT: and.b32 %r6, %r5, %r2;1191; SM30-NEXT: setp.ne.b32 %p2, %r16, %r6;1192; SM30-NEXT: mov.b32 %r16, %r6;1193; SM30-NEXT: @%p2 bra $L__BB9_1;1194; SM30-NEXT: $L__BB9_3: // %partword.cmpxchg.end1195; SM30-NEXT: membar.sys;1196; SM30-NEXT: st.param.b32 [func_retval0], %r12;1197; SM30-NEXT: ret;1198;1199; SM70-LABEL: seq_cst_sys_i16(1200; SM70: {1201; SM70-NEXT: .reg .pred %p<3>;1202; SM70-NEXT: .reg .b16 %rs<2>;1203; SM70-NEXT: .reg .b32 %r<17>;1204; SM70-NEXT: .reg .b64 %rd<3>;1205; SM70-EMPTY:1206; SM70-NEXT: // %bb.0:1207; SM70-NEXT: ld.param.b16 %rs1, [seq_cst_sys_i16_param_2];1208; SM70-NEXT: ld.param.b64 %rd2, [seq_cst_sys_i16_param_0];1209; SM70-NEXT: fence.sc.sys;1210; SM70-NEXT: ld.param.b16 %r7, [seq_cst_sys_i16_param_1];1211; SM70-NEXT: and.b64 %rd1, %rd2, -4;1212; SM70-NEXT: cvt.u32.u64 %r8, %rd2;1213; SM70-NEXT: and.b32 %r9, %r8, 3;1214; SM70-NEXT: shl.b32 %r1, %r9, 3;1215; SM70-NEXT: mov.b32 %r10, 65535;1216; SM70-NEXT: shl.b32 %r11, %r10, %r1;1217; SM70-NEXT: not.b32 %r2, %r11;1218; SM70-NEXT: cvt.u32.u16 %r12, %rs1;1219; SM70-NEXT: shl.b32 %r3, %r12, %r1;1220; SM70-NEXT: shl.b32 %r4, %r7, %r1;1221; SM70-NEXT: ld.b32 %r13, [%rd1];1222; SM70-NEXT: and.b32 %r16, %r13, %r2;1223; SM70-NEXT: $L__BB9_1: // %partword.cmpxchg.loop1224; SM70-NEXT: // =>This Inner Loop Header: Depth=11225; SM70-NEXT: or.b32 %r14, %r16, %r3;1226; SM70-NEXT: or.b32 %r15, %r16, %r4;1227; SM70-NEXT: atom.relaxed.sys.cas.b32 %r5, [%rd1], %r15, %r14;1228; SM70-NEXT: setp.eq.b32 %p1, %r5, %r15;1229; SM70-NEXT: @%p1 bra $L__BB9_3;1230; SM70-NEXT: // %bb.2: // %partword.cmpxchg.failure1231; SM70-NEXT: // in Loop: Header=BB9_1 Depth=11232; SM70-NEXT: and.b32 %r6, %r5, %r2;1233; SM70-NEXT: setp.ne.b32 %p2, %r16, %r6;1234; SM70-NEXT: mov.b32 %r16, %r6;1235; SM70-NEXT: @%p2 bra $L__BB9_1;1236; SM70-NEXT: $L__BB9_3: // %partword.cmpxchg.end1237; SM70-NEXT: fence.acq_rel.sys;1238; SM70-NEXT: st.param.b32 [func_retval0], %r12;1239; SM70-NEXT: ret;1240; SM90-LABEL: seq_cst_sys_i16(1241; SM90: {1242; SM90-NEXT: .reg .pred %p<3>;1243; SM90-NEXT: .reg .b16 %rs<2>;1244; SM90-NEXT: .reg .b32 %r<20>;1245; SM90-NEXT: .reg .b64 %rd<3>;1246; SM90-EMPTY:1247; SM90-NEXT: // %bb.0:1248; SM90-NEXT: ld.param.u16 %rs1, [seq_cst_sys_i16_param_2];1249; SM90-NEXT: ld.param.u64 %rd2, [seq_cst_sys_i16_param_0];1250; SM90-NEXT: fence.sc.sys;1251; SM90-NEXT: ld.param.u16 %r9, [seq_cst_sys_i16_param_1];1252; SM90-NEXT: and.b64 %rd1, %rd2, -4;1253; SM90-NEXT: cvt.u32.u64 %r10, %rd2;1254; SM90-NEXT: and.b32 %r11, %r10, 3;1255; SM90-NEXT: shl.b32 %r1, %r11, 3;1256; SM90-NEXT: mov.b32 %r12, 65535;1257; SM90-NEXT: shl.b32 %r13, %r12, %r1;1258; SM90-NEXT: not.b32 %r2, %r13;1259; SM90-NEXT: cvt.u32.u16 %r14, %rs1;1260; SM90-NEXT: shl.b32 %r3, %r14, %r1;1261; SM90-NEXT: shl.b32 %r4, %r9, %r1;1262; SM90-NEXT: ld.u32 %r15, [%rd1];1263; SM90-NEXT: and.b32 %r19, %r15, %r2;1264; SM90-NEXT: $L__BB9_1: // %partword.cmpxchg.loop1265; SM90-NEXT: // =>This Inner Loop Header: Depth=11266; SM90-NEXT: or.b32 %r16, %r19, %r3;1267; SM90-NEXT: or.b32 %r17, %r19, %r4;1268; SM90-NEXT: atom.cas.b32 %r7, [%rd1], %r17, %r16;1269; SM90-NEXT: setp.eq.s32 %p1, %r7, %r17;1270; SM90-NEXT: @%p1 bra $L__BB9_3;1271; SM90-NEXT: // %bb.2: // %partword.cmpxchg.failure1272; SM90-NEXT: // in Loop: Header=BB9_1 Depth=11273; SM90-NEXT: and.b32 %r8, %r7, %r2;1274; SM90-NEXT: setp.ne.s32 %p2, %r19, %r8;1275; SM90-NEXT: mov.u32 %r19, %r8;1276; SM90-NEXT: @%p2 bra $L__BB9_1;1277; SM90-NEXT: $L__BB9_3: // %partword.cmpxchg.end1278; SM90-NEXT: fence.acquire.sys;1279; SM90-NEXT: st.param.b32 [func_retval0], %r14;1280; SM90-NEXT: ret;1281 %pairold = cmpxchg ptr %addr, i16 %cmp, i16 %new seq_cst seq_cst1282 ret i16 %new1283}1284 1285; CHECK-LABEL: relaxed_sys_i321286define i32 @relaxed_sys_i32(ptr %addr, i32 %cmp, i32 %new) {1287; SM30-LABEL: relaxed_sys_i32(1288; SM30: {1289; SM30-NEXT: .reg .b32 %r<4>;1290; SM30-NEXT: .reg .b64 %rd<2>;1291; SM30-EMPTY:1292; SM30-NEXT: // %bb.0:1293; SM30-NEXT: ld.param.b64 %rd1, [relaxed_sys_i32_param_0];1294; SM30-NEXT: ld.param.b32 %r1, [relaxed_sys_i32_param_1];1295; SM30-NEXT: ld.param.b32 %r2, [relaxed_sys_i32_param_2];1296; SM30-NEXT: atom.cas.b32 %r3, [%rd1], %r1, %r2;1297; SM30-NEXT: st.param.b32 [func_retval0], %r2;1298; SM30-NEXT: ret;1299;1300; SM70-LABEL: relaxed_sys_i32(1301; SM70: {1302; SM70-NEXT: .reg .b32 %r<4>;1303; SM70-NEXT: .reg .b64 %rd<2>;1304; SM70-EMPTY:1305; SM70-NEXT: // %bb.0:1306; SM70-NEXT: ld.param.b64 %rd1, [relaxed_sys_i32_param_0];1307; SM70-NEXT: ld.param.b32 %r1, [relaxed_sys_i32_param_1];1308; SM70-NEXT: ld.param.b32 %r2, [relaxed_sys_i32_param_2];1309; SM70-NEXT: atom.relaxed.sys.cas.b32 %r3, [%rd1], %r1, %r2;1310; SM70-NEXT: st.param.b32 [func_retval0], %r2;1311; SM70-NEXT: ret;1312; SM90-LABEL: relaxed_sys_i32(1313; SM90: {1314; SM90-NEXT: .reg .b32 %r<4>;1315; SM90-NEXT: .reg .b64 %rd<2>;1316; SM90-EMPTY:1317; SM90-NEXT: // %bb.0:1318; SM90-NEXT: ld.param.u64 %rd1, [relaxed_sys_i32_param_0];1319; SM90-NEXT: ld.param.u32 %r1, [relaxed_sys_i32_param_1];1320; SM90-NEXT: ld.param.u32 %r2, [relaxed_sys_i32_param_2];1321; SM90-NEXT: atom.cas.b32 %r3, [%rd1], %r1, %r2;1322; SM90-NEXT: st.param.b32 [func_retval0], %r2;1323; SM90-NEXT: ret;1324 %pairold = cmpxchg ptr %addr, i32 %cmp, i32 %new monotonic monotonic1325 ret i32 %new1326}1327 1328define i32 @acq_rel_sys_i32(ptr %addr, i32 %cmp, i32 %new) {1329; SM30-LABEL: acq_rel_sys_i32(1330; SM30: {1331; SM30-NEXT: .reg .b32 %r<4>;1332; SM30-NEXT: .reg .b64 %rd<2>;1333; SM30-EMPTY:1334; SM30-NEXT: // %bb.0:1335; SM30-NEXT: ld.param.b64 %rd1, [acq_rel_sys_i32_param_0];1336; SM30-NEXT: ld.param.b32 %r1, [acq_rel_sys_i32_param_1];1337; SM30-NEXT: ld.param.b32 %r2, [acq_rel_sys_i32_param_2];1338; SM30-NEXT: atom.cas.b32 %r3, [%rd1], %r1, %r2;1339; SM30-NEXT: st.param.b32 [func_retval0], %r2;1340; SM30-NEXT: ret;1341;1342; SM70-LABEL: acq_rel_sys_i32(1343; SM70: {1344; SM70-NEXT: .reg .b32 %r<4>;1345; SM70-NEXT: .reg .b64 %rd<2>;1346; SM70-EMPTY:1347; SM70-NEXT: // %bb.0:1348; SM70-NEXT: ld.param.b64 %rd1, [acq_rel_sys_i32_param_0];1349; SM70-NEXT: ld.param.b32 %r1, [acq_rel_sys_i32_param_1];1350; SM70-NEXT: ld.param.b32 %r2, [acq_rel_sys_i32_param_2];1351; SM70-NEXT: atom.acq_rel.sys.cas.b32 %r3, [%rd1], %r1, %r2;1352; SM70-NEXT: st.param.b32 [func_retval0], %r2;1353; SM70-NEXT: ret;1354; SM90-LABEL: acq_rel_sys_i32(1355; SM90: {1356; SM90-NEXT: .reg .b32 %r<4>;1357; SM90-NEXT: .reg .b64 %rd<2>;1358; SM90-EMPTY:1359; SM90-NEXT: // %bb.0:1360; SM90-NEXT: ld.param.u64 %rd1, [acq_rel_sys_i32_param_0];1361; SM90-NEXT: ld.param.u32 %r1, [acq_rel_sys_i32_param_1];1362; SM90-NEXT: ld.param.u32 %r2, [acq_rel_sys_i32_param_2];1363; SM90-NEXT: atom.acq_rel.cas.b32 %r3, [%rd1], %r1, %r2;1364; SM90-NEXT: st.param.b32 [func_retval0], %r2;1365; SM90-NEXT: ret;1366 %pairold = cmpxchg ptr %addr, i32 %cmp, i32 %new acq_rel acquire1367 ret i32 %new1368}1369 1370define i32 @acquire_sys_i32(ptr %addr, i32 %cmp, i32 %new) {1371; SM30-LABEL: acquire_sys_i32(1372; SM30: {1373; SM30-NEXT: .reg .b32 %r<4>;1374; SM30-NEXT: .reg .b64 %rd<2>;1375; SM30-EMPTY:1376; SM30-NEXT: // %bb.0:1377; SM30-NEXT: ld.param.b64 %rd1, [acquire_sys_i32_param_0];1378; SM30-NEXT: ld.param.b32 %r1, [acquire_sys_i32_param_1];1379; SM30-NEXT: ld.param.b32 %r2, [acquire_sys_i32_param_2];1380; SM30-NEXT: atom.cas.b32 %r3, [%rd1], %r1, %r2;1381; SM30-NEXT: st.param.b32 [func_retval0], %r2;1382; SM30-NEXT: ret;1383;1384; SM70-LABEL: acquire_sys_i32(1385; SM70: {1386; SM70-NEXT: .reg .b32 %r<4>;1387; SM70-NEXT: .reg .b64 %rd<2>;1388; SM70-EMPTY:1389; SM70-NEXT: // %bb.0:1390; SM70-NEXT: ld.param.b64 %rd1, [acquire_sys_i32_param_0];1391; SM70-NEXT: ld.param.b32 %r1, [acquire_sys_i32_param_1];1392; SM70-NEXT: ld.param.b32 %r2, [acquire_sys_i32_param_2];1393; SM70-NEXT: atom.acquire.sys.cas.b32 %r3, [%rd1], %r1, %r2;1394; SM70-NEXT: st.param.b32 [func_retval0], %r2;1395; SM70-NEXT: ret;1396; SM90-LABEL: acquire_sys_i32(1397; SM90: {1398; SM90-NEXT: .reg .b32 %r<4>;1399; SM90-NEXT: .reg .b64 %rd<2>;1400; SM90-EMPTY:1401; SM90-NEXT: // %bb.0:1402; SM90-NEXT: ld.param.u64 %rd1, [acquire_sys_i32_param_0];1403; SM90-NEXT: ld.param.u32 %r1, [acquire_sys_i32_param_1];1404; SM90-NEXT: ld.param.u32 %r2, [acquire_sys_i32_param_2];1405; SM90-NEXT: atom.acquire.cas.b32 %r3, [%rd1], %r1, %r2;1406; SM90-NEXT: st.param.b32 [func_retval0], %r2;1407; SM90-NEXT: ret;1408 %pairold = cmpxchg ptr %addr, i32 %cmp, i32 %new acquire acquire1409 ret i32 %new1410}1411 1412define i32 @release_sys_i32(ptr %addr, i32 %cmp, i32 %new) {1413; SM30-LABEL: release_sys_i32(1414; SM30: {1415; SM30-NEXT: .reg .b32 %r<4>;1416; SM30-NEXT: .reg .b64 %rd<2>;1417; SM30-EMPTY:1418; SM30-NEXT: // %bb.0:1419; SM30-NEXT: ld.param.b64 %rd1, [release_sys_i32_param_0];1420; SM30-NEXT: ld.param.b32 %r1, [release_sys_i32_param_1];1421; SM30-NEXT: ld.param.b32 %r2, [release_sys_i32_param_2];1422; SM30-NEXT: atom.cas.b32 %r3, [%rd1], %r1, %r2;1423; SM30-NEXT: st.param.b32 [func_retval0], %r2;1424; SM30-NEXT: ret;1425;1426; SM70-LABEL: release_sys_i32(1427; SM70: {1428; SM70-NEXT: .reg .b32 %r<4>;1429; SM70-NEXT: .reg .b64 %rd<2>;1430; SM70-EMPTY:1431; SM70-NEXT: // %bb.0:1432; SM70-NEXT: ld.param.b64 %rd1, [release_sys_i32_param_0];1433; SM70-NEXT: ld.param.b32 %r1, [release_sys_i32_param_1];1434; SM70-NEXT: ld.param.b32 %r2, [release_sys_i32_param_2];1435; SM70-NEXT: atom.release.sys.cas.b32 %r3, [%rd1], %r1, %r2;1436; SM70-NEXT: st.param.b32 [func_retval0], %r2;1437; SM70-NEXT: ret;1438; SM90-LABEL: release_sys_i32(1439; SM90: {1440; SM90-NEXT: .reg .b32 %r<4>;1441; SM90-NEXT: .reg .b64 %rd<2>;1442; SM90-EMPTY:1443; SM90-NEXT: // %bb.0:1444; SM90-NEXT: ld.param.u64 %rd1, [release_sys_i32_param_0];1445; SM90-NEXT: ld.param.u32 %r1, [release_sys_i32_param_1];1446; SM90-NEXT: ld.param.u32 %r2, [release_sys_i32_param_2];1447; SM90-NEXT: atom.release.cas.b32 %r3, [%rd1], %r1, %r2;1448; SM90-NEXT: st.param.b32 [func_retval0], %r2;1449; SM90-NEXT: ret;1450 %pairold = cmpxchg ptr %addr, i32 %cmp, i32 %new release monotonic1451 ret i32 %new1452}1453 1454define i32 @seq_cst_sys_i32(ptr %addr, i32 %cmp, i32 %new) {1455; SM30-LABEL: seq_cst_sys_i32(1456; SM30: {1457; SM30-NEXT: .reg .b32 %r<4>;1458; SM30-NEXT: .reg .b64 %rd<2>;1459; SM30-EMPTY:1460; SM30-NEXT: // %bb.0:1461; SM30-NEXT: ld.param.b64 %rd1, [seq_cst_sys_i32_param_0];1462; SM30-NEXT: membar.sys;1463; SM30-NEXT: ld.param.b32 %r1, [seq_cst_sys_i32_param_1];1464; SM30-NEXT: ld.param.b32 %r2, [seq_cst_sys_i32_param_2];1465; SM30-NEXT: atom.cas.b32 %r3, [%rd1], %r1, %r2;1466; SM30-NEXT: st.param.b32 [func_retval0], %r2;1467; SM30-NEXT: ret;1468;1469; SM70-LABEL: seq_cst_sys_i32(1470; SM70: {1471; SM70-NEXT: .reg .b32 %r<4>;1472; SM70-NEXT: .reg .b64 %rd<2>;1473; SM70-EMPTY:1474; SM70-NEXT: // %bb.0:1475; SM70-NEXT: ld.param.b64 %rd1, [seq_cst_sys_i32_param_0];1476; SM70-NEXT: fence.sc.sys;1477; SM70-NEXT: ld.param.b32 %r1, [seq_cst_sys_i32_param_1];1478; SM70-NEXT: ld.param.b32 %r2, [seq_cst_sys_i32_param_2];1479; SM70-NEXT: atom.acquire.sys.cas.b32 %r3, [%rd1], %r1, %r2;1480; SM70-NEXT: st.param.b32 [func_retval0], %r2;1481; SM70-NEXT: ret;1482; SM90-LABEL: seq_cst_sys_i32(1483; SM90: {1484; SM90-NEXT: .reg .b32 %r<4>;1485; SM90-NEXT: .reg .b64 %rd<2>;1486; SM90-EMPTY:1487; SM90-NEXT: // %bb.0:1488; SM90-NEXT: ld.param.u64 %rd1, [seq_cst_sys_i32_param_0];1489; SM90-NEXT: fence.sc.sys;1490; SM90-NEXT: ld.param.u32 %r1, [seq_cst_sys_i32_param_1];1491; SM90-NEXT: ld.param.u32 %r2, [seq_cst_sys_i32_param_2];1492; SM90-NEXT: atom.cas.b32 %r3, [%rd1], %r1, %r2;1493; SM90-NEXT: fence.acquire.sys;1494; SM90-NEXT: st.param.b32 [func_retval0], %r2;1495; SM90-NEXT: ret;1496 %pairold = cmpxchg ptr %addr, i32 %cmp, i32 %new seq_cst seq_cst1497 ret i32 %new1498}1499 1500; CHECK-LABEL: relaxed_sys_i641501define i64 @relaxed_sys_i64(ptr %addr, i64 %cmp, i64 %new) {1502; SM30-LABEL: relaxed_sys_i64(1503; SM30: {1504; SM30-NEXT: .reg .b64 %rd<5>;1505; SM30-EMPTY:1506; SM30-NEXT: // %bb.0:1507; SM30-NEXT: ld.param.b64 %rd1, [relaxed_sys_i64_param_0];1508; SM30-NEXT: ld.param.b64 %rd2, [relaxed_sys_i64_param_1];1509; SM30-NEXT: ld.param.b64 %rd3, [relaxed_sys_i64_param_2];1510; SM30-NEXT: atom.cas.b64 %rd4, [%rd1], %rd2, %rd3;1511; SM30-NEXT: st.param.b64 [func_retval0], %rd3;1512; SM30-NEXT: ret;1513;1514; SM70-LABEL: relaxed_sys_i64(1515; SM70: {1516; SM70-NEXT: .reg .b64 %rd<5>;1517; SM70-EMPTY:1518; SM70-NEXT: // %bb.0:1519; SM70-NEXT: ld.param.b64 %rd1, [relaxed_sys_i64_param_0];1520; SM70-NEXT: ld.param.b64 %rd2, [relaxed_sys_i64_param_1];1521; SM70-NEXT: ld.param.b64 %rd3, [relaxed_sys_i64_param_2];1522; SM70-NEXT: atom.relaxed.sys.cas.b64 %rd4, [%rd1], %rd2, %rd3;1523; SM70-NEXT: st.param.b64 [func_retval0], %rd3;1524; SM70-NEXT: ret;1525; SM90-LABEL: relaxed_sys_i64(1526; SM90: {1527; SM90-NEXT: .reg .b64 %rd<5>;1528; SM90-EMPTY:1529; SM90-NEXT: // %bb.0:1530; SM90-NEXT: ld.param.u64 %rd1, [relaxed_sys_i64_param_0];1531; SM90-NEXT: ld.param.u64 %rd2, [relaxed_sys_i64_param_1];1532; SM90-NEXT: ld.param.u64 %rd3, [relaxed_sys_i64_param_2];1533; SM90-NEXT: atom.cas.b64 %rd4, [%rd1], %rd2, %rd3;1534; SM90-NEXT: st.param.b64 [func_retval0], %rd3;1535; SM90-NEXT: ret;1536 %pairold = cmpxchg ptr %addr, i64 %cmp, i64 %new monotonic monotonic1537 ret i64 %new1538}1539 1540define i64 @acquire_sys_i64(ptr %addr, i64 %cmp, i64 %new) {1541; SM30-LABEL: acquire_sys_i64(1542; SM30: {1543; SM30-NEXT: .reg .b64 %rd<5>;1544; SM30-EMPTY:1545; SM30-NEXT: // %bb.0:1546; SM30-NEXT: ld.param.b64 %rd1, [acquire_sys_i64_param_0];1547; SM30-NEXT: ld.param.b64 %rd2, [acquire_sys_i64_param_1];1548; SM30-NEXT: ld.param.b64 %rd3, [acquire_sys_i64_param_2];1549; SM30-NEXT: atom.cas.b64 %rd4, [%rd1], %rd2, %rd3;1550; SM30-NEXT: st.param.b64 [func_retval0], %rd3;1551; SM30-NEXT: ret;1552;1553; SM70-LABEL: acquire_sys_i64(1554; SM70: {1555; SM70-NEXT: .reg .b64 %rd<5>;1556; SM70-EMPTY:1557; SM70-NEXT: // %bb.0:1558; SM70-NEXT: ld.param.b64 %rd1, [acquire_sys_i64_param_0];1559; SM70-NEXT: ld.param.b64 %rd2, [acquire_sys_i64_param_1];1560; SM70-NEXT: ld.param.b64 %rd3, [acquire_sys_i64_param_2];1561; SM70-NEXT: atom.acquire.sys.cas.b64 %rd4, [%rd1], %rd2, %rd3;1562; SM70-NEXT: st.param.b64 [func_retval0], %rd3;1563; SM70-NEXT: ret;1564; SM90-LABEL: acquire_sys_i64(1565; SM90: {1566; SM90-NEXT: .reg .b64 %rd<5>;1567; SM90-EMPTY:1568; SM90-NEXT: // %bb.0:1569; SM90-NEXT: ld.param.u64 %rd1, [acquire_sys_i64_param_0];1570; SM90-NEXT: ld.param.u64 %rd2, [acquire_sys_i64_param_1];1571; SM90-NEXT: ld.param.u64 %rd3, [acquire_sys_i64_param_2];1572; SM90-NEXT: atom.acquire.cas.b64 %rd4, [%rd1], %rd2, %rd3;1573; SM90-NEXT: st.param.b64 [func_retval0], %rd3;1574; SM90-NEXT: ret;1575 %pairold = cmpxchg ptr %addr, i64 %cmp, i64 %new acquire acquire1576 ret i64 %new1577}1578 1579define i64 @acq_rel_sys_i64(ptr %addr, i64 %cmp, i64 %new) {1580; SM30-LABEL: acq_rel_sys_i64(1581; SM30: {1582; SM30-NEXT: .reg .b64 %rd<5>;1583; SM30-EMPTY:1584; SM30-NEXT: // %bb.0:1585; SM30-NEXT: ld.param.b64 %rd1, [acq_rel_sys_i64_param_0];1586; SM30-NEXT: ld.param.b64 %rd2, [acq_rel_sys_i64_param_1];1587; SM30-NEXT: ld.param.b64 %rd3, [acq_rel_sys_i64_param_2];1588; SM30-NEXT: atom.cas.b64 %rd4, [%rd1], %rd2, %rd3;1589; SM30-NEXT: st.param.b64 [func_retval0], %rd3;1590; SM30-NEXT: ret;1591;1592; SM70-LABEL: acq_rel_sys_i64(1593; SM70: {1594; SM70-NEXT: .reg .b64 %rd<5>;1595; SM70-EMPTY:1596; SM70-NEXT: // %bb.0:1597; SM70-NEXT: ld.param.b64 %rd1, [acq_rel_sys_i64_param_0];1598; SM70-NEXT: ld.param.b64 %rd2, [acq_rel_sys_i64_param_1];1599; SM70-NEXT: ld.param.b64 %rd3, [acq_rel_sys_i64_param_2];1600; SM70-NEXT: atom.acq_rel.sys.cas.b64 %rd4, [%rd1], %rd2, %rd3;1601; SM70-NEXT: st.param.b64 [func_retval0], %rd3;1602; SM70-NEXT: ret;1603; SM90-LABEL: acq_rel_sys_i64(1604; SM90: {1605; SM90-NEXT: .reg .b64 %rd<5>;1606; SM90-EMPTY:1607; SM90-NEXT: // %bb.0:1608; SM90-NEXT: ld.param.u64 %rd1, [acq_rel_sys_i64_param_0];1609; SM90-NEXT: ld.param.u64 %rd2, [acq_rel_sys_i64_param_1];1610; SM90-NEXT: ld.param.u64 %rd3, [acq_rel_sys_i64_param_2];1611; SM90-NEXT: atom.acq_rel.cas.b64 %rd4, [%rd1], %rd2, %rd3;1612; SM90-NEXT: st.param.b64 [func_retval0], %rd3;1613; SM90-NEXT: ret;1614 %pairold = cmpxchg ptr %addr, i64 %cmp, i64 %new acq_rel acquire1615 ret i64 %new1616}1617 1618define i64 @release_sys_i64(ptr %addr, i64 %cmp, i64 %new) {1619; SM30-LABEL: release_sys_i64(1620; SM30: {1621; SM30-NEXT: .reg .b64 %rd<5>;1622; SM30-EMPTY:1623; SM30-NEXT: // %bb.0:1624; SM30-NEXT: ld.param.b64 %rd1, [release_sys_i64_param_0];1625; SM30-NEXT: ld.param.b64 %rd2, [release_sys_i64_param_1];1626; SM30-NEXT: ld.param.b64 %rd3, [release_sys_i64_param_2];1627; SM30-NEXT: atom.cas.b64 %rd4, [%rd1], %rd2, %rd3;1628; SM30-NEXT: st.param.b64 [func_retval0], %rd3;1629; SM30-NEXT: ret;1630;1631; SM70-LABEL: release_sys_i64(1632; SM70: {1633; SM70-NEXT: .reg .b64 %rd<5>;1634; SM70-EMPTY:1635; SM70-NEXT: // %bb.0:1636; SM70-NEXT: ld.param.b64 %rd1, [release_sys_i64_param_0];1637; SM70-NEXT: ld.param.b64 %rd2, [release_sys_i64_param_1];1638; SM70-NEXT: ld.param.b64 %rd3, [release_sys_i64_param_2];1639; SM70-NEXT: atom.release.sys.cas.b64 %rd4, [%rd1], %rd2, %rd3;1640; SM70-NEXT: st.param.b64 [func_retval0], %rd3;1641; SM70-NEXT: ret;1642; SM90-LABEL: release_sys_i64(1643; SM90: {1644; SM90-NEXT: .reg .b64 %rd<5>;1645; SM90-EMPTY:1646; SM90-NEXT: // %bb.0:1647; SM90-NEXT: ld.param.u64 %rd1, [release_sys_i64_param_0];1648; SM90-NEXT: ld.param.u64 %rd2, [release_sys_i64_param_1];1649; SM90-NEXT: ld.param.u64 %rd3, [release_sys_i64_param_2];1650; SM90-NEXT: atom.release.cas.b64 %rd4, [%rd1], %rd2, %rd3;1651; SM90-NEXT: st.param.b64 [func_retval0], %rd3;1652; SM90-NEXT: ret;1653 %pairold = cmpxchg ptr %addr, i64 %cmp, i64 %new release monotonic1654 ret i64 %new1655}1656 1657define i64 @seq_cst_sys_i64(ptr %addr, i64 %cmp, i64 %new) {1658; SM30-LABEL: seq_cst_sys_i64(1659; SM30: {1660; SM30-NEXT: .reg .b64 %rd<5>;1661; SM30-EMPTY:1662; SM30-NEXT: // %bb.0:1663; SM30-NEXT: ld.param.b64 %rd1, [seq_cst_sys_i64_param_0];1664; SM30-NEXT: membar.sys;1665; SM30-NEXT: ld.param.b64 %rd2, [seq_cst_sys_i64_param_1];1666; SM30-NEXT: ld.param.b64 %rd3, [seq_cst_sys_i64_param_2];1667; SM30-NEXT: atom.cas.b64 %rd4, [%rd1], %rd2, %rd3;1668; SM30-NEXT: st.param.b64 [func_retval0], %rd3;1669; SM30-NEXT: ret;1670;1671; SM70-LABEL: seq_cst_sys_i64(1672; SM70: {1673; SM70-NEXT: .reg .b64 %rd<5>;1674; SM70-EMPTY:1675; SM70-NEXT: // %bb.0:1676; SM70-NEXT: ld.param.b64 %rd1, [seq_cst_sys_i64_param_0];1677; SM70-NEXT: fence.sc.sys;1678; SM70-NEXT: ld.param.b64 %rd2, [seq_cst_sys_i64_param_1];1679; SM70-NEXT: ld.param.b64 %rd3, [seq_cst_sys_i64_param_2];1680; SM70-NEXT: atom.acquire.sys.cas.b64 %rd4, [%rd1], %rd2, %rd3;1681; SM70-NEXT: st.param.b64 [func_retval0], %rd3;1682; SM70-NEXT: ret;1683; SM90-LABEL: seq_cst_sys_i64(1684; SM90: {1685; SM90-NEXT: .reg .b64 %rd<5>;1686; SM90-EMPTY:1687; SM90-NEXT: // %bb.0:1688; SM90-NEXT: ld.param.u64 %rd1, [seq_cst_sys_i64_param_0];1689; SM90-NEXT: fence.sc.sys;1690; SM90-NEXT: ld.param.u64 %rd2, [seq_cst_sys_i64_param_1];1691; SM90-NEXT: ld.param.u64 %rd3, [seq_cst_sys_i64_param_2];1692; SM90-NEXT: atom.cas.b64 %rd4, [%rd1], %rd2, %rd3;1693; SM90-NEXT: fence.acquire.sys;1694; SM90-NEXT: st.param.b64 [func_retval0], %rd3;1695; SM90-NEXT: ret;1696 %pairold = cmpxchg ptr %addr, i64 %cmp, i64 %new seq_cst seq_cst1697 ret i64 %new1698}1699 1700;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:1701; CHECK: {{.*}}1702