537 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_32 | FileCheck %s3; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_32 | %ptxas-verify %}4 5 6; CHECK-LABEL: atom07define i32 @atom0(ptr %addr, i32 %val) {8; CHECK-LABEL: atom0(9; CHECK: {10; CHECK-NEXT: .reg .b32 %r<3>;11; CHECK-NEXT: .reg .b64 %rd<2>;12; CHECK-EMPTY:13; CHECK-NEXT: // %bb.0:14; CHECK-NEXT: ld.param.b64 %rd1, [atom0_param_0];15; CHECK-NEXT: ld.param.b32 %r1, [atom0_param_1];16; CHECK-NEXT: atom.add.u32 %r2, [%rd1], %r1;17; CHECK-NEXT: st.param.b32 [func_retval0], %r2;18; CHECK-NEXT: ret;19 %ret = atomicrmw add ptr %addr, i32 %val seq_cst20 ret i32 %ret21}22 23; CHECK-LABEL: atom124define i64 @atom1(ptr %addr, i64 %val) {25; CHECK-LABEL: atom1(26; CHECK: {27; CHECK-NEXT: .reg .b64 %rd<4>;28; CHECK-EMPTY:29; CHECK-NEXT: // %bb.0:30; CHECK-NEXT: ld.param.b64 %rd1, [atom1_param_0];31; CHECK-NEXT: ld.param.b64 %rd2, [atom1_param_1];32; CHECK-NEXT: atom.add.u64 %rd3, [%rd1], %rd2;33; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;34; CHECK-NEXT: ret;35 %ret = atomicrmw add ptr %addr, i64 %val seq_cst36 ret i64 %ret37}38 39; CHECK-LABEL: atom240define i32 @atom2(ptr %subr, i32 %val) {41; CHECK-LABEL: atom2(42; CHECK: {43; CHECK-NEXT: .reg .b32 %r<4>;44; CHECK-NEXT: .reg .b64 %rd<2>;45; CHECK-EMPTY:46; CHECK-NEXT: // %bb.0:47; CHECK-NEXT: ld.param.b64 %rd1, [atom2_param_0];48; CHECK-NEXT: ld.param.b32 %r1, [atom2_param_1];49; CHECK-NEXT: neg.s32 %r2, %r1;50; CHECK-NEXT: atom.add.u32 %r3, [%rd1], %r2;51; CHECK-NEXT: st.param.b32 [func_retval0], %r3;52; CHECK-NEXT: ret;53 %ret = atomicrmw sub ptr %subr, i32 %val seq_cst54 ret i32 %ret55}56 57; CHECK-LABEL: atom358define i64 @atom3(ptr %subr, i64 %val) {59; CHECK-LABEL: atom3(60; CHECK: {61; CHECK-NEXT: .reg .b64 %rd<5>;62; CHECK-EMPTY:63; CHECK-NEXT: // %bb.0:64; CHECK-NEXT: ld.param.b64 %rd1, [atom3_param_0];65; CHECK-NEXT: ld.param.b64 %rd2, [atom3_param_1];66; CHECK-NEXT: neg.s64 %rd3, %rd2;67; CHECK-NEXT: atom.add.u64 %rd4, [%rd1], %rd3;68; CHECK-NEXT: st.param.b64 [func_retval0], %rd4;69; CHECK-NEXT: ret;70 %ret = atomicrmw sub ptr %subr, i64 %val seq_cst71 ret i64 %ret72}73 74; CHECK-LABEL: atom475define i32 @atom4(ptr %subr, i32 %val) {76; CHECK-LABEL: atom4(77; CHECK: {78; CHECK-NEXT: .reg .b32 %r<3>;79; CHECK-NEXT: .reg .b64 %rd<2>;80; CHECK-EMPTY:81; CHECK-NEXT: // %bb.0:82; CHECK-NEXT: ld.param.b64 %rd1, [atom4_param_0];83; CHECK-NEXT: ld.param.b32 %r1, [atom4_param_1];84; CHECK-NEXT: atom.and.b32 %r2, [%rd1], %r1;85; CHECK-NEXT: st.param.b32 [func_retval0], %r2;86; CHECK-NEXT: ret;87 %ret = atomicrmw and ptr %subr, i32 %val seq_cst88 ret i32 %ret89}90 91; CHECK-LABEL: atom592define i64 @atom5(ptr %subr, i64 %val) {93; CHECK-LABEL: atom5(94; CHECK: {95; CHECK-NEXT: .reg .b64 %rd<4>;96; CHECK-EMPTY:97; CHECK-NEXT: // %bb.0:98; CHECK-NEXT: ld.param.b64 %rd1, [atom5_param_0];99; CHECK-NEXT: ld.param.b64 %rd2, [atom5_param_1];100; CHECK-NEXT: atom.and.b64 %rd3, [%rd1], %rd2;101; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;102; CHECK-NEXT: ret;103 %ret = atomicrmw and ptr %subr, i64 %val seq_cst104 ret i64 %ret105}106 107;; NAND not yet supported108;define i32 @atom6(ptr %subr, i32 %val) {109; %ret = atomicrmw nand ptr %subr, i32 %val seq_cst110; ret i32 %ret111;}112 113;define i64 @atom7(ptr %subr, i64 %val) {114; %ret = atomicrmw nand ptr %subr, i64 %val seq_cst115; ret i64 %ret116;}117 118; CHECK-LABEL: atom8119define i32 @atom8(ptr %subr, i32 %val) {120; CHECK-LABEL: atom8(121; CHECK: {122; CHECK-NEXT: .reg .b32 %r<3>;123; CHECK-NEXT: .reg .b64 %rd<2>;124; CHECK-EMPTY:125; CHECK-NEXT: // %bb.0:126; CHECK-NEXT: ld.param.b64 %rd1, [atom8_param_0];127; CHECK-NEXT: ld.param.b32 %r1, [atom8_param_1];128; CHECK-NEXT: atom.or.b32 %r2, [%rd1], %r1;129; CHECK-NEXT: st.param.b32 [func_retval0], %r2;130; CHECK-NEXT: ret;131 %ret = atomicrmw or ptr %subr, i32 %val seq_cst132 ret i32 %ret133}134 135; CHECK-LABEL: atom9136define i64 @atom9(ptr %subr, i64 %val) {137; CHECK-LABEL: atom9(138; CHECK: {139; CHECK-NEXT: .reg .b64 %rd<4>;140; CHECK-EMPTY:141; CHECK-NEXT: // %bb.0:142; CHECK-NEXT: ld.param.b64 %rd1, [atom9_param_0];143; CHECK-NEXT: ld.param.b64 %rd2, [atom9_param_1];144; CHECK-NEXT: atom.or.b64 %rd3, [%rd1], %rd2;145; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;146; CHECK-NEXT: ret;147 %ret = atomicrmw or ptr %subr, i64 %val seq_cst148 ret i64 %ret149}150 151; CHECK-LABEL: atom10152define i32 @atom10(ptr %subr, i32 %val) {153; CHECK-LABEL: atom10(154; CHECK: {155; CHECK-NEXT: .reg .b32 %r<3>;156; CHECK-NEXT: .reg .b64 %rd<2>;157; CHECK-EMPTY:158; CHECK-NEXT: // %bb.0:159; CHECK-NEXT: ld.param.b64 %rd1, [atom10_param_0];160; CHECK-NEXT: ld.param.b32 %r1, [atom10_param_1];161; CHECK-NEXT: atom.xor.b32 %r2, [%rd1], %r1;162; CHECK-NEXT: st.param.b32 [func_retval0], %r2;163; CHECK-NEXT: ret;164 %ret = atomicrmw xor ptr %subr, i32 %val seq_cst165 ret i32 %ret166}167 168; CHECK-LABEL: atom11169define i64 @atom11(ptr %subr, i64 %val) {170; CHECK-LABEL: atom11(171; CHECK: {172; CHECK-NEXT: .reg .b64 %rd<4>;173; CHECK-EMPTY:174; CHECK-NEXT: // %bb.0:175; CHECK-NEXT: ld.param.b64 %rd1, [atom11_param_0];176; CHECK-NEXT: ld.param.b64 %rd2, [atom11_param_1];177; CHECK-NEXT: atom.xor.b64 %rd3, [%rd1], %rd2;178; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;179; CHECK-NEXT: ret;180 %ret = atomicrmw xor ptr %subr, i64 %val seq_cst181 ret i64 %ret182}183 184; CHECK-LABEL: atom12185define i32 @atom12(ptr %subr, i32 %val) {186; CHECK-LABEL: atom12(187; CHECK: {188; CHECK-NEXT: .reg .b32 %r<3>;189; CHECK-NEXT: .reg .b64 %rd<2>;190; CHECK-EMPTY:191; CHECK-NEXT: // %bb.0:192; CHECK-NEXT: ld.param.b64 %rd1, [atom12_param_0];193; CHECK-NEXT: ld.param.b32 %r1, [atom12_param_1];194; CHECK-NEXT: atom.max.s32 %r2, [%rd1], %r1;195; CHECK-NEXT: st.param.b32 [func_retval0], %r2;196; CHECK-NEXT: ret;197 %ret = atomicrmw max ptr %subr, i32 %val seq_cst198 ret i32 %ret199}200 201; CHECK-LABEL: atom13202define i64 @atom13(ptr %subr, i64 %val) {203; CHECK-LABEL: atom13(204; CHECK: {205; CHECK-NEXT: .reg .b64 %rd<4>;206; CHECK-EMPTY:207; CHECK-NEXT: // %bb.0:208; CHECK-NEXT: ld.param.b64 %rd1, [atom13_param_0];209; CHECK-NEXT: ld.param.b64 %rd2, [atom13_param_1];210; CHECK-NEXT: atom.max.s64 %rd3, [%rd1], %rd2;211; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;212; CHECK-NEXT: ret;213 %ret = atomicrmw max ptr %subr, i64 %val seq_cst214 ret i64 %ret215}216 217; CHECK-LABEL: atom14218define i32 @atom14(ptr %subr, i32 %val) {219; CHECK-LABEL: atom14(220; CHECK: {221; CHECK-NEXT: .reg .b32 %r<3>;222; CHECK-NEXT: .reg .b64 %rd<2>;223; CHECK-EMPTY:224; CHECK-NEXT: // %bb.0:225; CHECK-NEXT: ld.param.b64 %rd1, [atom14_param_0];226; CHECK-NEXT: ld.param.b32 %r1, [atom14_param_1];227; CHECK-NEXT: atom.min.s32 %r2, [%rd1], %r1;228; CHECK-NEXT: st.param.b32 [func_retval0], %r2;229; CHECK-NEXT: ret;230 %ret = atomicrmw min ptr %subr, i32 %val seq_cst231 ret i32 %ret232}233 234; CHECK-LABEL: atom15235define i64 @atom15(ptr %subr, i64 %val) {236; CHECK-LABEL: atom15(237; CHECK: {238; CHECK-NEXT: .reg .b64 %rd<4>;239; CHECK-EMPTY:240; CHECK-NEXT: // %bb.0:241; CHECK-NEXT: ld.param.b64 %rd1, [atom15_param_0];242; CHECK-NEXT: ld.param.b64 %rd2, [atom15_param_1];243; CHECK-NEXT: atom.min.s64 %rd3, [%rd1], %rd2;244; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;245; CHECK-NEXT: ret;246 %ret = atomicrmw min ptr %subr, i64 %val seq_cst247 ret i64 %ret248}249 250; CHECK-LABEL: atom16251define i32 @atom16(ptr %subr, i32 %val) {252; CHECK-LABEL: atom16(253; CHECK: {254; CHECK-NEXT: .reg .b32 %r<3>;255; CHECK-NEXT: .reg .b64 %rd<2>;256; CHECK-EMPTY:257; CHECK-NEXT: // %bb.0:258; CHECK-NEXT: ld.param.b64 %rd1, [atom16_param_0];259; CHECK-NEXT: ld.param.b32 %r1, [atom16_param_1];260; CHECK-NEXT: atom.max.u32 %r2, [%rd1], %r1;261; CHECK-NEXT: st.param.b32 [func_retval0], %r2;262; CHECK-NEXT: ret;263 %ret = atomicrmw umax ptr %subr, i32 %val seq_cst264 ret i32 %ret265}266 267; CHECK-LABEL: atom17268define i64 @atom17(ptr %subr, i64 %val) {269; CHECK-LABEL: atom17(270; CHECK: {271; CHECK-NEXT: .reg .b64 %rd<4>;272; CHECK-EMPTY:273; CHECK-NEXT: // %bb.0:274; CHECK-NEXT: ld.param.b64 %rd1, [atom17_param_0];275; CHECK-NEXT: ld.param.b64 %rd2, [atom17_param_1];276; CHECK-NEXT: atom.max.u64 %rd3, [%rd1], %rd2;277; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;278; CHECK-NEXT: ret;279 %ret = atomicrmw umax ptr %subr, i64 %val seq_cst280 ret i64 %ret281}282 283; CHECK-LABEL: atom18284define i32 @atom18(ptr %subr, i32 %val) {285; CHECK-LABEL: atom18(286; CHECK: {287; CHECK-NEXT: .reg .b32 %r<3>;288; CHECK-NEXT: .reg .b64 %rd<2>;289; CHECK-EMPTY:290; CHECK-NEXT: // %bb.0:291; CHECK-NEXT: ld.param.b64 %rd1, [atom18_param_0];292; CHECK-NEXT: ld.param.b32 %r1, [atom18_param_1];293; CHECK-NEXT: atom.min.u32 %r2, [%rd1], %r1;294; CHECK-NEXT: st.param.b32 [func_retval0], %r2;295; CHECK-NEXT: ret;296 %ret = atomicrmw umin ptr %subr, i32 %val seq_cst297 ret i32 %ret298}299 300; CHECK-LABEL: atom19301define i64 @atom19(ptr %subr, i64 %val) {302; CHECK-LABEL: atom19(303; CHECK: {304; CHECK-NEXT: .reg .b64 %rd<4>;305; CHECK-EMPTY:306; CHECK-NEXT: // %bb.0:307; CHECK-NEXT: ld.param.b64 %rd1, [atom19_param_0];308; CHECK-NEXT: ld.param.b64 %rd2, [atom19_param_1];309; CHECK-NEXT: atom.min.u64 %rd3, [%rd1], %rd2;310; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;311; CHECK-NEXT: ret;312 %ret = atomicrmw umin ptr %subr, i64 %val seq_cst313 ret i64 %ret314}315 316define i32 @atom20(ptr %subr, i32 %val) {317; CHECK-LABEL: atom20(318; CHECK: {319; CHECK-NEXT: .reg .b32 %r<3>;320; CHECK-NEXT: .reg .b64 %rd<2>;321; CHECK-EMPTY:322; CHECK-NEXT: // %bb.0:323; CHECK-NEXT: ld.param.b64 %rd1, [atom20_param_0];324; CHECK-NEXT: ld.param.b32 %r1, [atom20_param_1];325; CHECK-NEXT: atom.inc.u32 %r2, [%rd1], %r1;326; CHECK-NEXT: st.param.b32 [func_retval0], %r2;327; CHECK-NEXT: ret;328 %ret = atomicrmw uinc_wrap ptr %subr, i32 %val seq_cst329 ret i32 %ret330}331 332define i32 @atom21(ptr %subr, i32 %val) {333; CHECK-LABEL: atom21(334; CHECK: {335; CHECK-NEXT: .reg .b32 %r<3>;336; CHECK-NEXT: .reg .b64 %rd<2>;337; CHECK-EMPTY:338; CHECK-NEXT: // %bb.0:339; CHECK-NEXT: ld.param.b64 %rd1, [atom21_param_0];340; CHECK-NEXT: ld.param.b32 %r1, [atom21_param_1];341; CHECK-NEXT: atom.dec.u32 %r2, [%rd1], %r1;342; CHECK-NEXT: st.param.b32 [func_retval0], %r2;343; CHECK-NEXT: ret;344 %ret = atomicrmw udec_wrap ptr %subr, i32 %val seq_cst345 ret i32 %ret346}347 348declare float @llvm.nvvm.atomic.load.add.f32.p0(ptr %addr, float %val)349 350; CHECK-LABEL: atomic_add_f32_generic351define float @atomic_add_f32_generic(ptr %addr, float %val) {352; CHECK-LABEL: atomic_add_f32_generic(353; CHECK: {354; CHECK-NEXT: .reg .b32 %r<3>;355; CHECK-NEXT: .reg .b64 %rd<2>;356; CHECK-EMPTY:357; CHECK-NEXT: // %bb.0:358; CHECK-NEXT: ld.param.b64 %rd1, [atomic_add_f32_generic_param_0];359; CHECK-NEXT: ld.param.b32 %r1, [atomic_add_f32_generic_param_1];360; CHECK-NEXT: atom.add.f32 %r2, [%rd1], %r1;361; CHECK-NEXT: st.param.b32 [func_retval0], %r2;362; CHECK-NEXT: ret;363 %ret = call float @llvm.nvvm.atomic.load.add.f32.p0(ptr %addr, float %val)364 ret float %ret365}366 367declare float @llvm.nvvm.atomic.load.add.f32.p1(ptr addrspace(1) %addr, float %val)368 369; CHECK-LABEL: atomic_add_f32_addrspace1370define float @atomic_add_f32_addrspace1(ptr addrspace(1) %addr, float %val) {371; CHECK-LABEL: atomic_add_f32_addrspace1(372; CHECK: {373; CHECK-NEXT: .reg .b32 %r<3>;374; CHECK-NEXT: .reg .b64 %rd<2>;375; CHECK-EMPTY:376; CHECK-NEXT: // %bb.0:377; CHECK-NEXT: ld.param.b64 %rd1, [atomic_add_f32_addrspace1_param_0];378; CHECK-NEXT: ld.param.b32 %r1, [atomic_add_f32_addrspace1_param_1];379; CHECK-NEXT: atom.global.add.f32 %r2, [%rd1], %r1;380; CHECK-NEXT: st.param.b32 [func_retval0], %r2;381; CHECK-NEXT: ret;382 %ret = call float @llvm.nvvm.atomic.load.add.f32.p1(ptr addrspace(1) %addr, float %val)383 ret float %ret384}385 386declare float @llvm.nvvm.atomic.load.add.f32.p3(ptr addrspace(3) %addr, float %val)387 388; CHECK-LABEL: atomic_add_f32_addrspace3389define float @atomic_add_f32_addrspace3(ptr addrspace(3) %addr, float %val) {390; CHECK-LABEL: atomic_add_f32_addrspace3(391; CHECK: {392; CHECK-NEXT: .reg .b32 %r<3>;393; CHECK-NEXT: .reg .b64 %rd<2>;394; CHECK-EMPTY:395; CHECK-NEXT: // %bb.0:396; CHECK-NEXT: ld.param.b64 %rd1, [atomic_add_f32_addrspace3_param_0];397; CHECK-NEXT: ld.param.b32 %r1, [atomic_add_f32_addrspace3_param_1];398; CHECK-NEXT: atom.shared.add.f32 %r2, [%rd1], %r1;399; CHECK-NEXT: st.param.b32 [func_retval0], %r2;400; CHECK-NEXT: ret;401 %ret = call float @llvm.nvvm.atomic.load.add.f32.p3(ptr addrspace(3) %addr, float %val)402 ret float %ret403}404 405; CHECK-LABEL: atomicrmw_add_f32_generic406define float @atomicrmw_add_f32_generic(ptr %addr, float %val) {407; CHECK-LABEL: atomicrmw_add_f32_generic(408; CHECK: {409; CHECK-NEXT: .reg .b32 %r<3>;410; CHECK-NEXT: .reg .b64 %rd<2>;411; CHECK-EMPTY:412; CHECK-NEXT: // %bb.0:413; CHECK-NEXT: ld.param.b64 %rd1, [atomicrmw_add_f32_generic_param_0];414; CHECK-NEXT: ld.param.b32 %r1, [atomicrmw_add_f32_generic_param_1];415; CHECK-NEXT: atom.add.f32 %r2, [%rd1], %r1;416; CHECK-NEXT: st.param.b32 [func_retval0], %r2;417; CHECK-NEXT: ret;418 %ret = atomicrmw fadd ptr %addr, float %val seq_cst419 ret float %ret420}421 422; CHECK-LABEL: atomicrmw_add_f16_generic423define half @atomicrmw_add_f16_generic(ptr %addr, half %val) {424; CHECK-LABEL: atomicrmw_add_f16_generic(425; CHECK: {426; CHECK-NEXT: .reg .pred %p<2>;427; CHECK-NEXT: .reg .b16 %rs<4>;428; CHECK-NEXT: .reg .b32 %r<18>;429; CHECK-NEXT: .reg .b64 %rd<3>;430; CHECK-EMPTY:431; CHECK-NEXT: // %bb.0:432; CHECK-NEXT: ld.param.b16 %rs1, [atomicrmw_add_f16_generic_param_1];433; CHECK-NEXT: ld.param.b64 %rd2, [atomicrmw_add_f16_generic_param_0];434; CHECK-NEXT: and.b64 %rd1, %rd2, -4;435; CHECK-NEXT: cvt.u32.u64 %r4, %rd2;436; CHECK-NEXT: and.b32 %r5, %r4, 3;437; CHECK-NEXT: shl.b32 %r1, %r5, 3;438; CHECK-NEXT: mov.b32 %r6, 65535;439; CHECK-NEXT: shl.b32 %r7, %r6, %r1;440; CHECK-NEXT: not.b32 %r2, %r7;441; CHECK-NEXT: ld.b32 %r17, [%rd1];442; CHECK-NEXT: cvt.f32.f16 %r10, %rs1;443; CHECK-NEXT: $L__BB24_1: // %atomicrmw.start444; CHECK-NEXT: // =>This Inner Loop Header: Depth=1445; CHECK-NEXT: shr.u32 %r8, %r17, %r1;446; CHECK-NEXT: cvt.u16.u32 %rs2, %r8;447; CHECK-NEXT: cvt.f32.f16 %r9, %rs2;448; CHECK-NEXT: add.rn.f32 %r11, %r9, %r10;449; CHECK-NEXT: cvt.rn.f16.f32 %rs3, %r11;450; CHECK-NEXT: cvt.u32.u16 %r12, %rs3;451; CHECK-NEXT: shl.b32 %r13, %r12, %r1;452; CHECK-NEXT: and.b32 %r14, %r17, %r2;453; CHECK-NEXT: or.b32 %r15, %r14, %r13;454; CHECK-NEXT: membar.sys;455; CHECK-NEXT: atom.cas.b32 %r3, [%rd1], %r17, %r15;456; CHECK-NEXT: setp.ne.b32 %p1, %r3, %r17;457; CHECK-NEXT: mov.b32 %r17, %r3;458; CHECK-NEXT: @%p1 bra $L__BB24_1;459; CHECK-NEXT: // %bb.2: // %atomicrmw.end460; CHECK-NEXT: shr.u32 %r16, %r3, %r1;461; CHECK-NEXT: st.param.b16 [func_retval0], %r16;462; CHECK-NEXT: ret;463 %ret = atomicrmw fadd ptr %addr, half %val seq_cst464 ret half %ret465}466 467; CHECK-LABEL: atomicrmw_add_f32_addrspace1468define float @atomicrmw_add_f32_addrspace1(ptr addrspace(1) %addr, float %val) {469; CHECK-LABEL: atomicrmw_add_f32_addrspace1(470; CHECK: {471; CHECK-NEXT: .reg .b32 %r<3>;472; CHECK-NEXT: .reg .b64 %rd<2>;473; CHECK-EMPTY:474; CHECK-NEXT: // %bb.0:475; CHECK-NEXT: ld.param.b64 %rd1, [atomicrmw_add_f32_addrspace1_param_0];476; CHECK-NEXT: ld.param.b32 %r1, [atomicrmw_add_f32_addrspace1_param_1];477; CHECK-NEXT: atom.global.add.f32 %r2, [%rd1], %r1;478; CHECK-NEXT: st.param.b32 [func_retval0], %r2;479; CHECK-NEXT: ret;480 %ret = atomicrmw fadd ptr addrspace(1) %addr, float %val seq_cst481 ret float %ret482}483 484; CHECK-LABEL: atomicrmw_add_f32_addrspace3485define float @atomicrmw_add_f32_addrspace3(ptr addrspace(3) %addr, float %val) {486; CHECK-LABEL: atomicrmw_add_f32_addrspace3(487; CHECK: {488; CHECK-NEXT: .reg .b32 %r<3>;489; CHECK-NEXT: .reg .b64 %rd<2>;490; CHECK-EMPTY:491; CHECK-NEXT: // %bb.0:492; CHECK-NEXT: ld.param.b64 %rd1, [atomicrmw_add_f32_addrspace3_param_0];493; CHECK-NEXT: ld.param.b32 %r1, [atomicrmw_add_f32_addrspace3_param_1];494; CHECK-NEXT: atom.shared.add.f32 %r2, [%rd1], %r1;495; CHECK-NEXT: st.param.b32 [func_retval0], %r2;496; CHECK-NEXT: ret;497 %ret = atomicrmw fadd ptr addrspace(3) %addr, float %val seq_cst498 ret float %ret499}500 501; CHECK-LABEL: atomic_cmpxchg_i32502define i32 @atomic_cmpxchg_i32(ptr %addr, i32 %cmp, i32 %new) {503; CHECK-LABEL: atomic_cmpxchg_i32(504; CHECK: {505; CHECK-NEXT: .reg .b32 %r<4>;506; CHECK-NEXT: .reg .b64 %rd<2>;507; CHECK-EMPTY:508; CHECK-NEXT: // %bb.0:509; CHECK-NEXT: ld.param.b64 %rd1, [atomic_cmpxchg_i32_param_0];510; CHECK-NEXT: membar.sys;511; CHECK-NEXT: ld.param.b32 %r1, [atomic_cmpxchg_i32_param_1];512; CHECK-NEXT: ld.param.b32 %r2, [atomic_cmpxchg_i32_param_2];513; CHECK-NEXT: atom.cas.b32 %r3, [%rd1], %r1, %r2;514; CHECK-NEXT: st.param.b32 [func_retval0], %r2;515; CHECK-NEXT: ret;516 %pairold = cmpxchg ptr %addr, i32 %cmp, i32 %new seq_cst seq_cst517 ret i32 %new518}519 520; CHECK-LABEL: atomic_cmpxchg_i64521define i64 @atomic_cmpxchg_i64(ptr %addr, i64 %cmp, i64 %new) {522; CHECK-LABEL: atomic_cmpxchg_i64(523; CHECK: {524; CHECK-NEXT: .reg .b64 %rd<5>;525; CHECK-EMPTY:526; CHECK-NEXT: // %bb.0:527; CHECK-NEXT: ld.param.b64 %rd1, [atomic_cmpxchg_i64_param_0];528; CHECK-NEXT: membar.sys;529; CHECK-NEXT: ld.param.b64 %rd2, [atomic_cmpxchg_i64_param_1];530; CHECK-NEXT: ld.param.b64 %rd3, [atomic_cmpxchg_i64_param_2];531; CHECK-NEXT: atom.cas.b64 %rd4, [%rd1], %rd2, %rd3;532; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;533; CHECK-NEXT: ret;534 %pairold = cmpxchg ptr %addr, i64 %cmp, i64 %new seq_cst seq_cst535 ret i64 %new536}537