2033 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_90 -mattr=+ptx78 | FileCheck %s3; RUN: %if ptxas-sm_90 && ptxas-isa-7.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx78 | %ptxas-verify -arch=sm_90 %}4 5; TODO: fix "atomic load volatile acquire": generates "ld.acquire.sys;"6; but should generate "ld.mmio.relaxed.sys; fence.acq_rel.sys;"7; TODO: fix "atomic store volatile release": generates "st.release.sys;"8; but should generate "fence.acq_rel.sys; st.mmio.relaxed.sys;"9 10; TODO: fix "atomic load volatile seq_cst": generates "fence.sc.sys; ld.acquire.sys;"11; but should generate "fence.sc.sys; ld.relaxed.mmio.sys; fence.acq_rel.sys;"12; TODO: fix "atomic store volatile seq_cst": generates "fence.sc.sys; st.release.sys;"13; but should generate "fence.sc.sys; st.relaxed.mmio.sys;"14 15; TODO: add i1, <8 x i8>, and <6 x i8> vector tests.16 17; TODO: add test for vectors that exceed 128-bit length18; Per https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#vectors19; vectors cannot exceed 128-bit in length, i.e., .v4.u64 is not allowed.20 21; TODO: generate PTX that preserves Concurrent Forward Progress22; for atomic operations to local statespace23; by generating atomic or volatile operations.24 25; TODO: design exposure for atomic operations on vector types.26 27; TODO: implement and test thread scope.28 29; TODO: add weak,atomic,volatile,atomic volatile tests30; for .const and .param statespaces.31 32; TODO: optimize .shared.sys into .shared.cta or .shared.cluster .33 34;; generic statespace35 36; CHECK-LABEL: generic_unordered_cluster37define void @generic_unordered_cluster(ptr %a, ptr %b, ptr %c, ptr %d, ptr %e) local_unnamed_addr {38; CHECK-LABEL: generic_unordered_cluster(39; CHECK: {40; CHECK-NEXT: .reg .b16 %rs<5>;41; CHECK-NEXT: .reg .b32 %r<5>;42; CHECK-NEXT: .reg .b64 %rd<10>;43; CHECK-EMPTY:44; CHECK-NEXT: // %bb.0:45; CHECK-NEXT: ld.param.b64 %rd1, [generic_unordered_cluster_param_0];46; CHECK-NEXT: ld.relaxed.cluster.b8 %rs1, [%rd1];47; CHECK-NEXT: ld.param.b64 %rd2, [generic_unordered_cluster_param_1];48; CHECK-NEXT: add.s16 %rs2, %rs1, 1;49; CHECK-NEXT: ld.param.b64 %rd3, [generic_unordered_cluster_param_2];50; CHECK-NEXT: st.relaxed.cluster.b8 [%rd1], %rs2;51; CHECK-NEXT: ld.param.b64 %rd4, [generic_unordered_cluster_param_3];52; CHECK-NEXT: ld.relaxed.cluster.b16 %rs3, [%rd2];53; CHECK-NEXT: ld.param.b64 %rd5, [generic_unordered_cluster_param_4];54; CHECK-NEXT: add.s16 %rs4, %rs3, 1;55; CHECK-NEXT: st.relaxed.cluster.b16 [%rd2], %rs4;56; CHECK-NEXT: ld.relaxed.cluster.b32 %r1, [%rd3];57; CHECK-NEXT: add.s32 %r2, %r1, 1;58; CHECK-NEXT: st.relaxed.cluster.b32 [%rd3], %r2;59; CHECK-NEXT: ld.relaxed.cluster.b64 %rd6, [%rd4];60; CHECK-NEXT: add.s64 %rd7, %rd6, 1;61; CHECK-NEXT: st.relaxed.cluster.b64 [%rd4], %rd7;62; CHECK-NEXT: ld.relaxed.cluster.b32 %r3, [%rd5];63; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;64; CHECK-NEXT: st.relaxed.cluster.b32 [%rd5], %r4;65; CHECK-NEXT: ld.relaxed.cluster.b64 %rd8, [%rd5];66; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;67; CHECK-NEXT: st.relaxed.cluster.b64 [%rd5], %rd9;68; CHECK-NEXT: ret;69 %a.load = load atomic i8, ptr %a syncscope("cluster") unordered, align 170 %a.add = add i8 %a.load, 171 store atomic i8 %a.add, ptr %a syncscope("cluster") unordered, align 172 73 %b.load = load atomic i16, ptr %b syncscope("cluster") unordered, align 274 %b.add = add i16 %b.load, 175 store atomic i16 %b.add, ptr %b syncscope("cluster") unordered, align 276 77 %c.load = load atomic i32, ptr %c syncscope("cluster") unordered, align 478 %c.add = add i32 %c.load, 179 store atomic i32 %c.add, ptr %c syncscope("cluster") unordered, align 480 81 %d.load = load atomic i64, ptr %d syncscope("cluster") unordered, align 882 %d.add = add i64 %d.load, 183 store atomic i64 %d.add, ptr %d syncscope("cluster") unordered, align 884 85 %e.load = load atomic float, ptr %e syncscope("cluster") unordered, align 486 %e.add = fadd float %e.load, 1.87 store atomic float %e.add, ptr %e syncscope("cluster") unordered, align 488 89 %f.load = load atomic double, ptr %e syncscope("cluster") unordered, align 890 %f.add = fadd double %f.load, 1.91 store atomic double %f.add, ptr %e syncscope("cluster") unordered, align 892 93 ret void94}95 96; CHECK-LABEL: generic_unordered_volatile_cluster97define void @generic_unordered_volatile_cluster(ptr %a, ptr %b, ptr %c, ptr %d, ptr %e) local_unnamed_addr {98; CHECK-LABEL: generic_unordered_volatile_cluster(99; CHECK: {100; CHECK-NEXT: .reg .b16 %rs<5>;101; CHECK-NEXT: .reg .b32 %r<5>;102; CHECK-NEXT: .reg .b64 %rd<10>;103; CHECK-EMPTY:104; CHECK-NEXT: // %bb.0:105; CHECK-NEXT: ld.param.b64 %rd1, [generic_unordered_volatile_cluster_param_0];106; CHECK-NEXT: ld.volatile.b8 %rs1, [%rd1];107; CHECK-NEXT: ld.param.b64 %rd2, [generic_unordered_volatile_cluster_param_1];108; CHECK-NEXT: add.s16 %rs2, %rs1, 1;109; CHECK-NEXT: ld.param.b64 %rd3, [generic_unordered_volatile_cluster_param_2];110; CHECK-NEXT: st.volatile.b8 [%rd1], %rs2;111; CHECK-NEXT: ld.param.b64 %rd4, [generic_unordered_volatile_cluster_param_3];112; CHECK-NEXT: ld.volatile.b16 %rs3, [%rd2];113; CHECK-NEXT: ld.param.b64 %rd5, [generic_unordered_volatile_cluster_param_4];114; CHECK-NEXT: add.s16 %rs4, %rs3, 1;115; CHECK-NEXT: st.volatile.b16 [%rd2], %rs4;116; CHECK-NEXT: ld.volatile.b32 %r1, [%rd3];117; CHECK-NEXT: add.s32 %r2, %r1, 1;118; CHECK-NEXT: st.volatile.b32 [%rd3], %r2;119; CHECK-NEXT: ld.volatile.b64 %rd6, [%rd4];120; CHECK-NEXT: add.s64 %rd7, %rd6, 1;121; CHECK-NEXT: st.volatile.b64 [%rd4], %rd7;122; CHECK-NEXT: ld.volatile.b32 %r3, [%rd5];123; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;124; CHECK-NEXT: st.volatile.b32 [%rd5], %r4;125; CHECK-NEXT: ld.volatile.b64 %rd8, [%rd5];126; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;127; CHECK-NEXT: st.volatile.b64 [%rd5], %rd9;128; CHECK-NEXT: ret;129 %a.load = load atomic volatile i8, ptr %a syncscope("cluster") unordered, align 1130 %a.add = add i8 %a.load, 1131 store atomic volatile i8 %a.add, ptr %a syncscope("cluster") unordered, align 1132 133 %b.load = load atomic volatile i16, ptr %b syncscope("cluster") unordered, align 2134 %b.add = add i16 %b.load, 1135 store atomic volatile i16 %b.add, ptr %b syncscope("cluster") unordered, align 2136 137 %c.load = load atomic volatile i32, ptr %c syncscope("cluster") unordered, align 4138 %c.add = add i32 %c.load, 1139 store atomic volatile i32 %c.add, ptr %c syncscope("cluster") unordered, align 4140 141 %d.load = load atomic volatile i64, ptr %d syncscope("cluster") unordered, align 8142 %d.add = add i64 %d.load, 1143 store atomic volatile i64 %d.add, ptr %d syncscope("cluster") unordered, align 8144 145 %e.load = load atomic volatile float, ptr %e syncscope("cluster") unordered, align 4146 %e.add = fadd float %e.load, 1.147 store atomic volatile float %e.add, ptr %e syncscope("cluster") unordered, align 4148 149 %f.load = load atomic volatile double, ptr %e syncscope("cluster") unordered, align 8150 %f.add = fadd double %f.load, 1.151 store atomic volatile double %f.add, ptr %e syncscope("cluster") unordered, align 8152 153 ret void154}155 156; CHECK-LABEL: generic_monotonic_cluster157define void @generic_monotonic_cluster(ptr %a, ptr %b, ptr %c, ptr %d, ptr %e) local_unnamed_addr {158; CHECK-LABEL: generic_monotonic_cluster(159; CHECK: {160; CHECK-NEXT: .reg .b16 %rs<5>;161; CHECK-NEXT: .reg .b32 %r<5>;162; CHECK-NEXT: .reg .b64 %rd<10>;163; CHECK-EMPTY:164; CHECK-NEXT: // %bb.0:165; CHECK-NEXT: ld.param.b64 %rd1, [generic_monotonic_cluster_param_0];166; CHECK-NEXT: ld.relaxed.cluster.b8 %rs1, [%rd1];167; CHECK-NEXT: ld.param.b64 %rd2, [generic_monotonic_cluster_param_1];168; CHECK-NEXT: add.s16 %rs2, %rs1, 1;169; CHECK-NEXT: ld.param.b64 %rd3, [generic_monotonic_cluster_param_2];170; CHECK-NEXT: st.relaxed.cluster.b8 [%rd1], %rs2;171; CHECK-NEXT: ld.param.b64 %rd4, [generic_monotonic_cluster_param_3];172; CHECK-NEXT: ld.relaxed.cluster.b16 %rs3, [%rd2];173; CHECK-NEXT: ld.param.b64 %rd5, [generic_monotonic_cluster_param_4];174; CHECK-NEXT: add.s16 %rs4, %rs3, 1;175; CHECK-NEXT: st.relaxed.cluster.b16 [%rd2], %rs4;176; CHECK-NEXT: ld.relaxed.cluster.b32 %r1, [%rd3];177; CHECK-NEXT: add.s32 %r2, %r1, 1;178; CHECK-NEXT: st.relaxed.cluster.b32 [%rd3], %r2;179; CHECK-NEXT: ld.relaxed.cluster.b64 %rd6, [%rd4];180; CHECK-NEXT: add.s64 %rd7, %rd6, 1;181; CHECK-NEXT: st.relaxed.cluster.b64 [%rd4], %rd7;182; CHECK-NEXT: ld.relaxed.cluster.b32 %r3, [%rd5];183; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;184; CHECK-NEXT: st.relaxed.cluster.b32 [%rd5], %r4;185; CHECK-NEXT: ld.relaxed.cluster.b64 %rd8, [%rd5];186; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;187; CHECK-NEXT: st.relaxed.cluster.b64 [%rd5], %rd9;188; CHECK-NEXT: ret;189 %a.load = load atomic i8, ptr %a syncscope("cluster") monotonic, align 1190 %a.add = add i8 %a.load, 1191 store atomic i8 %a.add, ptr %a syncscope("cluster") monotonic, align 1192 193 %b.load = load atomic i16, ptr %b syncscope("cluster") monotonic, align 2194 %b.add = add i16 %b.load, 1195 store atomic i16 %b.add, ptr %b syncscope("cluster") monotonic, align 2196 197 %c.load = load atomic i32, ptr %c syncscope("cluster") monotonic, align 4198 %c.add = add i32 %c.load, 1199 store atomic i32 %c.add, ptr %c syncscope("cluster") monotonic, align 4200 201 %d.load = load atomic i64, ptr %d syncscope("cluster") monotonic, align 8202 %d.add = add i64 %d.load, 1203 store atomic i64 %d.add, ptr %d syncscope("cluster") monotonic, align 8204 205 %e.load = load atomic float, ptr %e syncscope("cluster") monotonic, align 4206 %e.add = fadd float %e.load, 1.207 store atomic float %e.add, ptr %e syncscope("cluster") monotonic, align 4208 209 %f.load = load atomic double, ptr %e syncscope("cluster") monotonic, align 8210 %f.add = fadd double %f.load, 1.211 store atomic double %f.add, ptr %e syncscope("cluster") monotonic, align 8212 213 ret void214}215 216; CHECK-LABEL: generic_monotonic_volatile_cluster217define void @generic_monotonic_volatile_cluster(ptr %a, ptr %b, ptr %c, ptr %d, ptr %e) local_unnamed_addr {218; CHECK-LABEL: generic_monotonic_volatile_cluster(219; CHECK: {220; CHECK-NEXT: .reg .b16 %rs<5>;221; CHECK-NEXT: .reg .b32 %r<5>;222; CHECK-NEXT: .reg .b64 %rd<10>;223; CHECK-EMPTY:224; CHECK-NEXT: // %bb.0:225; CHECK-NEXT: ld.param.b64 %rd1, [generic_monotonic_volatile_cluster_param_0];226; CHECK-NEXT: ld.volatile.b8 %rs1, [%rd1];227; CHECK-NEXT: ld.param.b64 %rd2, [generic_monotonic_volatile_cluster_param_1];228; CHECK-NEXT: add.s16 %rs2, %rs1, 1;229; CHECK-NEXT: ld.param.b64 %rd3, [generic_monotonic_volatile_cluster_param_2];230; CHECK-NEXT: st.volatile.b8 [%rd1], %rs2;231; CHECK-NEXT: ld.param.b64 %rd4, [generic_monotonic_volatile_cluster_param_3];232; CHECK-NEXT: ld.volatile.b16 %rs3, [%rd2];233; CHECK-NEXT: ld.param.b64 %rd5, [generic_monotonic_volatile_cluster_param_4];234; CHECK-NEXT: add.s16 %rs4, %rs3, 1;235; CHECK-NEXT: st.volatile.b16 [%rd2], %rs4;236; CHECK-NEXT: ld.volatile.b32 %r1, [%rd3];237; CHECK-NEXT: add.s32 %r2, %r1, 1;238; CHECK-NEXT: st.volatile.b32 [%rd3], %r2;239; CHECK-NEXT: ld.volatile.b64 %rd6, [%rd4];240; CHECK-NEXT: add.s64 %rd7, %rd6, 1;241; CHECK-NEXT: st.volatile.b64 [%rd4], %rd7;242; CHECK-NEXT: ld.volatile.b32 %r3, [%rd5];243; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;244; CHECK-NEXT: st.volatile.b32 [%rd5], %r4;245; CHECK-NEXT: ld.volatile.b64 %rd8, [%rd5];246; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;247; CHECK-NEXT: st.volatile.b64 [%rd5], %rd9;248; CHECK-NEXT: ret;249 %a.load = load atomic volatile i8, ptr %a syncscope("cluster") monotonic, align 1250 %a.add = add i8 %a.load, 1251 store atomic volatile i8 %a.add, ptr %a syncscope("cluster") monotonic, align 1252 253 %b.load = load atomic volatile i16, ptr %b syncscope("cluster") monotonic, align 2254 %b.add = add i16 %b.load, 1255 store atomic volatile i16 %b.add, ptr %b syncscope("cluster") monotonic, align 2256 257 %c.load = load atomic volatile i32, ptr %c syncscope("cluster") monotonic, align 4258 %c.add = add i32 %c.load, 1259 store atomic volatile i32 %c.add, ptr %c syncscope("cluster") monotonic, align 4260 261 %d.load = load atomic volatile i64, ptr %d syncscope("cluster") monotonic, align 8262 %d.add = add i64 %d.load, 1263 store atomic volatile i64 %d.add, ptr %d syncscope("cluster") monotonic, align 8264 265 %e.load = load atomic volatile float, ptr %e syncscope("cluster") monotonic, align 4266 %e.add = fadd float %e.load, 1.267 store atomic volatile float %e.add, ptr %e syncscope("cluster") monotonic, align 4268 269 %f.load = load atomic volatile double, ptr %e syncscope("cluster") monotonic, align 8270 %f.add = fadd double %f.load, 1.271 store atomic volatile double %f.add, ptr %e syncscope("cluster") monotonic, align 8272 273 ret void274}275 276; CHECK-LABEL: generic_acq_rel_cluster277define void @generic_acq_rel_cluster(ptr %a, ptr %b, ptr %c, ptr %d, ptr %e) local_unnamed_addr {278; CHECK-LABEL: generic_acq_rel_cluster(279; CHECK: {280; CHECK-NEXT: .reg .b16 %rs<5>;281; CHECK-NEXT: .reg .b32 %r<5>;282; CHECK-NEXT: .reg .b64 %rd<10>;283; CHECK-EMPTY:284; CHECK-NEXT: // %bb.0:285; CHECK-NEXT: ld.param.b64 %rd1, [generic_acq_rel_cluster_param_0];286; CHECK-NEXT: ld.acquire.cluster.b8 %rs1, [%rd1];287; CHECK-NEXT: ld.param.b64 %rd2, [generic_acq_rel_cluster_param_1];288; CHECK-NEXT: add.s16 %rs2, %rs1, 1;289; CHECK-NEXT: ld.param.b64 %rd3, [generic_acq_rel_cluster_param_2];290; CHECK-NEXT: st.release.cluster.b8 [%rd1], %rs2;291; CHECK-NEXT: ld.param.b64 %rd4, [generic_acq_rel_cluster_param_3];292; CHECK-NEXT: ld.acquire.cluster.b16 %rs3, [%rd2];293; CHECK-NEXT: ld.param.b64 %rd5, [generic_acq_rel_cluster_param_4];294; CHECK-NEXT: add.s16 %rs4, %rs3, 1;295; CHECK-NEXT: st.release.cluster.b16 [%rd2], %rs4;296; CHECK-NEXT: ld.acquire.cluster.b32 %r1, [%rd3];297; CHECK-NEXT: add.s32 %r2, %r1, 1;298; CHECK-NEXT: st.release.cluster.b32 [%rd3], %r2;299; CHECK-NEXT: ld.acquire.cluster.b64 %rd6, [%rd4];300; CHECK-NEXT: add.s64 %rd7, %rd6, 1;301; CHECK-NEXT: st.release.cluster.b64 [%rd4], %rd7;302; CHECK-NEXT: ld.acquire.cluster.b32 %r3, [%rd5];303; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;304; CHECK-NEXT: st.release.cluster.b32 [%rd5], %r4;305; CHECK-NEXT: ld.acquire.cluster.b64 %rd8, [%rd5];306; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;307; CHECK-NEXT: st.release.cluster.b64 [%rd5], %rd9;308; CHECK-NEXT: ret;309 %a.load = load atomic i8, ptr %a syncscope("cluster") acquire, align 1310 %a.add = add i8 %a.load, 1311 store atomic i8 %a.add, ptr %a syncscope("cluster") release, align 1312 313 %b.load = load atomic i16, ptr %b syncscope("cluster") acquire, align 2314 %b.add = add i16 %b.load, 1315 store atomic i16 %b.add, ptr %b syncscope("cluster") release, align 2316 317 %c.load = load atomic i32, ptr %c syncscope("cluster") acquire, align 4318 %c.add = add i32 %c.load, 1319 store atomic i32 %c.add, ptr %c syncscope("cluster") release, align 4320 321 %d.load = load atomic i64, ptr %d syncscope("cluster") acquire, align 8322 %d.add = add i64 %d.load, 1323 store atomic i64 %d.add, ptr %d syncscope("cluster") release, align 8324 325 %e.load = load atomic float, ptr %e syncscope("cluster") acquire, align 4326 %e.add = fadd float %e.load, 1.327 store atomic float %e.add, ptr %e syncscope("cluster") release, align 4328 329 %f.load = load atomic double, ptr %e syncscope("cluster") acquire, align 8330 %f.add = fadd double %f.load, 1.331 store atomic double %f.add, ptr %e syncscope("cluster") release, align 8332 333 ret void334}335 336; CHECK-LABEL: generic_acq_rel_volatile_cluster337define void @generic_acq_rel_volatile_cluster(ptr %a, ptr %b, ptr %c, ptr %d, ptr %e) local_unnamed_addr {338; CHECK-LABEL: generic_acq_rel_volatile_cluster(339; CHECK: {340; CHECK-NEXT: .reg .b16 %rs<5>;341; CHECK-NEXT: .reg .b32 %r<5>;342; CHECK-NEXT: .reg .b64 %rd<10>;343; CHECK-EMPTY:344; CHECK-NEXT: // %bb.0:345; CHECK-NEXT: ld.param.b64 %rd1, [generic_acq_rel_volatile_cluster_param_0];346; CHECK-NEXT: ld.acquire.sys.b8 %rs1, [%rd1];347; CHECK-NEXT: ld.param.b64 %rd2, [generic_acq_rel_volatile_cluster_param_1];348; CHECK-NEXT: add.s16 %rs2, %rs1, 1;349; CHECK-NEXT: ld.param.b64 %rd3, [generic_acq_rel_volatile_cluster_param_2];350; CHECK-NEXT: st.release.sys.b8 [%rd1], %rs2;351; CHECK-NEXT: ld.param.b64 %rd4, [generic_acq_rel_volatile_cluster_param_3];352; CHECK-NEXT: ld.acquire.sys.b16 %rs3, [%rd2];353; CHECK-NEXT: ld.param.b64 %rd5, [generic_acq_rel_volatile_cluster_param_4];354; CHECK-NEXT: add.s16 %rs4, %rs3, 1;355; CHECK-NEXT: st.release.sys.b16 [%rd2], %rs4;356; CHECK-NEXT: ld.acquire.sys.b32 %r1, [%rd3];357; CHECK-NEXT: add.s32 %r2, %r1, 1;358; CHECK-NEXT: st.release.sys.b32 [%rd3], %r2;359; CHECK-NEXT: ld.acquire.sys.b64 %rd6, [%rd4];360; CHECK-NEXT: add.s64 %rd7, %rd6, 1;361; CHECK-NEXT: st.release.sys.b64 [%rd4], %rd7;362; CHECK-NEXT: ld.acquire.sys.b32 %r3, [%rd5];363; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;364; CHECK-NEXT: st.release.sys.b32 [%rd5], %r4;365; CHECK-NEXT: ld.acquire.sys.b64 %rd8, [%rd5];366; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;367; CHECK-NEXT: st.release.sys.b64 [%rd5], %rd9;368; CHECK-NEXT: ret;369 %a.load = load atomic volatile i8, ptr %a syncscope("cluster") acquire, align 1370 %a.add = add i8 %a.load, 1371 store atomic volatile i8 %a.add, ptr %a syncscope("cluster") release, align 1372 373 %b.load = load atomic volatile i16, ptr %b syncscope("cluster") acquire, align 2374 %b.add = add i16 %b.load, 1375 store atomic volatile i16 %b.add, ptr %b syncscope("cluster") release, align 2376 377 %c.load = load atomic volatile i32, ptr %c syncscope("cluster") acquire, align 4378 %c.add = add i32 %c.load, 1379 store atomic volatile i32 %c.add, ptr %c syncscope("cluster") release, align 4380 381 %d.load = load atomic volatile i64, ptr %d syncscope("cluster") acquire, align 8382 %d.add = add i64 %d.load, 1383 store atomic volatile i64 %d.add, ptr %d syncscope("cluster") release, align 8384 385 %e.load = load atomic volatile float, ptr %e syncscope("cluster") acquire, align 4386 %e.add = fadd float %e.load, 1.387 store atomic volatile float %e.add, ptr %e syncscope("cluster") release, align 4388 389 %f.load = load atomic volatile double, ptr %e syncscope("cluster") acquire, align 8390 %f.add = fadd double %f.load, 1.391 store atomic volatile double %f.add, ptr %e syncscope("cluster") release, align 8392 393 ret void394}395 396; CHECK-LABEL: generic_sc_cluster397define void @generic_sc_cluster(ptr %a, ptr %b, ptr %c, ptr %d, ptr %e) local_unnamed_addr {398; CHECK-LABEL: generic_sc_cluster(399; CHECK: {400; CHECK-NEXT: .reg .b16 %rs<5>;401; CHECK-NEXT: .reg .b32 %r<5>;402; CHECK-NEXT: .reg .b64 %rd<10>;403; CHECK-EMPTY:404; CHECK-NEXT: // %bb.0:405; CHECK-NEXT: ld.param.b64 %rd1, [generic_sc_cluster_param_0];406; CHECK-NEXT: fence.sc.cluster;407; CHECK-NEXT: ld.acquire.cluster.b8 %rs1, [%rd1];408; CHECK-NEXT: ld.param.b64 %rd2, [generic_sc_cluster_param_1];409; CHECK-NEXT: add.s16 %rs2, %rs1, 1;410; CHECK-NEXT: ld.param.b64 %rd3, [generic_sc_cluster_param_2];411; CHECK-NEXT: fence.sc.cluster;412; CHECK-NEXT: st.release.cluster.b8 [%rd1], %rs2;413; CHECK-NEXT: ld.param.b64 %rd4, [generic_sc_cluster_param_3];414; CHECK-NEXT: fence.sc.cluster;415; CHECK-NEXT: ld.acquire.cluster.b16 %rs3, [%rd2];416; CHECK-NEXT: ld.param.b64 %rd5, [generic_sc_cluster_param_4];417; CHECK-NEXT: add.s16 %rs4, %rs3, 1;418; CHECK-NEXT: fence.sc.cluster;419; CHECK-NEXT: st.release.cluster.b16 [%rd2], %rs4;420; CHECK-NEXT: fence.sc.cluster;421; CHECK-NEXT: ld.acquire.cluster.b32 %r1, [%rd3];422; CHECK-NEXT: add.s32 %r2, %r1, 1;423; CHECK-NEXT: fence.sc.cluster;424; CHECK-NEXT: st.release.cluster.b32 [%rd3], %r2;425; CHECK-NEXT: fence.sc.cluster;426; CHECK-NEXT: ld.acquire.cluster.b64 %rd6, [%rd4];427; CHECK-NEXT: add.s64 %rd7, %rd6, 1;428; CHECK-NEXT: fence.sc.cluster;429; CHECK-NEXT: st.release.cluster.b64 [%rd4], %rd7;430; CHECK-NEXT: fence.sc.cluster;431; CHECK-NEXT: ld.acquire.cluster.b32 %r3, [%rd5];432; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;433; CHECK-NEXT: fence.sc.cluster;434; CHECK-NEXT: st.release.cluster.b32 [%rd5], %r4;435; CHECK-NEXT: fence.sc.cluster;436; CHECK-NEXT: ld.acquire.cluster.b64 %rd8, [%rd5];437; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;438; CHECK-NEXT: fence.sc.cluster;439; CHECK-NEXT: st.release.cluster.b64 [%rd5], %rd9;440; CHECK-NEXT: ret;441 %a.load = load atomic i8, ptr %a syncscope("cluster") seq_cst, align 1442 %a.add = add i8 %a.load, 1443 store atomic i8 %a.add, ptr %a syncscope("cluster") seq_cst, align 1444 445 %b.load = load atomic i16, ptr %b syncscope("cluster") seq_cst, align 2446 %b.add = add i16 %b.load, 1447 store atomic i16 %b.add, ptr %b syncscope("cluster") seq_cst, align 2448 449 %c.load = load atomic i32, ptr %c syncscope("cluster") seq_cst, align 4450 %c.add = add i32 %c.load, 1451 store atomic i32 %c.add, ptr %c syncscope("cluster") seq_cst, align 4452 453 %d.load = load atomic i64, ptr %d syncscope("cluster") seq_cst, align 8454 %d.add = add i64 %d.load, 1455 store atomic i64 %d.add, ptr %d syncscope("cluster") seq_cst, align 8456 457 %e.load = load atomic float, ptr %e syncscope("cluster") seq_cst, align 4458 %e.add = fadd float %e.load, 1.459 store atomic float %e.add, ptr %e syncscope("cluster") seq_cst, align 4460 461 %f.load = load atomic double, ptr %e syncscope("cluster") seq_cst, align 8462 %f.add = fadd double %f.load, 1.463 store atomic double %f.add, ptr %e syncscope("cluster") seq_cst, align 8464 465 ret void466}467 468; CHECK-LABEL: generic_sc_volatile_cluster469define void @generic_sc_volatile_cluster(ptr %a, ptr %b, ptr %c, ptr %d, ptr %e) local_unnamed_addr {470; CHECK-LABEL: generic_sc_volatile_cluster(471; CHECK: {472; CHECK-NEXT: .reg .b16 %rs<5>;473; CHECK-NEXT: .reg .b32 %r<5>;474; CHECK-NEXT: .reg .b64 %rd<10>;475; CHECK-EMPTY:476; CHECK-NEXT: // %bb.0:477; CHECK-NEXT: ld.param.b64 %rd1, [generic_sc_volatile_cluster_param_0];478; CHECK-NEXT: fence.sc.sys;479; CHECK-NEXT: ld.acquire.sys.b8 %rs1, [%rd1];480; CHECK-NEXT: ld.param.b64 %rd2, [generic_sc_volatile_cluster_param_1];481; CHECK-NEXT: add.s16 %rs2, %rs1, 1;482; CHECK-NEXT: ld.param.b64 %rd3, [generic_sc_volatile_cluster_param_2];483; CHECK-NEXT: fence.sc.sys;484; CHECK-NEXT: st.release.sys.b8 [%rd1], %rs2;485; CHECK-NEXT: ld.param.b64 %rd4, [generic_sc_volatile_cluster_param_3];486; CHECK-NEXT: fence.sc.sys;487; CHECK-NEXT: ld.acquire.sys.b16 %rs3, [%rd2];488; CHECK-NEXT: ld.param.b64 %rd5, [generic_sc_volatile_cluster_param_4];489; CHECK-NEXT: add.s16 %rs4, %rs3, 1;490; CHECK-NEXT: fence.sc.sys;491; CHECK-NEXT: st.release.sys.b16 [%rd2], %rs4;492; CHECK-NEXT: fence.sc.sys;493; CHECK-NEXT: ld.acquire.sys.b32 %r1, [%rd3];494; CHECK-NEXT: add.s32 %r2, %r1, 1;495; CHECK-NEXT: fence.sc.sys;496; CHECK-NEXT: st.release.sys.b32 [%rd3], %r2;497; CHECK-NEXT: fence.sc.sys;498; CHECK-NEXT: ld.acquire.sys.b64 %rd6, [%rd4];499; CHECK-NEXT: add.s64 %rd7, %rd6, 1;500; CHECK-NEXT: fence.sc.sys;501; CHECK-NEXT: st.release.sys.b64 [%rd4], %rd7;502; CHECK-NEXT: fence.sc.sys;503; CHECK-NEXT: ld.acquire.sys.b32 %r3, [%rd5];504; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;505; CHECK-NEXT: fence.sc.sys;506; CHECK-NEXT: st.release.sys.b32 [%rd5], %r4;507; CHECK-NEXT: fence.sc.sys;508; CHECK-NEXT: ld.acquire.sys.b64 %rd8, [%rd5];509; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;510; CHECK-NEXT: fence.sc.sys;511; CHECK-NEXT: st.release.sys.b64 [%rd5], %rd9;512; CHECK-NEXT: ret;513 %a.load = load atomic volatile i8, ptr %a syncscope("cluster") seq_cst, align 1514 %a.add = add i8 %a.load, 1515 store atomic volatile i8 %a.add, ptr %a syncscope("cluster") seq_cst, align 1516 517 %b.load = load atomic volatile i16, ptr %b syncscope("cluster") seq_cst, align 2518 %b.add = add i16 %b.load, 1519 store atomic volatile i16 %b.add, ptr %b syncscope("cluster") seq_cst, align 2520 521 %c.load = load atomic volatile i32, ptr %c syncscope("cluster") seq_cst, align 4522 %c.add = add i32 %c.load, 1523 store atomic volatile i32 %c.add, ptr %c syncscope("cluster") seq_cst, align 4524 525 %d.load = load atomic volatile i64, ptr %d syncscope("cluster") seq_cst, align 8526 %d.add = add i64 %d.load, 1527 store atomic volatile i64 %d.add, ptr %d syncscope("cluster") seq_cst, align 8528 529 %e.load = load atomic volatile float, ptr %e syncscope("cluster") seq_cst, align 4530 %e.add = fadd float %e.load, 1.531 store atomic volatile float %e.add, ptr %e syncscope("cluster") seq_cst, align 4532 533 %f.load = load atomic volatile double, ptr %e syncscope("cluster") seq_cst, align 8534 %f.add = fadd double %f.load, 1.535 store atomic volatile double %f.add, ptr %e syncscope("cluster") seq_cst, align 8536 537 ret void538}539 540;; global statespace541 542; CHECK-LABEL: global_unordered_cluster543define void @global_unordered_cluster(ptr addrspace(1) %a, ptr addrspace(1) %b, ptr addrspace(1) %c, ptr addrspace(1) %d, ptr addrspace(1) %e) local_unnamed_addr {544; CHECK-LABEL: global_unordered_cluster(545; CHECK: {546; CHECK-NEXT: .reg .b16 %rs<5>;547; CHECK-NEXT: .reg .b32 %r<5>;548; CHECK-NEXT: .reg .b64 %rd<10>;549; CHECK-EMPTY:550; CHECK-NEXT: // %bb.0:551; CHECK-NEXT: ld.param.b64 %rd1, [global_unordered_cluster_param_0];552; CHECK-NEXT: ld.relaxed.cluster.global.b8 %rs1, [%rd1];553; CHECK-NEXT: ld.param.b64 %rd2, [global_unordered_cluster_param_1];554; CHECK-NEXT: add.s16 %rs2, %rs1, 1;555; CHECK-NEXT: ld.param.b64 %rd3, [global_unordered_cluster_param_2];556; CHECK-NEXT: st.relaxed.cluster.global.b8 [%rd1], %rs2;557; CHECK-NEXT: ld.param.b64 %rd4, [global_unordered_cluster_param_3];558; CHECK-NEXT: ld.relaxed.cluster.global.b16 %rs3, [%rd2];559; CHECK-NEXT: ld.param.b64 %rd5, [global_unordered_cluster_param_4];560; CHECK-NEXT: add.s16 %rs4, %rs3, 1;561; CHECK-NEXT: st.relaxed.cluster.global.b16 [%rd2], %rs4;562; CHECK-NEXT: ld.relaxed.cluster.global.b32 %r1, [%rd3];563; CHECK-NEXT: add.s32 %r2, %r1, 1;564; CHECK-NEXT: st.relaxed.cluster.global.b32 [%rd3], %r2;565; CHECK-NEXT: ld.relaxed.cluster.global.b64 %rd6, [%rd4];566; CHECK-NEXT: add.s64 %rd7, %rd6, 1;567; CHECK-NEXT: st.relaxed.cluster.global.b64 [%rd4], %rd7;568; CHECK-NEXT: ld.relaxed.cluster.global.b32 %r3, [%rd5];569; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;570; CHECK-NEXT: st.relaxed.cluster.global.b32 [%rd5], %r4;571; CHECK-NEXT: ld.relaxed.cluster.global.b64 %rd8, [%rd5];572; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;573; CHECK-NEXT: st.relaxed.cluster.global.b64 [%rd5], %rd9;574; CHECK-NEXT: ret;575 %a.load = load atomic i8, ptr addrspace(1) %a syncscope("cluster") unordered, align 1576 %a.add = add i8 %a.load, 1577 store atomic i8 %a.add, ptr addrspace(1) %a syncscope("cluster") unordered, align 1578 579 %b.load = load atomic i16, ptr addrspace(1) %b syncscope("cluster") unordered, align 2580 %b.add = add i16 %b.load, 1581 store atomic i16 %b.add, ptr addrspace(1) %b syncscope("cluster") unordered, align 2582 583 %c.load = load atomic i32, ptr addrspace(1) %c syncscope("cluster") unordered, align 4584 %c.add = add i32 %c.load, 1585 store atomic i32 %c.add, ptr addrspace(1) %c syncscope("cluster") unordered, align 4586 587 %d.load = load atomic i64, ptr addrspace(1) %d syncscope("cluster") unordered, align 8588 %d.add = add i64 %d.load, 1589 store atomic i64 %d.add, ptr addrspace(1) %d syncscope("cluster") unordered, align 8590 591 %e.load = load atomic float, ptr addrspace(1) %e syncscope("cluster") unordered, align 4592 %e.add = fadd float %e.load, 1.593 store atomic float %e.add, ptr addrspace(1) %e syncscope("cluster") unordered, align 4594 595 %f.load = load atomic double, ptr addrspace(1) %e syncscope("cluster") unordered, align 8596 %f.add = fadd double %f.load, 1.597 store atomic double %f.add, ptr addrspace(1) %e syncscope("cluster") unordered, align 8598 599 ret void600}601 602; CHECK-LABEL: global_unordered_volatile_cluster603define void @global_unordered_volatile_cluster(ptr addrspace(1) %a, ptr addrspace(1) %b, ptr addrspace(1) %c, ptr addrspace(1) %d, ptr addrspace(1) %e) local_unnamed_addr {604; CHECK-LABEL: global_unordered_volatile_cluster(605; CHECK: {606; CHECK-NEXT: .reg .b16 %rs<5>;607; CHECK-NEXT: .reg .b32 %r<5>;608; CHECK-NEXT: .reg .b64 %rd<10>;609; CHECK-EMPTY:610; CHECK-NEXT: // %bb.0:611; CHECK-NEXT: ld.param.b64 %rd1, [global_unordered_volatile_cluster_param_0];612; CHECK-NEXT: ld.volatile.global.b8 %rs1, [%rd1];613; CHECK-NEXT: ld.param.b64 %rd2, [global_unordered_volatile_cluster_param_1];614; CHECK-NEXT: add.s16 %rs2, %rs1, 1;615; CHECK-NEXT: ld.param.b64 %rd3, [global_unordered_volatile_cluster_param_2];616; CHECK-NEXT: st.volatile.global.b8 [%rd1], %rs2;617; CHECK-NEXT: ld.param.b64 %rd4, [global_unordered_volatile_cluster_param_3];618; CHECK-NEXT: ld.volatile.global.b16 %rs3, [%rd2];619; CHECK-NEXT: ld.param.b64 %rd5, [global_unordered_volatile_cluster_param_4];620; CHECK-NEXT: add.s16 %rs4, %rs3, 1;621; CHECK-NEXT: st.volatile.global.b16 [%rd2], %rs4;622; CHECK-NEXT: ld.volatile.global.b32 %r1, [%rd3];623; CHECK-NEXT: add.s32 %r2, %r1, 1;624; CHECK-NEXT: st.volatile.global.b32 [%rd3], %r2;625; CHECK-NEXT: ld.volatile.global.b64 %rd6, [%rd4];626; CHECK-NEXT: add.s64 %rd7, %rd6, 1;627; CHECK-NEXT: st.volatile.global.b64 [%rd4], %rd7;628; CHECK-NEXT: ld.volatile.global.b32 %r3, [%rd5];629; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;630; CHECK-NEXT: st.volatile.global.b32 [%rd5], %r4;631; CHECK-NEXT: ld.volatile.global.b64 %rd8, [%rd5];632; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;633; CHECK-NEXT: st.volatile.global.b64 [%rd5], %rd9;634; CHECK-NEXT: ret;635 %a.load = load atomic volatile i8, ptr addrspace(1) %a syncscope("cluster") unordered, align 1636 %a.add = add i8 %a.load, 1637 store atomic volatile i8 %a.add, ptr addrspace(1) %a syncscope("cluster") unordered, align 1638 639 %b.load = load atomic volatile i16, ptr addrspace(1) %b syncscope("cluster") unordered, align 2640 %b.add = add i16 %b.load, 1641 store atomic volatile i16 %b.add, ptr addrspace(1) %b syncscope("cluster") unordered, align 2642 643 %c.load = load atomic volatile i32, ptr addrspace(1) %c syncscope("cluster") unordered, align 4644 %c.add = add i32 %c.load, 1645 store atomic volatile i32 %c.add, ptr addrspace(1) %c syncscope("cluster") unordered, align 4646 647 %d.load = load atomic volatile i64, ptr addrspace(1) %d syncscope("cluster") unordered, align 8648 %d.add = add i64 %d.load, 1649 store atomic volatile i64 %d.add, ptr addrspace(1) %d syncscope("cluster") unordered, align 8650 651 %e.load = load atomic volatile float, ptr addrspace(1) %e syncscope("cluster") unordered, align 4652 %e.add = fadd float %e.load, 1.653 store atomic volatile float %e.add, ptr addrspace(1) %e syncscope("cluster") unordered, align 4654 655 %f.load = load atomic volatile double, ptr addrspace(1) %e syncscope("cluster") unordered, align 8656 %f.add = fadd double %f.load, 1.657 store atomic volatile double %f.add, ptr addrspace(1) %e syncscope("cluster") unordered, align 8658 659 ret void660}661 662; CHECK-LABEL: global_monotonic_cluster663define void @global_monotonic_cluster(ptr addrspace(1) %a, ptr addrspace(1) %b, ptr addrspace(1) %c, ptr addrspace(1) %d, ptr addrspace(1) %e) local_unnamed_addr {664; CHECK-LABEL: global_monotonic_cluster(665; CHECK: {666; CHECK-NEXT: .reg .b16 %rs<5>;667; CHECK-NEXT: .reg .b32 %r<5>;668; CHECK-NEXT: .reg .b64 %rd<10>;669; CHECK-EMPTY:670; CHECK-NEXT: // %bb.0:671; CHECK-NEXT: ld.param.b64 %rd1, [global_monotonic_cluster_param_0];672; CHECK-NEXT: ld.relaxed.cluster.global.b8 %rs1, [%rd1];673; CHECK-NEXT: ld.param.b64 %rd2, [global_monotonic_cluster_param_1];674; CHECK-NEXT: add.s16 %rs2, %rs1, 1;675; CHECK-NEXT: ld.param.b64 %rd3, [global_monotonic_cluster_param_2];676; CHECK-NEXT: st.relaxed.cluster.global.b8 [%rd1], %rs2;677; CHECK-NEXT: ld.param.b64 %rd4, [global_monotonic_cluster_param_3];678; CHECK-NEXT: ld.relaxed.cluster.global.b16 %rs3, [%rd2];679; CHECK-NEXT: ld.param.b64 %rd5, [global_monotonic_cluster_param_4];680; CHECK-NEXT: add.s16 %rs4, %rs3, 1;681; CHECK-NEXT: st.relaxed.cluster.global.b16 [%rd2], %rs4;682; CHECK-NEXT: ld.relaxed.cluster.global.b32 %r1, [%rd3];683; CHECK-NEXT: add.s32 %r2, %r1, 1;684; CHECK-NEXT: st.relaxed.cluster.global.b32 [%rd3], %r2;685; CHECK-NEXT: ld.relaxed.cluster.global.b64 %rd6, [%rd4];686; CHECK-NEXT: add.s64 %rd7, %rd6, 1;687; CHECK-NEXT: st.relaxed.cluster.global.b64 [%rd4], %rd7;688; CHECK-NEXT: ld.relaxed.cluster.global.b32 %r3, [%rd5];689; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;690; CHECK-NEXT: st.relaxed.cluster.global.b32 [%rd5], %r4;691; CHECK-NEXT: ld.relaxed.cluster.global.b64 %rd8, [%rd5];692; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;693; CHECK-NEXT: st.relaxed.cluster.global.b64 [%rd5], %rd9;694; CHECK-NEXT: ret;695 %a.load = load atomic i8, ptr addrspace(1) %a syncscope("cluster") monotonic, align 1696 %a.add = add i8 %a.load, 1697 store atomic i8 %a.add, ptr addrspace(1) %a syncscope("cluster") monotonic, align 1698 699 %b.load = load atomic i16, ptr addrspace(1) %b syncscope("cluster") monotonic, align 2700 %b.add = add i16 %b.load, 1701 store atomic i16 %b.add, ptr addrspace(1) %b syncscope("cluster") monotonic, align 2702 703 %c.load = load atomic i32, ptr addrspace(1) %c syncscope("cluster") monotonic, align 4704 %c.add = add i32 %c.load, 1705 store atomic i32 %c.add, ptr addrspace(1) %c syncscope("cluster") monotonic, align 4706 707 %d.load = load atomic i64, ptr addrspace(1) %d syncscope("cluster") monotonic, align 8708 %d.add = add i64 %d.load, 1709 store atomic i64 %d.add, ptr addrspace(1) %d syncscope("cluster") monotonic, align 8710 711 %e.load = load atomic float, ptr addrspace(1) %e syncscope("cluster") monotonic, align 4712 %e.add = fadd float %e.load, 1.713 store atomic float %e.add, ptr addrspace(1) %e syncscope("cluster") monotonic, align 4714 715 %f.load = load atomic double, ptr addrspace(1) %e syncscope("cluster") monotonic, align 8716 %f.add = fadd double %f.load, 1.717 store atomic double %f.add, ptr addrspace(1) %e syncscope("cluster") monotonic, align 8718 719 ret void720}721 722; CHECK-LABEL: global_monotonic_volatile_cluster723define void @global_monotonic_volatile_cluster(ptr addrspace(1) %a, ptr addrspace(1) %b, ptr addrspace(1) %c, ptr addrspace(1) %d, ptr addrspace(1) %e) local_unnamed_addr {724; CHECK-LABEL: global_monotonic_volatile_cluster(725; CHECK: {726; CHECK-NEXT: .reg .b16 %rs<5>;727; CHECK-NEXT: .reg .b32 %r<5>;728; CHECK-NEXT: .reg .b64 %rd<10>;729; CHECK-EMPTY:730; CHECK-NEXT: // %bb.0:731; CHECK-NEXT: ld.param.b64 %rd1, [global_monotonic_volatile_cluster_param_0];732; CHECK-NEXT: ld.volatile.global.b8 %rs1, [%rd1];733; CHECK-NEXT: ld.param.b64 %rd2, [global_monotonic_volatile_cluster_param_1];734; CHECK-NEXT: add.s16 %rs2, %rs1, 1;735; CHECK-NEXT: ld.param.b64 %rd3, [global_monotonic_volatile_cluster_param_2];736; CHECK-NEXT: st.volatile.global.b8 [%rd1], %rs2;737; CHECK-NEXT: ld.param.b64 %rd4, [global_monotonic_volatile_cluster_param_3];738; CHECK-NEXT: ld.volatile.global.b16 %rs3, [%rd2];739; CHECK-NEXT: ld.param.b64 %rd5, [global_monotonic_volatile_cluster_param_4];740; CHECK-NEXT: add.s16 %rs4, %rs3, 1;741; CHECK-NEXT: st.volatile.global.b16 [%rd2], %rs4;742; CHECK-NEXT: ld.volatile.global.b32 %r1, [%rd3];743; CHECK-NEXT: add.s32 %r2, %r1, 1;744; CHECK-NEXT: st.volatile.global.b32 [%rd3], %r2;745; CHECK-NEXT: ld.volatile.global.b64 %rd6, [%rd4];746; CHECK-NEXT: add.s64 %rd7, %rd6, 1;747; CHECK-NEXT: st.volatile.global.b64 [%rd4], %rd7;748; CHECK-NEXT: ld.volatile.global.b32 %r3, [%rd5];749; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;750; CHECK-NEXT: st.volatile.global.b32 [%rd5], %r4;751; CHECK-NEXT: ld.volatile.global.b64 %rd8, [%rd5];752; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;753; CHECK-NEXT: st.volatile.global.b64 [%rd5], %rd9;754; CHECK-NEXT: ret;755 %a.load = load atomic volatile i8, ptr addrspace(1) %a syncscope("cluster") monotonic, align 1756 %a.add = add i8 %a.load, 1757 store atomic volatile i8 %a.add, ptr addrspace(1) %a syncscope("cluster") monotonic, align 1758 759 %b.load = load atomic volatile i16, ptr addrspace(1) %b syncscope("cluster") monotonic, align 2760 %b.add = add i16 %b.load, 1761 store atomic volatile i16 %b.add, ptr addrspace(1) %b syncscope("cluster") monotonic, align 2762 763 %c.load = load atomic volatile i32, ptr addrspace(1) %c syncscope("cluster") monotonic, align 4764 %c.add = add i32 %c.load, 1765 store atomic volatile i32 %c.add, ptr addrspace(1) %c syncscope("cluster") monotonic, align 4766 767 %d.load = load atomic volatile i64, ptr addrspace(1) %d syncscope("cluster") monotonic, align 8768 %d.add = add i64 %d.load, 1769 store atomic volatile i64 %d.add, ptr addrspace(1) %d syncscope("cluster") monotonic, align 8770 771 %e.load = load atomic volatile float, ptr addrspace(1) %e syncscope("cluster") monotonic, align 4772 %e.add = fadd float %e.load, 1.773 store atomic volatile float %e.add, ptr addrspace(1) %e syncscope("cluster") monotonic, align 4774 775 %f.load = load atomic volatile double, ptr addrspace(1) %e syncscope("cluster") monotonic, align 8776 %f.add = fadd double %f.load, 1.777 store atomic volatile double %f.add, ptr addrspace(1) %e syncscope("cluster") monotonic, align 8778 779 ret void780}781 782; CHECK-LABEL: global_acq_rel_cluster783define void @global_acq_rel_cluster(ptr addrspace(1) %a, ptr addrspace(1) %b, ptr addrspace(1) %c, ptr addrspace(1) %d, ptr addrspace(1) %e) local_unnamed_addr {784; CHECK-LABEL: global_acq_rel_cluster(785; CHECK: {786; CHECK-NEXT: .reg .b16 %rs<5>;787; CHECK-NEXT: .reg .b32 %r<5>;788; CHECK-NEXT: .reg .b64 %rd<10>;789; CHECK-EMPTY:790; CHECK-NEXT: // %bb.0:791; CHECK-NEXT: ld.param.b64 %rd1, [global_acq_rel_cluster_param_0];792; CHECK-NEXT: ld.acquire.cluster.global.b8 %rs1, [%rd1];793; CHECK-NEXT: ld.param.b64 %rd2, [global_acq_rel_cluster_param_1];794; CHECK-NEXT: add.s16 %rs2, %rs1, 1;795; CHECK-NEXT: ld.param.b64 %rd3, [global_acq_rel_cluster_param_2];796; CHECK-NEXT: st.release.cluster.global.b8 [%rd1], %rs2;797; CHECK-NEXT: ld.param.b64 %rd4, [global_acq_rel_cluster_param_3];798; CHECK-NEXT: ld.acquire.cluster.global.b16 %rs3, [%rd2];799; CHECK-NEXT: ld.param.b64 %rd5, [global_acq_rel_cluster_param_4];800; CHECK-NEXT: add.s16 %rs4, %rs3, 1;801; CHECK-NEXT: st.release.cluster.global.b16 [%rd2], %rs4;802; CHECK-NEXT: ld.acquire.cluster.global.b32 %r1, [%rd3];803; CHECK-NEXT: add.s32 %r2, %r1, 1;804; CHECK-NEXT: st.release.cluster.global.b32 [%rd3], %r2;805; CHECK-NEXT: ld.acquire.cluster.global.b64 %rd6, [%rd4];806; CHECK-NEXT: add.s64 %rd7, %rd6, 1;807; CHECK-NEXT: st.release.cluster.global.b64 [%rd4], %rd7;808; CHECK-NEXT: ld.acquire.cluster.global.b32 %r3, [%rd5];809; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;810; CHECK-NEXT: st.release.cluster.global.b32 [%rd5], %r4;811; CHECK-NEXT: ld.acquire.cluster.global.b64 %rd8, [%rd5];812; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;813; CHECK-NEXT: st.release.cluster.global.b64 [%rd5], %rd9;814; CHECK-NEXT: ret;815 %a.load = load atomic i8, ptr addrspace(1) %a syncscope("cluster") acquire, align 1816 %a.add = add i8 %a.load, 1817 store atomic i8 %a.add, ptr addrspace(1) %a syncscope("cluster") release, align 1818 819 %b.load = load atomic i16, ptr addrspace(1) %b syncscope("cluster") acquire, align 2820 %b.add = add i16 %b.load, 1821 store atomic i16 %b.add, ptr addrspace(1) %b syncscope("cluster") release, align 2822 823 %c.load = load atomic i32, ptr addrspace(1) %c syncscope("cluster") acquire, align 4824 %c.add = add i32 %c.load, 1825 store atomic i32 %c.add, ptr addrspace(1) %c syncscope("cluster") release, align 4826 827 %d.load = load atomic i64, ptr addrspace(1) %d syncscope("cluster") acquire, align 8828 %d.add = add i64 %d.load, 1829 store atomic i64 %d.add, ptr addrspace(1) %d syncscope("cluster") release, align 8830 831 %e.load = load atomic float, ptr addrspace(1) %e syncscope("cluster") acquire, align 4832 %e.add = fadd float %e.load, 1.833 store atomic float %e.add, ptr addrspace(1) %e syncscope("cluster") release, align 4834 835 %f.load = load atomic double, ptr addrspace(1) %e syncscope("cluster") acquire, align 8836 %f.add = fadd double %f.load, 1.837 store atomic double %f.add, ptr addrspace(1) %e syncscope("cluster") release, align 8838 839 ret void840}841 842; CHECK-LABEL: global_acq_rel_volatile_cluster843define void @global_acq_rel_volatile_cluster(ptr addrspace(1) %a, ptr addrspace(1) %b, ptr addrspace(1) %c, ptr addrspace(1) %d, ptr addrspace(1) %e) local_unnamed_addr {844; CHECK-LABEL: global_acq_rel_volatile_cluster(845; CHECK: {846; CHECK-NEXT: .reg .b16 %rs<5>;847; CHECK-NEXT: .reg .b32 %r<5>;848; CHECK-NEXT: .reg .b64 %rd<10>;849; CHECK-EMPTY:850; CHECK-NEXT: // %bb.0:851; CHECK-NEXT: ld.param.b64 %rd1, [global_acq_rel_volatile_cluster_param_0];852; CHECK-NEXT: ld.acquire.sys.global.b8 %rs1, [%rd1];853; CHECK-NEXT: ld.param.b64 %rd2, [global_acq_rel_volatile_cluster_param_1];854; CHECK-NEXT: add.s16 %rs2, %rs1, 1;855; CHECK-NEXT: ld.param.b64 %rd3, [global_acq_rel_volatile_cluster_param_2];856; CHECK-NEXT: st.release.sys.global.b8 [%rd1], %rs2;857; CHECK-NEXT: ld.param.b64 %rd4, [global_acq_rel_volatile_cluster_param_3];858; CHECK-NEXT: ld.acquire.sys.global.b16 %rs3, [%rd2];859; CHECK-NEXT: ld.param.b64 %rd5, [global_acq_rel_volatile_cluster_param_4];860; CHECK-NEXT: add.s16 %rs4, %rs3, 1;861; CHECK-NEXT: st.release.sys.global.b16 [%rd2], %rs4;862; CHECK-NEXT: ld.acquire.sys.global.b32 %r1, [%rd3];863; CHECK-NEXT: add.s32 %r2, %r1, 1;864; CHECK-NEXT: st.release.sys.global.b32 [%rd3], %r2;865; CHECK-NEXT: ld.acquire.sys.global.b64 %rd6, [%rd4];866; CHECK-NEXT: add.s64 %rd7, %rd6, 1;867; CHECK-NEXT: st.release.sys.global.b64 [%rd4], %rd7;868; CHECK-NEXT: ld.acquire.sys.global.b32 %r3, [%rd5];869; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;870; CHECK-NEXT: st.release.sys.global.b32 [%rd5], %r4;871; CHECK-NEXT: ld.acquire.sys.global.b64 %rd8, [%rd5];872; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;873; CHECK-NEXT: st.release.sys.global.b64 [%rd5], %rd9;874; CHECK-NEXT: ret;875 %a.load = load atomic volatile i8, ptr addrspace(1) %a syncscope("cluster") acquire, align 1876 %a.add = add i8 %a.load, 1877 store atomic volatile i8 %a.add, ptr addrspace(1) %a syncscope("cluster") release, align 1878 879 %b.load = load atomic volatile i16, ptr addrspace(1) %b syncscope("cluster") acquire, align 2880 %b.add = add i16 %b.load, 1881 store atomic volatile i16 %b.add, ptr addrspace(1) %b syncscope("cluster") release, align 2882 883 %c.load = load atomic volatile i32, ptr addrspace(1) %c syncscope("cluster") acquire, align 4884 %c.add = add i32 %c.load, 1885 store atomic volatile i32 %c.add, ptr addrspace(1) %c syncscope("cluster") release, align 4886 887 %d.load = load atomic volatile i64, ptr addrspace(1) %d syncscope("cluster") acquire, align 8888 %d.add = add i64 %d.load, 1889 store atomic volatile i64 %d.add, ptr addrspace(1) %d syncscope("cluster") release, align 8890 891 %e.load = load atomic volatile float, ptr addrspace(1) %e syncscope("cluster") acquire, align 4892 %e.add = fadd float %e.load, 1.893 store atomic volatile float %e.add, ptr addrspace(1) %e syncscope("cluster") release, align 4894 895 %f.load = load atomic volatile double, ptr addrspace(1) %e syncscope("cluster") acquire, align 8896 %f.add = fadd double %f.load, 1.897 store atomic volatile double %f.add, ptr addrspace(1) %e syncscope("cluster") release, align 8898 899 ret void900}901 902; CHECK-LABEL: global_seq_cst_cluster903define void @global_seq_cst_cluster(ptr addrspace(1) %a, ptr addrspace(1) %b, ptr addrspace(1) %c, ptr addrspace(1) %d, ptr addrspace(1) %e) local_unnamed_addr {904; CHECK-LABEL: global_seq_cst_cluster(905; CHECK: {906; CHECK-NEXT: .reg .b16 %rs<5>;907; CHECK-NEXT: .reg .b32 %r<5>;908; CHECK-NEXT: .reg .b64 %rd<10>;909; CHECK-EMPTY:910; CHECK-NEXT: // %bb.0:911; CHECK-NEXT: ld.param.b64 %rd1, [global_seq_cst_cluster_param_0];912; CHECK-NEXT: fence.sc.cluster;913; CHECK-NEXT: ld.acquire.cluster.global.b8 %rs1, [%rd1];914; CHECK-NEXT: ld.param.b64 %rd2, [global_seq_cst_cluster_param_1];915; CHECK-NEXT: add.s16 %rs2, %rs1, 1;916; CHECK-NEXT: ld.param.b64 %rd3, [global_seq_cst_cluster_param_2];917; CHECK-NEXT: fence.sc.cluster;918; CHECK-NEXT: st.release.cluster.global.b8 [%rd1], %rs2;919; CHECK-NEXT: ld.param.b64 %rd4, [global_seq_cst_cluster_param_3];920; CHECK-NEXT: fence.sc.cluster;921; CHECK-NEXT: ld.acquire.cluster.global.b16 %rs3, [%rd2];922; CHECK-NEXT: ld.param.b64 %rd5, [global_seq_cst_cluster_param_4];923; CHECK-NEXT: add.s16 %rs4, %rs3, 1;924; CHECK-NEXT: fence.sc.cluster;925; CHECK-NEXT: st.release.cluster.global.b16 [%rd2], %rs4;926; CHECK-NEXT: fence.sc.cluster;927; CHECK-NEXT: ld.acquire.cluster.global.b32 %r1, [%rd3];928; CHECK-NEXT: add.s32 %r2, %r1, 1;929; CHECK-NEXT: fence.sc.cluster;930; CHECK-NEXT: st.release.cluster.global.b32 [%rd3], %r2;931; CHECK-NEXT: fence.sc.cluster;932; CHECK-NEXT: ld.acquire.cluster.global.b64 %rd6, [%rd4];933; CHECK-NEXT: add.s64 %rd7, %rd6, 1;934; CHECK-NEXT: fence.sc.cluster;935; CHECK-NEXT: st.release.cluster.global.b64 [%rd4], %rd7;936; CHECK-NEXT: fence.sc.cluster;937; CHECK-NEXT: ld.acquire.cluster.global.b32 %r3, [%rd5];938; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;939; CHECK-NEXT: fence.sc.cluster;940; CHECK-NEXT: st.release.cluster.global.b32 [%rd5], %r4;941; CHECK-NEXT: fence.sc.cluster;942; CHECK-NEXT: ld.acquire.cluster.global.b64 %rd8, [%rd5];943; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;944; CHECK-NEXT: fence.sc.cluster;945; CHECK-NEXT: st.release.cluster.global.b64 [%rd5], %rd9;946; CHECK-NEXT: ret;947 %a.load = load atomic i8, ptr addrspace(1) %a syncscope("cluster") seq_cst, align 1948 %a.add = add i8 %a.load, 1949 store atomic i8 %a.add, ptr addrspace(1) %a syncscope("cluster") seq_cst, align 1950 951 %b.load = load atomic i16, ptr addrspace(1) %b syncscope("cluster") seq_cst, align 2952 %b.add = add i16 %b.load, 1953 store atomic i16 %b.add, ptr addrspace(1) %b syncscope("cluster") seq_cst, align 2954 955 %c.load = load atomic i32, ptr addrspace(1) %c syncscope("cluster") seq_cst, align 4956 %c.add = add i32 %c.load, 1957 store atomic i32 %c.add, ptr addrspace(1) %c syncscope("cluster") seq_cst, align 4958 959 %d.load = load atomic i64, ptr addrspace(1) %d syncscope("cluster") seq_cst, align 8960 %d.add = add i64 %d.load, 1961 store atomic i64 %d.add, ptr addrspace(1) %d syncscope("cluster") seq_cst, align 8962 963 %e.load = load atomic float, ptr addrspace(1) %e syncscope("cluster") seq_cst, align 4964 %e.add = fadd float %e.load, 1.965 store atomic float %e.add, ptr addrspace(1) %e syncscope("cluster") seq_cst, align 4966 967 %f.load = load atomic double, ptr addrspace(1) %e syncscope("cluster") seq_cst, align 8968 %f.add = fadd double %f.load, 1.969 store atomic double %f.add, ptr addrspace(1) %e syncscope("cluster") seq_cst, align 8970 971 ret void972}973 974; CHECK-LABEL: global_seq_cst_volatile_cluster975define void @global_seq_cst_volatile_cluster(ptr addrspace(1) %a, ptr addrspace(1) %b, ptr addrspace(1) %c, ptr addrspace(1) %d, ptr addrspace(1) %e) local_unnamed_addr {976; CHECK-LABEL: global_seq_cst_volatile_cluster(977; CHECK: {978; CHECK-NEXT: .reg .b16 %rs<5>;979; CHECK-NEXT: .reg .b32 %r<5>;980; CHECK-NEXT: .reg .b64 %rd<10>;981; CHECK-EMPTY:982; CHECK-NEXT: // %bb.0:983; CHECK-NEXT: ld.param.b64 %rd1, [global_seq_cst_volatile_cluster_param_0];984; CHECK-NEXT: fence.sc.sys;985; CHECK-NEXT: ld.acquire.sys.global.b8 %rs1, [%rd1];986; CHECK-NEXT: ld.param.b64 %rd2, [global_seq_cst_volatile_cluster_param_1];987; CHECK-NEXT: add.s16 %rs2, %rs1, 1;988; CHECK-NEXT: ld.param.b64 %rd3, [global_seq_cst_volatile_cluster_param_2];989; CHECK-NEXT: fence.sc.sys;990; CHECK-NEXT: st.release.sys.global.b8 [%rd1], %rs2;991; CHECK-NEXT: ld.param.b64 %rd4, [global_seq_cst_volatile_cluster_param_3];992; CHECK-NEXT: fence.sc.sys;993; CHECK-NEXT: ld.acquire.sys.global.b16 %rs3, [%rd2];994; CHECK-NEXT: ld.param.b64 %rd5, [global_seq_cst_volatile_cluster_param_4];995; CHECK-NEXT: add.s16 %rs4, %rs3, 1;996; CHECK-NEXT: fence.sc.sys;997; CHECK-NEXT: st.release.sys.global.b16 [%rd2], %rs4;998; CHECK-NEXT: fence.sc.sys;999; CHECK-NEXT: ld.acquire.sys.global.b32 %r1, [%rd3];1000; CHECK-NEXT: add.s32 %r2, %r1, 1;1001; CHECK-NEXT: fence.sc.sys;1002; CHECK-NEXT: st.release.sys.global.b32 [%rd3], %r2;1003; CHECK-NEXT: fence.sc.sys;1004; CHECK-NEXT: ld.acquire.sys.global.b64 %rd6, [%rd4];1005; CHECK-NEXT: add.s64 %rd7, %rd6, 1;1006; CHECK-NEXT: fence.sc.sys;1007; CHECK-NEXT: st.release.sys.global.b64 [%rd4], %rd7;1008; CHECK-NEXT: fence.sc.sys;1009; CHECK-NEXT: ld.acquire.sys.global.b32 %r3, [%rd5];1010; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;1011; CHECK-NEXT: fence.sc.sys;1012; CHECK-NEXT: st.release.sys.global.b32 [%rd5], %r4;1013; CHECK-NEXT: fence.sc.sys;1014; CHECK-NEXT: ld.acquire.sys.global.b64 %rd8, [%rd5];1015; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;1016; CHECK-NEXT: fence.sc.sys;1017; CHECK-NEXT: st.release.sys.global.b64 [%rd5], %rd9;1018; CHECK-NEXT: ret;1019 %a.load = load atomic volatile i8, ptr addrspace(1) %a syncscope("cluster") seq_cst, align 11020 %a.add = add i8 %a.load, 11021 store atomic volatile i8 %a.add, ptr addrspace(1) %a syncscope("cluster") seq_cst, align 11022 1023 %b.load = load atomic volatile i16, ptr addrspace(1) %b syncscope("cluster") seq_cst, align 21024 %b.add = add i16 %b.load, 11025 store atomic volatile i16 %b.add, ptr addrspace(1) %b syncscope("cluster") seq_cst, align 21026 1027 %c.load = load atomic volatile i32, ptr addrspace(1) %c syncscope("cluster") seq_cst, align 41028 %c.add = add i32 %c.load, 11029 store atomic volatile i32 %c.add, ptr addrspace(1) %c syncscope("cluster") seq_cst, align 41030 1031 %d.load = load atomic volatile i64, ptr addrspace(1) %d syncscope("cluster") seq_cst, align 81032 %d.add = add i64 %d.load, 11033 store atomic volatile i64 %d.add, ptr addrspace(1) %d syncscope("cluster") seq_cst, align 81034 1035 %e.load = load atomic volatile float, ptr addrspace(1) %e syncscope("cluster") seq_cst, align 41036 %e.add = fadd float %e.load, 1.1037 store atomic volatile float %e.add, ptr addrspace(1) %e syncscope("cluster") seq_cst, align 41038 1039 %f.load = load atomic volatile double, ptr addrspace(1) %e syncscope("cluster") seq_cst, align 81040 %f.add = fadd double %f.load, 1.1041 store atomic volatile double %f.add, ptr addrspace(1) %e syncscope("cluster") seq_cst, align 81042 1043 ret void1044}1045 1046;; shared1047 1048; CHECK-LABEL: shared_unordered_cluster1049define void @shared_unordered_cluster(ptr addrspace(3) %a, ptr addrspace(3) %b, ptr addrspace(3) %c, ptr addrspace(3) %d, ptr addrspace(3) %e) local_unnamed_addr {1050; CHECK-LABEL: shared_unordered_cluster(1051; CHECK: {1052; CHECK-NEXT: .reg .b16 %rs<5>;1053; CHECK-NEXT: .reg .b32 %r<5>;1054; CHECK-NEXT: .reg .b64 %rd<10>;1055; CHECK-EMPTY:1056; CHECK-NEXT: // %bb.0:1057; CHECK-NEXT: ld.param.b64 %rd1, [shared_unordered_cluster_param_0];1058; CHECK-NEXT: ld.relaxed.cluster.shared.b8 %rs1, [%rd1];1059; CHECK-NEXT: ld.param.b64 %rd2, [shared_unordered_cluster_param_1];1060; CHECK-NEXT: add.s16 %rs2, %rs1, 1;1061; CHECK-NEXT: ld.param.b64 %rd3, [shared_unordered_cluster_param_2];1062; CHECK-NEXT: st.relaxed.cluster.shared.b8 [%rd1], %rs2;1063; CHECK-NEXT: ld.param.b64 %rd4, [shared_unordered_cluster_param_3];1064; CHECK-NEXT: ld.relaxed.cluster.shared.b16 %rs3, [%rd2];1065; CHECK-NEXT: ld.param.b64 %rd5, [shared_unordered_cluster_param_4];1066; CHECK-NEXT: add.s16 %rs4, %rs3, 1;1067; CHECK-NEXT: st.relaxed.cluster.shared.b16 [%rd2], %rs4;1068; CHECK-NEXT: ld.relaxed.cluster.shared.b32 %r1, [%rd3];1069; CHECK-NEXT: add.s32 %r2, %r1, 1;1070; CHECK-NEXT: st.relaxed.cluster.shared.b32 [%rd3], %r2;1071; CHECK-NEXT: ld.relaxed.cluster.shared.b64 %rd6, [%rd4];1072; CHECK-NEXT: add.s64 %rd7, %rd6, 1;1073; CHECK-NEXT: st.relaxed.cluster.shared.b64 [%rd4], %rd7;1074; CHECK-NEXT: ld.relaxed.cluster.shared.b32 %r3, [%rd5];1075; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;1076; CHECK-NEXT: st.relaxed.cluster.shared.b32 [%rd5], %r4;1077; CHECK-NEXT: ld.relaxed.cluster.shared.b64 %rd8, [%rd5];1078; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;1079; CHECK-NEXT: st.relaxed.cluster.shared.b64 [%rd5], %rd9;1080; CHECK-NEXT: ret;1081 %a.load = load atomic i8, ptr addrspace(3) %a syncscope("cluster") unordered, align 11082 %a.add = add i8 %a.load, 11083 store atomic i8 %a.add, ptr addrspace(3) %a syncscope("cluster") unordered, align 11084 1085 %b.load = load atomic i16, ptr addrspace(3) %b syncscope("cluster") unordered, align 21086 %b.add = add i16 %b.load, 11087 store atomic i16 %b.add, ptr addrspace(3) %b syncscope("cluster") unordered, align 21088 1089 %c.load = load atomic i32, ptr addrspace(3) %c syncscope("cluster") unordered, align 41090 %c.add = add i32 %c.load, 11091 store atomic i32 %c.add, ptr addrspace(3) %c syncscope("cluster") unordered, align 41092 1093 %d.load = load atomic i64, ptr addrspace(3) %d syncscope("cluster") unordered, align 81094 %d.add = add i64 %d.load, 11095 store atomic i64 %d.add, ptr addrspace(3) %d syncscope("cluster") unordered, align 81096 1097 %e.load = load atomic float, ptr addrspace(3) %e syncscope("cluster") unordered, align 41098 %e.add = fadd float %e.load, 1.1099 store atomic float %e.add, ptr addrspace(3) %e syncscope("cluster") unordered, align 41100 1101 %f.load = load atomic double, ptr addrspace(3) %e syncscope("cluster") unordered, align 81102 %f.add = fadd double %f.load, 1.1103 store atomic double %f.add, ptr addrspace(3) %e syncscope("cluster") unordered, align 81104 1105 ret void1106}1107 1108; CHECK-LABEL: shared_unordered_volatile_cluster1109define void @shared_unordered_volatile_cluster(ptr addrspace(3) %a, ptr addrspace(3) %b, ptr addrspace(3) %c, ptr addrspace(3) %d, ptr addrspace(3) %e) local_unnamed_addr {1110; CHECK-LABEL: shared_unordered_volatile_cluster(1111; CHECK: {1112; CHECK-NEXT: .reg .b16 %rs<5>;1113; CHECK-NEXT: .reg .b32 %r<5>;1114; CHECK-NEXT: .reg .b64 %rd<10>;1115; CHECK-EMPTY:1116; CHECK-NEXT: // %bb.0:1117; CHECK-NEXT: ld.param.b64 %rd1, [shared_unordered_volatile_cluster_param_0];1118; CHECK-NEXT: ld.volatile.shared.b8 %rs1, [%rd1];1119; CHECK-NEXT: ld.param.b64 %rd2, [shared_unordered_volatile_cluster_param_1];1120; CHECK-NEXT: add.s16 %rs2, %rs1, 1;1121; CHECK-NEXT: ld.param.b64 %rd3, [shared_unordered_volatile_cluster_param_2];1122; CHECK-NEXT: st.volatile.shared.b8 [%rd1], %rs2;1123; CHECK-NEXT: ld.param.b64 %rd4, [shared_unordered_volatile_cluster_param_3];1124; CHECK-NEXT: ld.volatile.shared.b16 %rs3, [%rd2];1125; CHECK-NEXT: ld.param.b64 %rd5, [shared_unordered_volatile_cluster_param_4];1126; CHECK-NEXT: add.s16 %rs4, %rs3, 1;1127; CHECK-NEXT: st.volatile.shared.b16 [%rd2], %rs4;1128; CHECK-NEXT: ld.volatile.shared.b32 %r1, [%rd3];1129; CHECK-NEXT: add.s32 %r2, %r1, 1;1130; CHECK-NEXT: st.volatile.shared.b32 [%rd3], %r2;1131; CHECK-NEXT: ld.volatile.shared.b64 %rd6, [%rd4];1132; CHECK-NEXT: add.s64 %rd7, %rd6, 1;1133; CHECK-NEXT: st.volatile.shared.b64 [%rd4], %rd7;1134; CHECK-NEXT: ld.volatile.shared.b32 %r3, [%rd5];1135; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;1136; CHECK-NEXT: st.volatile.shared.b32 [%rd5], %r4;1137; CHECK-NEXT: ld.volatile.shared.b64 %rd8, [%rd5];1138; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;1139; CHECK-NEXT: st.volatile.shared.b64 [%rd5], %rd9;1140; CHECK-NEXT: ret;1141 %a.load = load atomic volatile i8, ptr addrspace(3) %a syncscope("cluster") unordered, align 11142 %a.add = add i8 %a.load, 11143 store atomic volatile i8 %a.add, ptr addrspace(3) %a syncscope("cluster") unordered, align 11144 1145 %b.load = load atomic volatile i16, ptr addrspace(3) %b syncscope("cluster") unordered, align 21146 %b.add = add i16 %b.load, 11147 store atomic volatile i16 %b.add, ptr addrspace(3) %b syncscope("cluster") unordered, align 21148 1149 %c.load = load atomic volatile i32, ptr addrspace(3) %c syncscope("cluster") unordered, align 41150 %c.add = add i32 %c.load, 11151 store atomic volatile i32 %c.add, ptr addrspace(3) %c syncscope("cluster") unordered, align 41152 1153 %d.load = load atomic volatile i64, ptr addrspace(3) %d syncscope("cluster") unordered, align 81154 %d.add = add i64 %d.load, 11155 store atomic volatile i64 %d.add, ptr addrspace(3) %d syncscope("cluster") unordered, align 81156 1157 %e.load = load atomic volatile float, ptr addrspace(3) %e syncscope("cluster") unordered, align 41158 %e.add = fadd float %e.load, 1.1159 store atomic volatile float %e.add, ptr addrspace(3) %e syncscope("cluster") unordered, align 41160 1161 %f.load = load atomic volatile double, ptr addrspace(3) %e syncscope("cluster") unordered, align 81162 %f.add = fadd double %f.load, 1.1163 store atomic volatile double %f.add, ptr addrspace(3) %e syncscope("cluster") unordered, align 81164 1165 ret void1166}1167 1168; CHECK-LABEL: shared_monotonic_cluster1169define void @shared_monotonic_cluster(ptr addrspace(3) %a, ptr addrspace(3) %b, ptr addrspace(3) %c, ptr addrspace(3) %d, ptr addrspace(3) %e) local_unnamed_addr {1170; CHECK-LABEL: shared_monotonic_cluster(1171; CHECK: {1172; CHECK-NEXT: .reg .b16 %rs<5>;1173; CHECK-NEXT: .reg .b32 %r<5>;1174; CHECK-NEXT: .reg .b64 %rd<10>;1175; CHECK-EMPTY:1176; CHECK-NEXT: // %bb.0:1177; CHECK-NEXT: ld.param.b64 %rd1, [shared_monotonic_cluster_param_0];1178; CHECK-NEXT: ld.relaxed.cluster.shared.b8 %rs1, [%rd1];1179; CHECK-NEXT: ld.param.b64 %rd2, [shared_monotonic_cluster_param_1];1180; CHECK-NEXT: add.s16 %rs2, %rs1, 1;1181; CHECK-NEXT: ld.param.b64 %rd3, [shared_monotonic_cluster_param_2];1182; CHECK-NEXT: st.relaxed.cluster.shared.b8 [%rd1], %rs2;1183; CHECK-NEXT: ld.param.b64 %rd4, [shared_monotonic_cluster_param_3];1184; CHECK-NEXT: ld.relaxed.cluster.shared.b16 %rs3, [%rd2];1185; CHECK-NEXT: ld.param.b64 %rd5, [shared_monotonic_cluster_param_4];1186; CHECK-NEXT: add.s16 %rs4, %rs3, 1;1187; CHECK-NEXT: st.relaxed.cluster.shared.b16 [%rd2], %rs4;1188; CHECK-NEXT: ld.relaxed.cluster.shared.b32 %r1, [%rd3];1189; CHECK-NEXT: add.s32 %r2, %r1, 1;1190; CHECK-NEXT: st.relaxed.cluster.shared.b32 [%rd3], %r2;1191; CHECK-NEXT: ld.relaxed.cluster.shared.b64 %rd6, [%rd4];1192; CHECK-NEXT: add.s64 %rd7, %rd6, 1;1193; CHECK-NEXT: st.relaxed.cluster.shared.b64 [%rd4], %rd7;1194; CHECK-NEXT: ld.relaxed.cluster.shared.b32 %r3, [%rd5];1195; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;1196; CHECK-NEXT: st.relaxed.cluster.shared.b32 [%rd5], %r4;1197; CHECK-NEXT: ld.relaxed.cluster.shared.b64 %rd8, [%rd5];1198; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;1199; CHECK-NEXT: st.relaxed.cluster.shared.b64 [%rd5], %rd9;1200; CHECK-NEXT: ret;1201 %a.load = load atomic i8, ptr addrspace(3) %a syncscope("cluster") monotonic, align 11202 %a.add = add i8 %a.load, 11203 store atomic i8 %a.add, ptr addrspace(3) %a syncscope("cluster") monotonic, align 11204 1205 %b.load = load atomic i16, ptr addrspace(3) %b syncscope("cluster") monotonic, align 21206 %b.add = add i16 %b.load, 11207 store atomic i16 %b.add, ptr addrspace(3) %b syncscope("cluster") monotonic, align 21208 1209 %c.load = load atomic i32, ptr addrspace(3) %c syncscope("cluster") monotonic, align 41210 %c.add = add i32 %c.load, 11211 store atomic i32 %c.add, ptr addrspace(3) %c syncscope("cluster") monotonic, align 41212 1213 %d.load = load atomic i64, ptr addrspace(3) %d syncscope("cluster") monotonic, align 81214 %d.add = add i64 %d.load, 11215 store atomic i64 %d.add, ptr addrspace(3) %d syncscope("cluster") monotonic, align 81216 1217 %e.load = load atomic float, ptr addrspace(3) %e syncscope("cluster") monotonic, align 41218 %e.add = fadd float %e.load, 1.1219 store atomic float %e.add, ptr addrspace(3) %e syncscope("cluster") monotonic, align 41220 1221 %f.load = load atomic double, ptr addrspace(3) %e syncscope("cluster") monotonic, align 81222 %f.add = fadd double %f.load, 1.1223 store atomic double %f.add, ptr addrspace(3) %e syncscope("cluster") monotonic, align 81224 1225 ret void1226}1227 1228; CHECK-LABEL: shared_monotonic_volatile_cluster1229define void @shared_monotonic_volatile_cluster(ptr addrspace(3) %a, ptr addrspace(3) %b, ptr addrspace(3) %c, ptr addrspace(3) %d, ptr addrspace(3) %e) local_unnamed_addr {1230; CHECK-LABEL: shared_monotonic_volatile_cluster(1231; CHECK: {1232; CHECK-NEXT: .reg .b16 %rs<5>;1233; CHECK-NEXT: .reg .b32 %r<5>;1234; CHECK-NEXT: .reg .b64 %rd<10>;1235; CHECK-EMPTY:1236; CHECK-NEXT: // %bb.0:1237; CHECK-NEXT: ld.param.b64 %rd1, [shared_monotonic_volatile_cluster_param_0];1238; CHECK-NEXT: ld.volatile.shared.b8 %rs1, [%rd1];1239; CHECK-NEXT: ld.param.b64 %rd2, [shared_monotonic_volatile_cluster_param_1];1240; CHECK-NEXT: add.s16 %rs2, %rs1, 1;1241; CHECK-NEXT: ld.param.b64 %rd3, [shared_monotonic_volatile_cluster_param_2];1242; CHECK-NEXT: st.volatile.shared.b8 [%rd1], %rs2;1243; CHECK-NEXT: ld.param.b64 %rd4, [shared_monotonic_volatile_cluster_param_3];1244; CHECK-NEXT: ld.volatile.shared.b16 %rs3, [%rd2];1245; CHECK-NEXT: ld.param.b64 %rd5, [shared_monotonic_volatile_cluster_param_4];1246; CHECK-NEXT: add.s16 %rs4, %rs3, 1;1247; CHECK-NEXT: st.volatile.shared.b16 [%rd2], %rs4;1248; CHECK-NEXT: ld.volatile.shared.b32 %r1, [%rd3];1249; CHECK-NEXT: add.s32 %r2, %r1, 1;1250; CHECK-NEXT: st.volatile.shared.b32 [%rd3], %r2;1251; CHECK-NEXT: ld.volatile.shared.b64 %rd6, [%rd4];1252; CHECK-NEXT: add.s64 %rd7, %rd6, 1;1253; CHECK-NEXT: st.volatile.shared.b64 [%rd4], %rd7;1254; CHECK-NEXT: ld.volatile.shared.b32 %r3, [%rd5];1255; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;1256; CHECK-NEXT: st.volatile.shared.b32 [%rd5], %r4;1257; CHECK-NEXT: ld.volatile.shared.b64 %rd8, [%rd5];1258; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;1259; CHECK-NEXT: st.volatile.shared.b64 [%rd5], %rd9;1260; CHECK-NEXT: ret;1261 %a.load = load atomic volatile i8, ptr addrspace(3) %a syncscope("cluster") monotonic, align 11262 %a.add = add i8 %a.load, 11263 store atomic volatile i8 %a.add, ptr addrspace(3) %a syncscope("cluster") monotonic, align 11264 1265 %b.load = load atomic volatile i16, ptr addrspace(3) %b syncscope("cluster") monotonic, align 21266 %b.add = add i16 %b.load, 11267 store atomic volatile i16 %b.add, ptr addrspace(3) %b syncscope("cluster") monotonic, align 21268 1269 %c.load = load atomic volatile i32, ptr addrspace(3) %c syncscope("cluster") monotonic, align 41270 %c.add = add i32 %c.load, 11271 store atomic volatile i32 %c.add, ptr addrspace(3) %c syncscope("cluster") monotonic, align 41272 1273 %d.load = load atomic volatile i64, ptr addrspace(3) %d syncscope("cluster") monotonic, align 81274 %d.add = add i64 %d.load, 11275 store atomic volatile i64 %d.add, ptr addrspace(3) %d syncscope("cluster") monotonic, align 81276 1277 %e.load = load atomic volatile float, ptr addrspace(3) %e syncscope("cluster") monotonic, align 41278 %e.add = fadd float %e.load, 1.1279 store atomic volatile float %e.add, ptr addrspace(3) %e syncscope("cluster") monotonic, align 41280 1281 %f.load = load atomic volatile double, ptr addrspace(3) %e syncscope("cluster") monotonic, align 81282 %f.add = fadd double %f.load, 1.1283 store atomic volatile double %f.add, ptr addrspace(3) %e syncscope("cluster") monotonic, align 81284 1285 ret void1286}1287 1288; CHECK-LABEL: shared_acq_rel_cluster1289define void @shared_acq_rel_cluster(ptr addrspace(3) %a, ptr addrspace(3) %b, ptr addrspace(3) %c, ptr addrspace(3) %d, ptr addrspace(3) %e) local_unnamed_addr {1290; CHECK-LABEL: shared_acq_rel_cluster(1291; CHECK: {1292; CHECK-NEXT: .reg .b16 %rs<5>;1293; CHECK-NEXT: .reg .b32 %r<5>;1294; CHECK-NEXT: .reg .b64 %rd<10>;1295; CHECK-EMPTY:1296; CHECK-NEXT: // %bb.0:1297; CHECK-NEXT: ld.param.b64 %rd1, [shared_acq_rel_cluster_param_0];1298; CHECK-NEXT: ld.acquire.cluster.shared.b8 %rs1, [%rd1];1299; CHECK-NEXT: ld.param.b64 %rd2, [shared_acq_rel_cluster_param_1];1300; CHECK-NEXT: add.s16 %rs2, %rs1, 1;1301; CHECK-NEXT: ld.param.b64 %rd3, [shared_acq_rel_cluster_param_2];1302; CHECK-NEXT: st.release.cluster.shared.b8 [%rd1], %rs2;1303; CHECK-NEXT: ld.param.b64 %rd4, [shared_acq_rel_cluster_param_3];1304; CHECK-NEXT: ld.acquire.cluster.shared.b16 %rs3, [%rd2];1305; CHECK-NEXT: ld.param.b64 %rd5, [shared_acq_rel_cluster_param_4];1306; CHECK-NEXT: add.s16 %rs4, %rs3, 1;1307; CHECK-NEXT: st.release.cluster.shared.b16 [%rd2], %rs4;1308; CHECK-NEXT: ld.acquire.cluster.shared.b32 %r1, [%rd3];1309; CHECK-NEXT: add.s32 %r2, %r1, 1;1310; CHECK-NEXT: st.release.cluster.shared.b32 [%rd3], %r2;1311; CHECK-NEXT: ld.acquire.cluster.shared.b64 %rd6, [%rd4];1312; CHECK-NEXT: add.s64 %rd7, %rd6, 1;1313; CHECK-NEXT: st.release.cluster.shared.b64 [%rd4], %rd7;1314; CHECK-NEXT: ld.acquire.cluster.shared.b32 %r3, [%rd5];1315; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;1316; CHECK-NEXT: st.release.cluster.shared.b32 [%rd5], %r4;1317; CHECK-NEXT: ld.acquire.cluster.shared.b64 %rd8, [%rd5];1318; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;1319; CHECK-NEXT: st.release.cluster.shared.b64 [%rd5], %rd9;1320; CHECK-NEXT: ret;1321 %a.load = load atomic i8, ptr addrspace(3) %a syncscope("cluster") acquire, align 11322 %a.add = add i8 %a.load, 11323 store atomic i8 %a.add, ptr addrspace(3) %a syncscope("cluster") release, align 11324 1325 %b.load = load atomic i16, ptr addrspace(3) %b syncscope("cluster") acquire, align 21326 %b.add = add i16 %b.load, 11327 store atomic i16 %b.add, ptr addrspace(3) %b syncscope("cluster") release, align 21328 1329 %c.load = load atomic i32, ptr addrspace(3) %c syncscope("cluster") acquire, align 41330 %c.add = add i32 %c.load, 11331 store atomic i32 %c.add, ptr addrspace(3) %c syncscope("cluster") release, align 41332 1333 %d.load = load atomic i64, ptr addrspace(3) %d syncscope("cluster") acquire, align 81334 %d.add = add i64 %d.load, 11335 store atomic i64 %d.add, ptr addrspace(3) %d syncscope("cluster") release, align 81336 1337 %e.load = load atomic float, ptr addrspace(3) %e syncscope("cluster") acquire, align 41338 %e.add = fadd float %e.load, 1.1339 store atomic float %e.add, ptr addrspace(3) %e syncscope("cluster") release, align 41340 1341 %f.load = load atomic double, ptr addrspace(3) %e syncscope("cluster") acquire, align 81342 %f.add = fadd double %f.load, 1.1343 store atomic double %f.add, ptr addrspace(3) %e syncscope("cluster") release, align 81344 1345 ret void1346}1347 1348; CHECK-LABEL: shared_acq_rel_volatile_cluster1349define void @shared_acq_rel_volatile_cluster(ptr addrspace(3) %a, ptr addrspace(3) %b, ptr addrspace(3) %c, ptr addrspace(3) %d, ptr addrspace(3) %e) local_unnamed_addr {1350; CHECK-LABEL: shared_acq_rel_volatile_cluster(1351; CHECK: {1352; CHECK-NEXT: .reg .b16 %rs<5>;1353; CHECK-NEXT: .reg .b32 %r<5>;1354; CHECK-NEXT: .reg .b64 %rd<10>;1355; CHECK-EMPTY:1356; CHECK-NEXT: // %bb.0:1357; CHECK-NEXT: ld.param.b64 %rd1, [shared_acq_rel_volatile_cluster_param_0];1358; CHECK-NEXT: ld.acquire.sys.shared.b8 %rs1, [%rd1];1359; CHECK-NEXT: ld.param.b64 %rd2, [shared_acq_rel_volatile_cluster_param_1];1360; CHECK-NEXT: add.s16 %rs2, %rs1, 1;1361; CHECK-NEXT: ld.param.b64 %rd3, [shared_acq_rel_volatile_cluster_param_2];1362; CHECK-NEXT: st.release.sys.shared.b8 [%rd1], %rs2;1363; CHECK-NEXT: ld.param.b64 %rd4, [shared_acq_rel_volatile_cluster_param_3];1364; CHECK-NEXT: ld.acquire.sys.shared.b16 %rs3, [%rd2];1365; CHECK-NEXT: ld.param.b64 %rd5, [shared_acq_rel_volatile_cluster_param_4];1366; CHECK-NEXT: add.s16 %rs4, %rs3, 1;1367; CHECK-NEXT: st.release.sys.shared.b16 [%rd2], %rs4;1368; CHECK-NEXT: ld.acquire.sys.shared.b32 %r1, [%rd3];1369; CHECK-NEXT: add.s32 %r2, %r1, 1;1370; CHECK-NEXT: st.release.sys.shared.b32 [%rd3], %r2;1371; CHECK-NEXT: ld.acquire.sys.shared.b64 %rd6, [%rd4];1372; CHECK-NEXT: add.s64 %rd7, %rd6, 1;1373; CHECK-NEXT: st.release.sys.shared.b64 [%rd4], %rd7;1374; CHECK-NEXT: ld.acquire.sys.shared.b32 %r3, [%rd5];1375; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;1376; CHECK-NEXT: st.release.sys.shared.b32 [%rd5], %r4;1377; CHECK-NEXT: ld.acquire.sys.shared.b64 %rd8, [%rd5];1378; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;1379; CHECK-NEXT: st.release.sys.shared.b64 [%rd5], %rd9;1380; CHECK-NEXT: ret;1381 %a.load = load atomic volatile i8, ptr addrspace(3) %a syncscope("cluster") acquire, align 11382 %a.add = add i8 %a.load, 11383 store atomic volatile i8 %a.add, ptr addrspace(3) %a syncscope("cluster") release, align 11384 1385 %b.load = load atomic volatile i16, ptr addrspace(3) %b syncscope("cluster") acquire, align 21386 %b.add = add i16 %b.load, 11387 store atomic volatile i16 %b.add, ptr addrspace(3) %b syncscope("cluster") release, align 21388 1389 %c.load = load atomic volatile i32, ptr addrspace(3) %c syncscope("cluster") acquire, align 41390 %c.add = add i32 %c.load, 11391 store atomic volatile i32 %c.add, ptr addrspace(3) %c syncscope("cluster") release, align 41392 1393 %d.load = load atomic volatile i64, ptr addrspace(3) %d syncscope("cluster") acquire, align 81394 %d.add = add i64 %d.load, 11395 store atomic volatile i64 %d.add, ptr addrspace(3) %d syncscope("cluster") release, align 81396 1397 %e.load = load atomic volatile float, ptr addrspace(3) %e syncscope("cluster") acquire, align 41398 %e.add = fadd float %e.load, 1.1399 store atomic volatile float %e.add, ptr addrspace(3) %e syncscope("cluster") release, align 41400 1401 %f.load = load atomic volatile double, ptr addrspace(3) %e syncscope("cluster") acquire, align 81402 %f.add = fadd double %f.load, 1.1403 store atomic volatile double %f.add, ptr addrspace(3) %e syncscope("cluster") release, align 81404 1405 ret void1406}1407 1408; CHECK-LABEL: shared_seq_cst_cluster1409define void @shared_seq_cst_cluster(ptr addrspace(3) %a, ptr addrspace(3) %b, ptr addrspace(3) %c, ptr addrspace(3) %d, ptr addrspace(3) %e) local_unnamed_addr {1410; CHECK-LABEL: shared_seq_cst_cluster(1411; CHECK: {1412; CHECK-NEXT: .reg .b16 %rs<5>;1413; CHECK-NEXT: .reg .b32 %r<5>;1414; CHECK-NEXT: .reg .b64 %rd<10>;1415; CHECK-EMPTY:1416; CHECK-NEXT: // %bb.0:1417; CHECK-NEXT: ld.param.b64 %rd1, [shared_seq_cst_cluster_param_0];1418; CHECK-NEXT: fence.sc.cluster;1419; CHECK-NEXT: ld.acquire.cluster.shared.b8 %rs1, [%rd1];1420; CHECK-NEXT: ld.param.b64 %rd2, [shared_seq_cst_cluster_param_1];1421; CHECK-NEXT: add.s16 %rs2, %rs1, 1;1422; CHECK-NEXT: ld.param.b64 %rd3, [shared_seq_cst_cluster_param_2];1423; CHECK-NEXT: fence.sc.cluster;1424; CHECK-NEXT: st.release.cluster.shared.b8 [%rd1], %rs2;1425; CHECK-NEXT: ld.param.b64 %rd4, [shared_seq_cst_cluster_param_3];1426; CHECK-NEXT: fence.sc.cluster;1427; CHECK-NEXT: ld.acquire.cluster.shared.b16 %rs3, [%rd2];1428; CHECK-NEXT: ld.param.b64 %rd5, [shared_seq_cst_cluster_param_4];1429; CHECK-NEXT: add.s16 %rs4, %rs3, 1;1430; CHECK-NEXT: fence.sc.cluster;1431; CHECK-NEXT: st.release.cluster.shared.b16 [%rd2], %rs4;1432; CHECK-NEXT: fence.sc.cluster;1433; CHECK-NEXT: ld.acquire.cluster.shared.b32 %r1, [%rd3];1434; CHECK-NEXT: add.s32 %r2, %r1, 1;1435; CHECK-NEXT: fence.sc.cluster;1436; CHECK-NEXT: st.release.cluster.shared.b32 [%rd3], %r2;1437; CHECK-NEXT: fence.sc.cluster;1438; CHECK-NEXT: ld.acquire.cluster.shared.b64 %rd6, [%rd4];1439; CHECK-NEXT: add.s64 %rd7, %rd6, 1;1440; CHECK-NEXT: fence.sc.cluster;1441; CHECK-NEXT: st.release.cluster.shared.b64 [%rd4], %rd7;1442; CHECK-NEXT: fence.sc.cluster;1443; CHECK-NEXT: ld.acquire.cluster.shared.b32 %r3, [%rd5];1444; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;1445; CHECK-NEXT: fence.sc.cluster;1446; CHECK-NEXT: st.release.cluster.shared.b32 [%rd5], %r4;1447; CHECK-NEXT: fence.sc.cluster;1448; CHECK-NEXT: ld.acquire.cluster.shared.b64 %rd8, [%rd5];1449; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;1450; CHECK-NEXT: fence.sc.cluster;1451; CHECK-NEXT: st.release.cluster.shared.b64 [%rd5], %rd9;1452; CHECK-NEXT: ret;1453 %a.load = load atomic i8, ptr addrspace(3) %a syncscope("cluster") seq_cst, align 11454 %a.add = add i8 %a.load, 11455 store atomic i8 %a.add, ptr addrspace(3) %a syncscope("cluster") seq_cst, align 11456 1457 %b.load = load atomic i16, ptr addrspace(3) %b syncscope("cluster") seq_cst, align 21458 %b.add = add i16 %b.load, 11459 store atomic i16 %b.add, ptr addrspace(3) %b syncscope("cluster") seq_cst, align 21460 1461 %c.load = load atomic i32, ptr addrspace(3) %c syncscope("cluster") seq_cst, align 41462 %c.add = add i32 %c.load, 11463 store atomic i32 %c.add, ptr addrspace(3) %c syncscope("cluster") seq_cst, align 41464 1465 %d.load = load atomic i64, ptr addrspace(3) %d syncscope("cluster") seq_cst, align 81466 %d.add = add i64 %d.load, 11467 store atomic i64 %d.add, ptr addrspace(3) %d syncscope("cluster") seq_cst, align 81468 1469 %e.load = load atomic float, ptr addrspace(3) %e syncscope("cluster") seq_cst, align 41470 %e.add = fadd float %e.load, 1.1471 store atomic float %e.add, ptr addrspace(3) %e syncscope("cluster") seq_cst, align 41472 1473 %f.load = load atomic double, ptr addrspace(3) %e syncscope("cluster") seq_cst, align 81474 %f.add = fadd double %f.load, 1.1475 store atomic double %f.add, ptr addrspace(3) %e syncscope("cluster") seq_cst, align 81476 1477 ret void1478}1479 1480; CHECK-LABEL: shared_seq_cst_volatile_cluster1481define void @shared_seq_cst_volatile_cluster(ptr addrspace(3) %a, ptr addrspace(3) %b, ptr addrspace(3) %c, ptr addrspace(3) %d, ptr addrspace(3) %e) local_unnamed_addr {1482; CHECK-LABEL: shared_seq_cst_volatile_cluster(1483; CHECK: {1484; CHECK-NEXT: .reg .b16 %rs<5>;1485; CHECK-NEXT: .reg .b32 %r<5>;1486; CHECK-NEXT: .reg .b64 %rd<10>;1487; CHECK-EMPTY:1488; CHECK-NEXT: // %bb.0:1489; CHECK-NEXT: ld.param.b64 %rd1, [shared_seq_cst_volatile_cluster_param_0];1490; CHECK-NEXT: fence.sc.sys;1491; CHECK-NEXT: ld.acquire.sys.shared.b8 %rs1, [%rd1];1492; CHECK-NEXT: ld.param.b64 %rd2, [shared_seq_cst_volatile_cluster_param_1];1493; CHECK-NEXT: add.s16 %rs2, %rs1, 1;1494; CHECK-NEXT: ld.param.b64 %rd3, [shared_seq_cst_volatile_cluster_param_2];1495; CHECK-NEXT: fence.sc.sys;1496; CHECK-NEXT: st.release.sys.shared.b8 [%rd1], %rs2;1497; CHECK-NEXT: ld.param.b64 %rd4, [shared_seq_cst_volatile_cluster_param_3];1498; CHECK-NEXT: fence.sc.sys;1499; CHECK-NEXT: ld.acquire.sys.shared.b16 %rs3, [%rd2];1500; CHECK-NEXT: ld.param.b64 %rd5, [shared_seq_cst_volatile_cluster_param_4];1501; CHECK-NEXT: add.s16 %rs4, %rs3, 1;1502; CHECK-NEXT: fence.sc.sys;1503; CHECK-NEXT: st.release.sys.shared.b16 [%rd2], %rs4;1504; CHECK-NEXT: fence.sc.sys;1505; CHECK-NEXT: ld.acquire.sys.shared.b32 %r1, [%rd3];1506; CHECK-NEXT: add.s32 %r2, %r1, 1;1507; CHECK-NEXT: fence.sc.sys;1508; CHECK-NEXT: st.release.sys.shared.b32 [%rd3], %r2;1509; CHECK-NEXT: fence.sc.sys;1510; CHECK-NEXT: ld.acquire.sys.shared.b64 %rd6, [%rd4];1511; CHECK-NEXT: add.s64 %rd7, %rd6, 1;1512; CHECK-NEXT: fence.sc.sys;1513; CHECK-NEXT: st.release.sys.shared.b64 [%rd4], %rd7;1514; CHECK-NEXT: fence.sc.sys;1515; CHECK-NEXT: ld.acquire.sys.shared.b32 %r3, [%rd5];1516; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;1517; CHECK-NEXT: fence.sc.sys;1518; CHECK-NEXT: st.release.sys.shared.b32 [%rd5], %r4;1519; CHECK-NEXT: fence.sc.sys;1520; CHECK-NEXT: ld.acquire.sys.shared.b64 %rd8, [%rd5];1521; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;1522; CHECK-NEXT: fence.sc.sys;1523; CHECK-NEXT: st.release.sys.shared.b64 [%rd5], %rd9;1524; CHECK-NEXT: ret;1525 %a.load = load atomic volatile i8, ptr addrspace(3) %a syncscope("cluster") seq_cst, align 11526 %a.add = add i8 %a.load, 11527 store atomic volatile i8 %a.add, ptr addrspace(3) %a syncscope("cluster") seq_cst, align 11528 1529 %b.load = load atomic volatile i16, ptr addrspace(3) %b syncscope("cluster") seq_cst, align 21530 %b.add = add i16 %b.load, 11531 store atomic volatile i16 %b.add, ptr addrspace(3) %b syncscope("cluster") seq_cst, align 21532 1533 %c.load = load atomic volatile i32, ptr addrspace(3) %c syncscope("cluster") seq_cst, align 41534 %c.add = add i32 %c.load, 11535 store atomic volatile i32 %c.add, ptr addrspace(3) %c syncscope("cluster") seq_cst, align 41536 1537 %d.load = load atomic volatile i64, ptr addrspace(3) %d syncscope("cluster") seq_cst, align 81538 %d.add = add i64 %d.load, 11539 store atomic volatile i64 %d.add, ptr addrspace(3) %d syncscope("cluster") seq_cst, align 81540 1541 %e.load = load atomic volatile float, ptr addrspace(3) %e syncscope("cluster") seq_cst, align 41542 %e.add = fadd float %e.load, 1.1543 store atomic volatile float %e.add, ptr addrspace(3) %e syncscope("cluster") seq_cst, align 41544 1545 %f.load = load atomic volatile double, ptr addrspace(3) %e syncscope("cluster") seq_cst, align 81546 %f.add = fadd double %f.load, 1.1547 store atomic volatile double %f.add, ptr addrspace(3) %e syncscope("cluster") seq_cst, align 81548 1549 ret void1550}1551 1552;; local statespace1553 1554; CHECK-LABEL: local_unordered_cluster1555define void @local_unordered_cluster(ptr addrspace(5) %a, ptr addrspace(5) %b, ptr addrspace(5) %c, ptr addrspace(5) %d, ptr addrspace(5) %e) local_unnamed_addr {1556; CHECK-LABEL: local_unordered_cluster(1557; CHECK: {1558; CHECK-NEXT: .reg .b16 %rs<5>;1559; CHECK-NEXT: .reg .b32 %r<5>;1560; CHECK-NEXT: .reg .b64 %rd<10>;1561; CHECK-EMPTY:1562; CHECK-NEXT: // %bb.0:1563; CHECK-NEXT: ld.param.b64 %rd1, [local_unordered_cluster_param_0];1564; CHECK-NEXT: ld.local.b8 %rs1, [%rd1];1565; CHECK-NEXT: ld.param.b64 %rd2, [local_unordered_cluster_param_1];1566; CHECK-NEXT: add.s16 %rs2, %rs1, 1;1567; CHECK-NEXT: ld.param.b64 %rd3, [local_unordered_cluster_param_2];1568; CHECK-NEXT: st.local.b8 [%rd1], %rs2;1569; CHECK-NEXT: ld.param.b64 %rd4, [local_unordered_cluster_param_3];1570; CHECK-NEXT: ld.local.b16 %rs3, [%rd2];1571; CHECK-NEXT: ld.param.b64 %rd5, [local_unordered_cluster_param_4];1572; CHECK-NEXT: add.s16 %rs4, %rs3, 1;1573; CHECK-NEXT: st.local.b16 [%rd2], %rs4;1574; CHECK-NEXT: ld.local.b32 %r1, [%rd3];1575; CHECK-NEXT: add.s32 %r2, %r1, 1;1576; CHECK-NEXT: st.local.b32 [%rd3], %r2;1577; CHECK-NEXT: ld.local.b64 %rd6, [%rd4];1578; CHECK-NEXT: add.s64 %rd7, %rd6, 1;1579; CHECK-NEXT: st.local.b64 [%rd4], %rd7;1580; CHECK-NEXT: ld.local.b32 %r3, [%rd5];1581; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;1582; CHECK-NEXT: st.local.b32 [%rd5], %r4;1583; CHECK-NEXT: ld.local.b64 %rd8, [%rd5];1584; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;1585; CHECK-NEXT: st.local.b64 [%rd5], %rd9;1586; CHECK-NEXT: ret;1587 %a.load = load atomic i8, ptr addrspace(5) %a syncscope("cluster") unordered, align 11588 %a.add = add i8 %a.load, 11589 store atomic i8 %a.add, ptr addrspace(5) %a syncscope("cluster") unordered, align 11590 1591 %b.load = load atomic i16, ptr addrspace(5) %b syncscope("cluster") unordered, align 21592 %b.add = add i16 %b.load, 11593 store atomic i16 %b.add, ptr addrspace(5) %b syncscope("cluster") unordered, align 21594 1595 %c.load = load atomic i32, ptr addrspace(5) %c syncscope("cluster") unordered, align 41596 %c.add = add i32 %c.load, 11597 store atomic i32 %c.add, ptr addrspace(5) %c syncscope("cluster") unordered, align 41598 1599 %d.load = load atomic i64, ptr addrspace(5) %d syncscope("cluster") unordered, align 81600 %d.add = add i64 %d.load, 11601 store atomic i64 %d.add, ptr addrspace(5) %d syncscope("cluster") unordered, align 81602 1603 %e.load = load atomic float, ptr addrspace(5) %e syncscope("cluster") unordered, align 41604 %e.add = fadd float %e.load, 1.1605 store atomic float %e.add, ptr addrspace(5) %e syncscope("cluster") unordered, align 41606 1607 %f.load = load atomic double, ptr addrspace(5) %e syncscope("cluster") unordered, align 81608 %f.add = fadd double %f.load, 1.1609 store atomic double %f.add, ptr addrspace(5) %e syncscope("cluster") unordered, align 81610 1611 ret void1612}1613 1614; CHECK-LABEL: local_unordered_volatile_cluster1615define void @local_unordered_volatile_cluster(ptr addrspace(5) %a, ptr addrspace(5) %b, ptr addrspace(5) %c, ptr addrspace(5) %d, ptr addrspace(5) %e) local_unnamed_addr {1616; CHECK-LABEL: local_unordered_volatile_cluster(1617; CHECK: {1618; CHECK-NEXT: .reg .b16 %rs<5>;1619; CHECK-NEXT: .reg .b32 %r<5>;1620; CHECK-NEXT: .reg .b64 %rd<10>;1621; CHECK-EMPTY:1622; CHECK-NEXT: // %bb.0:1623; CHECK-NEXT: ld.param.b64 %rd1, [local_unordered_volatile_cluster_param_0];1624; CHECK-NEXT: ld.local.b8 %rs1, [%rd1];1625; CHECK-NEXT: ld.param.b64 %rd2, [local_unordered_volatile_cluster_param_1];1626; CHECK-NEXT: add.s16 %rs2, %rs1, 1;1627; CHECK-NEXT: ld.param.b64 %rd3, [local_unordered_volatile_cluster_param_2];1628; CHECK-NEXT: st.local.b8 [%rd1], %rs2;1629; CHECK-NEXT: ld.param.b64 %rd4, [local_unordered_volatile_cluster_param_3];1630; CHECK-NEXT: ld.local.b16 %rs3, [%rd2];1631; CHECK-NEXT: ld.param.b64 %rd5, [local_unordered_volatile_cluster_param_4];1632; CHECK-NEXT: add.s16 %rs4, %rs3, 1;1633; CHECK-NEXT: st.local.b16 [%rd2], %rs4;1634; CHECK-NEXT: ld.local.b32 %r1, [%rd3];1635; CHECK-NEXT: add.s32 %r2, %r1, 1;1636; CHECK-NEXT: st.local.b32 [%rd3], %r2;1637; CHECK-NEXT: ld.local.b64 %rd6, [%rd4];1638; CHECK-NEXT: add.s64 %rd7, %rd6, 1;1639; CHECK-NEXT: st.local.b64 [%rd4], %rd7;1640; CHECK-NEXT: ld.local.b32 %r3, [%rd5];1641; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;1642; CHECK-NEXT: st.local.b32 [%rd5], %r4;1643; CHECK-NEXT: ld.local.b64 %rd8, [%rd5];1644; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;1645; CHECK-NEXT: st.local.b64 [%rd5], %rd9;1646; CHECK-NEXT: ret;1647 %a.load = load atomic volatile i8, ptr addrspace(5) %a syncscope("cluster") unordered, align 11648 %a.add = add i8 %a.load, 11649 store atomic volatile i8 %a.add, ptr addrspace(5) %a syncscope("cluster") unordered, align 11650 1651 %b.load = load atomic volatile i16, ptr addrspace(5) %b syncscope("cluster") unordered, align 21652 %b.add = add i16 %b.load, 11653 store atomic volatile i16 %b.add, ptr addrspace(5) %b syncscope("cluster") unordered, align 21654 1655 %c.load = load atomic volatile i32, ptr addrspace(5) %c syncscope("cluster") unordered, align 41656 %c.add = add i32 %c.load, 11657 store atomic volatile i32 %c.add, ptr addrspace(5) %c syncscope("cluster") unordered, align 41658 1659 %d.load = load atomic volatile i64, ptr addrspace(5) %d syncscope("cluster") unordered, align 81660 %d.add = add i64 %d.load, 11661 store atomic volatile i64 %d.add, ptr addrspace(5) %d syncscope("cluster") unordered, align 81662 1663 %e.load = load atomic volatile float, ptr addrspace(5) %e syncscope("cluster") unordered, align 41664 %e.add = fadd float %e.load, 1.1665 store atomic volatile float %e.add, ptr addrspace(5) %e syncscope("cluster") unordered, align 41666 1667 %f.load = load atomic volatile double, ptr addrspace(5) %e syncscope("cluster") unordered, align 81668 %f.add = fadd double %f.load, 1.1669 store atomic volatile double %f.add, ptr addrspace(5) %e syncscope("cluster") unordered, align 81670 1671 ret void1672}1673 1674; CHECK-LABEL: local_monotonic_cluster1675define void @local_monotonic_cluster(ptr addrspace(5) %a, ptr addrspace(5) %b, ptr addrspace(5) %c, ptr addrspace(5) %d, ptr addrspace(5) %e) local_unnamed_addr {1676; CHECK-LABEL: local_monotonic_cluster(1677; CHECK: {1678; CHECK-NEXT: .reg .b16 %rs<5>;1679; CHECK-NEXT: .reg .b32 %r<5>;1680; CHECK-NEXT: .reg .b64 %rd<10>;1681; CHECK-EMPTY:1682; CHECK-NEXT: // %bb.0:1683; CHECK-NEXT: ld.param.b64 %rd1, [local_monotonic_cluster_param_0];1684; CHECK-NEXT: ld.local.b8 %rs1, [%rd1];1685; CHECK-NEXT: ld.param.b64 %rd2, [local_monotonic_cluster_param_1];1686; CHECK-NEXT: add.s16 %rs2, %rs1, 1;1687; CHECK-NEXT: ld.param.b64 %rd3, [local_monotonic_cluster_param_2];1688; CHECK-NEXT: st.local.b8 [%rd1], %rs2;1689; CHECK-NEXT: ld.param.b64 %rd4, [local_monotonic_cluster_param_3];1690; CHECK-NEXT: ld.local.b16 %rs3, [%rd2];1691; CHECK-NEXT: ld.param.b64 %rd5, [local_monotonic_cluster_param_4];1692; CHECK-NEXT: add.s16 %rs4, %rs3, 1;1693; CHECK-NEXT: st.local.b16 [%rd2], %rs4;1694; CHECK-NEXT: ld.local.b32 %r1, [%rd3];1695; CHECK-NEXT: add.s32 %r2, %r1, 1;1696; CHECK-NEXT: st.local.b32 [%rd3], %r2;1697; CHECK-NEXT: ld.local.b64 %rd6, [%rd4];1698; CHECK-NEXT: add.s64 %rd7, %rd6, 1;1699; CHECK-NEXT: st.local.b64 [%rd4], %rd7;1700; CHECK-NEXT: ld.local.b32 %r3, [%rd5];1701; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;1702; CHECK-NEXT: st.local.b32 [%rd5], %r4;1703; CHECK-NEXT: ld.local.b64 %rd8, [%rd5];1704; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;1705; CHECK-NEXT: st.local.b64 [%rd5], %rd9;1706; CHECK-NEXT: ret;1707 %a.load = load atomic i8, ptr addrspace(5) %a syncscope("cluster") monotonic, align 11708 %a.add = add i8 %a.load, 11709 store atomic i8 %a.add, ptr addrspace(5) %a syncscope("cluster") monotonic, align 11710 1711 %b.load = load atomic i16, ptr addrspace(5) %b syncscope("cluster") monotonic, align 21712 %b.add = add i16 %b.load, 11713 store atomic i16 %b.add, ptr addrspace(5) %b syncscope("cluster") monotonic, align 21714 1715 %c.load = load atomic i32, ptr addrspace(5) %c syncscope("cluster") monotonic, align 41716 %c.add = add i32 %c.load, 11717 store atomic i32 %c.add, ptr addrspace(5) %c syncscope("cluster") monotonic, align 41718 1719 %d.load = load atomic i64, ptr addrspace(5) %d syncscope("cluster") monotonic, align 81720 %d.add = add i64 %d.load, 11721 store atomic i64 %d.add, ptr addrspace(5) %d syncscope("cluster") monotonic, align 81722 1723 %e.load = load atomic float, ptr addrspace(5) %e syncscope("cluster") monotonic, align 41724 %e.add = fadd float %e.load, 1.1725 store atomic float %e.add, ptr addrspace(5) %e syncscope("cluster") monotonic, align 41726 1727 %f.load = load atomic double, ptr addrspace(5) %e syncscope("cluster") monotonic, align 81728 %f.add = fadd double %f.load, 1.1729 store atomic double %f.add, ptr addrspace(5) %e syncscope("cluster") monotonic, align 81730 1731 ret void1732}1733 1734; CHECK-LABEL: local_monotonic_volatile_cluster1735define void @local_monotonic_volatile_cluster(ptr addrspace(5) %a, ptr addrspace(5) %b, ptr addrspace(5) %c, ptr addrspace(5) %d, ptr addrspace(5) %e) local_unnamed_addr {1736; CHECK-LABEL: local_monotonic_volatile_cluster(1737; CHECK: {1738; CHECK-NEXT: .reg .b16 %rs<5>;1739; CHECK-NEXT: .reg .b32 %r<5>;1740; CHECK-NEXT: .reg .b64 %rd<10>;1741; CHECK-EMPTY:1742; CHECK-NEXT: // %bb.0:1743; CHECK-NEXT: ld.param.b64 %rd1, [local_monotonic_volatile_cluster_param_0];1744; CHECK-NEXT: ld.local.b8 %rs1, [%rd1];1745; CHECK-NEXT: ld.param.b64 %rd2, [local_monotonic_volatile_cluster_param_1];1746; CHECK-NEXT: add.s16 %rs2, %rs1, 1;1747; CHECK-NEXT: ld.param.b64 %rd3, [local_monotonic_volatile_cluster_param_2];1748; CHECK-NEXT: st.local.b8 [%rd1], %rs2;1749; CHECK-NEXT: ld.param.b64 %rd4, [local_monotonic_volatile_cluster_param_3];1750; CHECK-NEXT: ld.local.b16 %rs3, [%rd2];1751; CHECK-NEXT: ld.param.b64 %rd5, [local_monotonic_volatile_cluster_param_4];1752; CHECK-NEXT: add.s16 %rs4, %rs3, 1;1753; CHECK-NEXT: st.local.b16 [%rd2], %rs4;1754; CHECK-NEXT: ld.local.b32 %r1, [%rd3];1755; CHECK-NEXT: add.s32 %r2, %r1, 1;1756; CHECK-NEXT: st.local.b32 [%rd3], %r2;1757; CHECK-NEXT: ld.local.b64 %rd6, [%rd4];1758; CHECK-NEXT: add.s64 %rd7, %rd6, 1;1759; CHECK-NEXT: st.local.b64 [%rd4], %rd7;1760; CHECK-NEXT: ld.local.b32 %r3, [%rd5];1761; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;1762; CHECK-NEXT: st.local.b32 [%rd5], %r4;1763; CHECK-NEXT: ld.local.b64 %rd8, [%rd5];1764; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;1765; CHECK-NEXT: st.local.b64 [%rd5], %rd9;1766; CHECK-NEXT: ret;1767 %a.load = load atomic volatile i8, ptr addrspace(5) %a syncscope("cluster") monotonic, align 11768 %a.add = add i8 %a.load, 11769 store atomic volatile i8 %a.add, ptr addrspace(5) %a syncscope("cluster") monotonic, align 11770 1771 %b.load = load atomic volatile i16, ptr addrspace(5) %b syncscope("cluster") monotonic, align 21772 %b.add = add i16 %b.load, 11773 store atomic volatile i16 %b.add, ptr addrspace(5) %b syncscope("cluster") monotonic, align 21774 1775 %c.load = load atomic volatile i32, ptr addrspace(5) %c syncscope("cluster") monotonic, align 41776 %c.add = add i32 %c.load, 11777 store atomic volatile i32 %c.add, ptr addrspace(5) %c syncscope("cluster") monotonic, align 41778 1779 %d.load = load atomic volatile i64, ptr addrspace(5) %d syncscope("cluster") monotonic, align 81780 %d.add = add i64 %d.load, 11781 store atomic volatile i64 %d.add, ptr addrspace(5) %d syncscope("cluster") monotonic, align 81782 1783 %e.load = load atomic volatile float, ptr addrspace(5) %e syncscope("cluster") monotonic, align 41784 %e.add = fadd float %e.load, 1.1785 store atomic volatile float %e.add, ptr addrspace(5) %e syncscope("cluster") monotonic, align 41786 1787 %f.load = load atomic volatile double, ptr addrspace(5) %e syncscope("cluster") monotonic, align 81788 %f.add = fadd double %f.load, 1.1789 store atomic volatile double %f.add, ptr addrspace(5) %e syncscope("cluster") monotonic, align 81790 1791 ret void1792}1793 1794; CHECK-LABEL: local_acq_rel_cluster1795define void @local_acq_rel_cluster(ptr addrspace(5) %a, ptr addrspace(5) %b, ptr addrspace(5) %c, ptr addrspace(5) %d, ptr addrspace(5) %e) local_unnamed_addr {1796; CHECK-LABEL: local_acq_rel_cluster(1797; CHECK: {1798; CHECK-NEXT: .reg .b16 %rs<5>;1799; CHECK-NEXT: .reg .b32 %r<5>;1800; CHECK-NEXT: .reg .b64 %rd<10>;1801; CHECK-EMPTY:1802; CHECK-NEXT: // %bb.0:1803; CHECK-NEXT: ld.param.b64 %rd1, [local_acq_rel_cluster_param_0];1804; CHECK-NEXT: ld.local.b8 %rs1, [%rd1];1805; CHECK-NEXT: ld.param.b64 %rd2, [local_acq_rel_cluster_param_1];1806; CHECK-NEXT: add.s16 %rs2, %rs1, 1;1807; CHECK-NEXT: ld.param.b64 %rd3, [local_acq_rel_cluster_param_2];1808; CHECK-NEXT: st.local.b8 [%rd1], %rs2;1809; CHECK-NEXT: ld.param.b64 %rd4, [local_acq_rel_cluster_param_3];1810; CHECK-NEXT: ld.local.b16 %rs3, [%rd2];1811; CHECK-NEXT: ld.param.b64 %rd5, [local_acq_rel_cluster_param_4];1812; CHECK-NEXT: add.s16 %rs4, %rs3, 1;1813; CHECK-NEXT: st.local.b16 [%rd2], %rs4;1814; CHECK-NEXT: ld.local.b32 %r1, [%rd3];1815; CHECK-NEXT: add.s32 %r2, %r1, 1;1816; CHECK-NEXT: st.local.b32 [%rd3], %r2;1817; CHECK-NEXT: ld.local.b64 %rd6, [%rd4];1818; CHECK-NEXT: add.s64 %rd7, %rd6, 1;1819; CHECK-NEXT: st.local.b64 [%rd4], %rd7;1820; CHECK-NEXT: ld.local.b32 %r3, [%rd5];1821; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;1822; CHECK-NEXT: st.local.b32 [%rd5], %r4;1823; CHECK-NEXT: ld.local.b64 %rd8, [%rd5];1824; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;1825; CHECK-NEXT: st.local.b64 [%rd5], %rd9;1826; CHECK-NEXT: ret;1827 %a.load = load atomic i8, ptr addrspace(5) %a syncscope("cluster") acquire, align 11828 %a.add = add i8 %a.load, 11829 store atomic i8 %a.add, ptr addrspace(5) %a syncscope("cluster") release, align 11830 1831 %b.load = load atomic i16, ptr addrspace(5) %b syncscope("cluster") acquire, align 21832 %b.add = add i16 %b.load, 11833 store atomic i16 %b.add, ptr addrspace(5) %b syncscope("cluster") release, align 21834 1835 %c.load = load atomic i32, ptr addrspace(5) %c syncscope("cluster") acquire, align 41836 %c.add = add i32 %c.load, 11837 store atomic i32 %c.add, ptr addrspace(5) %c syncscope("cluster") release, align 41838 1839 %d.load = load atomic i64, ptr addrspace(5) %d syncscope("cluster") acquire, align 81840 %d.add = add i64 %d.load, 11841 store atomic i64 %d.add, ptr addrspace(5) %d syncscope("cluster") release, align 81842 1843 %e.load = load atomic float, ptr addrspace(5) %e syncscope("cluster") acquire, align 41844 %e.add = fadd float %e.load, 1.1845 store atomic float %e.add, ptr addrspace(5) %e syncscope("cluster") release, align 41846 1847 %f.load = load atomic double, ptr addrspace(5) %e syncscope("cluster") acquire, align 81848 %f.add = fadd double %f.load, 1.1849 store atomic double %f.add, ptr addrspace(5) %e syncscope("cluster") release, align 81850 1851 ret void1852}1853 1854; CHECK-LABEL: local_acq_rel_volatile_cluster1855define void @local_acq_rel_volatile_cluster(ptr addrspace(5) %a, ptr addrspace(5) %b, ptr addrspace(5) %c, ptr addrspace(5) %d, ptr addrspace(5) %e) local_unnamed_addr {1856; CHECK-LABEL: local_acq_rel_volatile_cluster(1857; CHECK: {1858; CHECK-NEXT: .reg .b16 %rs<5>;1859; CHECK-NEXT: .reg .b32 %r<5>;1860; CHECK-NEXT: .reg .b64 %rd<10>;1861; CHECK-EMPTY:1862; CHECK-NEXT: // %bb.0:1863; CHECK-NEXT: ld.param.b64 %rd1, [local_acq_rel_volatile_cluster_param_0];1864; CHECK-NEXT: ld.local.b8 %rs1, [%rd1];1865; CHECK-NEXT: ld.param.b64 %rd2, [local_acq_rel_volatile_cluster_param_1];1866; CHECK-NEXT: add.s16 %rs2, %rs1, 1;1867; CHECK-NEXT: ld.param.b64 %rd3, [local_acq_rel_volatile_cluster_param_2];1868; CHECK-NEXT: st.local.b8 [%rd1], %rs2;1869; CHECK-NEXT: ld.param.b64 %rd4, [local_acq_rel_volatile_cluster_param_3];1870; CHECK-NEXT: ld.local.b16 %rs3, [%rd2];1871; CHECK-NEXT: ld.param.b64 %rd5, [local_acq_rel_volatile_cluster_param_4];1872; CHECK-NEXT: add.s16 %rs4, %rs3, 1;1873; CHECK-NEXT: st.local.b16 [%rd2], %rs4;1874; CHECK-NEXT: ld.local.b32 %r1, [%rd3];1875; CHECK-NEXT: add.s32 %r2, %r1, 1;1876; CHECK-NEXT: st.local.b32 [%rd3], %r2;1877; CHECK-NEXT: ld.local.b64 %rd6, [%rd4];1878; CHECK-NEXT: add.s64 %rd7, %rd6, 1;1879; CHECK-NEXT: st.local.b64 [%rd4], %rd7;1880; CHECK-NEXT: ld.local.b32 %r3, [%rd5];1881; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;1882; CHECK-NEXT: st.local.b32 [%rd5], %r4;1883; CHECK-NEXT: ld.local.b64 %rd8, [%rd5];1884; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;1885; CHECK-NEXT: st.local.b64 [%rd5], %rd9;1886; CHECK-NEXT: ret;1887 %a.load = load atomic volatile i8, ptr addrspace(5) %a syncscope("cluster") acquire, align 11888 %a.add = add i8 %a.load, 11889 store atomic volatile i8 %a.add, ptr addrspace(5) %a syncscope("cluster") release, align 11890 1891 %b.load = load atomic volatile i16, ptr addrspace(5) %b syncscope("cluster") acquire, align 21892 %b.add = add i16 %b.load, 11893 store atomic volatile i16 %b.add, ptr addrspace(5) %b syncscope("cluster") release, align 21894 1895 %c.load = load atomic volatile i32, ptr addrspace(5) %c syncscope("cluster") acquire, align 41896 %c.add = add i32 %c.load, 11897 store atomic volatile i32 %c.add, ptr addrspace(5) %c syncscope("cluster") release, align 41898 1899 %d.load = load atomic volatile i64, ptr addrspace(5) %d syncscope("cluster") acquire, align 81900 %d.add = add i64 %d.load, 11901 store atomic volatile i64 %d.add, ptr addrspace(5) %d syncscope("cluster") release, align 81902 1903 %e.load = load atomic volatile float, ptr addrspace(5) %e syncscope("cluster") acquire, align 41904 %e.add = fadd float %e.load, 1.1905 store atomic volatile float %e.add, ptr addrspace(5) %e syncscope("cluster") release, align 41906 1907 %f.load = load atomic volatile double, ptr addrspace(5) %e syncscope("cluster") acquire, align 81908 %f.add = fadd double %f.load, 1.1909 store atomic volatile double %f.add, ptr addrspace(5) %e syncscope("cluster") release, align 81910 1911 ret void1912}1913 1914; CHECK-LABEL: local_seq_cst_cluster1915define void @local_seq_cst_cluster(ptr addrspace(5) %a, ptr addrspace(5) %b, ptr addrspace(5) %c, ptr addrspace(5) %d, ptr addrspace(5) %e) local_unnamed_addr {1916; CHECK-LABEL: local_seq_cst_cluster(1917; CHECK: {1918; CHECK-NEXT: .reg .b16 %rs<5>;1919; CHECK-NEXT: .reg .b32 %r<5>;1920; CHECK-NEXT: .reg .b64 %rd<10>;1921; CHECK-EMPTY:1922; CHECK-NEXT: // %bb.0:1923; CHECK-NEXT: ld.param.b64 %rd1, [local_seq_cst_cluster_param_0];1924; CHECK-NEXT: ld.local.b8 %rs1, [%rd1];1925; CHECK-NEXT: ld.param.b64 %rd2, [local_seq_cst_cluster_param_1];1926; CHECK-NEXT: add.s16 %rs2, %rs1, 1;1927; CHECK-NEXT: ld.param.b64 %rd3, [local_seq_cst_cluster_param_2];1928; CHECK-NEXT: st.local.b8 [%rd1], %rs2;1929; CHECK-NEXT: ld.param.b64 %rd4, [local_seq_cst_cluster_param_3];1930; CHECK-NEXT: ld.local.b16 %rs3, [%rd2];1931; CHECK-NEXT: ld.param.b64 %rd5, [local_seq_cst_cluster_param_4];1932; CHECK-NEXT: add.s16 %rs4, %rs3, 1;1933; CHECK-NEXT: st.local.b16 [%rd2], %rs4;1934; CHECK-NEXT: ld.local.b32 %r1, [%rd3];1935; CHECK-NEXT: add.s32 %r2, %r1, 1;1936; CHECK-NEXT: st.local.b32 [%rd3], %r2;1937; CHECK-NEXT: ld.local.b64 %rd6, [%rd4];1938; CHECK-NEXT: add.s64 %rd7, %rd6, 1;1939; CHECK-NEXT: st.local.b64 [%rd4], %rd7;1940; CHECK-NEXT: ld.local.b32 %r3, [%rd5];1941; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;1942; CHECK-NEXT: st.local.b32 [%rd5], %r4;1943; CHECK-NEXT: ld.local.b64 %rd8, [%rd5];1944; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;1945; CHECK-NEXT: st.local.b64 [%rd5], %rd9;1946; CHECK-NEXT: ret;1947 %a.load = load atomic i8, ptr addrspace(5) %a syncscope("cluster") seq_cst, align 11948 %a.add = add i8 %a.load, 11949 store atomic i8 %a.add, ptr addrspace(5) %a syncscope("cluster") seq_cst, align 11950 1951 %b.load = load atomic i16, ptr addrspace(5) %b syncscope("cluster") seq_cst, align 21952 %b.add = add i16 %b.load, 11953 store atomic i16 %b.add, ptr addrspace(5) %b syncscope("cluster") seq_cst, align 21954 1955 %c.load = load atomic i32, ptr addrspace(5) %c syncscope("cluster") seq_cst, align 41956 %c.add = add i32 %c.load, 11957 store atomic i32 %c.add, ptr addrspace(5) %c syncscope("cluster") seq_cst, align 41958 1959 %d.load = load atomic i64, ptr addrspace(5) %d syncscope("cluster") seq_cst, align 81960 %d.add = add i64 %d.load, 11961 store atomic i64 %d.add, ptr addrspace(5) %d syncscope("cluster") seq_cst, align 81962 1963 %e.load = load atomic float, ptr addrspace(5) %e syncscope("cluster") seq_cst, align 41964 %e.add = fadd float %e.load, 1.1965 store atomic float %e.add, ptr addrspace(5) %e syncscope("cluster") seq_cst, align 41966 1967 %f.load = load atomic double, ptr addrspace(5) %e syncscope("cluster") seq_cst, align 81968 %f.add = fadd double %f.load, 1.1969 store atomic double %f.add, ptr addrspace(5) %e syncscope("cluster") seq_cst, align 81970 1971 ret void1972}1973 1974; CHECK-LABEL: local_seq_cst_volatile_cluster1975define void @local_seq_cst_volatile_cluster(ptr addrspace(5) %a, ptr addrspace(5) %b, ptr addrspace(5) %c, ptr addrspace(5) %d, ptr addrspace(5) %e) local_unnamed_addr {1976; CHECK-LABEL: local_seq_cst_volatile_cluster(1977; CHECK: {1978; CHECK-NEXT: .reg .b16 %rs<5>;1979; CHECK-NEXT: .reg .b32 %r<5>;1980; CHECK-NEXT: .reg .b64 %rd<10>;1981; CHECK-EMPTY:1982; CHECK-NEXT: // %bb.0:1983; CHECK-NEXT: ld.param.b64 %rd1, [local_seq_cst_volatile_cluster_param_0];1984; CHECK-NEXT: ld.local.b8 %rs1, [%rd1];1985; CHECK-NEXT: ld.param.b64 %rd2, [local_seq_cst_volatile_cluster_param_1];1986; CHECK-NEXT: add.s16 %rs2, %rs1, 1;1987; CHECK-NEXT: ld.param.b64 %rd3, [local_seq_cst_volatile_cluster_param_2];1988; CHECK-NEXT: st.local.b8 [%rd1], %rs2;1989; CHECK-NEXT: ld.param.b64 %rd4, [local_seq_cst_volatile_cluster_param_3];1990; CHECK-NEXT: ld.local.b16 %rs3, [%rd2];1991; CHECK-NEXT: ld.param.b64 %rd5, [local_seq_cst_volatile_cluster_param_4];1992; CHECK-NEXT: add.s16 %rs4, %rs3, 1;1993; CHECK-NEXT: st.local.b16 [%rd2], %rs4;1994; CHECK-NEXT: ld.local.b32 %r1, [%rd3];1995; CHECK-NEXT: add.s32 %r2, %r1, 1;1996; CHECK-NEXT: st.local.b32 [%rd3], %r2;1997; CHECK-NEXT: ld.local.b64 %rd6, [%rd4];1998; CHECK-NEXT: add.s64 %rd7, %rd6, 1;1999; CHECK-NEXT: st.local.b64 [%rd4], %rd7;2000; CHECK-NEXT: ld.local.b32 %r3, [%rd5];2001; CHECK-NEXT: add.rn.f32 %r4, %r3, 0f3F800000;2002; CHECK-NEXT: st.local.b32 [%rd5], %r4;2003; CHECK-NEXT: ld.local.b64 %rd8, [%rd5];2004; CHECK-NEXT: add.rn.f64 %rd9, %rd8, 0d3FF0000000000000;2005; CHECK-NEXT: st.local.b64 [%rd5], %rd9;2006; CHECK-NEXT: ret;2007 %a.load = load atomic volatile i8, ptr addrspace(5) %a syncscope("cluster") seq_cst, align 12008 %a.add = add i8 %a.load, 12009 store atomic volatile i8 %a.add, ptr addrspace(5) %a syncscope("cluster") seq_cst, align 12010 2011 %b.load = load atomic volatile i16, ptr addrspace(5) %b syncscope("cluster") seq_cst, align 22012 %b.add = add i16 %b.load, 12013 store atomic volatile i16 %b.add, ptr addrspace(5) %b syncscope("cluster") seq_cst, align 22014 2015 %c.load = load atomic volatile i32, ptr addrspace(5) %c syncscope("cluster") seq_cst, align 42016 %c.add = add i32 %c.load, 12017 store atomic volatile i32 %c.add, ptr addrspace(5) %c syncscope("cluster") seq_cst, align 42018 2019 %d.load = load atomic volatile i64, ptr addrspace(5) %d syncscope("cluster") seq_cst, align 82020 %d.add = add i64 %d.load, 12021 store atomic volatile i64 %d.add, ptr addrspace(5) %d syncscope("cluster") seq_cst, align 82022 2023 %e.load = load atomic volatile float, ptr addrspace(5) %e syncscope("cluster") seq_cst, align 42024 %e.add = fadd float %e.load, 1.2025 store atomic volatile float %e.add, ptr addrspace(5) %e syncscope("cluster") seq_cst, align 42026 2027 %f.load = load atomic volatile double, ptr addrspace(5) %e syncscope("cluster") seq_cst, align 82028 %f.add = fadd double %f.load, 1.2029 store atomic volatile double %f.add, ptr addrspace(5) %e syncscope("cluster") seq_cst, align 82030 2031 ret void2032}2033