3175 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_20 | FileCheck -check-prefixes=CHECK,SM60 %s3; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}4; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_70 -mattr=+ptx82 | FileCheck %s -check-prefixes=CHECK,SM705; RUN: %if ptxas-sm_70 && ptxas-isa-8.2 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_70 -mattr=+ptx82 | %ptxas-verify -arch=sm_70 %}6 7; TODO: generate PTX that preserves Concurrent Forward Progress8; for atomic operations to local statespace9; by generating atomic or volatile operations.10 11; TODO: add weak,atomic,volatile,atomic volatile tests12; for .const and .param statespaces.13 14; TODO: optimize .sys.shared into .cta.shared or .cluster.shared .15 16;; generic statespace17 18; generic19 20define void @generic_i8(ptr %a) {21; CHECK-LABEL: generic_i8(22; CHECK: {23; CHECK-NEXT: .reg .b16 %rs<3>;24; CHECK-NEXT: .reg .b64 %rd<2>;25; CHECK-EMPTY:26; CHECK-NEXT: // %bb.0:27; CHECK-NEXT: ld.param.b64 %rd1, [generic_i8_param_0];28; CHECK-NEXT: ld.b8 %rs1, [%rd1];29; CHECK-NEXT: add.s16 %rs2, %rs1, 1;30; CHECK-NEXT: st.b8 [%rd1], %rs2;31; CHECK-NEXT: ret;32 %a.load = load i8, ptr %a33 %a.add = add i8 %a.load, 134 store i8 %a.add, ptr %a35 ret void36}37 38define void @generic_i16(ptr %a) {39; CHECK-LABEL: generic_i16(40; CHECK: {41; CHECK-NEXT: .reg .b16 %rs<3>;42; CHECK-NEXT: .reg .b64 %rd<2>;43; CHECK-EMPTY:44; CHECK-NEXT: // %bb.0:45; CHECK-NEXT: ld.param.b64 %rd1, [generic_i16_param_0];46; CHECK-NEXT: ld.b16 %rs1, [%rd1];47; CHECK-NEXT: add.s16 %rs2, %rs1, 1;48; CHECK-NEXT: st.b16 [%rd1], %rs2;49; CHECK-NEXT: ret;50 %a.load = load i16, ptr %a51 %a.add = add i16 %a.load, 152 store i16 %a.add, ptr %a53 ret void54}55 56define void @generic_i32(ptr %a) {57; CHECK-LABEL: generic_i32(58; CHECK: {59; CHECK-NEXT: .reg .b32 %r<3>;60; CHECK-NEXT: .reg .b64 %rd<2>;61; CHECK-EMPTY:62; CHECK-NEXT: // %bb.0:63; CHECK-NEXT: ld.param.b64 %rd1, [generic_i32_param_0];64; CHECK-NEXT: ld.b32 %r1, [%rd1];65; CHECK-NEXT: add.s32 %r2, %r1, 1;66; CHECK-NEXT: st.b32 [%rd1], %r2;67; CHECK-NEXT: ret;68 %a.load = load i32, ptr %a69 %a.add = add i32 %a.load, 170 store i32 %a.add, ptr %a71 ret void72}73 74define void @generic_i64(ptr %a) {75; CHECK-LABEL: generic_i64(76; CHECK: {77; CHECK-NEXT: .reg .b64 %rd<4>;78; CHECK-EMPTY:79; CHECK-NEXT: // %bb.0:80; CHECK-NEXT: ld.param.b64 %rd1, [generic_i64_param_0];81; CHECK-NEXT: ld.b64 %rd2, [%rd1];82; CHECK-NEXT: add.s64 %rd3, %rd2, 1;83; CHECK-NEXT: st.b64 [%rd1], %rd3;84; CHECK-NEXT: ret;85 %a.load = load i64, ptr %a86 %a.add = add i64 %a.load, 187 store i64 %a.add, ptr %a88 ret void89}90 91define void @generic_float(ptr %a) {92; CHECK-LABEL: generic_float(93; CHECK: {94; CHECK-NEXT: .reg .b32 %r<3>;95; CHECK-NEXT: .reg .b64 %rd<2>;96; CHECK-EMPTY:97; CHECK-NEXT: // %bb.0:98; CHECK-NEXT: ld.param.b64 %rd1, [generic_float_param_0];99; CHECK-NEXT: ld.b32 %r1, [%rd1];100; CHECK-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;101; CHECK-NEXT: st.b32 [%rd1], %r2;102; CHECK-NEXT: ret;103 %a.load = load float, ptr %a104 %a.add = fadd float %a.load, 1.105 store float %a.add, ptr %a106 ret void107}108 109define void @generic_double(ptr %a) {110; CHECK-LABEL: generic_double(111; CHECK: {112; CHECK-NEXT: .reg .b64 %rd<4>;113; CHECK-EMPTY:114; CHECK-NEXT: // %bb.0:115; CHECK-NEXT: ld.param.b64 %rd1, [generic_double_param_0];116; CHECK-NEXT: ld.b64 %rd2, [%rd1];117; CHECK-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;118; CHECK-NEXT: st.b64 [%rd1], %rd3;119; CHECK-NEXT: ret;120 %a.load = load double, ptr %a121 %a.add = fadd double %a.load, 1.122 store double %a.add, ptr %a123 ret void124}125 126; generic_volatile127 128define void @generic_volatile_i8(ptr %a) {129; CHECK-LABEL: generic_volatile_i8(130; CHECK: {131; CHECK-NEXT: .reg .b16 %rs<3>;132; CHECK-NEXT: .reg .b64 %rd<2>;133; CHECK-EMPTY:134; CHECK-NEXT: // %bb.0:135; CHECK-NEXT: ld.param.b64 %rd1, [generic_volatile_i8_param_0];136; CHECK-NEXT: ld.volatile.b8 %rs1, [%rd1];137; CHECK-NEXT: add.s16 %rs2, %rs1, 1;138; CHECK-NEXT: st.volatile.b8 [%rd1], %rs2;139; CHECK-NEXT: ret;140 %a.load = load volatile i8, ptr %a141 %a.add = add i8 %a.load, 1142 store volatile i8 %a.add, ptr %a143 ret void144}145 146define void @generic_volatile_i16(ptr %a) {147; CHECK-LABEL: generic_volatile_i16(148; CHECK: {149; CHECK-NEXT: .reg .b16 %rs<3>;150; CHECK-NEXT: .reg .b64 %rd<2>;151; CHECK-EMPTY:152; CHECK-NEXT: // %bb.0:153; CHECK-NEXT: ld.param.b64 %rd1, [generic_volatile_i16_param_0];154; CHECK-NEXT: ld.volatile.b16 %rs1, [%rd1];155; CHECK-NEXT: add.s16 %rs2, %rs1, 1;156; CHECK-NEXT: st.volatile.b16 [%rd1], %rs2;157; CHECK-NEXT: ret;158 %a.load = load volatile i16, ptr %a159 %a.add = add i16 %a.load, 1160 store volatile i16 %a.add, ptr %a161 ret void162}163 164define void @generic_volatile_i32(ptr %a) {165; CHECK-LABEL: generic_volatile_i32(166; CHECK: {167; CHECK-NEXT: .reg .b32 %r<3>;168; CHECK-NEXT: .reg .b64 %rd<2>;169; CHECK-EMPTY:170; CHECK-NEXT: // %bb.0:171; CHECK-NEXT: ld.param.b64 %rd1, [generic_volatile_i32_param_0];172; CHECK-NEXT: ld.volatile.b32 %r1, [%rd1];173; CHECK-NEXT: add.s32 %r2, %r1, 1;174; CHECK-NEXT: st.volatile.b32 [%rd1], %r2;175; CHECK-NEXT: ret;176 %a.load = load volatile i32, ptr %a177 %a.add = add i32 %a.load, 1178 store volatile i32 %a.add, ptr %a179 ret void180}181 182define void @generic_volatile_i64(ptr %a) {183; CHECK-LABEL: generic_volatile_i64(184; CHECK: {185; CHECK-NEXT: .reg .b64 %rd<4>;186; CHECK-EMPTY:187; CHECK-NEXT: // %bb.0:188; CHECK-NEXT: ld.param.b64 %rd1, [generic_volatile_i64_param_0];189; CHECK-NEXT: ld.volatile.b64 %rd2, [%rd1];190; CHECK-NEXT: add.s64 %rd3, %rd2, 1;191; CHECK-NEXT: st.volatile.b64 [%rd1], %rd3;192; CHECK-NEXT: ret;193 %a.load = load volatile i64, ptr %a194 %a.add = add i64 %a.load, 1195 store volatile i64 %a.add, ptr %a196 ret void197}198 199define void @generic_volatile_float(ptr %a) {200; CHECK-LABEL: generic_volatile_float(201; CHECK: {202; CHECK-NEXT: .reg .b32 %r<3>;203; CHECK-NEXT: .reg .b64 %rd<2>;204; CHECK-EMPTY:205; CHECK-NEXT: // %bb.0:206; CHECK-NEXT: ld.param.b64 %rd1, [generic_volatile_float_param_0];207; CHECK-NEXT: ld.volatile.b32 %r1, [%rd1];208; CHECK-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;209; CHECK-NEXT: st.volatile.b32 [%rd1], %r2;210; CHECK-NEXT: ret;211 %a.load = load volatile float, ptr %a212 %a.add = fadd float %a.load, 1.213 store volatile float %a.add, ptr %a214 ret void215}216 217define void @generic_volatile_double(ptr %a) {218; CHECK-LABEL: generic_volatile_double(219; CHECK: {220; CHECK-NEXT: .reg .b64 %rd<4>;221; CHECK-EMPTY:222; CHECK-NEXT: // %bb.0:223; CHECK-NEXT: ld.param.b64 %rd1, [generic_volatile_double_param_0];224; CHECK-NEXT: ld.volatile.b64 %rd2, [%rd1];225; CHECK-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;226; CHECK-NEXT: st.volatile.b64 [%rd1], %rd3;227; CHECK-NEXT: ret;228 %a.load = load volatile double, ptr %a229 %a.add = fadd double %a.load, 1.230 store volatile double %a.add, ptr %a231 ret void232}233 234; generic_unordered_sys235 236define void @generic_unordered_sys_i8(ptr %a) {237; SM60-LABEL: generic_unordered_sys_i8(238; SM60: {239; SM60-NEXT: .reg .b16 %rs<3>;240; SM60-NEXT: .reg .b64 %rd<2>;241; SM60-EMPTY:242; SM60-NEXT: // %bb.0:243; SM60-NEXT: ld.param.b64 %rd1, [generic_unordered_sys_i8_param_0];244; SM60-NEXT: ld.volatile.b8 %rs1, [%rd1];245; SM60-NEXT: add.s16 %rs2, %rs1, 1;246; SM60-NEXT: st.volatile.b8 [%rd1], %rs2;247; SM60-NEXT: ret;248;249; SM70-LABEL: generic_unordered_sys_i8(250; SM70: {251; SM70-NEXT: .reg .b16 %rs<3>;252; SM70-NEXT: .reg .b64 %rd<2>;253; SM70-EMPTY:254; SM70-NEXT: // %bb.0:255; SM70-NEXT: ld.param.b64 %rd1, [generic_unordered_sys_i8_param_0];256; SM70-NEXT: ld.relaxed.sys.b8 %rs1, [%rd1];257; SM70-NEXT: add.s16 %rs2, %rs1, 1;258; SM70-NEXT: st.relaxed.sys.b8 [%rd1], %rs2;259; SM70-NEXT: ret;260 %a.load = load atomic i8, ptr %a unordered, align 1261 %a.add = add i8 %a.load, 1262 store atomic i8 %a.add, ptr %a unordered, align 1263 ret void264}265 266define void @generic_unordered_sys_i16(ptr %a) {267; SM60-LABEL: generic_unordered_sys_i16(268; SM60: {269; SM60-NEXT: .reg .b16 %rs<3>;270; SM60-NEXT: .reg .b64 %rd<2>;271; SM60-EMPTY:272; SM60-NEXT: // %bb.0:273; SM60-NEXT: ld.param.b64 %rd1, [generic_unordered_sys_i16_param_0];274; SM60-NEXT: ld.volatile.b16 %rs1, [%rd1];275; SM60-NEXT: add.s16 %rs2, %rs1, 1;276; SM60-NEXT: st.volatile.b16 [%rd1], %rs2;277; SM60-NEXT: ret;278;279; SM70-LABEL: generic_unordered_sys_i16(280; SM70: {281; SM70-NEXT: .reg .b16 %rs<3>;282; SM70-NEXT: .reg .b64 %rd<2>;283; SM70-EMPTY:284; SM70-NEXT: // %bb.0:285; SM70-NEXT: ld.param.b64 %rd1, [generic_unordered_sys_i16_param_0];286; SM70-NEXT: ld.relaxed.sys.b16 %rs1, [%rd1];287; SM70-NEXT: add.s16 %rs2, %rs1, 1;288; SM70-NEXT: st.relaxed.sys.b16 [%rd1], %rs2;289; SM70-NEXT: ret;290 %a.load = load atomic i16, ptr %a unordered, align 2291 %a.add = add i16 %a.load, 1292 store atomic i16 %a.add, ptr %a unordered, align 2293 ret void294}295 296define void @generic_unordered_sys_i32(ptr %a) {297; SM60-LABEL: generic_unordered_sys_i32(298; SM60: {299; SM60-NEXT: .reg .b32 %r<3>;300; SM60-NEXT: .reg .b64 %rd<2>;301; SM60-EMPTY:302; SM60-NEXT: // %bb.0:303; SM60-NEXT: ld.param.b64 %rd1, [generic_unordered_sys_i32_param_0];304; SM60-NEXT: ld.volatile.b32 %r1, [%rd1];305; SM60-NEXT: add.s32 %r2, %r1, 1;306; SM60-NEXT: st.volatile.b32 [%rd1], %r2;307; SM60-NEXT: ret;308;309; SM70-LABEL: generic_unordered_sys_i32(310; SM70: {311; SM70-NEXT: .reg .b32 %r<3>;312; SM70-NEXT: .reg .b64 %rd<2>;313; SM70-EMPTY:314; SM70-NEXT: // %bb.0:315; SM70-NEXT: ld.param.b64 %rd1, [generic_unordered_sys_i32_param_0];316; SM70-NEXT: ld.relaxed.sys.b32 %r1, [%rd1];317; SM70-NEXT: add.s32 %r2, %r1, 1;318; SM70-NEXT: st.relaxed.sys.b32 [%rd1], %r2;319; SM70-NEXT: ret;320 %a.load = load atomic i32, ptr %a unordered, align 4321 %a.add = add i32 %a.load, 1322 store atomic i32 %a.add, ptr %a unordered, align 4323 ret void324}325 326define void @generic_unordered_sys_i64(ptr %a) {327; SM60-LABEL: generic_unordered_sys_i64(328; SM60: {329; SM60-NEXT: .reg .b64 %rd<4>;330; SM60-EMPTY:331; SM60-NEXT: // %bb.0:332; SM60-NEXT: ld.param.b64 %rd1, [generic_unordered_sys_i64_param_0];333; SM60-NEXT: ld.volatile.b64 %rd2, [%rd1];334; SM60-NEXT: add.s64 %rd3, %rd2, 1;335; SM60-NEXT: st.volatile.b64 [%rd1], %rd3;336; SM60-NEXT: ret;337;338; SM70-LABEL: generic_unordered_sys_i64(339; SM70: {340; SM70-NEXT: .reg .b64 %rd<4>;341; SM70-EMPTY:342; SM70-NEXT: // %bb.0:343; SM70-NEXT: ld.param.b64 %rd1, [generic_unordered_sys_i64_param_0];344; SM70-NEXT: ld.relaxed.sys.b64 %rd2, [%rd1];345; SM70-NEXT: add.s64 %rd3, %rd2, 1;346; SM70-NEXT: st.relaxed.sys.b64 [%rd1], %rd3;347; SM70-NEXT: ret;348 %a.load = load atomic i64, ptr %a unordered, align 8349 %a.add = add i64 %a.load, 1350 store atomic i64 %a.add, ptr %a unordered, align 8351 ret void352}353 354define void @generic_unordered_sys_float(ptr %a) {355; SM60-LABEL: generic_unordered_sys_float(356; SM60: {357; SM60-NEXT: .reg .b32 %r<3>;358; SM60-NEXT: .reg .b64 %rd<2>;359; SM60-EMPTY:360; SM60-NEXT: // %bb.0:361; SM60-NEXT: ld.param.b64 %rd1, [generic_unordered_sys_float_param_0];362; SM60-NEXT: ld.volatile.b32 %r1, [%rd1];363; SM60-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;364; SM60-NEXT: st.volatile.b32 [%rd1], %r2;365; SM60-NEXT: ret;366;367; SM70-LABEL: generic_unordered_sys_float(368; SM70: {369; SM70-NEXT: .reg .b32 %r<3>;370; SM70-NEXT: .reg .b64 %rd<2>;371; SM70-EMPTY:372; SM70-NEXT: // %bb.0:373; SM70-NEXT: ld.param.b64 %rd1, [generic_unordered_sys_float_param_0];374; SM70-NEXT: ld.relaxed.sys.b32 %r1, [%rd1];375; SM70-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;376; SM70-NEXT: st.relaxed.sys.b32 [%rd1], %r2;377; SM70-NEXT: ret;378 %a.load = load atomic float, ptr %a unordered, align 4379 %a.add = fadd float %a.load, 1.380 store atomic float %a.add, ptr %a unordered, align 4381 ret void382}383 384define void @generic_unordered_sys_double(ptr %a) {385; SM60-LABEL: generic_unordered_sys_double(386; SM60: {387; SM60-NEXT: .reg .b64 %rd<4>;388; SM60-EMPTY:389; SM60-NEXT: // %bb.0:390; SM60-NEXT: ld.param.b64 %rd1, [generic_unordered_sys_double_param_0];391; SM60-NEXT: ld.volatile.b64 %rd2, [%rd1];392; SM60-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;393; SM60-NEXT: st.volatile.b64 [%rd1], %rd3;394; SM60-NEXT: ret;395;396; SM70-LABEL: generic_unordered_sys_double(397; SM70: {398; SM70-NEXT: .reg .b64 %rd<4>;399; SM70-EMPTY:400; SM70-NEXT: // %bb.0:401; SM70-NEXT: ld.param.b64 %rd1, [generic_unordered_sys_double_param_0];402; SM70-NEXT: ld.relaxed.sys.b64 %rd2, [%rd1];403; SM70-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;404; SM70-NEXT: st.relaxed.sys.b64 [%rd1], %rd3;405; SM70-NEXT: ret;406 %a.load = load atomic double, ptr %a unordered, align 8407 %a.add = fadd double %a.load, 1.408 store atomic double %a.add, ptr %a unordered, align 8409 ret void410}411 412; generic_unordered_volatile_sys413 414define void @generic_unordered_volatile_sys_i8(ptr %a) {415; CHECK-LABEL: generic_unordered_volatile_sys_i8(416; CHECK: {417; CHECK-NEXT: .reg .b16 %rs<3>;418; CHECK-NEXT: .reg .b64 %rd<2>;419; CHECK-EMPTY:420; CHECK-NEXT: // %bb.0:421; CHECK-NEXT: ld.param.b64 %rd1, [generic_unordered_volatile_sys_i8_param_0];422; CHECK-NEXT: ld.volatile.b8 %rs1, [%rd1];423; CHECK-NEXT: add.s16 %rs2, %rs1, 1;424; CHECK-NEXT: st.volatile.b8 [%rd1], %rs2;425; CHECK-NEXT: ret;426 %a.load = load atomic volatile i8, ptr %a unordered, align 1427 %a.add = add i8 %a.load, 1428 store atomic volatile i8 %a.add, ptr %a unordered, align 1429 ret void430}431 432define void @generic_unordered_volatile_sys_i16(ptr %a) {433; CHECK-LABEL: generic_unordered_volatile_sys_i16(434; CHECK: {435; CHECK-NEXT: .reg .b16 %rs<3>;436; CHECK-NEXT: .reg .b64 %rd<2>;437; CHECK-EMPTY:438; CHECK-NEXT: // %bb.0:439; CHECK-NEXT: ld.param.b64 %rd1, [generic_unordered_volatile_sys_i16_param_0];440; CHECK-NEXT: ld.volatile.b16 %rs1, [%rd1];441; CHECK-NEXT: add.s16 %rs2, %rs1, 1;442; CHECK-NEXT: st.volatile.b16 [%rd1], %rs2;443; CHECK-NEXT: ret;444 %a.load = load atomic volatile i16, ptr %a unordered, align 2445 %a.add = add i16 %a.load, 1446 store atomic volatile i16 %a.add, ptr %a unordered, align 2447 ret void448}449 450define void @generic_unordered_volatile_sys_i32(ptr %a) {451; CHECK-LABEL: generic_unordered_volatile_sys_i32(452; CHECK: {453; CHECK-NEXT: .reg .b32 %r<3>;454; CHECK-NEXT: .reg .b64 %rd<2>;455; CHECK-EMPTY:456; CHECK-NEXT: // %bb.0:457; CHECK-NEXT: ld.param.b64 %rd1, [generic_unordered_volatile_sys_i32_param_0];458; CHECK-NEXT: ld.volatile.b32 %r1, [%rd1];459; CHECK-NEXT: add.s32 %r2, %r1, 1;460; CHECK-NEXT: st.volatile.b32 [%rd1], %r2;461; CHECK-NEXT: ret;462 %a.load = load atomic volatile i32, ptr %a unordered, align 4463 %a.add = add i32 %a.load, 1464 store atomic volatile i32 %a.add, ptr %a unordered, align 4465 ret void466}467 468define void @generic_unordered_volatile_sys_i64(ptr %a) {469; CHECK-LABEL: generic_unordered_volatile_sys_i64(470; CHECK: {471; CHECK-NEXT: .reg .b64 %rd<4>;472; CHECK-EMPTY:473; CHECK-NEXT: // %bb.0:474; CHECK-NEXT: ld.param.b64 %rd1, [generic_unordered_volatile_sys_i64_param_0];475; CHECK-NEXT: ld.volatile.b64 %rd2, [%rd1];476; CHECK-NEXT: add.s64 %rd3, %rd2, 1;477; CHECK-NEXT: st.volatile.b64 [%rd1], %rd3;478; CHECK-NEXT: ret;479 %a.load = load atomic volatile i64, ptr %a unordered, align 8480 %a.add = add i64 %a.load, 1481 store atomic volatile i64 %a.add, ptr %a unordered, align 8482 ret void483}484 485define void @generic_unordered_volatile_sys_float(ptr %a) {486; CHECK-LABEL: generic_unordered_volatile_sys_float(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, [generic_unordered_volatile_sys_float_param_0];493; CHECK-NEXT: ld.volatile.b32 %r1, [%rd1];494; CHECK-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;495; CHECK-NEXT: st.volatile.b32 [%rd1], %r2;496; CHECK-NEXT: ret;497 %a.load = load atomic volatile float, ptr %a unordered, align 4498 %a.add = fadd float %a.load, 1.499 store atomic volatile float %a.add, ptr %a unordered, align 4500 ret void501}502 503define void @generic_unordered_volatile_sys_double(ptr %a) {504; CHECK-LABEL: generic_unordered_volatile_sys_double(505; CHECK: {506; CHECK-NEXT: .reg .b64 %rd<4>;507; CHECK-EMPTY:508; CHECK-NEXT: // %bb.0:509; CHECK-NEXT: ld.param.b64 %rd1, [generic_unordered_volatile_sys_double_param_0];510; CHECK-NEXT: ld.volatile.b64 %rd2, [%rd1];511; CHECK-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;512; CHECK-NEXT: st.volatile.b64 [%rd1], %rd3;513; CHECK-NEXT: ret;514 %a.load = load atomic volatile double, ptr %a unordered, align 8515 %a.add = fadd double %a.load, 1.516 store atomic volatile double %a.add, ptr %a unordered, align 8517 ret void518}519 520; generic_monotonic_sys521 522define void @generic_monotonic_sys_i8(ptr %a) {523; SM60-LABEL: generic_monotonic_sys_i8(524; SM60: {525; SM60-NEXT: .reg .b16 %rs<3>;526; SM60-NEXT: .reg .b64 %rd<2>;527; SM60-EMPTY:528; SM60-NEXT: // %bb.0:529; SM60-NEXT: ld.param.b64 %rd1, [generic_monotonic_sys_i8_param_0];530; SM60-NEXT: ld.volatile.b8 %rs1, [%rd1];531; SM60-NEXT: add.s16 %rs2, %rs1, 1;532; SM60-NEXT: st.volatile.b8 [%rd1], %rs2;533; SM60-NEXT: ret;534;535; SM70-LABEL: generic_monotonic_sys_i8(536; SM70: {537; SM70-NEXT: .reg .b16 %rs<3>;538; SM70-NEXT: .reg .b64 %rd<2>;539; SM70-EMPTY:540; SM70-NEXT: // %bb.0:541; SM70-NEXT: ld.param.b64 %rd1, [generic_monotonic_sys_i8_param_0];542; SM70-NEXT: ld.relaxed.sys.b8 %rs1, [%rd1];543; SM70-NEXT: add.s16 %rs2, %rs1, 1;544; SM70-NEXT: st.relaxed.sys.b8 [%rd1], %rs2;545; SM70-NEXT: ret;546 %a.load = load atomic i8, ptr %a monotonic, align 1547 %a.add = add i8 %a.load, 1548 store atomic i8 %a.add, ptr %a monotonic, align 1549 ret void550}551 552define void @generic_monotonic_sys_i16(ptr %a) {553; SM60-LABEL: generic_monotonic_sys_i16(554; SM60: {555; SM60-NEXT: .reg .b16 %rs<3>;556; SM60-NEXT: .reg .b64 %rd<2>;557; SM60-EMPTY:558; SM60-NEXT: // %bb.0:559; SM60-NEXT: ld.param.b64 %rd1, [generic_monotonic_sys_i16_param_0];560; SM60-NEXT: ld.volatile.b16 %rs1, [%rd1];561; SM60-NEXT: add.s16 %rs2, %rs1, 1;562; SM60-NEXT: st.volatile.b16 [%rd1], %rs2;563; SM60-NEXT: ret;564;565; SM70-LABEL: generic_monotonic_sys_i16(566; SM70: {567; SM70-NEXT: .reg .b16 %rs<3>;568; SM70-NEXT: .reg .b64 %rd<2>;569; SM70-EMPTY:570; SM70-NEXT: // %bb.0:571; SM70-NEXT: ld.param.b64 %rd1, [generic_monotonic_sys_i16_param_0];572; SM70-NEXT: ld.relaxed.sys.b16 %rs1, [%rd1];573; SM70-NEXT: add.s16 %rs2, %rs1, 1;574; SM70-NEXT: st.relaxed.sys.b16 [%rd1], %rs2;575; SM70-NEXT: ret;576 %a.load = load atomic i16, ptr %a monotonic, align 2577 %a.add = add i16 %a.load, 1578 store atomic i16 %a.add, ptr %a monotonic, align 2579 ret void580}581 582define void @generic_monotonic_sys_i32(ptr %a) {583; SM60-LABEL: generic_monotonic_sys_i32(584; SM60: {585; SM60-NEXT: .reg .b32 %r<3>;586; SM60-NEXT: .reg .b64 %rd<2>;587; SM60-EMPTY:588; SM60-NEXT: // %bb.0:589; SM60-NEXT: ld.param.b64 %rd1, [generic_monotonic_sys_i32_param_0];590; SM60-NEXT: ld.volatile.b32 %r1, [%rd1];591; SM60-NEXT: add.s32 %r2, %r1, 1;592; SM60-NEXT: st.volatile.b32 [%rd1], %r2;593; SM60-NEXT: ret;594;595; SM70-LABEL: generic_monotonic_sys_i32(596; SM70: {597; SM70-NEXT: .reg .b32 %r<3>;598; SM70-NEXT: .reg .b64 %rd<2>;599; SM70-EMPTY:600; SM70-NEXT: // %bb.0:601; SM70-NEXT: ld.param.b64 %rd1, [generic_monotonic_sys_i32_param_0];602; SM70-NEXT: ld.relaxed.sys.b32 %r1, [%rd1];603; SM70-NEXT: add.s32 %r2, %r1, 1;604; SM70-NEXT: st.relaxed.sys.b32 [%rd1], %r2;605; SM70-NEXT: ret;606 %a.load = load atomic i32, ptr %a monotonic, align 4607 %a.add = add i32 %a.load, 1608 store atomic i32 %a.add, ptr %a monotonic, align 4609 ret void610}611 612define void @generic_monotonic_sys_i64(ptr %a) {613; SM60-LABEL: generic_monotonic_sys_i64(614; SM60: {615; SM60-NEXT: .reg .b64 %rd<4>;616; SM60-EMPTY:617; SM60-NEXT: // %bb.0:618; SM60-NEXT: ld.param.b64 %rd1, [generic_monotonic_sys_i64_param_0];619; SM60-NEXT: ld.volatile.b64 %rd2, [%rd1];620; SM60-NEXT: add.s64 %rd3, %rd2, 1;621; SM60-NEXT: st.volatile.b64 [%rd1], %rd3;622; SM60-NEXT: ret;623;624; SM70-LABEL: generic_monotonic_sys_i64(625; SM70: {626; SM70-NEXT: .reg .b64 %rd<4>;627; SM70-EMPTY:628; SM70-NEXT: // %bb.0:629; SM70-NEXT: ld.param.b64 %rd1, [generic_monotonic_sys_i64_param_0];630; SM70-NEXT: ld.relaxed.sys.b64 %rd2, [%rd1];631; SM70-NEXT: add.s64 %rd3, %rd2, 1;632; SM70-NEXT: st.relaxed.sys.b64 [%rd1], %rd3;633; SM70-NEXT: ret;634 %a.load = load atomic i64, ptr %a monotonic, align 8635 %a.add = add i64 %a.load, 1636 store atomic i64 %a.add, ptr %a monotonic, align 8637 ret void638}639 640define void @generic_monotonic_sys_float(ptr %a) {641; SM60-LABEL: generic_monotonic_sys_float(642; SM60: {643; SM60-NEXT: .reg .b32 %r<3>;644; SM60-NEXT: .reg .b64 %rd<2>;645; SM60-EMPTY:646; SM60-NEXT: // %bb.0:647; SM60-NEXT: ld.param.b64 %rd1, [generic_monotonic_sys_float_param_0];648; SM60-NEXT: ld.volatile.b32 %r1, [%rd1];649; SM60-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;650; SM60-NEXT: st.volatile.b32 [%rd1], %r2;651; SM60-NEXT: ret;652;653; SM70-LABEL: generic_monotonic_sys_float(654; SM70: {655; SM70-NEXT: .reg .b32 %r<3>;656; SM70-NEXT: .reg .b64 %rd<2>;657; SM70-EMPTY:658; SM70-NEXT: // %bb.0:659; SM70-NEXT: ld.param.b64 %rd1, [generic_monotonic_sys_float_param_0];660; SM70-NEXT: ld.relaxed.sys.b32 %r1, [%rd1];661; SM70-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;662; SM70-NEXT: st.relaxed.sys.b32 [%rd1], %r2;663; SM70-NEXT: ret;664 %a.load = load atomic float, ptr %a monotonic, align 4665 %a.add = fadd float %a.load, 1.666 store atomic float %a.add, ptr %a monotonic, align 4667 ret void668}669 670define void @generic_monotonic_sys_double(ptr %a) {671; SM60-LABEL: generic_monotonic_sys_double(672; SM60: {673; SM60-NEXT: .reg .b64 %rd<4>;674; SM60-EMPTY:675; SM60-NEXT: // %bb.0:676; SM60-NEXT: ld.param.b64 %rd1, [generic_monotonic_sys_double_param_0];677; SM60-NEXT: ld.volatile.b64 %rd2, [%rd1];678; SM60-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;679; SM60-NEXT: st.volatile.b64 [%rd1], %rd3;680; SM60-NEXT: ret;681;682; SM70-LABEL: generic_monotonic_sys_double(683; SM70: {684; SM70-NEXT: .reg .b64 %rd<4>;685; SM70-EMPTY:686; SM70-NEXT: // %bb.0:687; SM70-NEXT: ld.param.b64 %rd1, [generic_monotonic_sys_double_param_0];688; SM70-NEXT: ld.relaxed.sys.b64 %rd2, [%rd1];689; SM70-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;690; SM70-NEXT: st.relaxed.sys.b64 [%rd1], %rd3;691; SM70-NEXT: ret;692 %a.load = load atomic double, ptr %a monotonic, align 8693 %a.add = fadd double %a.load, 1.694 store atomic double %a.add, ptr %a monotonic, align 8695 ret void696}697 698; generic_monotonic_volatile_sys699 700define void @generic_monotonic_volatile_sys_i8(ptr %a) {701; CHECK-LABEL: generic_monotonic_volatile_sys_i8(702; CHECK: {703; CHECK-NEXT: .reg .b16 %rs<3>;704; CHECK-NEXT: .reg .b64 %rd<2>;705; CHECK-EMPTY:706; CHECK-NEXT: // %bb.0:707; CHECK-NEXT: ld.param.b64 %rd1, [generic_monotonic_volatile_sys_i8_param_0];708; CHECK-NEXT: ld.volatile.b8 %rs1, [%rd1];709; CHECK-NEXT: add.s16 %rs2, %rs1, 1;710; CHECK-NEXT: st.volatile.b8 [%rd1], %rs2;711; CHECK-NEXT: ret;712 %a.load = load atomic volatile i8, ptr %a monotonic, align 1713 %a.add = add i8 %a.load, 1714 store atomic volatile i8 %a.add, ptr %a monotonic, align 1715 ret void716}717 718define void @generic_monotonic_volatile_sys_i16(ptr %a) {719; CHECK-LABEL: generic_monotonic_volatile_sys_i16(720; CHECK: {721; CHECK-NEXT: .reg .b16 %rs<3>;722; CHECK-NEXT: .reg .b64 %rd<2>;723; CHECK-EMPTY:724; CHECK-NEXT: // %bb.0:725; CHECK-NEXT: ld.param.b64 %rd1, [generic_monotonic_volatile_sys_i16_param_0];726; CHECK-NEXT: ld.volatile.b16 %rs1, [%rd1];727; CHECK-NEXT: add.s16 %rs2, %rs1, 1;728; CHECK-NEXT: st.volatile.b16 [%rd1], %rs2;729; CHECK-NEXT: ret;730 %a.load = load atomic volatile i16, ptr %a monotonic, align 2731 %a.add = add i16 %a.load, 1732 store atomic volatile i16 %a.add, ptr %a monotonic, align 2733 ret void734}735 736define void @generic_monotonic_volatile_sys_i32(ptr %a) {737; CHECK-LABEL: generic_monotonic_volatile_sys_i32(738; CHECK: {739; CHECK-NEXT: .reg .b32 %r<3>;740; CHECK-NEXT: .reg .b64 %rd<2>;741; CHECK-EMPTY:742; CHECK-NEXT: // %bb.0:743; CHECK-NEXT: ld.param.b64 %rd1, [generic_monotonic_volatile_sys_i32_param_0];744; CHECK-NEXT: ld.volatile.b32 %r1, [%rd1];745; CHECK-NEXT: add.s32 %r2, %r1, 1;746; CHECK-NEXT: st.volatile.b32 [%rd1], %r2;747; CHECK-NEXT: ret;748 %a.load = load atomic volatile i32, ptr %a monotonic, align 4749 %a.add = add i32 %a.load, 1750 store atomic volatile i32 %a.add, ptr %a monotonic, align 4751 ret void752}753 754define void @generic_monotonic_volatile_sys_i64(ptr %a) {755; CHECK-LABEL: generic_monotonic_volatile_sys_i64(756; CHECK: {757; CHECK-NEXT: .reg .b64 %rd<4>;758; CHECK-EMPTY:759; CHECK-NEXT: // %bb.0:760; CHECK-NEXT: ld.param.b64 %rd1, [generic_monotonic_volatile_sys_i64_param_0];761; CHECK-NEXT: ld.volatile.b64 %rd2, [%rd1];762; CHECK-NEXT: add.s64 %rd3, %rd2, 1;763; CHECK-NEXT: st.volatile.b64 [%rd1], %rd3;764; CHECK-NEXT: ret;765 %a.load = load atomic volatile i64, ptr %a monotonic, align 8766 %a.add = add i64 %a.load, 1767 store atomic volatile i64 %a.add, ptr %a monotonic, align 8768 ret void769}770 771define void @generic_monotonic_volatile_sys_float(ptr %a) {772; CHECK-LABEL: generic_monotonic_volatile_sys_float(773; CHECK: {774; CHECK-NEXT: .reg .b32 %r<3>;775; CHECK-NEXT: .reg .b64 %rd<2>;776; CHECK-EMPTY:777; CHECK-NEXT: // %bb.0:778; CHECK-NEXT: ld.param.b64 %rd1, [generic_monotonic_volatile_sys_float_param_0];779; CHECK-NEXT: ld.volatile.b32 %r1, [%rd1];780; CHECK-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;781; CHECK-NEXT: st.volatile.b32 [%rd1], %r2;782; CHECK-NEXT: ret;783 %a.load = load atomic volatile float, ptr %a monotonic, align 4784 %a.add = fadd float %a.load, 1.785 store atomic volatile float %a.add, ptr %a monotonic, align 4786 ret void787}788 789define void @generic_monotonic_volatile_sys_double(ptr %a) {790; CHECK-LABEL: generic_monotonic_volatile_sys_double(791; CHECK: {792; CHECK-NEXT: .reg .b64 %rd<4>;793; CHECK-EMPTY:794; CHECK-NEXT: // %bb.0:795; CHECK-NEXT: ld.param.b64 %rd1, [generic_monotonic_volatile_sys_double_param_0];796; CHECK-NEXT: ld.volatile.b64 %rd2, [%rd1];797; CHECK-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;798; CHECK-NEXT: st.volatile.b64 [%rd1], %rd3;799; CHECK-NEXT: ret;800 %a.load = load atomic volatile double, ptr %a monotonic, align 8801 %a.add = fadd double %a.load, 1.802 store atomic volatile double %a.add, ptr %a monotonic, align 8803 ret void804}805 806;; global statespace807 808; global809 810define void @global_i8(ptr addrspace(1) %a) {811; CHECK-LABEL: global_i8(812; CHECK: {813; CHECK-NEXT: .reg .b16 %rs<3>;814; CHECK-NEXT: .reg .b64 %rd<2>;815; CHECK-EMPTY:816; CHECK-NEXT: // %bb.0:817; CHECK-NEXT: ld.param.b64 %rd1, [global_i8_param_0];818; CHECK-NEXT: ld.global.b8 %rs1, [%rd1];819; CHECK-NEXT: add.s16 %rs2, %rs1, 1;820; CHECK-NEXT: st.global.b8 [%rd1], %rs2;821; CHECK-NEXT: ret;822 %a.load = load i8, ptr addrspace(1) %a823 %a.add = add i8 %a.load, 1824 store i8 %a.add, ptr addrspace(1) %a825 ret void826}827 828define void @global_i16(ptr addrspace(1) %a) {829; CHECK-LABEL: global_i16(830; CHECK: {831; CHECK-NEXT: .reg .b16 %rs<3>;832; CHECK-NEXT: .reg .b64 %rd<2>;833; CHECK-EMPTY:834; CHECK-NEXT: // %bb.0:835; CHECK-NEXT: ld.param.b64 %rd1, [global_i16_param_0];836; CHECK-NEXT: ld.global.b16 %rs1, [%rd1];837; CHECK-NEXT: add.s16 %rs2, %rs1, 1;838; CHECK-NEXT: st.global.b16 [%rd1], %rs2;839; CHECK-NEXT: ret;840 %a.load = load i16, ptr addrspace(1) %a841 %a.add = add i16 %a.load, 1842 store i16 %a.add, ptr addrspace(1) %a843 ret void844}845 846define void @global_i32(ptr addrspace(1) %a) {847; CHECK-LABEL: global_i32(848; CHECK: {849; CHECK-NEXT: .reg .b32 %r<3>;850; CHECK-NEXT: .reg .b64 %rd<2>;851; CHECK-EMPTY:852; CHECK-NEXT: // %bb.0:853; CHECK-NEXT: ld.param.b64 %rd1, [global_i32_param_0];854; CHECK-NEXT: ld.global.b32 %r1, [%rd1];855; CHECK-NEXT: add.s32 %r2, %r1, 1;856; CHECK-NEXT: st.global.b32 [%rd1], %r2;857; CHECK-NEXT: ret;858 %a.load = load i32, ptr addrspace(1) %a859 %a.add = add i32 %a.load, 1860 store i32 %a.add, ptr addrspace(1) %a861 ret void862}863 864define void @global_i64(ptr addrspace(1) %a) {865; CHECK-LABEL: global_i64(866; CHECK: {867; CHECK-NEXT: .reg .b64 %rd<4>;868; CHECK-EMPTY:869; CHECK-NEXT: // %bb.0:870; CHECK-NEXT: ld.param.b64 %rd1, [global_i64_param_0];871; CHECK-NEXT: ld.global.b64 %rd2, [%rd1];872; CHECK-NEXT: add.s64 %rd3, %rd2, 1;873; CHECK-NEXT: st.global.b64 [%rd1], %rd3;874; CHECK-NEXT: ret;875 %a.load = load i64, ptr addrspace(1) %a876 %a.add = add i64 %a.load, 1877 store i64 %a.add, ptr addrspace(1) %a878 ret void879}880 881define void @global_float(ptr addrspace(1) %a) {882; CHECK-LABEL: global_float(883; CHECK: {884; CHECK-NEXT: .reg .b32 %r<3>;885; CHECK-NEXT: .reg .b64 %rd<2>;886; CHECK-EMPTY:887; CHECK-NEXT: // %bb.0:888; CHECK-NEXT: ld.param.b64 %rd1, [global_float_param_0];889; CHECK-NEXT: ld.global.b32 %r1, [%rd1];890; CHECK-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;891; CHECK-NEXT: st.global.b32 [%rd1], %r2;892; CHECK-NEXT: ret;893 %a.load = load float, ptr addrspace(1) %a894 %a.add = fadd float %a.load, 1.895 store float %a.add, ptr addrspace(1) %a896 ret void897}898 899define void @global_double(ptr addrspace(1) %a) {900; CHECK-LABEL: global_double(901; CHECK: {902; CHECK-NEXT: .reg .b64 %rd<4>;903; CHECK-EMPTY:904; CHECK-NEXT: // %bb.0:905; CHECK-NEXT: ld.param.b64 %rd1, [global_double_param_0];906; CHECK-NEXT: ld.global.b64 %rd2, [%rd1];907; CHECK-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;908; CHECK-NEXT: st.global.b64 [%rd1], %rd3;909; CHECK-NEXT: ret;910 %a.load = load double, ptr addrspace(1) %a911 %a.add = fadd double %a.load, 1.912 store double %a.add, ptr addrspace(1) %a913 ret void914}915 916; global_volatile917 918define void @global_volatile_i8(ptr addrspace(1) %a) {919; CHECK-LABEL: global_volatile_i8(920; CHECK: {921; CHECK-NEXT: .reg .b16 %rs<3>;922; CHECK-NEXT: .reg .b64 %rd<2>;923; CHECK-EMPTY:924; CHECK-NEXT: // %bb.0:925; CHECK-NEXT: ld.param.b64 %rd1, [global_volatile_i8_param_0];926; CHECK-NEXT: ld.volatile.global.b8 %rs1, [%rd1];927; CHECK-NEXT: add.s16 %rs2, %rs1, 1;928; CHECK-NEXT: st.volatile.global.b8 [%rd1], %rs2;929; CHECK-NEXT: ret;930 %a.load = load volatile i8, ptr addrspace(1) %a931 %a.add = add i8 %a.load, 1932 store volatile i8 %a.add, ptr addrspace(1) %a933 ret void934}935 936define void @global_volatile_i16(ptr addrspace(1) %a) {937; CHECK-LABEL: global_volatile_i16(938; CHECK: {939; CHECK-NEXT: .reg .b16 %rs<3>;940; CHECK-NEXT: .reg .b64 %rd<2>;941; CHECK-EMPTY:942; CHECK-NEXT: // %bb.0:943; CHECK-NEXT: ld.param.b64 %rd1, [global_volatile_i16_param_0];944; CHECK-NEXT: ld.volatile.global.b16 %rs1, [%rd1];945; CHECK-NEXT: add.s16 %rs2, %rs1, 1;946; CHECK-NEXT: st.volatile.global.b16 [%rd1], %rs2;947; CHECK-NEXT: ret;948 %a.load = load volatile i16, ptr addrspace(1) %a949 %a.add = add i16 %a.load, 1950 store volatile i16 %a.add, ptr addrspace(1) %a951 ret void952}953 954define void @global_volatile_i32(ptr addrspace(1) %a) {955; CHECK-LABEL: global_volatile_i32(956; CHECK: {957; CHECK-NEXT: .reg .b32 %r<3>;958; CHECK-NEXT: .reg .b64 %rd<2>;959; CHECK-EMPTY:960; CHECK-NEXT: // %bb.0:961; CHECK-NEXT: ld.param.b64 %rd1, [global_volatile_i32_param_0];962; CHECK-NEXT: ld.volatile.global.b32 %r1, [%rd1];963; CHECK-NEXT: add.s32 %r2, %r1, 1;964; CHECK-NEXT: st.volatile.global.b32 [%rd1], %r2;965; CHECK-NEXT: ret;966 %a.load = load volatile i32, ptr addrspace(1) %a967 %a.add = add i32 %a.load, 1968 store volatile i32 %a.add, ptr addrspace(1) %a969 ret void970}971 972define void @global_volatile_i64(ptr addrspace(1) %a) {973; CHECK-LABEL: global_volatile_i64(974; CHECK: {975; CHECK-NEXT: .reg .b64 %rd<4>;976; CHECK-EMPTY:977; CHECK-NEXT: // %bb.0:978; CHECK-NEXT: ld.param.b64 %rd1, [global_volatile_i64_param_0];979; CHECK-NEXT: ld.volatile.global.b64 %rd2, [%rd1];980; CHECK-NEXT: add.s64 %rd3, %rd2, 1;981; CHECK-NEXT: st.volatile.global.b64 [%rd1], %rd3;982; CHECK-NEXT: ret;983 %a.load = load volatile i64, ptr addrspace(1) %a984 %a.add = add i64 %a.load, 1985 store volatile i64 %a.add, ptr addrspace(1) %a986 ret void987}988 989define void @global_volatile_float(ptr addrspace(1) %a) {990; CHECK-LABEL: global_volatile_float(991; CHECK: {992; CHECK-NEXT: .reg .b32 %r<3>;993; CHECK-NEXT: .reg .b64 %rd<2>;994; CHECK-EMPTY:995; CHECK-NEXT: // %bb.0:996; CHECK-NEXT: ld.param.b64 %rd1, [global_volatile_float_param_0];997; CHECK-NEXT: ld.volatile.global.b32 %r1, [%rd1];998; CHECK-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;999; CHECK-NEXT: st.volatile.global.b32 [%rd1], %r2;1000; CHECK-NEXT: ret;1001 %a.load = load volatile float, ptr addrspace(1) %a1002 %a.add = fadd float %a.load, 1.1003 store volatile float %a.add, ptr addrspace(1) %a1004 ret void1005}1006 1007define void @global_volatile_double(ptr addrspace(1) %a) {1008; CHECK-LABEL: global_volatile_double(1009; CHECK: {1010; CHECK-NEXT: .reg .b64 %rd<4>;1011; CHECK-EMPTY:1012; CHECK-NEXT: // %bb.0:1013; CHECK-NEXT: ld.param.b64 %rd1, [global_volatile_double_param_0];1014; CHECK-NEXT: ld.volatile.global.b64 %rd2, [%rd1];1015; CHECK-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;1016; CHECK-NEXT: st.volatile.global.b64 [%rd1], %rd3;1017; CHECK-NEXT: ret;1018 %a.load = load volatile double, ptr addrspace(1) %a1019 %a.add = fadd double %a.load, 1.1020 store volatile double %a.add, ptr addrspace(1) %a1021 ret void1022}1023 1024; global_unordered_sys1025 1026define void @global_unordered_sys_i8(ptr addrspace(1) %a) {1027; SM60-LABEL: global_unordered_sys_i8(1028; SM60: {1029; SM60-NEXT: .reg .b16 %rs<3>;1030; SM60-NEXT: .reg .b64 %rd<2>;1031; SM60-EMPTY:1032; SM60-NEXT: // %bb.0:1033; SM60-NEXT: ld.param.b64 %rd1, [global_unordered_sys_i8_param_0];1034; SM60-NEXT: ld.volatile.global.b8 %rs1, [%rd1];1035; SM60-NEXT: add.s16 %rs2, %rs1, 1;1036; SM60-NEXT: st.volatile.global.b8 [%rd1], %rs2;1037; SM60-NEXT: ret;1038;1039; SM70-LABEL: global_unordered_sys_i8(1040; SM70: {1041; SM70-NEXT: .reg .b16 %rs<3>;1042; SM70-NEXT: .reg .b64 %rd<2>;1043; SM70-EMPTY:1044; SM70-NEXT: // %bb.0:1045; SM70-NEXT: ld.param.b64 %rd1, [global_unordered_sys_i8_param_0];1046; SM70-NEXT: ld.relaxed.sys.global.b8 %rs1, [%rd1];1047; SM70-NEXT: add.s16 %rs2, %rs1, 1;1048; SM70-NEXT: st.relaxed.sys.global.b8 [%rd1], %rs2;1049; SM70-NEXT: ret;1050 %a.load = load atomic i8, ptr addrspace(1) %a unordered, align 11051 %a.add = add i8 %a.load, 11052 store atomic i8 %a.add, ptr addrspace(1) %a unordered, align 11053 ret void1054}1055 1056define void @global_unordered_sys_i16(ptr addrspace(1) %a) {1057; SM60-LABEL: global_unordered_sys_i16(1058; SM60: {1059; SM60-NEXT: .reg .b16 %rs<3>;1060; SM60-NEXT: .reg .b64 %rd<2>;1061; SM60-EMPTY:1062; SM60-NEXT: // %bb.0:1063; SM60-NEXT: ld.param.b64 %rd1, [global_unordered_sys_i16_param_0];1064; SM60-NEXT: ld.volatile.global.b16 %rs1, [%rd1];1065; SM60-NEXT: add.s16 %rs2, %rs1, 1;1066; SM60-NEXT: st.volatile.global.b16 [%rd1], %rs2;1067; SM60-NEXT: ret;1068;1069; SM70-LABEL: global_unordered_sys_i16(1070; SM70: {1071; SM70-NEXT: .reg .b16 %rs<3>;1072; SM70-NEXT: .reg .b64 %rd<2>;1073; SM70-EMPTY:1074; SM70-NEXT: // %bb.0:1075; SM70-NEXT: ld.param.b64 %rd1, [global_unordered_sys_i16_param_0];1076; SM70-NEXT: ld.relaxed.sys.global.b16 %rs1, [%rd1];1077; SM70-NEXT: add.s16 %rs2, %rs1, 1;1078; SM70-NEXT: st.relaxed.sys.global.b16 [%rd1], %rs2;1079; SM70-NEXT: ret;1080 %a.load = load atomic i16, ptr addrspace(1) %a unordered, align 21081 %a.add = add i16 %a.load, 11082 store atomic i16 %a.add, ptr addrspace(1) %a unordered, align 21083 ret void1084}1085 1086define void @global_unordered_sys_i32(ptr addrspace(1) %a) {1087; SM60-LABEL: global_unordered_sys_i32(1088; SM60: {1089; SM60-NEXT: .reg .b32 %r<3>;1090; SM60-NEXT: .reg .b64 %rd<2>;1091; SM60-EMPTY:1092; SM60-NEXT: // %bb.0:1093; SM60-NEXT: ld.param.b64 %rd1, [global_unordered_sys_i32_param_0];1094; SM60-NEXT: ld.volatile.global.b32 %r1, [%rd1];1095; SM60-NEXT: add.s32 %r2, %r1, 1;1096; SM60-NEXT: st.volatile.global.b32 [%rd1], %r2;1097; SM60-NEXT: ret;1098;1099; SM70-LABEL: global_unordered_sys_i32(1100; SM70: {1101; SM70-NEXT: .reg .b32 %r<3>;1102; SM70-NEXT: .reg .b64 %rd<2>;1103; SM70-EMPTY:1104; SM70-NEXT: // %bb.0:1105; SM70-NEXT: ld.param.b64 %rd1, [global_unordered_sys_i32_param_0];1106; SM70-NEXT: ld.relaxed.sys.global.b32 %r1, [%rd1];1107; SM70-NEXT: add.s32 %r2, %r1, 1;1108; SM70-NEXT: st.relaxed.sys.global.b32 [%rd1], %r2;1109; SM70-NEXT: ret;1110 %a.load = load atomic i32, ptr addrspace(1) %a unordered, align 41111 %a.add = add i32 %a.load, 11112 store atomic i32 %a.add, ptr addrspace(1) %a unordered, align 41113 ret void1114}1115 1116define void @global_unordered_sys_i64(ptr addrspace(1) %a) {1117; SM60-LABEL: global_unordered_sys_i64(1118; SM60: {1119; SM60-NEXT: .reg .b64 %rd<4>;1120; SM60-EMPTY:1121; SM60-NEXT: // %bb.0:1122; SM60-NEXT: ld.param.b64 %rd1, [global_unordered_sys_i64_param_0];1123; SM60-NEXT: ld.volatile.global.b64 %rd2, [%rd1];1124; SM60-NEXT: add.s64 %rd3, %rd2, 1;1125; SM60-NEXT: st.volatile.global.b64 [%rd1], %rd3;1126; SM60-NEXT: ret;1127;1128; SM70-LABEL: global_unordered_sys_i64(1129; SM70: {1130; SM70-NEXT: .reg .b64 %rd<4>;1131; SM70-EMPTY:1132; SM70-NEXT: // %bb.0:1133; SM70-NEXT: ld.param.b64 %rd1, [global_unordered_sys_i64_param_0];1134; SM70-NEXT: ld.relaxed.sys.global.b64 %rd2, [%rd1];1135; SM70-NEXT: add.s64 %rd3, %rd2, 1;1136; SM70-NEXT: st.relaxed.sys.global.b64 [%rd1], %rd3;1137; SM70-NEXT: ret;1138 %a.load = load atomic i64, ptr addrspace(1) %a unordered, align 81139 %a.add = add i64 %a.load, 11140 store atomic i64 %a.add, ptr addrspace(1) %a unordered, align 81141 ret void1142}1143 1144define void @global_unordered_sys_float(ptr addrspace(1) %a) {1145; SM60-LABEL: global_unordered_sys_float(1146; SM60: {1147; SM60-NEXT: .reg .b32 %r<3>;1148; SM60-NEXT: .reg .b64 %rd<2>;1149; SM60-EMPTY:1150; SM60-NEXT: // %bb.0:1151; SM60-NEXT: ld.param.b64 %rd1, [global_unordered_sys_float_param_0];1152; SM60-NEXT: ld.volatile.global.b32 %r1, [%rd1];1153; SM60-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;1154; SM60-NEXT: st.volatile.global.b32 [%rd1], %r2;1155; SM60-NEXT: ret;1156;1157; SM70-LABEL: global_unordered_sys_float(1158; SM70: {1159; SM70-NEXT: .reg .b32 %r<3>;1160; SM70-NEXT: .reg .b64 %rd<2>;1161; SM70-EMPTY:1162; SM70-NEXT: // %bb.0:1163; SM70-NEXT: ld.param.b64 %rd1, [global_unordered_sys_float_param_0];1164; SM70-NEXT: ld.relaxed.sys.global.b32 %r1, [%rd1];1165; SM70-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;1166; SM70-NEXT: st.relaxed.sys.global.b32 [%rd1], %r2;1167; SM70-NEXT: ret;1168 %a.load = load atomic float, ptr addrspace(1) %a unordered, align 41169 %a.add = fadd float %a.load, 1.1170 store atomic float %a.add, ptr addrspace(1) %a unordered, align 41171 ret void1172}1173 1174define void @global_unordered_sys_double(ptr addrspace(1) %a) {1175; SM60-LABEL: global_unordered_sys_double(1176; SM60: {1177; SM60-NEXT: .reg .b64 %rd<4>;1178; SM60-EMPTY:1179; SM60-NEXT: // %bb.0:1180; SM60-NEXT: ld.param.b64 %rd1, [global_unordered_sys_double_param_0];1181; SM60-NEXT: ld.volatile.global.b64 %rd2, [%rd1];1182; SM60-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;1183; SM60-NEXT: st.volatile.global.b64 [%rd1], %rd3;1184; SM60-NEXT: ret;1185;1186; SM70-LABEL: global_unordered_sys_double(1187; SM70: {1188; SM70-NEXT: .reg .b64 %rd<4>;1189; SM70-EMPTY:1190; SM70-NEXT: // %bb.0:1191; SM70-NEXT: ld.param.b64 %rd1, [global_unordered_sys_double_param_0];1192; SM70-NEXT: ld.relaxed.sys.global.b64 %rd2, [%rd1];1193; SM70-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;1194; SM70-NEXT: st.relaxed.sys.global.b64 [%rd1], %rd3;1195; SM70-NEXT: ret;1196 %a.load = load atomic double, ptr addrspace(1) %a unordered, align 81197 %a.add = fadd double %a.load, 1.1198 store atomic double %a.add, ptr addrspace(1) %a unordered, align 81199 ret void1200}1201 1202; global_unordered_volatile_sys1203 1204define void @global_unordered_volatile_sys_i8(ptr addrspace(1) %a) {1205; SM60-LABEL: global_unordered_volatile_sys_i8(1206; SM60: {1207; SM60-NEXT: .reg .b16 %rs<3>;1208; SM60-NEXT: .reg .b64 %rd<2>;1209; SM60-EMPTY:1210; SM60-NEXT: // %bb.0:1211; SM60-NEXT: ld.param.b64 %rd1, [global_unordered_volatile_sys_i8_param_0];1212; SM60-NEXT: ld.volatile.global.b8 %rs1, [%rd1];1213; SM60-NEXT: add.s16 %rs2, %rs1, 1;1214; SM60-NEXT: st.volatile.global.b8 [%rd1], %rs2;1215; SM60-NEXT: ret;1216;1217; SM70-LABEL: global_unordered_volatile_sys_i8(1218; SM70: {1219; SM70-NEXT: .reg .b16 %rs<3>;1220; SM70-NEXT: .reg .b64 %rd<2>;1221; SM70-EMPTY:1222; SM70-NEXT: // %bb.0:1223; SM70-NEXT: ld.param.b64 %rd1, [global_unordered_volatile_sys_i8_param_0];1224; SM70-NEXT: ld.mmio.relaxed.sys.global.b8 %rs1, [%rd1];1225; SM70-NEXT: add.s16 %rs2, %rs1, 1;1226; SM70-NEXT: st.mmio.relaxed.sys.global.b8 [%rd1], %rs2;1227; SM70-NEXT: ret;1228 %a.load = load atomic volatile i8, ptr addrspace(1) %a unordered, align 11229 %a.add = add i8 %a.load, 11230 store atomic volatile i8 %a.add, ptr addrspace(1) %a unordered, align 11231 ret void1232}1233 1234define void @global_unordered_volatile_sys_i16(ptr addrspace(1) %a) {1235; SM60-LABEL: global_unordered_volatile_sys_i16(1236; SM60: {1237; SM60-NEXT: .reg .b16 %rs<3>;1238; SM60-NEXT: .reg .b64 %rd<2>;1239; SM60-EMPTY:1240; SM60-NEXT: // %bb.0:1241; SM60-NEXT: ld.param.b64 %rd1, [global_unordered_volatile_sys_i16_param_0];1242; SM60-NEXT: ld.volatile.global.b16 %rs1, [%rd1];1243; SM60-NEXT: add.s16 %rs2, %rs1, 1;1244; SM60-NEXT: st.volatile.global.b16 [%rd1], %rs2;1245; SM60-NEXT: ret;1246;1247; SM70-LABEL: global_unordered_volatile_sys_i16(1248; SM70: {1249; SM70-NEXT: .reg .b16 %rs<3>;1250; SM70-NEXT: .reg .b64 %rd<2>;1251; SM70-EMPTY:1252; SM70-NEXT: // %bb.0:1253; SM70-NEXT: ld.param.b64 %rd1, [global_unordered_volatile_sys_i16_param_0];1254; SM70-NEXT: ld.mmio.relaxed.sys.global.b16 %rs1, [%rd1];1255; SM70-NEXT: add.s16 %rs2, %rs1, 1;1256; SM70-NEXT: st.mmio.relaxed.sys.global.b16 [%rd1], %rs2;1257; SM70-NEXT: ret;1258 %a.load = load atomic volatile i16, ptr addrspace(1) %a unordered, align 21259 %a.add = add i16 %a.load, 11260 store atomic volatile i16 %a.add, ptr addrspace(1) %a unordered, align 21261 ret void1262}1263 1264define void @global_unordered_volatile_sys_i32(ptr addrspace(1) %a) {1265; SM60-LABEL: global_unordered_volatile_sys_i32(1266; SM60: {1267; SM60-NEXT: .reg .b32 %r<3>;1268; SM60-NEXT: .reg .b64 %rd<2>;1269; SM60-EMPTY:1270; SM60-NEXT: // %bb.0:1271; SM60-NEXT: ld.param.b64 %rd1, [global_unordered_volatile_sys_i32_param_0];1272; SM60-NEXT: ld.volatile.global.b32 %r1, [%rd1];1273; SM60-NEXT: add.s32 %r2, %r1, 1;1274; SM60-NEXT: st.volatile.global.b32 [%rd1], %r2;1275; SM60-NEXT: ret;1276;1277; SM70-LABEL: global_unordered_volatile_sys_i32(1278; SM70: {1279; SM70-NEXT: .reg .b32 %r<3>;1280; SM70-NEXT: .reg .b64 %rd<2>;1281; SM70-EMPTY:1282; SM70-NEXT: // %bb.0:1283; SM70-NEXT: ld.param.b64 %rd1, [global_unordered_volatile_sys_i32_param_0];1284; SM70-NEXT: ld.mmio.relaxed.sys.global.b32 %r1, [%rd1];1285; SM70-NEXT: add.s32 %r2, %r1, 1;1286; SM70-NEXT: st.mmio.relaxed.sys.global.b32 [%rd1], %r2;1287; SM70-NEXT: ret;1288 %a.load = load atomic volatile i32, ptr addrspace(1) %a unordered, align 41289 %a.add = add i32 %a.load, 11290 store atomic volatile i32 %a.add, ptr addrspace(1) %a unordered, align 41291 ret void1292}1293 1294define void @global_unordered_volatile_sys_i64(ptr addrspace(1) %a) {1295; SM60-LABEL: global_unordered_volatile_sys_i64(1296; SM60: {1297; SM60-NEXT: .reg .b64 %rd<4>;1298; SM60-EMPTY:1299; SM60-NEXT: // %bb.0:1300; SM60-NEXT: ld.param.b64 %rd1, [global_unordered_volatile_sys_i64_param_0];1301; SM60-NEXT: ld.volatile.global.b64 %rd2, [%rd1];1302; SM60-NEXT: add.s64 %rd3, %rd2, 1;1303; SM60-NEXT: st.volatile.global.b64 [%rd1], %rd3;1304; SM60-NEXT: ret;1305;1306; SM70-LABEL: global_unordered_volatile_sys_i64(1307; SM70: {1308; SM70-NEXT: .reg .b64 %rd<4>;1309; SM70-EMPTY:1310; SM70-NEXT: // %bb.0:1311; SM70-NEXT: ld.param.b64 %rd1, [global_unordered_volatile_sys_i64_param_0];1312; SM70-NEXT: ld.mmio.relaxed.sys.global.b64 %rd2, [%rd1];1313; SM70-NEXT: add.s64 %rd3, %rd2, 1;1314; SM70-NEXT: st.mmio.relaxed.sys.global.b64 [%rd1], %rd3;1315; SM70-NEXT: ret;1316 %a.load = load atomic volatile i64, ptr addrspace(1) %a unordered, align 81317 %a.add = add i64 %a.load, 11318 store atomic volatile i64 %a.add, ptr addrspace(1) %a unordered, align 81319 ret void1320}1321 1322define void @global_unordered_volatile_sys_float(ptr addrspace(1) %a) {1323; SM60-LABEL: global_unordered_volatile_sys_float(1324; SM60: {1325; SM60-NEXT: .reg .b32 %r<3>;1326; SM60-NEXT: .reg .b64 %rd<2>;1327; SM60-EMPTY:1328; SM60-NEXT: // %bb.0:1329; SM60-NEXT: ld.param.b64 %rd1, [global_unordered_volatile_sys_float_param_0];1330; SM60-NEXT: ld.volatile.global.b32 %r1, [%rd1];1331; SM60-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;1332; SM60-NEXT: st.volatile.global.b32 [%rd1], %r2;1333; SM60-NEXT: ret;1334;1335; SM70-LABEL: global_unordered_volatile_sys_float(1336; SM70: {1337; SM70-NEXT: .reg .b32 %r<3>;1338; SM70-NEXT: .reg .b64 %rd<2>;1339; SM70-EMPTY:1340; SM70-NEXT: // %bb.0:1341; SM70-NEXT: ld.param.b64 %rd1, [global_unordered_volatile_sys_float_param_0];1342; SM70-NEXT: ld.mmio.relaxed.sys.global.b32 %r1, [%rd1];1343; SM70-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;1344; SM70-NEXT: st.mmio.relaxed.sys.global.b32 [%rd1], %r2;1345; SM70-NEXT: ret;1346 %a.load = load atomic volatile float, ptr addrspace(1) %a unordered, align 41347 %a.add = fadd float %a.load, 1.1348 store atomic volatile float %a.add, ptr addrspace(1) %a unordered, align 41349 ret void1350}1351 1352define void @global_unordered_volatile_sys_double(ptr addrspace(1) %a) {1353; SM60-LABEL: global_unordered_volatile_sys_double(1354; SM60: {1355; SM60-NEXT: .reg .b64 %rd<4>;1356; SM60-EMPTY:1357; SM60-NEXT: // %bb.0:1358; SM60-NEXT: ld.param.b64 %rd1, [global_unordered_volatile_sys_double_param_0];1359; SM60-NEXT: ld.volatile.global.b64 %rd2, [%rd1];1360; SM60-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;1361; SM60-NEXT: st.volatile.global.b64 [%rd1], %rd3;1362; SM60-NEXT: ret;1363;1364; SM70-LABEL: global_unordered_volatile_sys_double(1365; SM70: {1366; SM70-NEXT: .reg .b64 %rd<4>;1367; SM70-EMPTY:1368; SM70-NEXT: // %bb.0:1369; SM70-NEXT: ld.param.b64 %rd1, [global_unordered_volatile_sys_double_param_0];1370; SM70-NEXT: ld.mmio.relaxed.sys.global.b64 %rd2, [%rd1];1371; SM70-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;1372; SM70-NEXT: st.mmio.relaxed.sys.global.b64 [%rd1], %rd3;1373; SM70-NEXT: ret;1374 %a.load = load atomic volatile double, ptr addrspace(1) %a unordered, align 81375 %a.add = fadd double %a.load, 1.1376 store atomic volatile double %a.add, ptr addrspace(1) %a unordered, align 81377 ret void1378}1379 1380; global_monotonic_sys1381 1382define void @global_monotonic_sys_i8(ptr addrspace(1) %a) {1383; SM60-LABEL: global_monotonic_sys_i8(1384; SM60: {1385; SM60-NEXT: .reg .b16 %rs<3>;1386; SM60-NEXT: .reg .b64 %rd<2>;1387; SM60-EMPTY:1388; SM60-NEXT: // %bb.0:1389; SM60-NEXT: ld.param.b64 %rd1, [global_monotonic_sys_i8_param_0];1390; SM60-NEXT: ld.volatile.global.b8 %rs1, [%rd1];1391; SM60-NEXT: add.s16 %rs2, %rs1, 1;1392; SM60-NEXT: st.volatile.global.b8 [%rd1], %rs2;1393; SM60-NEXT: ret;1394;1395; SM70-LABEL: global_monotonic_sys_i8(1396; SM70: {1397; SM70-NEXT: .reg .b16 %rs<3>;1398; SM70-NEXT: .reg .b64 %rd<2>;1399; SM70-EMPTY:1400; SM70-NEXT: // %bb.0:1401; SM70-NEXT: ld.param.b64 %rd1, [global_monotonic_sys_i8_param_0];1402; SM70-NEXT: ld.relaxed.sys.global.b8 %rs1, [%rd1];1403; SM70-NEXT: add.s16 %rs2, %rs1, 1;1404; SM70-NEXT: st.relaxed.sys.global.b8 [%rd1], %rs2;1405; SM70-NEXT: ret;1406 %a.load = load atomic i8, ptr addrspace(1) %a monotonic, align 11407 %a.add = add i8 %a.load, 11408 store atomic i8 %a.add, ptr addrspace(1) %a monotonic, align 11409 ret void1410}1411 1412define void @global_monotonic_sys_i16(ptr addrspace(1) %a) {1413; SM60-LABEL: global_monotonic_sys_i16(1414; SM60: {1415; SM60-NEXT: .reg .b16 %rs<3>;1416; SM60-NEXT: .reg .b64 %rd<2>;1417; SM60-EMPTY:1418; SM60-NEXT: // %bb.0:1419; SM60-NEXT: ld.param.b64 %rd1, [global_monotonic_sys_i16_param_0];1420; SM60-NEXT: ld.volatile.global.b16 %rs1, [%rd1];1421; SM60-NEXT: add.s16 %rs2, %rs1, 1;1422; SM60-NEXT: st.volatile.global.b16 [%rd1], %rs2;1423; SM60-NEXT: ret;1424;1425; SM70-LABEL: global_monotonic_sys_i16(1426; SM70: {1427; SM70-NEXT: .reg .b16 %rs<3>;1428; SM70-NEXT: .reg .b64 %rd<2>;1429; SM70-EMPTY:1430; SM70-NEXT: // %bb.0:1431; SM70-NEXT: ld.param.b64 %rd1, [global_monotonic_sys_i16_param_0];1432; SM70-NEXT: ld.relaxed.sys.global.b16 %rs1, [%rd1];1433; SM70-NEXT: add.s16 %rs2, %rs1, 1;1434; SM70-NEXT: st.relaxed.sys.global.b16 [%rd1], %rs2;1435; SM70-NEXT: ret;1436 %a.load = load atomic i16, ptr addrspace(1) %a monotonic, align 21437 %a.add = add i16 %a.load, 11438 store atomic i16 %a.add, ptr addrspace(1) %a monotonic, align 21439 ret void1440}1441 1442define void @global_monotonic_sys_i32(ptr addrspace(1) %a) {1443; SM60-LABEL: global_monotonic_sys_i32(1444; SM60: {1445; SM60-NEXT: .reg .b32 %r<3>;1446; SM60-NEXT: .reg .b64 %rd<2>;1447; SM60-EMPTY:1448; SM60-NEXT: // %bb.0:1449; SM60-NEXT: ld.param.b64 %rd1, [global_monotonic_sys_i32_param_0];1450; SM60-NEXT: ld.volatile.global.b32 %r1, [%rd1];1451; SM60-NEXT: add.s32 %r2, %r1, 1;1452; SM60-NEXT: st.volatile.global.b32 [%rd1], %r2;1453; SM60-NEXT: ret;1454;1455; SM70-LABEL: global_monotonic_sys_i32(1456; SM70: {1457; SM70-NEXT: .reg .b32 %r<3>;1458; SM70-NEXT: .reg .b64 %rd<2>;1459; SM70-EMPTY:1460; SM70-NEXT: // %bb.0:1461; SM70-NEXT: ld.param.b64 %rd1, [global_monotonic_sys_i32_param_0];1462; SM70-NEXT: ld.relaxed.sys.global.b32 %r1, [%rd1];1463; SM70-NEXT: add.s32 %r2, %r1, 1;1464; SM70-NEXT: st.relaxed.sys.global.b32 [%rd1], %r2;1465; SM70-NEXT: ret;1466 %a.load = load atomic i32, ptr addrspace(1) %a monotonic, align 41467 %a.add = add i32 %a.load, 11468 store atomic i32 %a.add, ptr addrspace(1) %a monotonic, align 41469 ret void1470}1471 1472define void @global_monotonic_sys_i64(ptr addrspace(1) %a) {1473; SM60-LABEL: global_monotonic_sys_i64(1474; SM60: {1475; SM60-NEXT: .reg .b64 %rd<4>;1476; SM60-EMPTY:1477; SM60-NEXT: // %bb.0:1478; SM60-NEXT: ld.param.b64 %rd1, [global_monotonic_sys_i64_param_0];1479; SM60-NEXT: ld.volatile.global.b64 %rd2, [%rd1];1480; SM60-NEXT: add.s64 %rd3, %rd2, 1;1481; SM60-NEXT: st.volatile.global.b64 [%rd1], %rd3;1482; SM60-NEXT: ret;1483;1484; SM70-LABEL: global_monotonic_sys_i64(1485; SM70: {1486; SM70-NEXT: .reg .b64 %rd<4>;1487; SM70-EMPTY:1488; SM70-NEXT: // %bb.0:1489; SM70-NEXT: ld.param.b64 %rd1, [global_monotonic_sys_i64_param_0];1490; SM70-NEXT: ld.relaxed.sys.global.b64 %rd2, [%rd1];1491; SM70-NEXT: add.s64 %rd3, %rd2, 1;1492; SM70-NEXT: st.relaxed.sys.global.b64 [%rd1], %rd3;1493; SM70-NEXT: ret;1494 %a.load = load atomic i64, ptr addrspace(1) %a monotonic, align 81495 %a.add = add i64 %a.load, 11496 store atomic i64 %a.add, ptr addrspace(1) %a monotonic, align 81497 ret void1498}1499 1500define void @global_monotonic_sys_float(ptr addrspace(1) %a) {1501; SM60-LABEL: global_monotonic_sys_float(1502; SM60: {1503; SM60-NEXT: .reg .b32 %r<3>;1504; SM60-NEXT: .reg .b64 %rd<2>;1505; SM60-EMPTY:1506; SM60-NEXT: // %bb.0:1507; SM60-NEXT: ld.param.b64 %rd1, [global_monotonic_sys_float_param_0];1508; SM60-NEXT: ld.volatile.global.b32 %r1, [%rd1];1509; SM60-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;1510; SM60-NEXT: st.volatile.global.b32 [%rd1], %r2;1511; SM60-NEXT: ret;1512;1513; SM70-LABEL: global_monotonic_sys_float(1514; SM70: {1515; SM70-NEXT: .reg .b32 %r<3>;1516; SM70-NEXT: .reg .b64 %rd<2>;1517; SM70-EMPTY:1518; SM70-NEXT: // %bb.0:1519; SM70-NEXT: ld.param.b64 %rd1, [global_monotonic_sys_float_param_0];1520; SM70-NEXT: ld.relaxed.sys.global.b32 %r1, [%rd1];1521; SM70-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;1522; SM70-NEXT: st.relaxed.sys.global.b32 [%rd1], %r2;1523; SM70-NEXT: ret;1524 %a.load = load atomic float, ptr addrspace(1) %a monotonic, align 41525 %a.add = fadd float %a.load, 1.1526 store atomic float %a.add, ptr addrspace(1) %a monotonic, align 41527 ret void1528}1529 1530define void @global_monotonic_sys_double(ptr addrspace(1) %a) {1531; SM60-LABEL: global_monotonic_sys_double(1532; SM60: {1533; SM60-NEXT: .reg .b64 %rd<4>;1534; SM60-EMPTY:1535; SM60-NEXT: // %bb.0:1536; SM60-NEXT: ld.param.b64 %rd1, [global_monotonic_sys_double_param_0];1537; SM60-NEXT: ld.volatile.global.b64 %rd2, [%rd1];1538; SM60-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;1539; SM60-NEXT: st.volatile.global.b64 [%rd1], %rd3;1540; SM60-NEXT: ret;1541;1542; SM70-LABEL: global_monotonic_sys_double(1543; SM70: {1544; SM70-NEXT: .reg .b64 %rd<4>;1545; SM70-EMPTY:1546; SM70-NEXT: // %bb.0:1547; SM70-NEXT: ld.param.b64 %rd1, [global_monotonic_sys_double_param_0];1548; SM70-NEXT: ld.relaxed.sys.global.b64 %rd2, [%rd1];1549; SM70-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;1550; SM70-NEXT: st.relaxed.sys.global.b64 [%rd1], %rd3;1551; SM70-NEXT: ret;1552 %a.load = load atomic double, ptr addrspace(1) %a monotonic, align 81553 %a.add = fadd double %a.load, 1.1554 store atomic double %a.add, ptr addrspace(1) %a monotonic, align 81555 ret void1556}1557 1558; global_monotonic_volatile_sys1559 1560define void @global_monotonic_volatile_sys_i8(ptr addrspace(1) %a) {1561; SM60-LABEL: global_monotonic_volatile_sys_i8(1562; SM60: {1563; SM60-NEXT: .reg .b16 %rs<3>;1564; SM60-NEXT: .reg .b64 %rd<2>;1565; SM60-EMPTY:1566; SM60-NEXT: // %bb.0:1567; SM60-NEXT: ld.param.b64 %rd1, [global_monotonic_volatile_sys_i8_param_0];1568; SM60-NEXT: ld.volatile.global.b8 %rs1, [%rd1];1569; SM60-NEXT: add.s16 %rs2, %rs1, 1;1570; SM60-NEXT: st.volatile.global.b8 [%rd1], %rs2;1571; SM60-NEXT: ret;1572;1573; SM70-LABEL: global_monotonic_volatile_sys_i8(1574; SM70: {1575; SM70-NEXT: .reg .b16 %rs<3>;1576; SM70-NEXT: .reg .b64 %rd<2>;1577; SM70-EMPTY:1578; SM70-NEXT: // %bb.0:1579; SM70-NEXT: ld.param.b64 %rd1, [global_monotonic_volatile_sys_i8_param_0];1580; SM70-NEXT: ld.mmio.relaxed.sys.global.b8 %rs1, [%rd1];1581; SM70-NEXT: add.s16 %rs2, %rs1, 1;1582; SM70-NEXT: st.mmio.relaxed.sys.global.b8 [%rd1], %rs2;1583; SM70-NEXT: ret;1584 %a.load = load atomic volatile i8, ptr addrspace(1) %a monotonic, align 11585 %a.add = add i8 %a.load, 11586 store atomic volatile i8 %a.add, ptr addrspace(1) %a monotonic, align 11587 ret void1588}1589 1590define void @global_monotonic_volatile_sys_i16(ptr addrspace(1) %a) {1591; SM60-LABEL: global_monotonic_volatile_sys_i16(1592; SM60: {1593; SM60-NEXT: .reg .b16 %rs<3>;1594; SM60-NEXT: .reg .b64 %rd<2>;1595; SM60-EMPTY:1596; SM60-NEXT: // %bb.0:1597; SM60-NEXT: ld.param.b64 %rd1, [global_monotonic_volatile_sys_i16_param_0];1598; SM60-NEXT: ld.volatile.global.b16 %rs1, [%rd1];1599; SM60-NEXT: add.s16 %rs2, %rs1, 1;1600; SM60-NEXT: st.volatile.global.b16 [%rd1], %rs2;1601; SM60-NEXT: ret;1602;1603; SM70-LABEL: global_monotonic_volatile_sys_i16(1604; SM70: {1605; SM70-NEXT: .reg .b16 %rs<3>;1606; SM70-NEXT: .reg .b64 %rd<2>;1607; SM70-EMPTY:1608; SM70-NEXT: // %bb.0:1609; SM70-NEXT: ld.param.b64 %rd1, [global_monotonic_volatile_sys_i16_param_0];1610; SM70-NEXT: ld.mmio.relaxed.sys.global.b16 %rs1, [%rd1];1611; SM70-NEXT: add.s16 %rs2, %rs1, 1;1612; SM70-NEXT: st.mmio.relaxed.sys.global.b16 [%rd1], %rs2;1613; SM70-NEXT: ret;1614 %a.load = load atomic volatile i16, ptr addrspace(1) %a monotonic, align 21615 %a.add = add i16 %a.load, 11616 store atomic volatile i16 %a.add, ptr addrspace(1) %a monotonic, align 21617 ret void1618}1619 1620define void @global_monotonic_volatile_sys_i32(ptr addrspace(1) %a) {1621; SM60-LABEL: global_monotonic_volatile_sys_i32(1622; SM60: {1623; SM60-NEXT: .reg .b32 %r<3>;1624; SM60-NEXT: .reg .b64 %rd<2>;1625; SM60-EMPTY:1626; SM60-NEXT: // %bb.0:1627; SM60-NEXT: ld.param.b64 %rd1, [global_monotonic_volatile_sys_i32_param_0];1628; SM60-NEXT: ld.volatile.global.b32 %r1, [%rd1];1629; SM60-NEXT: add.s32 %r2, %r1, 1;1630; SM60-NEXT: st.volatile.global.b32 [%rd1], %r2;1631; SM60-NEXT: ret;1632;1633; SM70-LABEL: global_monotonic_volatile_sys_i32(1634; SM70: {1635; SM70-NEXT: .reg .b32 %r<3>;1636; SM70-NEXT: .reg .b64 %rd<2>;1637; SM70-EMPTY:1638; SM70-NEXT: // %bb.0:1639; SM70-NEXT: ld.param.b64 %rd1, [global_monotonic_volatile_sys_i32_param_0];1640; SM70-NEXT: ld.mmio.relaxed.sys.global.b32 %r1, [%rd1];1641; SM70-NEXT: add.s32 %r2, %r1, 1;1642; SM70-NEXT: st.mmio.relaxed.sys.global.b32 [%rd1], %r2;1643; SM70-NEXT: ret;1644 %a.load = load atomic volatile i32, ptr addrspace(1) %a monotonic, align 41645 %a.add = add i32 %a.load, 11646 store atomic volatile i32 %a.add, ptr addrspace(1) %a monotonic, align 41647 ret void1648}1649 1650define void @global_monotonic_volatile_sys_i64(ptr addrspace(1) %a) {1651; SM60-LABEL: global_monotonic_volatile_sys_i64(1652; SM60: {1653; SM60-NEXT: .reg .b64 %rd<4>;1654; SM60-EMPTY:1655; SM60-NEXT: // %bb.0:1656; SM60-NEXT: ld.param.b64 %rd1, [global_monotonic_volatile_sys_i64_param_0];1657; SM60-NEXT: ld.volatile.global.b64 %rd2, [%rd1];1658; SM60-NEXT: add.s64 %rd3, %rd2, 1;1659; SM60-NEXT: st.volatile.global.b64 [%rd1], %rd3;1660; SM60-NEXT: ret;1661;1662; SM70-LABEL: global_monotonic_volatile_sys_i64(1663; SM70: {1664; SM70-NEXT: .reg .b64 %rd<4>;1665; SM70-EMPTY:1666; SM70-NEXT: // %bb.0:1667; SM70-NEXT: ld.param.b64 %rd1, [global_monotonic_volatile_sys_i64_param_0];1668; SM70-NEXT: ld.mmio.relaxed.sys.global.b64 %rd2, [%rd1];1669; SM70-NEXT: add.s64 %rd3, %rd2, 1;1670; SM70-NEXT: st.mmio.relaxed.sys.global.b64 [%rd1], %rd3;1671; SM70-NEXT: ret;1672 %a.load = load atomic volatile i64, ptr addrspace(1) %a monotonic, align 81673 %a.add = add i64 %a.load, 11674 store atomic volatile i64 %a.add, ptr addrspace(1) %a monotonic, align 81675 ret void1676}1677 1678define void @global_monotonic_volatile_sys_float(ptr addrspace(1) %a) {1679; SM60-LABEL: global_monotonic_volatile_sys_float(1680; SM60: {1681; SM60-NEXT: .reg .b32 %r<3>;1682; SM60-NEXT: .reg .b64 %rd<2>;1683; SM60-EMPTY:1684; SM60-NEXT: // %bb.0:1685; SM60-NEXT: ld.param.b64 %rd1, [global_monotonic_volatile_sys_float_param_0];1686; SM60-NEXT: ld.volatile.global.b32 %r1, [%rd1];1687; SM60-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;1688; SM60-NEXT: st.volatile.global.b32 [%rd1], %r2;1689; SM60-NEXT: ret;1690;1691; SM70-LABEL: global_monotonic_volatile_sys_float(1692; SM70: {1693; SM70-NEXT: .reg .b32 %r<3>;1694; SM70-NEXT: .reg .b64 %rd<2>;1695; SM70-EMPTY:1696; SM70-NEXT: // %bb.0:1697; SM70-NEXT: ld.param.b64 %rd1, [global_monotonic_volatile_sys_float_param_0];1698; SM70-NEXT: ld.mmio.relaxed.sys.global.b32 %r1, [%rd1];1699; SM70-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;1700; SM70-NEXT: st.mmio.relaxed.sys.global.b32 [%rd1], %r2;1701; SM70-NEXT: ret;1702 %a.load = load atomic volatile float, ptr addrspace(1) %a monotonic, align 41703 %a.add = fadd float %a.load, 1.1704 store atomic volatile float %a.add, ptr addrspace(1) %a monotonic, align 41705 ret void1706}1707 1708define void @global_monotonic_volatile_sys_double(ptr addrspace(1) %a) {1709; SM60-LABEL: global_monotonic_volatile_sys_double(1710; SM60: {1711; SM60-NEXT: .reg .b64 %rd<4>;1712; SM60-EMPTY:1713; SM60-NEXT: // %bb.0:1714; SM60-NEXT: ld.param.b64 %rd1, [global_monotonic_volatile_sys_double_param_0];1715; SM60-NEXT: ld.volatile.global.b64 %rd2, [%rd1];1716; SM60-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;1717; SM60-NEXT: st.volatile.global.b64 [%rd1], %rd3;1718; SM60-NEXT: ret;1719;1720; SM70-LABEL: global_monotonic_volatile_sys_double(1721; SM70: {1722; SM70-NEXT: .reg .b64 %rd<4>;1723; SM70-EMPTY:1724; SM70-NEXT: // %bb.0:1725; SM70-NEXT: ld.param.b64 %rd1, [global_monotonic_volatile_sys_double_param_0];1726; SM70-NEXT: ld.mmio.relaxed.sys.global.b64 %rd2, [%rd1];1727; SM70-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;1728; SM70-NEXT: st.mmio.relaxed.sys.global.b64 [%rd1], %rd3;1729; SM70-NEXT: ret;1730 %a.load = load atomic volatile double, ptr addrspace(1) %a monotonic, align 81731 %a.add = fadd double %a.load, 1.1732 store atomic volatile double %a.add, ptr addrspace(1) %a monotonic, align 81733 ret void1734}1735 1736;; shared statespace1737 1738; shared1739 1740define void @shared_i8(ptr addrspace(3) %a) {1741; CHECK-LABEL: shared_i8(1742; CHECK: {1743; CHECK-NEXT: .reg .b16 %rs<3>;1744; CHECK-NEXT: .reg .b64 %rd<2>;1745; CHECK-EMPTY:1746; CHECK-NEXT: // %bb.0:1747; CHECK-NEXT: ld.param.b64 %rd1, [shared_i8_param_0];1748; CHECK-NEXT: ld.shared.b8 %rs1, [%rd1];1749; CHECK-NEXT: add.s16 %rs2, %rs1, 1;1750; CHECK-NEXT: st.shared.b8 [%rd1], %rs2;1751; CHECK-NEXT: ret;1752 %a.load = load i8, ptr addrspace(3) %a1753 %a.add = add i8 %a.load, 11754 store i8 %a.add, ptr addrspace(3) %a1755 ret void1756}1757 1758define void @shared_i16(ptr addrspace(3) %a) {1759; CHECK-LABEL: shared_i16(1760; CHECK: {1761; CHECK-NEXT: .reg .b16 %rs<3>;1762; CHECK-NEXT: .reg .b64 %rd<2>;1763; CHECK-EMPTY:1764; CHECK-NEXT: // %bb.0:1765; CHECK-NEXT: ld.param.b64 %rd1, [shared_i16_param_0];1766; CHECK-NEXT: ld.shared.b16 %rs1, [%rd1];1767; CHECK-NEXT: add.s16 %rs2, %rs1, 1;1768; CHECK-NEXT: st.shared.b16 [%rd1], %rs2;1769; CHECK-NEXT: ret;1770 %a.load = load i16, ptr addrspace(3) %a1771 %a.add = add i16 %a.load, 11772 store i16 %a.add, ptr addrspace(3) %a1773 ret void1774}1775 1776define void @shared_i32(ptr addrspace(3) %a) {1777; CHECK-LABEL: shared_i32(1778; CHECK: {1779; CHECK-NEXT: .reg .b32 %r<3>;1780; CHECK-NEXT: .reg .b64 %rd<2>;1781; CHECK-EMPTY:1782; CHECK-NEXT: // %bb.0:1783; CHECK-NEXT: ld.param.b64 %rd1, [shared_i32_param_0];1784; CHECK-NEXT: ld.shared.b32 %r1, [%rd1];1785; CHECK-NEXT: add.s32 %r2, %r1, 1;1786; CHECK-NEXT: st.shared.b32 [%rd1], %r2;1787; CHECK-NEXT: ret;1788 %a.load = load i32, ptr addrspace(3) %a1789 %a.add = add i32 %a.load, 11790 store i32 %a.add, ptr addrspace(3) %a1791 ret void1792}1793 1794define void @shared_i64(ptr addrspace(3) %a) {1795; CHECK-LABEL: shared_i64(1796; CHECK: {1797; CHECK-NEXT: .reg .b64 %rd<4>;1798; CHECK-EMPTY:1799; CHECK-NEXT: // %bb.0:1800; CHECK-NEXT: ld.param.b64 %rd1, [shared_i64_param_0];1801; CHECK-NEXT: ld.shared.b64 %rd2, [%rd1];1802; CHECK-NEXT: add.s64 %rd3, %rd2, 1;1803; CHECK-NEXT: st.shared.b64 [%rd1], %rd3;1804; CHECK-NEXT: ret;1805 %a.load = load i64, ptr addrspace(3) %a1806 %a.add = add i64 %a.load, 11807 store i64 %a.add, ptr addrspace(3) %a1808 ret void1809}1810 1811define void @shared_float(ptr addrspace(3) %a) {1812; CHECK-LABEL: shared_float(1813; CHECK: {1814; CHECK-NEXT: .reg .b32 %r<3>;1815; CHECK-NEXT: .reg .b64 %rd<2>;1816; CHECK-EMPTY:1817; CHECK-NEXT: // %bb.0:1818; CHECK-NEXT: ld.param.b64 %rd1, [shared_float_param_0];1819; CHECK-NEXT: ld.shared.b32 %r1, [%rd1];1820; CHECK-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;1821; CHECK-NEXT: st.shared.b32 [%rd1], %r2;1822; CHECK-NEXT: ret;1823 %a.load = load float, ptr addrspace(3) %a1824 %a.add = fadd float %a.load, 1.1825 store float %a.add, ptr addrspace(3) %a1826 ret void1827}1828 1829define void @shared_double(ptr addrspace(3) %a) {1830; CHECK-LABEL: shared_double(1831; CHECK: {1832; CHECK-NEXT: .reg .b64 %rd<4>;1833; CHECK-EMPTY:1834; CHECK-NEXT: // %bb.0:1835; CHECK-NEXT: ld.param.b64 %rd1, [shared_double_param_0];1836; CHECK-NEXT: ld.shared.b64 %rd2, [%rd1];1837; CHECK-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;1838; CHECK-NEXT: st.shared.b64 [%rd1], %rd3;1839; CHECK-NEXT: ret;1840 %a.load = load double, ptr addrspace(3) %a1841 %a.add = fadd double %a.load, 1.1842 store double %a.add, ptr addrspace(3) %a1843 ret void1844}1845 1846; shared_volatile1847 1848define void @shared_volatile_i8(ptr addrspace(3) %a) {1849; CHECK-LABEL: shared_volatile_i8(1850; CHECK: {1851; CHECK-NEXT: .reg .b16 %rs<3>;1852; CHECK-NEXT: .reg .b64 %rd<2>;1853; CHECK-EMPTY:1854; CHECK-NEXT: // %bb.0:1855; CHECK-NEXT: ld.param.b64 %rd1, [shared_volatile_i8_param_0];1856; CHECK-NEXT: ld.volatile.shared.b8 %rs1, [%rd1];1857; CHECK-NEXT: add.s16 %rs2, %rs1, 1;1858; CHECK-NEXT: st.volatile.shared.b8 [%rd1], %rs2;1859; CHECK-NEXT: ret;1860 %a.load = load volatile i8, ptr addrspace(3) %a1861 %a.add = add i8 %a.load, 11862 store volatile i8 %a.add, ptr addrspace(3) %a1863 ret void1864}1865 1866define void @shared_volatile_i16(ptr addrspace(3) %a) {1867; CHECK-LABEL: shared_volatile_i16(1868; CHECK: {1869; CHECK-NEXT: .reg .b16 %rs<3>;1870; CHECK-NEXT: .reg .b64 %rd<2>;1871; CHECK-EMPTY:1872; CHECK-NEXT: // %bb.0:1873; CHECK-NEXT: ld.param.b64 %rd1, [shared_volatile_i16_param_0];1874; CHECK-NEXT: ld.volatile.shared.b16 %rs1, [%rd1];1875; CHECK-NEXT: add.s16 %rs2, %rs1, 1;1876; CHECK-NEXT: st.volatile.shared.b16 [%rd1], %rs2;1877; CHECK-NEXT: ret;1878 %a.load = load volatile i16, ptr addrspace(3) %a1879 %a.add = add i16 %a.load, 11880 store volatile i16 %a.add, ptr addrspace(3) %a1881 ret void1882}1883 1884define void @shared_volatile_i32(ptr addrspace(3) %a) {1885; CHECK-LABEL: shared_volatile_i32(1886; CHECK: {1887; CHECK-NEXT: .reg .b32 %r<3>;1888; CHECK-NEXT: .reg .b64 %rd<2>;1889; CHECK-EMPTY:1890; CHECK-NEXT: // %bb.0:1891; CHECK-NEXT: ld.param.b64 %rd1, [shared_volatile_i32_param_0];1892; CHECK-NEXT: ld.volatile.shared.b32 %r1, [%rd1];1893; CHECK-NEXT: add.s32 %r2, %r1, 1;1894; CHECK-NEXT: st.volatile.shared.b32 [%rd1], %r2;1895; CHECK-NEXT: ret;1896 %a.load = load volatile i32, ptr addrspace(3) %a1897 %a.add = add i32 %a.load, 11898 store volatile i32 %a.add, ptr addrspace(3) %a1899 ret void1900}1901 1902define void @shared_volatile_i64(ptr addrspace(3) %a) {1903; CHECK-LABEL: shared_volatile_i64(1904; CHECK: {1905; CHECK-NEXT: .reg .b64 %rd<4>;1906; CHECK-EMPTY:1907; CHECK-NEXT: // %bb.0:1908; CHECK-NEXT: ld.param.b64 %rd1, [shared_volatile_i64_param_0];1909; CHECK-NEXT: ld.volatile.shared.b64 %rd2, [%rd1];1910; CHECK-NEXT: add.s64 %rd3, %rd2, 1;1911; CHECK-NEXT: st.volatile.shared.b64 [%rd1], %rd3;1912; CHECK-NEXT: ret;1913 %a.load = load volatile i64, ptr addrspace(3) %a1914 %a.add = add i64 %a.load, 11915 store volatile i64 %a.add, ptr addrspace(3) %a1916 ret void1917}1918 1919define void @shared_volatile_float(ptr addrspace(3) %a) {1920; CHECK-LABEL: shared_volatile_float(1921; CHECK: {1922; CHECK-NEXT: .reg .b32 %r<3>;1923; CHECK-NEXT: .reg .b64 %rd<2>;1924; CHECK-EMPTY:1925; CHECK-NEXT: // %bb.0:1926; CHECK-NEXT: ld.param.b64 %rd1, [shared_volatile_float_param_0];1927; CHECK-NEXT: ld.volatile.shared.b32 %r1, [%rd1];1928; CHECK-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;1929; CHECK-NEXT: st.volatile.shared.b32 [%rd1], %r2;1930; CHECK-NEXT: ret;1931 %a.load = load volatile float, ptr addrspace(3) %a1932 %a.add = fadd float %a.load, 1.1933 store volatile float %a.add, ptr addrspace(3) %a1934 ret void1935}1936 1937define void @shared_volatile_double(ptr addrspace(3) %a) {1938; CHECK-LABEL: shared_volatile_double(1939; CHECK: {1940; CHECK-NEXT: .reg .b64 %rd<4>;1941; CHECK-EMPTY:1942; CHECK-NEXT: // %bb.0:1943; CHECK-NEXT: ld.param.b64 %rd1, [shared_volatile_double_param_0];1944; CHECK-NEXT: ld.volatile.shared.b64 %rd2, [%rd1];1945; CHECK-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;1946; CHECK-NEXT: st.volatile.shared.b64 [%rd1], %rd3;1947; CHECK-NEXT: ret;1948 %a.load = load volatile double, ptr addrspace(3) %a1949 %a.add = fadd double %a.load, 1.1950 store volatile double %a.add, ptr addrspace(3) %a1951 ret void1952}1953 1954; shared_unordered_sys1955 1956define void @shared_unordered_sys_i8(ptr addrspace(3) %a) {1957; SM60-LABEL: shared_unordered_sys_i8(1958; SM60: {1959; SM60-NEXT: .reg .b16 %rs<3>;1960; SM60-NEXT: .reg .b64 %rd<2>;1961; SM60-EMPTY:1962; SM60-NEXT: // %bb.0:1963; SM60-NEXT: ld.param.b64 %rd1, [shared_unordered_sys_i8_param_0];1964; SM60-NEXT: ld.volatile.shared.b8 %rs1, [%rd1];1965; SM60-NEXT: add.s16 %rs2, %rs1, 1;1966; SM60-NEXT: st.volatile.shared.b8 [%rd1], %rs2;1967; SM60-NEXT: ret;1968;1969; SM70-LABEL: shared_unordered_sys_i8(1970; SM70: {1971; SM70-NEXT: .reg .b16 %rs<3>;1972; SM70-NEXT: .reg .b64 %rd<2>;1973; SM70-EMPTY:1974; SM70-NEXT: // %bb.0:1975; SM70-NEXT: ld.param.b64 %rd1, [shared_unordered_sys_i8_param_0];1976; SM70-NEXT: ld.relaxed.sys.shared.b8 %rs1, [%rd1];1977; SM70-NEXT: add.s16 %rs2, %rs1, 1;1978; SM70-NEXT: st.relaxed.sys.shared.b8 [%rd1], %rs2;1979; SM70-NEXT: ret;1980 %a.load = load atomic i8, ptr addrspace(3) %a unordered, align 11981 %a.add = add i8 %a.load, 11982 store atomic i8 %a.add, ptr addrspace(3) %a unordered, align 11983 ret void1984}1985 1986define void @shared_unordered_sys_i16(ptr addrspace(3) %a) {1987; SM60-LABEL: shared_unordered_sys_i16(1988; SM60: {1989; SM60-NEXT: .reg .b16 %rs<3>;1990; SM60-NEXT: .reg .b64 %rd<2>;1991; SM60-EMPTY:1992; SM60-NEXT: // %bb.0:1993; SM60-NEXT: ld.param.b64 %rd1, [shared_unordered_sys_i16_param_0];1994; SM60-NEXT: ld.volatile.shared.b16 %rs1, [%rd1];1995; SM60-NEXT: add.s16 %rs2, %rs1, 1;1996; SM60-NEXT: st.volatile.shared.b16 [%rd1], %rs2;1997; SM60-NEXT: ret;1998;1999; SM70-LABEL: shared_unordered_sys_i16(2000; SM70: {2001; SM70-NEXT: .reg .b16 %rs<3>;2002; SM70-NEXT: .reg .b64 %rd<2>;2003; SM70-EMPTY:2004; SM70-NEXT: // %bb.0:2005; SM70-NEXT: ld.param.b64 %rd1, [shared_unordered_sys_i16_param_0];2006; SM70-NEXT: ld.relaxed.sys.shared.b16 %rs1, [%rd1];2007; SM70-NEXT: add.s16 %rs2, %rs1, 1;2008; SM70-NEXT: st.relaxed.sys.shared.b16 [%rd1], %rs2;2009; SM70-NEXT: ret;2010 %a.load = load atomic i16, ptr addrspace(3) %a unordered, align 22011 %a.add = add i16 %a.load, 12012 store atomic i16 %a.add, ptr addrspace(3) %a unordered, align 22013 ret void2014}2015 2016define void @shared_unordered_sys_i32(ptr addrspace(3) %a) {2017; SM60-LABEL: shared_unordered_sys_i32(2018; SM60: {2019; SM60-NEXT: .reg .b32 %r<3>;2020; SM60-NEXT: .reg .b64 %rd<2>;2021; SM60-EMPTY:2022; SM60-NEXT: // %bb.0:2023; SM60-NEXT: ld.param.b64 %rd1, [shared_unordered_sys_i32_param_0];2024; SM60-NEXT: ld.volatile.shared.b32 %r1, [%rd1];2025; SM60-NEXT: add.s32 %r2, %r1, 1;2026; SM60-NEXT: st.volatile.shared.b32 [%rd1], %r2;2027; SM60-NEXT: ret;2028;2029; SM70-LABEL: shared_unordered_sys_i32(2030; SM70: {2031; SM70-NEXT: .reg .b32 %r<3>;2032; SM70-NEXT: .reg .b64 %rd<2>;2033; SM70-EMPTY:2034; SM70-NEXT: // %bb.0:2035; SM70-NEXT: ld.param.b64 %rd1, [shared_unordered_sys_i32_param_0];2036; SM70-NEXT: ld.relaxed.sys.shared.b32 %r1, [%rd1];2037; SM70-NEXT: add.s32 %r2, %r1, 1;2038; SM70-NEXT: st.relaxed.sys.shared.b32 [%rd1], %r2;2039; SM70-NEXT: ret;2040 %a.load = load atomic i32, ptr addrspace(3) %a unordered, align 42041 %a.add = add i32 %a.load, 12042 store atomic i32 %a.add, ptr addrspace(3) %a unordered, align 42043 ret void2044}2045 2046define void @shared_unordered_sys_i64(ptr addrspace(3) %a) {2047; SM60-LABEL: shared_unordered_sys_i64(2048; SM60: {2049; SM60-NEXT: .reg .b64 %rd<4>;2050; SM60-EMPTY:2051; SM60-NEXT: // %bb.0:2052; SM60-NEXT: ld.param.b64 %rd1, [shared_unordered_sys_i64_param_0];2053; SM60-NEXT: ld.volatile.shared.b64 %rd2, [%rd1];2054; SM60-NEXT: add.s64 %rd3, %rd2, 1;2055; SM60-NEXT: st.volatile.shared.b64 [%rd1], %rd3;2056; SM60-NEXT: ret;2057;2058; SM70-LABEL: shared_unordered_sys_i64(2059; SM70: {2060; SM70-NEXT: .reg .b64 %rd<4>;2061; SM70-EMPTY:2062; SM70-NEXT: // %bb.0:2063; SM70-NEXT: ld.param.b64 %rd1, [shared_unordered_sys_i64_param_0];2064; SM70-NEXT: ld.relaxed.sys.shared.b64 %rd2, [%rd1];2065; SM70-NEXT: add.s64 %rd3, %rd2, 1;2066; SM70-NEXT: st.relaxed.sys.shared.b64 [%rd1], %rd3;2067; SM70-NEXT: ret;2068 %a.load = load atomic i64, ptr addrspace(3) %a unordered, align 82069 %a.add = add i64 %a.load, 12070 store atomic i64 %a.add, ptr addrspace(3) %a unordered, align 82071 ret void2072}2073 2074define void @shared_unordered_sys_float(ptr addrspace(3) %a) {2075; SM60-LABEL: shared_unordered_sys_float(2076; SM60: {2077; SM60-NEXT: .reg .b32 %r<3>;2078; SM60-NEXT: .reg .b64 %rd<2>;2079; SM60-EMPTY:2080; SM60-NEXT: // %bb.0:2081; SM60-NEXT: ld.param.b64 %rd1, [shared_unordered_sys_float_param_0];2082; SM60-NEXT: ld.volatile.shared.b32 %r1, [%rd1];2083; SM60-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;2084; SM60-NEXT: st.volatile.shared.b32 [%rd1], %r2;2085; SM60-NEXT: ret;2086;2087; SM70-LABEL: shared_unordered_sys_float(2088; SM70: {2089; SM70-NEXT: .reg .b32 %r<3>;2090; SM70-NEXT: .reg .b64 %rd<2>;2091; SM70-EMPTY:2092; SM70-NEXT: // %bb.0:2093; SM70-NEXT: ld.param.b64 %rd1, [shared_unordered_sys_float_param_0];2094; SM70-NEXT: ld.relaxed.sys.shared.b32 %r1, [%rd1];2095; SM70-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;2096; SM70-NEXT: st.relaxed.sys.shared.b32 [%rd1], %r2;2097; SM70-NEXT: ret;2098 %a.load = load atomic float, ptr addrspace(3) %a unordered, align 42099 %a.add = fadd float %a.load, 1.2100 store atomic float %a.add, ptr addrspace(3) %a unordered, align 42101 ret void2102}2103 2104define void @shared_unordered_sys_double(ptr addrspace(3) %a) {2105; SM60-LABEL: shared_unordered_sys_double(2106; SM60: {2107; SM60-NEXT: .reg .b64 %rd<4>;2108; SM60-EMPTY:2109; SM60-NEXT: // %bb.0:2110; SM60-NEXT: ld.param.b64 %rd1, [shared_unordered_sys_double_param_0];2111; SM60-NEXT: ld.volatile.shared.b64 %rd2, [%rd1];2112; SM60-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;2113; SM60-NEXT: st.volatile.shared.b64 [%rd1], %rd3;2114; SM60-NEXT: ret;2115;2116; SM70-LABEL: shared_unordered_sys_double(2117; SM70: {2118; SM70-NEXT: .reg .b64 %rd<4>;2119; SM70-EMPTY:2120; SM70-NEXT: // %bb.0:2121; SM70-NEXT: ld.param.b64 %rd1, [shared_unordered_sys_double_param_0];2122; SM70-NEXT: ld.relaxed.sys.shared.b64 %rd2, [%rd1];2123; SM70-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;2124; SM70-NEXT: st.relaxed.sys.shared.b64 [%rd1], %rd3;2125; SM70-NEXT: ret;2126 %a.load = load atomic double, ptr addrspace(3) %a unordered, align 82127 %a.add = fadd double %a.load, 1.2128 store atomic double %a.add, ptr addrspace(3) %a unordered, align 82129 ret void2130}2131 2132; shared_unordered_volatile_sys2133 2134define void @shared_unordered_volatile_sys_i8(ptr addrspace(3) %a) {2135; CHECK-LABEL: shared_unordered_volatile_sys_i8(2136; CHECK: {2137; CHECK-NEXT: .reg .b16 %rs<3>;2138; CHECK-NEXT: .reg .b64 %rd<2>;2139; CHECK-EMPTY:2140; CHECK-NEXT: // %bb.0:2141; CHECK-NEXT: ld.param.b64 %rd1, [shared_unordered_volatile_sys_i8_param_0];2142; CHECK-NEXT: ld.volatile.shared.b8 %rs1, [%rd1];2143; CHECK-NEXT: add.s16 %rs2, %rs1, 1;2144; CHECK-NEXT: st.volatile.shared.b8 [%rd1], %rs2;2145; CHECK-NEXT: ret;2146 %a.load = load atomic volatile i8, ptr addrspace(3) %a unordered, align 12147 %a.add = add i8 %a.load, 12148 store atomic volatile i8 %a.add, ptr addrspace(3) %a unordered, align 12149 ret void2150}2151 2152define void @shared_unordered_volatile_sys_i16(ptr addrspace(3) %a) {2153; CHECK-LABEL: shared_unordered_volatile_sys_i16(2154; CHECK: {2155; CHECK-NEXT: .reg .b16 %rs<3>;2156; CHECK-NEXT: .reg .b64 %rd<2>;2157; CHECK-EMPTY:2158; CHECK-NEXT: // %bb.0:2159; CHECK-NEXT: ld.param.b64 %rd1, [shared_unordered_volatile_sys_i16_param_0];2160; CHECK-NEXT: ld.volatile.shared.b16 %rs1, [%rd1];2161; CHECK-NEXT: add.s16 %rs2, %rs1, 1;2162; CHECK-NEXT: st.volatile.shared.b16 [%rd1], %rs2;2163; CHECK-NEXT: ret;2164 %a.load = load atomic volatile i16, ptr addrspace(3) %a unordered, align 22165 %a.add = add i16 %a.load, 12166 store atomic volatile i16 %a.add, ptr addrspace(3) %a unordered, align 22167 ret void2168}2169 2170define void @shared_unordered_volatile_sys_i32(ptr addrspace(3) %a) {2171; CHECK-LABEL: shared_unordered_volatile_sys_i32(2172; CHECK: {2173; CHECK-NEXT: .reg .b32 %r<3>;2174; CHECK-NEXT: .reg .b64 %rd<2>;2175; CHECK-EMPTY:2176; CHECK-NEXT: // %bb.0:2177; CHECK-NEXT: ld.param.b64 %rd1, [shared_unordered_volatile_sys_i32_param_0];2178; CHECK-NEXT: ld.volatile.shared.b32 %r1, [%rd1];2179; CHECK-NEXT: add.s32 %r2, %r1, 1;2180; CHECK-NEXT: st.volatile.shared.b32 [%rd1], %r2;2181; CHECK-NEXT: ret;2182 %a.load = load atomic volatile i32, ptr addrspace(3) %a unordered, align 42183 %a.add = add i32 %a.load, 12184 store atomic volatile i32 %a.add, ptr addrspace(3) %a unordered, align 42185 ret void2186}2187 2188define void @shared_unordered_volatile_sys_i64(ptr addrspace(3) %a) {2189; CHECK-LABEL: shared_unordered_volatile_sys_i64(2190; CHECK: {2191; CHECK-NEXT: .reg .b64 %rd<4>;2192; CHECK-EMPTY:2193; CHECK-NEXT: // %bb.0:2194; CHECK-NEXT: ld.param.b64 %rd1, [shared_unordered_volatile_sys_i64_param_0];2195; CHECK-NEXT: ld.volatile.shared.b64 %rd2, [%rd1];2196; CHECK-NEXT: add.s64 %rd3, %rd2, 1;2197; CHECK-NEXT: st.volatile.shared.b64 [%rd1], %rd3;2198; CHECK-NEXT: ret;2199 %a.load = load atomic volatile i64, ptr addrspace(3) %a unordered, align 82200 %a.add = add i64 %a.load, 12201 store atomic volatile i64 %a.add, ptr addrspace(3) %a unordered, align 82202 ret void2203}2204 2205define void @shared_unordered_volatile_sys_float(ptr addrspace(3) %a) {2206; CHECK-LABEL: shared_unordered_volatile_sys_float(2207; CHECK: {2208; CHECK-NEXT: .reg .b32 %r<3>;2209; CHECK-NEXT: .reg .b64 %rd<2>;2210; CHECK-EMPTY:2211; CHECK-NEXT: // %bb.0:2212; CHECK-NEXT: ld.param.b64 %rd1, [shared_unordered_volatile_sys_float_param_0];2213; CHECK-NEXT: ld.volatile.shared.b32 %r1, [%rd1];2214; CHECK-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;2215; CHECK-NEXT: st.volatile.shared.b32 [%rd1], %r2;2216; CHECK-NEXT: ret;2217 %a.load = load atomic volatile float, ptr addrspace(3) %a unordered, align 42218 %a.add = fadd float %a.load, 1.2219 store atomic volatile float %a.add, ptr addrspace(3) %a unordered, align 42220 ret void2221}2222 2223define void @shared_unordered_volatile_sys_double(ptr addrspace(3) %a) {2224; CHECK-LABEL: shared_unordered_volatile_sys_double(2225; CHECK: {2226; CHECK-NEXT: .reg .b64 %rd<4>;2227; CHECK-EMPTY:2228; CHECK-NEXT: // %bb.0:2229; CHECK-NEXT: ld.param.b64 %rd1, [shared_unordered_volatile_sys_double_param_0];2230; CHECK-NEXT: ld.volatile.shared.b64 %rd2, [%rd1];2231; CHECK-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;2232; CHECK-NEXT: st.volatile.shared.b64 [%rd1], %rd3;2233; CHECK-NEXT: ret;2234 %a.load = load atomic volatile double, ptr addrspace(3) %a unordered, align 82235 %a.add = fadd double %a.load, 1.2236 store atomic volatile double %a.add, ptr addrspace(3) %a unordered, align 82237 ret void2238}2239 2240; shared_monotonic_sys2241 2242define void @shared_monotonic_sys_i8(ptr addrspace(3) %a) {2243; SM60-LABEL: shared_monotonic_sys_i8(2244; SM60: {2245; SM60-NEXT: .reg .b16 %rs<3>;2246; SM60-NEXT: .reg .b64 %rd<2>;2247; SM60-EMPTY:2248; SM60-NEXT: // %bb.0:2249; SM60-NEXT: ld.param.b64 %rd1, [shared_monotonic_sys_i8_param_0];2250; SM60-NEXT: ld.volatile.shared.b8 %rs1, [%rd1];2251; SM60-NEXT: add.s16 %rs2, %rs1, 1;2252; SM60-NEXT: st.volatile.shared.b8 [%rd1], %rs2;2253; SM60-NEXT: ret;2254;2255; SM70-LABEL: shared_monotonic_sys_i8(2256; SM70: {2257; SM70-NEXT: .reg .b16 %rs<3>;2258; SM70-NEXT: .reg .b64 %rd<2>;2259; SM70-EMPTY:2260; SM70-NEXT: // %bb.0:2261; SM70-NEXT: ld.param.b64 %rd1, [shared_monotonic_sys_i8_param_0];2262; SM70-NEXT: ld.relaxed.sys.shared.b8 %rs1, [%rd1];2263; SM70-NEXT: add.s16 %rs2, %rs1, 1;2264; SM70-NEXT: st.relaxed.sys.shared.b8 [%rd1], %rs2;2265; SM70-NEXT: ret;2266 %a.load = load atomic i8, ptr addrspace(3) %a monotonic, align 12267 %a.add = add i8 %a.load, 12268 store atomic i8 %a.add, ptr addrspace(3) %a monotonic, align 12269 ret void2270}2271 2272define void @shared_monotonic_sys_i16(ptr addrspace(3) %a) {2273; SM60-LABEL: shared_monotonic_sys_i16(2274; SM60: {2275; SM60-NEXT: .reg .b16 %rs<3>;2276; SM60-NEXT: .reg .b64 %rd<2>;2277; SM60-EMPTY:2278; SM60-NEXT: // %bb.0:2279; SM60-NEXT: ld.param.b64 %rd1, [shared_monotonic_sys_i16_param_0];2280; SM60-NEXT: ld.volatile.shared.b16 %rs1, [%rd1];2281; SM60-NEXT: add.s16 %rs2, %rs1, 1;2282; SM60-NEXT: st.volatile.shared.b16 [%rd1], %rs2;2283; SM60-NEXT: ret;2284;2285; SM70-LABEL: shared_monotonic_sys_i16(2286; SM70: {2287; SM70-NEXT: .reg .b16 %rs<3>;2288; SM70-NEXT: .reg .b64 %rd<2>;2289; SM70-EMPTY:2290; SM70-NEXT: // %bb.0:2291; SM70-NEXT: ld.param.b64 %rd1, [shared_monotonic_sys_i16_param_0];2292; SM70-NEXT: ld.relaxed.sys.shared.b16 %rs1, [%rd1];2293; SM70-NEXT: add.s16 %rs2, %rs1, 1;2294; SM70-NEXT: st.relaxed.sys.shared.b16 [%rd1], %rs2;2295; SM70-NEXT: ret;2296 %a.load = load atomic i16, ptr addrspace(3) %a monotonic, align 22297 %a.add = add i16 %a.load, 12298 store atomic i16 %a.add, ptr addrspace(3) %a monotonic, align 22299 ret void2300}2301 2302define void @shared_monotonic_sys_i32(ptr addrspace(3) %a) {2303; SM60-LABEL: shared_monotonic_sys_i32(2304; SM60: {2305; SM60-NEXT: .reg .b32 %r<3>;2306; SM60-NEXT: .reg .b64 %rd<2>;2307; SM60-EMPTY:2308; SM60-NEXT: // %bb.0:2309; SM60-NEXT: ld.param.b64 %rd1, [shared_monotonic_sys_i32_param_0];2310; SM60-NEXT: ld.volatile.shared.b32 %r1, [%rd1];2311; SM60-NEXT: add.s32 %r2, %r1, 1;2312; SM60-NEXT: st.volatile.shared.b32 [%rd1], %r2;2313; SM60-NEXT: ret;2314;2315; SM70-LABEL: shared_monotonic_sys_i32(2316; SM70: {2317; SM70-NEXT: .reg .b32 %r<3>;2318; SM70-NEXT: .reg .b64 %rd<2>;2319; SM70-EMPTY:2320; SM70-NEXT: // %bb.0:2321; SM70-NEXT: ld.param.b64 %rd1, [shared_monotonic_sys_i32_param_0];2322; SM70-NEXT: ld.relaxed.sys.shared.b32 %r1, [%rd1];2323; SM70-NEXT: add.s32 %r2, %r1, 1;2324; SM70-NEXT: st.relaxed.sys.shared.b32 [%rd1], %r2;2325; SM70-NEXT: ret;2326 %a.load = load atomic i32, ptr addrspace(3) %a monotonic, align 42327 %a.add = add i32 %a.load, 12328 store atomic i32 %a.add, ptr addrspace(3) %a monotonic, align 42329 ret void2330}2331 2332define void @shared_monotonic_sys_i64(ptr addrspace(3) %a) {2333; SM60-LABEL: shared_monotonic_sys_i64(2334; SM60: {2335; SM60-NEXT: .reg .b64 %rd<4>;2336; SM60-EMPTY:2337; SM60-NEXT: // %bb.0:2338; SM60-NEXT: ld.param.b64 %rd1, [shared_monotonic_sys_i64_param_0];2339; SM60-NEXT: ld.volatile.shared.b64 %rd2, [%rd1];2340; SM60-NEXT: add.s64 %rd3, %rd2, 1;2341; SM60-NEXT: st.volatile.shared.b64 [%rd1], %rd3;2342; SM60-NEXT: ret;2343;2344; SM70-LABEL: shared_monotonic_sys_i64(2345; SM70: {2346; SM70-NEXT: .reg .b64 %rd<4>;2347; SM70-EMPTY:2348; SM70-NEXT: // %bb.0:2349; SM70-NEXT: ld.param.b64 %rd1, [shared_monotonic_sys_i64_param_0];2350; SM70-NEXT: ld.relaxed.sys.shared.b64 %rd2, [%rd1];2351; SM70-NEXT: add.s64 %rd3, %rd2, 1;2352; SM70-NEXT: st.relaxed.sys.shared.b64 [%rd1], %rd3;2353; SM70-NEXT: ret;2354 %a.load = load atomic i64, ptr addrspace(3) %a monotonic, align 82355 %a.add = add i64 %a.load, 12356 store atomic i64 %a.add, ptr addrspace(3) %a monotonic, align 82357 ret void2358}2359 2360define void @shared_monotonic_sys_float(ptr addrspace(3) %a) {2361; SM60-LABEL: shared_monotonic_sys_float(2362; SM60: {2363; SM60-NEXT: .reg .b32 %r<3>;2364; SM60-NEXT: .reg .b64 %rd<2>;2365; SM60-EMPTY:2366; SM60-NEXT: // %bb.0:2367; SM60-NEXT: ld.param.b64 %rd1, [shared_monotonic_sys_float_param_0];2368; SM60-NEXT: ld.volatile.shared.b32 %r1, [%rd1];2369; SM60-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;2370; SM60-NEXT: st.volatile.shared.b32 [%rd1], %r2;2371; SM60-NEXT: ret;2372;2373; SM70-LABEL: shared_monotonic_sys_float(2374; SM70: {2375; SM70-NEXT: .reg .b32 %r<3>;2376; SM70-NEXT: .reg .b64 %rd<2>;2377; SM70-EMPTY:2378; SM70-NEXT: // %bb.0:2379; SM70-NEXT: ld.param.b64 %rd1, [shared_monotonic_sys_float_param_0];2380; SM70-NEXT: ld.relaxed.sys.shared.b32 %r1, [%rd1];2381; SM70-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;2382; SM70-NEXT: st.relaxed.sys.shared.b32 [%rd1], %r2;2383; SM70-NEXT: ret;2384 %a.load = load atomic float, ptr addrspace(3) %a monotonic, align 42385 %a.add = fadd float %a.load, 1.2386 store atomic float %a.add, ptr addrspace(3) %a monotonic, align 42387 ret void2388}2389 2390define void @shared_monotonic_sys_double(ptr addrspace(3) %a) {2391; SM60-LABEL: shared_monotonic_sys_double(2392; SM60: {2393; SM60-NEXT: .reg .b64 %rd<4>;2394; SM60-EMPTY:2395; SM60-NEXT: // %bb.0:2396; SM60-NEXT: ld.param.b64 %rd1, [shared_monotonic_sys_double_param_0];2397; SM60-NEXT: ld.volatile.shared.b64 %rd2, [%rd1];2398; SM60-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;2399; SM60-NEXT: st.volatile.shared.b64 [%rd1], %rd3;2400; SM60-NEXT: ret;2401;2402; SM70-LABEL: shared_monotonic_sys_double(2403; SM70: {2404; SM70-NEXT: .reg .b64 %rd<4>;2405; SM70-EMPTY:2406; SM70-NEXT: // %bb.0:2407; SM70-NEXT: ld.param.b64 %rd1, [shared_monotonic_sys_double_param_0];2408; SM70-NEXT: ld.relaxed.sys.shared.b64 %rd2, [%rd1];2409; SM70-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;2410; SM70-NEXT: st.relaxed.sys.shared.b64 [%rd1], %rd3;2411; SM70-NEXT: ret;2412 %a.load = load atomic double, ptr addrspace(3) %a monotonic, align 82413 %a.add = fadd double %a.load, 1.2414 store atomic double %a.add, ptr addrspace(3) %a monotonic, align 82415 ret void2416}2417 2418; shared_monotonic_volatile_sys2419 2420define void @shared_monotonic_volatile_sys_i8(ptr addrspace(3) %a) {2421; CHECK-LABEL: shared_monotonic_volatile_sys_i8(2422; CHECK: {2423; CHECK-NEXT: .reg .b16 %rs<3>;2424; CHECK-NEXT: .reg .b64 %rd<2>;2425; CHECK-EMPTY:2426; CHECK-NEXT: // %bb.0:2427; CHECK-NEXT: ld.param.b64 %rd1, [shared_monotonic_volatile_sys_i8_param_0];2428; CHECK-NEXT: ld.volatile.shared.b8 %rs1, [%rd1];2429; CHECK-NEXT: add.s16 %rs2, %rs1, 1;2430; CHECK-NEXT: st.volatile.shared.b8 [%rd1], %rs2;2431; CHECK-NEXT: ret;2432 %a.load = load atomic volatile i8, ptr addrspace(3) %a monotonic, align 12433 %a.add = add i8 %a.load, 12434 store atomic volatile i8 %a.add, ptr addrspace(3) %a monotonic, align 12435 ret void2436}2437 2438define void @shared_monotonic_volatile_sys_i16(ptr addrspace(3) %a) {2439; CHECK-LABEL: shared_monotonic_volatile_sys_i16(2440; CHECK: {2441; CHECK-NEXT: .reg .b16 %rs<3>;2442; CHECK-NEXT: .reg .b64 %rd<2>;2443; CHECK-EMPTY:2444; CHECK-NEXT: // %bb.0:2445; CHECK-NEXT: ld.param.b64 %rd1, [shared_monotonic_volatile_sys_i16_param_0];2446; CHECK-NEXT: ld.volatile.shared.b16 %rs1, [%rd1];2447; CHECK-NEXT: add.s16 %rs2, %rs1, 1;2448; CHECK-NEXT: st.volatile.shared.b16 [%rd1], %rs2;2449; CHECK-NEXT: ret;2450 %a.load = load atomic volatile i16, ptr addrspace(3) %a monotonic, align 22451 %a.add = add i16 %a.load, 12452 store atomic volatile i16 %a.add, ptr addrspace(3) %a monotonic, align 22453 ret void2454}2455 2456define void @shared_monotonic_volatile_sys_i32(ptr addrspace(3) %a) {2457; CHECK-LABEL: shared_monotonic_volatile_sys_i32(2458; CHECK: {2459; CHECK-NEXT: .reg .b32 %r<3>;2460; CHECK-NEXT: .reg .b64 %rd<2>;2461; CHECK-EMPTY:2462; CHECK-NEXT: // %bb.0:2463; CHECK-NEXT: ld.param.b64 %rd1, [shared_monotonic_volatile_sys_i32_param_0];2464; CHECK-NEXT: ld.volatile.shared.b32 %r1, [%rd1];2465; CHECK-NEXT: add.s32 %r2, %r1, 1;2466; CHECK-NEXT: st.volatile.shared.b32 [%rd1], %r2;2467; CHECK-NEXT: ret;2468 %a.load = load atomic volatile i32, ptr addrspace(3) %a monotonic, align 42469 %a.add = add i32 %a.load, 12470 store atomic volatile i32 %a.add, ptr addrspace(3) %a monotonic, align 42471 ret void2472}2473 2474define void @shared_monotonic_volatile_sys_i64(ptr addrspace(3) %a) {2475; CHECK-LABEL: shared_monotonic_volatile_sys_i64(2476; CHECK: {2477; CHECK-NEXT: .reg .b64 %rd<4>;2478; CHECK-EMPTY:2479; CHECK-NEXT: // %bb.0:2480; CHECK-NEXT: ld.param.b64 %rd1, [shared_monotonic_volatile_sys_i64_param_0];2481; CHECK-NEXT: ld.volatile.shared.b64 %rd2, [%rd1];2482; CHECK-NEXT: add.s64 %rd3, %rd2, 1;2483; CHECK-NEXT: st.volatile.shared.b64 [%rd1], %rd3;2484; CHECK-NEXT: ret;2485 %a.load = load atomic volatile i64, ptr addrspace(3) %a monotonic, align 82486 %a.add = add i64 %a.load, 12487 store atomic volatile i64 %a.add, ptr addrspace(3) %a monotonic, align 82488 ret void2489}2490 2491define void @shared_monotonic_volatile_sys_float(ptr addrspace(3) %a) {2492; CHECK-LABEL: shared_monotonic_volatile_sys_float(2493; CHECK: {2494; CHECK-NEXT: .reg .b32 %r<3>;2495; CHECK-NEXT: .reg .b64 %rd<2>;2496; CHECK-EMPTY:2497; CHECK-NEXT: // %bb.0:2498; CHECK-NEXT: ld.param.b64 %rd1, [shared_monotonic_volatile_sys_float_param_0];2499; CHECK-NEXT: ld.volatile.shared.b32 %r1, [%rd1];2500; CHECK-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;2501; CHECK-NEXT: st.volatile.shared.b32 [%rd1], %r2;2502; CHECK-NEXT: ret;2503 %a.load = load atomic volatile float, ptr addrspace(3) %a monotonic, align 42504 %a.add = fadd float %a.load, 1.2505 store atomic volatile float %a.add, ptr addrspace(3) %a monotonic, align 42506 ret void2507}2508 2509define void @shared_monotonic_volatile_sys_double(ptr addrspace(3) %a) {2510; CHECK-LABEL: shared_monotonic_volatile_sys_double(2511; CHECK: {2512; CHECK-NEXT: .reg .b64 %rd<4>;2513; CHECK-EMPTY:2514; CHECK-NEXT: // %bb.0:2515; CHECK-NEXT: ld.param.b64 %rd1, [shared_monotonic_volatile_sys_double_param_0];2516; CHECK-NEXT: ld.volatile.shared.b64 %rd2, [%rd1];2517; CHECK-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;2518; CHECK-NEXT: st.volatile.shared.b64 [%rd1], %rd3;2519; CHECK-NEXT: ret;2520 %a.load = load atomic volatile double, ptr addrspace(3) %a monotonic, align 82521 %a.add = fadd double %a.load, 1.2522 store atomic volatile double %a.add, ptr addrspace(3) %a monotonic, align 82523 ret void2524}2525 2526;; local statespace2527 2528; local2529 2530define void @local_i8(ptr addrspace(5) %a) {2531; CHECK-LABEL: local_i8(2532; CHECK: {2533; CHECK-NEXT: .reg .b16 %rs<3>;2534; CHECK-NEXT: .reg .b64 %rd<2>;2535; CHECK-EMPTY:2536; CHECK-NEXT: // %bb.0:2537; CHECK-NEXT: ld.param.b64 %rd1, [local_i8_param_0];2538; CHECK-NEXT: ld.local.b8 %rs1, [%rd1];2539; CHECK-NEXT: add.s16 %rs2, %rs1, 1;2540; CHECK-NEXT: st.local.b8 [%rd1], %rs2;2541; CHECK-NEXT: ret;2542 %a.load = load i8, ptr addrspace(5) %a2543 %a.add = add i8 %a.load, 12544 store i8 %a.add, ptr addrspace(5) %a2545 ret void2546}2547 2548define void @local_i16(ptr addrspace(5) %a) {2549; CHECK-LABEL: local_i16(2550; CHECK: {2551; CHECK-NEXT: .reg .b16 %rs<3>;2552; CHECK-NEXT: .reg .b64 %rd<2>;2553; CHECK-EMPTY:2554; CHECK-NEXT: // %bb.0:2555; CHECK-NEXT: ld.param.b64 %rd1, [local_i16_param_0];2556; CHECK-NEXT: ld.local.b16 %rs1, [%rd1];2557; CHECK-NEXT: add.s16 %rs2, %rs1, 1;2558; CHECK-NEXT: st.local.b16 [%rd1], %rs2;2559; CHECK-NEXT: ret;2560 %a.load = load i16, ptr addrspace(5) %a2561 %a.add = add i16 %a.load, 12562 store i16 %a.add, ptr addrspace(5) %a2563 ret void2564}2565 2566define void @local_i32(ptr addrspace(5) %a) {2567; CHECK-LABEL: local_i32(2568; CHECK: {2569; CHECK-NEXT: .reg .b32 %r<3>;2570; CHECK-NEXT: .reg .b64 %rd<2>;2571; CHECK-EMPTY:2572; CHECK-NEXT: // %bb.0:2573; CHECK-NEXT: ld.param.b64 %rd1, [local_i32_param_0];2574; CHECK-NEXT: ld.local.b32 %r1, [%rd1];2575; CHECK-NEXT: add.s32 %r2, %r1, 1;2576; CHECK-NEXT: st.local.b32 [%rd1], %r2;2577; CHECK-NEXT: ret;2578 %a.load = load i32, ptr addrspace(5) %a2579 %a.add = add i32 %a.load, 12580 store i32 %a.add, ptr addrspace(5) %a2581 ret void2582}2583 2584define void @local_i64(ptr addrspace(5) %a) {2585; CHECK-LABEL: local_i64(2586; CHECK: {2587; CHECK-NEXT: .reg .b64 %rd<4>;2588; CHECK-EMPTY:2589; CHECK-NEXT: // %bb.0:2590; CHECK-NEXT: ld.param.b64 %rd1, [local_i64_param_0];2591; CHECK-NEXT: ld.local.b64 %rd2, [%rd1];2592; CHECK-NEXT: add.s64 %rd3, %rd2, 1;2593; CHECK-NEXT: st.local.b64 [%rd1], %rd3;2594; CHECK-NEXT: ret;2595 %a.load = load i64, ptr addrspace(5) %a2596 %a.add = add i64 %a.load, 12597 store i64 %a.add, ptr addrspace(5) %a2598 ret void2599}2600 2601define void @local_float(ptr addrspace(5) %a) {2602; CHECK-LABEL: local_float(2603; CHECK: {2604; CHECK-NEXT: .reg .b32 %r<3>;2605; CHECK-NEXT: .reg .b64 %rd<2>;2606; CHECK-EMPTY:2607; CHECK-NEXT: // %bb.0:2608; CHECK-NEXT: ld.param.b64 %rd1, [local_float_param_0];2609; CHECK-NEXT: ld.local.b32 %r1, [%rd1];2610; CHECK-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;2611; CHECK-NEXT: st.local.b32 [%rd1], %r2;2612; CHECK-NEXT: ret;2613 %a.load = load float, ptr addrspace(5) %a2614 %a.add = fadd float %a.load, 1.2615 store float %a.add, ptr addrspace(5) %a2616 ret void2617}2618 2619define void @local_double(ptr addrspace(5) %a) {2620; CHECK-LABEL: local_double(2621; CHECK: {2622; CHECK-NEXT: .reg .b64 %rd<4>;2623; CHECK-EMPTY:2624; CHECK-NEXT: // %bb.0:2625; CHECK-NEXT: ld.param.b64 %rd1, [local_double_param_0];2626; CHECK-NEXT: ld.local.b64 %rd2, [%rd1];2627; CHECK-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;2628; CHECK-NEXT: st.local.b64 [%rd1], %rd3;2629; CHECK-NEXT: ret;2630 %a.load = load double, ptr addrspace(5) %a2631 %a.add = fadd double %a.load, 1.2632 store double %a.add, ptr addrspace(5) %a2633 ret void2634}2635 2636; local_volatile2637 2638define void @local_volatile_i8(ptr addrspace(5) %a) {2639; CHECK-LABEL: local_volatile_i8(2640; CHECK: {2641; CHECK-NEXT: .reg .b16 %rs<3>;2642; CHECK-NEXT: .reg .b64 %rd<2>;2643; CHECK-EMPTY:2644; CHECK-NEXT: // %bb.0:2645; CHECK-NEXT: ld.param.b64 %rd1, [local_volatile_i8_param_0];2646; CHECK-NEXT: ld.local.b8 %rs1, [%rd1];2647; CHECK-NEXT: add.s16 %rs2, %rs1, 1;2648; CHECK-NEXT: st.local.b8 [%rd1], %rs2;2649; CHECK-NEXT: ret;2650 %a.load = load volatile i8, ptr addrspace(5) %a2651 %a.add = add i8 %a.load, 12652 store volatile i8 %a.add, ptr addrspace(5) %a2653 ret void2654}2655 2656define void @local_volatile_i16(ptr addrspace(5) %a) {2657; CHECK-LABEL: local_volatile_i16(2658; CHECK: {2659; CHECK-NEXT: .reg .b16 %rs<3>;2660; CHECK-NEXT: .reg .b64 %rd<2>;2661; CHECK-EMPTY:2662; CHECK-NEXT: // %bb.0:2663; CHECK-NEXT: ld.param.b64 %rd1, [local_volatile_i16_param_0];2664; CHECK-NEXT: ld.local.b16 %rs1, [%rd1];2665; CHECK-NEXT: add.s16 %rs2, %rs1, 1;2666; CHECK-NEXT: st.local.b16 [%rd1], %rs2;2667; CHECK-NEXT: ret;2668 %a.load = load volatile i16, ptr addrspace(5) %a2669 %a.add = add i16 %a.load, 12670 store volatile i16 %a.add, ptr addrspace(5) %a2671 ret void2672}2673 2674define void @local_volatile_i32(ptr addrspace(5) %a) {2675; CHECK-LABEL: local_volatile_i32(2676; CHECK: {2677; CHECK-NEXT: .reg .b32 %r<3>;2678; CHECK-NEXT: .reg .b64 %rd<2>;2679; CHECK-EMPTY:2680; CHECK-NEXT: // %bb.0:2681; CHECK-NEXT: ld.param.b64 %rd1, [local_volatile_i32_param_0];2682; CHECK-NEXT: ld.local.b32 %r1, [%rd1];2683; CHECK-NEXT: add.s32 %r2, %r1, 1;2684; CHECK-NEXT: st.local.b32 [%rd1], %r2;2685; CHECK-NEXT: ret;2686 %a.load = load volatile i32, ptr addrspace(5) %a2687 %a.add = add i32 %a.load, 12688 store volatile i32 %a.add, ptr addrspace(5) %a2689 ret void2690}2691 2692define void @local_volatile_i64(ptr addrspace(5) %a) {2693; CHECK-LABEL: local_volatile_i64(2694; CHECK: {2695; CHECK-NEXT: .reg .b64 %rd<4>;2696; CHECK-EMPTY:2697; CHECK-NEXT: // %bb.0:2698; CHECK-NEXT: ld.param.b64 %rd1, [local_volatile_i64_param_0];2699; CHECK-NEXT: ld.local.b64 %rd2, [%rd1];2700; CHECK-NEXT: add.s64 %rd3, %rd2, 1;2701; CHECK-NEXT: st.local.b64 [%rd1], %rd3;2702; CHECK-NEXT: ret;2703 %a.load = load volatile i64, ptr addrspace(5) %a2704 %a.add = add i64 %a.load, 12705 store volatile i64 %a.add, ptr addrspace(5) %a2706 ret void2707}2708 2709define void @local_volatile_float(ptr addrspace(5) %a) {2710; CHECK-LABEL: local_volatile_float(2711; CHECK: {2712; CHECK-NEXT: .reg .b32 %r<3>;2713; CHECK-NEXT: .reg .b64 %rd<2>;2714; CHECK-EMPTY:2715; CHECK-NEXT: // %bb.0:2716; CHECK-NEXT: ld.param.b64 %rd1, [local_volatile_float_param_0];2717; CHECK-NEXT: ld.local.b32 %r1, [%rd1];2718; CHECK-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;2719; CHECK-NEXT: st.local.b32 [%rd1], %r2;2720; CHECK-NEXT: ret;2721 %a.load = load volatile float, ptr addrspace(5) %a2722 %a.add = fadd float %a.load, 1.2723 store volatile float %a.add, ptr addrspace(5) %a2724 ret void2725}2726 2727define void @local_volatile_double(ptr addrspace(5) %a) {2728; CHECK-LABEL: local_volatile_double(2729; CHECK: {2730; CHECK-NEXT: .reg .b64 %rd<4>;2731; CHECK-EMPTY:2732; CHECK-NEXT: // %bb.0:2733; CHECK-NEXT: ld.param.b64 %rd1, [local_volatile_double_param_0];2734; CHECK-NEXT: ld.local.b64 %rd2, [%rd1];2735; CHECK-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;2736; CHECK-NEXT: st.local.b64 [%rd1], %rd3;2737; CHECK-NEXT: ret;2738 %a.load = load volatile double, ptr addrspace(5) %a2739 %a.add = fadd double %a.load, 1.2740 store volatile double %a.add, ptr addrspace(5) %a2741 ret void2742}2743 2744; local_unordered_sys2745 2746define void @local_unordered_sys_i8(ptr addrspace(5) %a) {2747; CHECK-LABEL: local_unordered_sys_i8(2748; CHECK: {2749; CHECK-NEXT: .reg .b16 %rs<3>;2750; CHECK-NEXT: .reg .b64 %rd<2>;2751; CHECK-EMPTY:2752; CHECK-NEXT: // %bb.0:2753; CHECK-NEXT: ld.param.b64 %rd1, [local_unordered_sys_i8_param_0];2754; CHECK-NEXT: ld.local.b8 %rs1, [%rd1];2755; CHECK-NEXT: add.s16 %rs2, %rs1, 1;2756; CHECK-NEXT: st.local.b8 [%rd1], %rs2;2757; CHECK-NEXT: ret;2758 %a.load = load atomic i8, ptr addrspace(5) %a unordered, align 12759 %a.add = add i8 %a.load, 12760 store atomic i8 %a.add, ptr addrspace(5) %a unordered, align 12761 ret void2762}2763 2764define void @local_unordered_sys_i16(ptr addrspace(5) %a) {2765; CHECK-LABEL: local_unordered_sys_i16(2766; CHECK: {2767; CHECK-NEXT: .reg .b16 %rs<3>;2768; CHECK-NEXT: .reg .b64 %rd<2>;2769; CHECK-EMPTY:2770; CHECK-NEXT: // %bb.0:2771; CHECK-NEXT: ld.param.b64 %rd1, [local_unordered_sys_i16_param_0];2772; CHECK-NEXT: ld.local.b16 %rs1, [%rd1];2773; CHECK-NEXT: add.s16 %rs2, %rs1, 1;2774; CHECK-NEXT: st.local.b16 [%rd1], %rs2;2775; CHECK-NEXT: ret;2776 %a.load = load atomic i16, ptr addrspace(5) %a unordered, align 22777 %a.add = add i16 %a.load, 12778 store atomic i16 %a.add, ptr addrspace(5) %a unordered, align 22779 ret void2780}2781 2782define void @local_unordered_sys_i32(ptr addrspace(5) %a) {2783; CHECK-LABEL: local_unordered_sys_i32(2784; CHECK: {2785; CHECK-NEXT: .reg .b32 %r<3>;2786; CHECK-NEXT: .reg .b64 %rd<2>;2787; CHECK-EMPTY:2788; CHECK-NEXT: // %bb.0:2789; CHECK-NEXT: ld.param.b64 %rd1, [local_unordered_sys_i32_param_0];2790; CHECK-NEXT: ld.local.b32 %r1, [%rd1];2791; CHECK-NEXT: add.s32 %r2, %r1, 1;2792; CHECK-NEXT: st.local.b32 [%rd1], %r2;2793; CHECK-NEXT: ret;2794 %a.load = load atomic i32, ptr addrspace(5) %a unordered, align 42795 %a.add = add i32 %a.load, 12796 store atomic i32 %a.add, ptr addrspace(5) %a unordered, align 42797 ret void2798}2799 2800define void @local_unordered_sys_i64(ptr addrspace(5) %a) {2801; CHECK-LABEL: local_unordered_sys_i64(2802; CHECK: {2803; CHECK-NEXT: .reg .b64 %rd<4>;2804; CHECK-EMPTY:2805; CHECK-NEXT: // %bb.0:2806; CHECK-NEXT: ld.param.b64 %rd1, [local_unordered_sys_i64_param_0];2807; CHECK-NEXT: ld.local.b64 %rd2, [%rd1];2808; CHECK-NEXT: add.s64 %rd3, %rd2, 1;2809; CHECK-NEXT: st.local.b64 [%rd1], %rd3;2810; CHECK-NEXT: ret;2811 %a.load = load atomic i64, ptr addrspace(5) %a unordered, align 82812 %a.add = add i64 %a.load, 12813 store atomic i64 %a.add, ptr addrspace(5) %a unordered, align 82814 ret void2815}2816 2817define void @local_unordered_sys_float(ptr addrspace(5) %a) {2818; CHECK-LABEL: local_unordered_sys_float(2819; CHECK: {2820; CHECK-NEXT: .reg .b32 %r<3>;2821; CHECK-NEXT: .reg .b64 %rd<2>;2822; CHECK-EMPTY:2823; CHECK-NEXT: // %bb.0:2824; CHECK-NEXT: ld.param.b64 %rd1, [local_unordered_sys_float_param_0];2825; CHECK-NEXT: ld.local.b32 %r1, [%rd1];2826; CHECK-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;2827; CHECK-NEXT: st.local.b32 [%rd1], %r2;2828; CHECK-NEXT: ret;2829 %a.load = load atomic float, ptr addrspace(5) %a unordered, align 42830 %a.add = fadd float %a.load, 1.2831 store atomic float %a.add, ptr addrspace(5) %a unordered, align 42832 ret void2833}2834 2835define void @local_unordered_sys_double(ptr addrspace(5) %a) {2836; CHECK-LABEL: local_unordered_sys_double(2837; CHECK: {2838; CHECK-NEXT: .reg .b64 %rd<4>;2839; CHECK-EMPTY:2840; CHECK-NEXT: // %bb.0:2841; CHECK-NEXT: ld.param.b64 %rd1, [local_unordered_sys_double_param_0];2842; CHECK-NEXT: ld.local.b64 %rd2, [%rd1];2843; CHECK-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;2844; CHECK-NEXT: st.local.b64 [%rd1], %rd3;2845; CHECK-NEXT: ret;2846 %a.load = load atomic double, ptr addrspace(5) %a unordered, align 82847 %a.add = fadd double %a.load, 1.2848 store atomic double %a.add, ptr addrspace(5) %a unordered, align 82849 ret void2850}2851 2852; local_unordered_volatile_sys2853 2854define void @local_unordered_volatile_sys_i8(ptr addrspace(5) %a) {2855; CHECK-LABEL: local_unordered_volatile_sys_i8(2856; CHECK: {2857; CHECK-NEXT: .reg .b16 %rs<3>;2858; CHECK-NEXT: .reg .b64 %rd<2>;2859; CHECK-EMPTY:2860; CHECK-NEXT: // %bb.0:2861; CHECK-NEXT: ld.param.b64 %rd1, [local_unordered_volatile_sys_i8_param_0];2862; CHECK-NEXT: ld.local.b8 %rs1, [%rd1];2863; CHECK-NEXT: add.s16 %rs2, %rs1, 1;2864; CHECK-NEXT: st.local.b8 [%rd1], %rs2;2865; CHECK-NEXT: ret;2866 %a.load = load atomic volatile i8, ptr addrspace(5) %a unordered, align 12867 %a.add = add i8 %a.load, 12868 store atomic volatile i8 %a.add, ptr addrspace(5) %a unordered, align 12869 ret void2870}2871 2872define void @local_unordered_volatile_sys_i16(ptr addrspace(5) %a) {2873; CHECK-LABEL: local_unordered_volatile_sys_i16(2874; CHECK: {2875; CHECK-NEXT: .reg .b16 %rs<3>;2876; CHECK-NEXT: .reg .b64 %rd<2>;2877; CHECK-EMPTY:2878; CHECK-NEXT: // %bb.0:2879; CHECK-NEXT: ld.param.b64 %rd1, [local_unordered_volatile_sys_i16_param_0];2880; CHECK-NEXT: ld.local.b16 %rs1, [%rd1];2881; CHECK-NEXT: add.s16 %rs2, %rs1, 1;2882; CHECK-NEXT: st.local.b16 [%rd1], %rs2;2883; CHECK-NEXT: ret;2884 %a.load = load atomic volatile i16, ptr addrspace(5) %a unordered, align 22885 %a.add = add i16 %a.load, 12886 store atomic volatile i16 %a.add, ptr addrspace(5) %a unordered, align 22887 ret void2888}2889 2890define void @local_unordered_volatile_sys_i32(ptr addrspace(5) %a) {2891; CHECK-LABEL: local_unordered_volatile_sys_i32(2892; CHECK: {2893; CHECK-NEXT: .reg .b32 %r<3>;2894; CHECK-NEXT: .reg .b64 %rd<2>;2895; CHECK-EMPTY:2896; CHECK-NEXT: // %bb.0:2897; CHECK-NEXT: ld.param.b64 %rd1, [local_unordered_volatile_sys_i32_param_0];2898; CHECK-NEXT: ld.local.b32 %r1, [%rd1];2899; CHECK-NEXT: add.s32 %r2, %r1, 1;2900; CHECK-NEXT: st.local.b32 [%rd1], %r2;2901; CHECK-NEXT: ret;2902 %a.load = load atomic volatile i32, ptr addrspace(5) %a unordered, align 42903 %a.add = add i32 %a.load, 12904 store atomic volatile i32 %a.add, ptr addrspace(5) %a unordered, align 42905 ret void2906}2907 2908define void @local_unordered_volatile_sys_i64(ptr addrspace(5) %a) {2909; CHECK-LABEL: local_unordered_volatile_sys_i64(2910; CHECK: {2911; CHECK-NEXT: .reg .b64 %rd<4>;2912; CHECK-EMPTY:2913; CHECK-NEXT: // %bb.0:2914; CHECK-NEXT: ld.param.b64 %rd1, [local_unordered_volatile_sys_i64_param_0];2915; CHECK-NEXT: ld.local.b64 %rd2, [%rd1];2916; CHECK-NEXT: add.s64 %rd3, %rd2, 1;2917; CHECK-NEXT: st.local.b64 [%rd1], %rd3;2918; CHECK-NEXT: ret;2919 %a.load = load atomic volatile i64, ptr addrspace(5) %a unordered, align 82920 %a.add = add i64 %a.load, 12921 store atomic volatile i64 %a.add, ptr addrspace(5) %a unordered, align 82922 ret void2923}2924 2925define void @local_unordered_volatile_sys_float(ptr addrspace(5) %a) {2926; CHECK-LABEL: local_unordered_volatile_sys_float(2927; CHECK: {2928; CHECK-NEXT: .reg .b32 %r<3>;2929; CHECK-NEXT: .reg .b64 %rd<2>;2930; CHECK-EMPTY:2931; CHECK-NEXT: // %bb.0:2932; CHECK-NEXT: ld.param.b64 %rd1, [local_unordered_volatile_sys_float_param_0];2933; CHECK-NEXT: ld.local.b32 %r1, [%rd1];2934; CHECK-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;2935; CHECK-NEXT: st.local.b32 [%rd1], %r2;2936; CHECK-NEXT: ret;2937 %a.load = load atomic volatile float, ptr addrspace(5) %a unordered, align 42938 %a.add = fadd float %a.load, 1.2939 store atomic volatile float %a.add, ptr addrspace(5) %a unordered, align 42940 ret void2941}2942 2943define void @local_unordered_volatile_sys_double(ptr addrspace(5) %a) {2944; CHECK-LABEL: local_unordered_volatile_sys_double(2945; CHECK: {2946; CHECK-NEXT: .reg .b64 %rd<4>;2947; CHECK-EMPTY:2948; CHECK-NEXT: // %bb.0:2949; CHECK-NEXT: ld.param.b64 %rd1, [local_unordered_volatile_sys_double_param_0];2950; CHECK-NEXT: ld.local.b64 %rd2, [%rd1];2951; CHECK-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;2952; CHECK-NEXT: st.local.b64 [%rd1], %rd3;2953; CHECK-NEXT: ret;2954 %a.load = load atomic volatile double, ptr addrspace(5) %a unordered, align 82955 %a.add = fadd double %a.load, 1.2956 store atomic volatile double %a.add, ptr addrspace(5) %a unordered, align 82957 ret void2958}2959 2960; local_monotonic_sys2961 2962define void @local_monotonic_sys_i8(ptr addrspace(5) %a) {2963; CHECK-LABEL: local_monotonic_sys_i8(2964; CHECK: {2965; CHECK-NEXT: .reg .b16 %rs<3>;2966; CHECK-NEXT: .reg .b64 %rd<2>;2967; CHECK-EMPTY:2968; CHECK-NEXT: // %bb.0:2969; CHECK-NEXT: ld.param.b64 %rd1, [local_monotonic_sys_i8_param_0];2970; CHECK-NEXT: ld.local.b8 %rs1, [%rd1];2971; CHECK-NEXT: add.s16 %rs2, %rs1, 1;2972; CHECK-NEXT: st.local.b8 [%rd1], %rs2;2973; CHECK-NEXT: ret;2974 %a.load = load atomic i8, ptr addrspace(5) %a monotonic, align 12975 %a.add = add i8 %a.load, 12976 store atomic i8 %a.add, ptr addrspace(5) %a monotonic, align 12977 ret void2978}2979 2980define void @local_monotonic_sys_i16(ptr addrspace(5) %a) {2981; CHECK-LABEL: local_monotonic_sys_i16(2982; CHECK: {2983; CHECK-NEXT: .reg .b16 %rs<3>;2984; CHECK-NEXT: .reg .b64 %rd<2>;2985; CHECK-EMPTY:2986; CHECK-NEXT: // %bb.0:2987; CHECK-NEXT: ld.param.b64 %rd1, [local_monotonic_sys_i16_param_0];2988; CHECK-NEXT: ld.local.b16 %rs1, [%rd1];2989; CHECK-NEXT: add.s16 %rs2, %rs1, 1;2990; CHECK-NEXT: st.local.b16 [%rd1], %rs2;2991; CHECK-NEXT: ret;2992 %a.load = load atomic i16, ptr addrspace(5) %a monotonic, align 22993 %a.add = add i16 %a.load, 12994 store atomic i16 %a.add, ptr addrspace(5) %a monotonic, align 22995 ret void2996}2997 2998define void @local_monotonic_sys_i32(ptr addrspace(5) %a) {2999; CHECK-LABEL: local_monotonic_sys_i32(3000; CHECK: {3001; CHECK-NEXT: .reg .b32 %r<3>;3002; CHECK-NEXT: .reg .b64 %rd<2>;3003; CHECK-EMPTY:3004; CHECK-NEXT: // %bb.0:3005; CHECK-NEXT: ld.param.b64 %rd1, [local_monotonic_sys_i32_param_0];3006; CHECK-NEXT: ld.local.b32 %r1, [%rd1];3007; CHECK-NEXT: add.s32 %r2, %r1, 1;3008; CHECK-NEXT: st.local.b32 [%rd1], %r2;3009; CHECK-NEXT: ret;3010 %a.load = load atomic i32, ptr addrspace(5) %a monotonic, align 43011 %a.add = add i32 %a.load, 13012 store atomic i32 %a.add, ptr addrspace(5) %a monotonic, align 43013 ret void3014}3015 3016define void @local_monotonic_sys_i64(ptr addrspace(5) %a) {3017; CHECK-LABEL: local_monotonic_sys_i64(3018; CHECK: {3019; CHECK-NEXT: .reg .b64 %rd<4>;3020; CHECK-EMPTY:3021; CHECK-NEXT: // %bb.0:3022; CHECK-NEXT: ld.param.b64 %rd1, [local_monotonic_sys_i64_param_0];3023; CHECK-NEXT: ld.local.b64 %rd2, [%rd1];3024; CHECK-NEXT: add.s64 %rd3, %rd2, 1;3025; CHECK-NEXT: st.local.b64 [%rd1], %rd3;3026; CHECK-NEXT: ret;3027 %a.load = load atomic i64, ptr addrspace(5) %a monotonic, align 83028 %a.add = add i64 %a.load, 13029 store atomic i64 %a.add, ptr addrspace(5) %a monotonic, align 83030 ret void3031}3032 3033define void @local_monotonic_sys_float(ptr addrspace(5) %a) {3034; CHECK-LABEL: local_monotonic_sys_float(3035; CHECK: {3036; CHECK-NEXT: .reg .b32 %r<3>;3037; CHECK-NEXT: .reg .b64 %rd<2>;3038; CHECK-EMPTY:3039; CHECK-NEXT: // %bb.0:3040; CHECK-NEXT: ld.param.b64 %rd1, [local_monotonic_sys_float_param_0];3041; CHECK-NEXT: ld.local.b32 %r1, [%rd1];3042; CHECK-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;3043; CHECK-NEXT: st.local.b32 [%rd1], %r2;3044; CHECK-NEXT: ret;3045 %a.load = load atomic float, ptr addrspace(5) %a monotonic, align 43046 %a.add = fadd float %a.load, 1.3047 store atomic float %a.add, ptr addrspace(5) %a monotonic, align 43048 ret void3049}3050 3051define void @local_monotonic_sys_double(ptr addrspace(5) %a) {3052; CHECK-LABEL: local_monotonic_sys_double(3053; CHECK: {3054; CHECK-NEXT: .reg .b64 %rd<4>;3055; CHECK-EMPTY:3056; CHECK-NEXT: // %bb.0:3057; CHECK-NEXT: ld.param.b64 %rd1, [local_monotonic_sys_double_param_0];3058; CHECK-NEXT: ld.local.b64 %rd2, [%rd1];3059; CHECK-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;3060; CHECK-NEXT: st.local.b64 [%rd1], %rd3;3061; CHECK-NEXT: ret;3062 %a.load = load atomic double, ptr addrspace(5) %a monotonic, align 83063 %a.add = fadd double %a.load, 1.3064 store atomic double %a.add, ptr addrspace(5) %a monotonic, align 83065 ret void3066}3067 3068; local_monotonic_volatile_sys3069 3070define void @local_monotonic_volatile_sys_i8(ptr addrspace(5) %a) {3071; CHECK-LABEL: local_monotonic_volatile_sys_i8(3072; CHECK: {3073; CHECK-NEXT: .reg .b16 %rs<3>;3074; CHECK-NEXT: .reg .b64 %rd<2>;3075; CHECK-EMPTY:3076; CHECK-NEXT: // %bb.0:3077; CHECK-NEXT: ld.param.b64 %rd1, [local_monotonic_volatile_sys_i8_param_0];3078; CHECK-NEXT: ld.local.b8 %rs1, [%rd1];3079; CHECK-NEXT: add.s16 %rs2, %rs1, 1;3080; CHECK-NEXT: st.local.b8 [%rd1], %rs2;3081; CHECK-NEXT: ret;3082 %a.load = load atomic volatile i8, ptr addrspace(5) %a monotonic, align 13083 %a.add = add i8 %a.load, 13084 store atomic volatile i8 %a.add, ptr addrspace(5) %a monotonic, align 13085 ret void3086}3087 3088define void @local_monotonic_volatile_sys_i16(ptr addrspace(5) %a) {3089; CHECK-LABEL: local_monotonic_volatile_sys_i16(3090; CHECK: {3091; CHECK-NEXT: .reg .b16 %rs<3>;3092; CHECK-NEXT: .reg .b64 %rd<2>;3093; CHECK-EMPTY:3094; CHECK-NEXT: // %bb.0:3095; CHECK-NEXT: ld.param.b64 %rd1, [local_monotonic_volatile_sys_i16_param_0];3096; CHECK-NEXT: ld.local.b16 %rs1, [%rd1];3097; CHECK-NEXT: add.s16 %rs2, %rs1, 1;3098; CHECK-NEXT: st.local.b16 [%rd1], %rs2;3099; CHECK-NEXT: ret;3100 %a.load = load atomic volatile i16, ptr addrspace(5) %a monotonic, align 23101 %a.add = add i16 %a.load, 13102 store atomic volatile i16 %a.add, ptr addrspace(5) %a monotonic, align 23103 ret void3104}3105 3106define void @local_monotonic_volatile_sys_i32(ptr addrspace(5) %a) {3107; CHECK-LABEL: local_monotonic_volatile_sys_i32(3108; CHECK: {3109; CHECK-NEXT: .reg .b32 %r<3>;3110; CHECK-NEXT: .reg .b64 %rd<2>;3111; CHECK-EMPTY:3112; CHECK-NEXT: // %bb.0:3113; CHECK-NEXT: ld.param.b64 %rd1, [local_monotonic_volatile_sys_i32_param_0];3114; CHECK-NEXT: ld.local.b32 %r1, [%rd1];3115; CHECK-NEXT: add.s32 %r2, %r1, 1;3116; CHECK-NEXT: st.local.b32 [%rd1], %r2;3117; CHECK-NEXT: ret;3118 %a.load = load atomic volatile i32, ptr addrspace(5) %a monotonic, align 43119 %a.add = add i32 %a.load, 13120 store atomic volatile i32 %a.add, ptr addrspace(5) %a monotonic, align 43121 ret void3122}3123 3124define void @local_monotonic_volatile_sys_i64(ptr addrspace(5) %a) {3125; CHECK-LABEL: local_monotonic_volatile_sys_i64(3126; CHECK: {3127; CHECK-NEXT: .reg .b64 %rd<4>;3128; CHECK-EMPTY:3129; CHECK-NEXT: // %bb.0:3130; CHECK-NEXT: ld.param.b64 %rd1, [local_monotonic_volatile_sys_i64_param_0];3131; CHECK-NEXT: ld.local.b64 %rd2, [%rd1];3132; CHECK-NEXT: add.s64 %rd3, %rd2, 1;3133; CHECK-NEXT: st.local.b64 [%rd1], %rd3;3134; CHECK-NEXT: ret;3135 %a.load = load atomic volatile i64, ptr addrspace(5) %a monotonic, align 83136 %a.add = add i64 %a.load, 13137 store atomic volatile i64 %a.add, ptr addrspace(5) %a monotonic, align 83138 ret void3139}3140 3141define void @local_monotonic_volatile_sys_float(ptr addrspace(5) %a) {3142; CHECK-LABEL: local_monotonic_volatile_sys_float(3143; CHECK: {3144; CHECK-NEXT: .reg .b32 %r<3>;3145; CHECK-NEXT: .reg .b64 %rd<2>;3146; CHECK-EMPTY:3147; CHECK-NEXT: // %bb.0:3148; CHECK-NEXT: ld.param.b64 %rd1, [local_monotonic_volatile_sys_float_param_0];3149; CHECK-NEXT: ld.local.b32 %r1, [%rd1];3150; CHECK-NEXT: add.rn.f32 %r2, %r1, 0f3F800000;3151; CHECK-NEXT: st.local.b32 [%rd1], %r2;3152; CHECK-NEXT: ret;3153 %a.load = load atomic volatile float, ptr addrspace(5) %a monotonic, align 43154 %a.add = fadd float %a.load, 1.3155 store atomic volatile float %a.add, ptr addrspace(5) %a monotonic, align 43156 ret void3157}3158 3159define void @local_monotonic_volatile_sys_double(ptr addrspace(5) %a) {3160; CHECK-LABEL: local_monotonic_volatile_sys_double(3161; CHECK: {3162; CHECK-NEXT: .reg .b64 %rd<4>;3163; CHECK-EMPTY:3164; CHECK-NEXT: // %bb.0:3165; CHECK-NEXT: ld.param.b64 %rd1, [local_monotonic_volatile_sys_double_param_0];3166; CHECK-NEXT: ld.local.b64 %rd2, [%rd1];3167; CHECK-NEXT: add.rn.f64 %rd3, %rd2, 0d3FF0000000000000;3168; CHECK-NEXT: st.local.b64 [%rd1], %rd3;3169; CHECK-NEXT: ret;3170 %a.load = load atomic volatile double, ptr addrspace(5) %a monotonic, align 83171 %a.add = fadd double %a.load, 1.3172 store atomic volatile double %a.add, ptr addrspace(5) %a monotonic, align 83173 ret void3174}3175