1034 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: not llc < %s -mcpu=sm_90 -mattr=+ptx82 2>&1 | FileCheck %s --check-prefix=ERROR3; RUN: not llc < %s -mcpu=sm_80 -mattr=+ptx84 2>&1 | FileCheck %s --check-prefix=ERROR4; RUN: llc < %s -mcpu=sm_90 -mattr=+ptx84 | FileCheck %s --check-prefix=CHECK5; RUN: %if ptxas-sm_90 && ptxas-isa-8.4 %{ llc < %s -mcpu=sm_90 -mattr=+ptx84 | %ptxas-verify -arch=sm_90 %}6 7;; TODO: Update cmpxchg.py so that it can automatically generate the IR for8;; these test cases.9 10target triple = "nvptx64-nvidia-cuda"11 12;; Check that the first couple of error messages are correct.13; ERROR: error: unsupported cmpxchg14; ERROR: error: unsupported cmpxchg15 16define i128 @test_xchg_generic(ptr %addr, i128 %amt) {17; CHECK-LABEL: test_xchg_generic(18; CHECK: {19; CHECK-NEXT: .reg .b64 %rd<6>;20; CHECK-EMPTY:21; CHECK-NEXT: // %bb.0:22; CHECK-NEXT: ld.param.b64 %rd1, [test_xchg_generic_param_0];23; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_xchg_generic_param_1];24; CHECK-NEXT: {25; CHECK-NEXT: .reg .b128 amt, dst;26; CHECK-NEXT: mov.b128 amt, {%rd2, %rd3};27; CHECK-NEXT: atom.release.sys.exch.b128 dst, [%rd1], amt;28; CHECK-NEXT: mov.b128 {%rd4, %rd5}, dst;29; CHECK-NEXT: }30; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};31; CHECK-NEXT: ret;32 %old = atomicrmw xchg ptr %addr, i128 %amt release33 ret i128 %old34}35 36define i128 @test_xchg_global(ptr addrspace(1) %addr, i128 %amt) {37; CHECK-LABEL: test_xchg_global(38; CHECK: {39; CHECK-NEXT: .reg .b64 %rd<6>;40; CHECK-EMPTY:41; CHECK-NEXT: // %bb.0:42; CHECK-NEXT: ld.param.b64 %rd1, [test_xchg_global_param_0];43; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_xchg_global_param_1];44; CHECK-NEXT: {45; CHECK-NEXT: .reg .b128 amt, dst;46; CHECK-NEXT: mov.b128 amt, {%rd2, %rd3};47; CHECK-NEXT: atom.release.sys.global.exch.b128 dst, [%rd1], amt;48; CHECK-NEXT: mov.b128 {%rd4, %rd5}, dst;49; CHECK-NEXT: }50; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};51; CHECK-NEXT: ret;52 %old = atomicrmw xchg ptr addrspace(1) %addr, i128 %amt release53 ret i128 %old54}55 56define i128 @test_xchg_shared(ptr addrspace(3) %addr, i128 %amt) {57; CHECK-LABEL: test_xchg_shared(58; CHECK: {59; CHECK-NEXT: .reg .b64 %rd<6>;60; CHECK-EMPTY:61; CHECK-NEXT: // %bb.0:62; CHECK-NEXT: ld.param.b64 %rd1, [test_xchg_shared_param_0];63; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_xchg_shared_param_1];64; CHECK-NEXT: {65; CHECK-NEXT: .reg .b128 amt, dst;66; CHECK-NEXT: mov.b128 amt, {%rd2, %rd3};67; CHECK-NEXT: atom.release.sys.shared.exch.b128 dst, [%rd1], amt;68; CHECK-NEXT: mov.b128 {%rd4, %rd5}, dst;69; CHECK-NEXT: }70; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};71; CHECK-NEXT: ret;72 %old = atomicrmw xchg ptr addrspace(3) %addr, i128 %amt release73 ret i128 %old74}75 76define i128 @test_xchg_shared_cluster(ptr addrspace(7) %addr, i128 %amt) {77; CHECK-LABEL: test_xchg_shared_cluster(78; CHECK: {79; CHECK-NEXT: .reg .b64 %rd<6>;80; CHECK-EMPTY:81; CHECK-NEXT: // %bb.0:82; CHECK-NEXT: ld.param.b64 %rd1, [test_xchg_shared_cluster_param_0];83; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_xchg_shared_cluster_param_1];84; CHECK-NEXT: {85; CHECK-NEXT: .reg .b128 amt, dst;86; CHECK-NEXT: mov.b128 amt, {%rd2, %rd3};87; CHECK-NEXT: atom.release.sys.shared::cluster.exch.b128 dst, [%rd1], amt;88; CHECK-NEXT: mov.b128 {%rd4, %rd5}, dst;89; CHECK-NEXT: }90; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};91; CHECK-NEXT: ret;92 %old = atomicrmw xchg ptr addrspace(7) %addr, i128 %amt release93 ret i128 %old94}95 96define i128 @test_xchg_block(ptr %addr, i128 %amt) {97; CHECK-LABEL: test_xchg_block(98; CHECK: {99; CHECK-NEXT: .reg .b64 %rd<6>;100; CHECK-EMPTY:101; CHECK-NEXT: // %bb.0:102; CHECK-NEXT: ld.param.b64 %rd1, [test_xchg_block_param_0];103; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_xchg_block_param_1];104; CHECK-NEXT: {105; CHECK-NEXT: .reg .b128 amt, dst;106; CHECK-NEXT: mov.b128 amt, {%rd2, %rd3};107; CHECK-NEXT: atom.release.cta.exch.b128 dst, [%rd1], amt;108; CHECK-NEXT: mov.b128 {%rd4, %rd5}, dst;109; CHECK-NEXT: }110; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};111; CHECK-NEXT: ret;112 %old = atomicrmw xchg ptr %addr, i128 %amt syncscope("block") release113 ret i128 %old114}115 116define i128 @test_xchg_cluster(ptr %addr, i128 %amt) {117; CHECK-LABEL: test_xchg_cluster(118; CHECK: {119; CHECK-NEXT: .reg .b64 %rd<6>;120; CHECK-EMPTY:121; CHECK-NEXT: // %bb.0:122; CHECK-NEXT: ld.param.b64 %rd1, [test_xchg_cluster_param_0];123; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_xchg_cluster_param_1];124; CHECK-NEXT: {125; CHECK-NEXT: .reg .b128 amt, dst;126; CHECK-NEXT: mov.b128 amt, {%rd2, %rd3};127; CHECK-NEXT: atom.release.cluster.exch.b128 dst, [%rd1], amt;128; CHECK-NEXT: mov.b128 {%rd4, %rd5}, dst;129; CHECK-NEXT: }130; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};131; CHECK-NEXT: ret;132 %old = atomicrmw xchg ptr %addr, i128 %amt syncscope("cluster") release133 ret i128 %old134}135 136define i128 @test_xchg_gpu(ptr %addr, i128 %amt) {137; CHECK-LABEL: test_xchg_gpu(138; CHECK: {139; CHECK-NEXT: .reg .b64 %rd<6>;140; CHECK-EMPTY:141; CHECK-NEXT: // %bb.0:142; CHECK-NEXT: ld.param.b64 %rd1, [test_xchg_gpu_param_0];143; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_xchg_gpu_param_1];144; CHECK-NEXT: {145; CHECK-NEXT: .reg .b128 amt, dst;146; CHECK-NEXT: mov.b128 amt, {%rd2, %rd3};147; CHECK-NEXT: atom.release.gpu.exch.b128 dst, [%rd1], amt;148; CHECK-NEXT: mov.b128 {%rd4, %rd5}, dst;149; CHECK-NEXT: }150; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};151; CHECK-NEXT: ret;152 %old = atomicrmw xchg ptr %addr, i128 %amt syncscope("device") release153 ret i128 %old154}155 156define i128 @test_xchg_sys(ptr %addr, i128 %amt) {157; CHECK-LABEL: test_xchg_sys(158; CHECK: {159; CHECK-NEXT: .reg .b64 %rd<6>;160; CHECK-EMPTY:161; CHECK-NEXT: // %bb.0:162; CHECK-NEXT: ld.param.b64 %rd1, [test_xchg_sys_param_0];163; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_xchg_sys_param_1];164; CHECK-NEXT: {165; CHECK-NEXT: .reg .b128 amt, dst;166; CHECK-NEXT: mov.b128 amt, {%rd2, %rd3};167; CHECK-NEXT: atom.release.sys.exch.b128 dst, [%rd1], amt;168; CHECK-NEXT: mov.b128 {%rd4, %rd5}, dst;169; CHECK-NEXT: }170; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};171; CHECK-NEXT: ret;172 %old = atomicrmw xchg ptr %addr, i128 %amt release173 ret i128 %old174}175 176define i128 @test_xchg_relaxed(ptr %addr, i128 %amt) {177; CHECK-LABEL: test_xchg_relaxed(178; CHECK: {179; CHECK-NEXT: .reg .b64 %rd<6>;180; CHECK-EMPTY:181; CHECK-NEXT: // %bb.0:182; CHECK-NEXT: ld.param.b64 %rd1, [test_xchg_relaxed_param_0];183; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_xchg_relaxed_param_1];184; CHECK-NEXT: {185; CHECK-NEXT: .reg .b128 amt, dst;186; CHECK-NEXT: mov.b128 amt, {%rd2, %rd3};187; CHECK-NEXT: atom.relaxed.sys.exch.b128 dst, [%rd1], amt;188; CHECK-NEXT: mov.b128 {%rd4, %rd5}, dst;189; CHECK-NEXT: }190; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};191; CHECK-NEXT: ret;192 %old = atomicrmw xchg ptr %addr, i128 %amt monotonic193 ret i128 %old194}195 196define i128 @test_xchg_acquire(ptr %addr, i128 %amt) {197; CHECK-LABEL: test_xchg_acquire(198; CHECK: {199; CHECK-NEXT: .reg .b64 %rd<6>;200; CHECK-EMPTY:201; CHECK-NEXT: // %bb.0:202; CHECK-NEXT: ld.param.b64 %rd1, [test_xchg_acquire_param_0];203; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_xchg_acquire_param_1];204; CHECK-NEXT: {205; CHECK-NEXT: .reg .b128 amt, dst;206; CHECK-NEXT: mov.b128 amt, {%rd2, %rd3};207; CHECK-NEXT: atom.acquire.sys.exch.b128 dst, [%rd1], amt;208; CHECK-NEXT: mov.b128 {%rd4, %rd5}, dst;209; CHECK-NEXT: }210; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};211; CHECK-NEXT: ret;212 %old = atomicrmw xchg ptr %addr, i128 %amt acquire213 ret i128 %old214}215 216define i128 @test_xchg_release(ptr %addr, i128 %amt) {217; CHECK-LABEL: test_xchg_release(218; CHECK: {219; CHECK-NEXT: .reg .b64 %rd<6>;220; CHECK-EMPTY:221; CHECK-NEXT: // %bb.0:222; CHECK-NEXT: ld.param.b64 %rd1, [test_xchg_release_param_0];223; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_xchg_release_param_1];224; CHECK-NEXT: {225; CHECK-NEXT: .reg .b128 amt, dst;226; CHECK-NEXT: mov.b128 amt, {%rd2, %rd3};227; CHECK-NEXT: atom.release.sys.exch.b128 dst, [%rd1], amt;228; CHECK-NEXT: mov.b128 {%rd4, %rd5}, dst;229; CHECK-NEXT: }230; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};231; CHECK-NEXT: ret;232 %old = atomicrmw xchg ptr %addr, i128 %amt release233 ret i128 %old234}235 236define i128 @test_xchg_acq_rel(ptr %addr, i128 %amt) {237; CHECK-LABEL: test_xchg_acq_rel(238; CHECK: {239; CHECK-NEXT: .reg .b64 %rd<6>;240; CHECK-EMPTY:241; CHECK-NEXT: // %bb.0:242; CHECK-NEXT: ld.param.b64 %rd1, [test_xchg_acq_rel_param_0];243; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_xchg_acq_rel_param_1];244; CHECK-NEXT: {245; CHECK-NEXT: .reg .b128 amt, dst;246; CHECK-NEXT: mov.b128 amt, {%rd2, %rd3};247; CHECK-NEXT: atom.acq_rel.sys.exch.b128 dst, [%rd1], amt;248; CHECK-NEXT: mov.b128 {%rd4, %rd5}, dst;249; CHECK-NEXT: }250; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};251; CHECK-NEXT: ret;252 %old = atomicrmw xchg ptr %addr, i128 %amt acq_rel253 ret i128 %old254}255 256define i128 @test_cmpxchg_generic(ptr %addr, i128 %cmp, i128 %new) {257; CHECK-LABEL: test_cmpxchg_generic(258; CHECK: {259; CHECK-NEXT: .reg .b64 %rd<8>;260; CHECK-EMPTY:261; CHECK-NEXT: // %bb.0:262; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_generic_param_0];263; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_generic_param_1];264; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_generic_param_2];265; CHECK-NEXT: {266; CHECK-NEXT: .reg .b128 cmp, swap, dst;267; CHECK-NEXT: mov.b128 cmp, {%rd2, %rd3};268; CHECK-NEXT: mov.b128 swap, {%rd4, %rd5};269; CHECK-NEXT: atom.relaxed.sys.cas.b128 dst, [%rd1], cmp, swap;270; CHECK-NEXT: mov.b128 {%rd6, %rd7}, dst;271; CHECK-NEXT: }272; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};273; CHECK-NEXT: ret;274 %pairold = cmpxchg ptr %addr, i128 %cmp, i128 %new monotonic monotonic275 ret i128 %new276}277 278define i128 @test_cmpxchg_global(ptr addrspace(1) %addr, i128 %cmp, i128 %new) {279; CHECK-LABEL: test_cmpxchg_global(280; CHECK: {281; CHECK-NEXT: .reg .b64 %rd<8>;282; CHECK-EMPTY:283; CHECK-NEXT: // %bb.0:284; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_global_param_0];285; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_global_param_1];286; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_global_param_2];287; CHECK-NEXT: {288; CHECK-NEXT: .reg .b128 cmp, swap, dst;289; CHECK-NEXT: mov.b128 cmp, {%rd2, %rd3};290; CHECK-NEXT: mov.b128 swap, {%rd4, %rd5};291; CHECK-NEXT: atom.relaxed.sys.global.cas.b128 dst, [%rd1], cmp, swap;292; CHECK-NEXT: mov.b128 {%rd6, %rd7}, dst;293; CHECK-NEXT: }294; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};295; CHECK-NEXT: ret;296 %pairold = cmpxchg ptr addrspace(1) %addr, i128 %cmp, i128 %new monotonic monotonic297 ret i128 %new298}299 300define i128 @test_cmpxchg_shared(ptr addrspace(3) %addr, i128 %cmp, i128 %new) {301; CHECK-LABEL: test_cmpxchg_shared(302; CHECK: {303; CHECK-NEXT: .reg .b64 %rd<8>;304; CHECK-EMPTY:305; CHECK-NEXT: // %bb.0:306; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_shared_param_0];307; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_shared_param_1];308; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_shared_param_2];309; CHECK-NEXT: {310; CHECK-NEXT: .reg .b128 cmp, swap, dst;311; CHECK-NEXT: mov.b128 cmp, {%rd2, %rd3};312; CHECK-NEXT: mov.b128 swap, {%rd4, %rd5};313; CHECK-NEXT: atom.relaxed.sys.shared.cas.b128 dst, [%rd1], cmp, swap;314; CHECK-NEXT: mov.b128 {%rd6, %rd7}, dst;315; CHECK-NEXT: }316; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};317; CHECK-NEXT: ret;318 %pairold = cmpxchg ptr addrspace(3) %addr, i128 %cmp, i128 %new monotonic monotonic319 ret i128 %new320}321 322define i128 @test_cmpxchg_block(ptr %addr, i128 %cmp, i128 %new) {323; CHECK-LABEL: test_cmpxchg_block(324; CHECK: {325; CHECK-NEXT: .reg .b64 %rd<8>;326; CHECK-EMPTY:327; CHECK-NEXT: // %bb.0:328; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_block_param_0];329; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_block_param_1];330; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_block_param_2];331; CHECK-NEXT: {332; CHECK-NEXT: .reg .b128 cmp, swap, dst;333; CHECK-NEXT: mov.b128 cmp, {%rd2, %rd3};334; CHECK-NEXT: mov.b128 swap, {%rd4, %rd5};335; CHECK-NEXT: atom.relaxed.cta.cas.b128 dst, [%rd1], cmp, swap;336; CHECK-NEXT: mov.b128 {%rd6, %rd7}, dst;337; CHECK-NEXT: }338; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};339; CHECK-NEXT: ret;340 %pairold = cmpxchg ptr %addr, i128 %cmp, i128 %new syncscope("block") monotonic monotonic341 ret i128 %new342}343 344define i128 @test_cmpxchg_cluster(ptr %addr, i128 %cmp, i128 %new) {345; CHECK-LABEL: test_cmpxchg_cluster(346; CHECK: {347; CHECK-NEXT: .reg .b64 %rd<8>;348; CHECK-EMPTY:349; CHECK-NEXT: // %bb.0:350; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_cluster_param_0];351; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_cluster_param_1];352; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_cluster_param_2];353; CHECK-NEXT: {354; CHECK-NEXT: .reg .b128 cmp, swap, dst;355; CHECK-NEXT: mov.b128 cmp, {%rd2, %rd3};356; CHECK-NEXT: mov.b128 swap, {%rd4, %rd5};357; CHECK-NEXT: atom.relaxed.cluster.cas.b128 dst, [%rd1], cmp, swap;358; CHECK-NEXT: mov.b128 {%rd6, %rd7}, dst;359; CHECK-NEXT: }360; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};361; CHECK-NEXT: ret;362 %pairold = cmpxchg ptr %addr, i128 %cmp, i128 %new syncscope("cluster") monotonic monotonic363 ret i128 %new364}365 366define i128 @test_cmpxchg_gpu(ptr %addr, i128 %cmp, i128 %new) {367; CHECK-LABEL: test_cmpxchg_gpu(368; CHECK: {369; CHECK-NEXT: .reg .b64 %rd<8>;370; CHECK-EMPTY:371; CHECK-NEXT: // %bb.0:372; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_gpu_param_0];373; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_gpu_param_1];374; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_gpu_param_2];375; CHECK-NEXT: {376; CHECK-NEXT: .reg .b128 cmp, swap, dst;377; CHECK-NEXT: mov.b128 cmp, {%rd2, %rd3};378; CHECK-NEXT: mov.b128 swap, {%rd4, %rd5};379; CHECK-NEXT: atom.relaxed.gpu.cas.b128 dst, [%rd1], cmp, swap;380; CHECK-NEXT: mov.b128 {%rd6, %rd7}, dst;381; CHECK-NEXT: }382; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};383; CHECK-NEXT: ret;384 %pairold = cmpxchg ptr %addr, i128 %cmp, i128 %new syncscope("device") monotonic monotonic385 ret i128 %new386}387 388define i128 @test_cmpxchg_shared_cluster(ptr addrspace(7) %addr, i128 %cmp, i128 %new) {389; CHECK-LABEL: test_cmpxchg_shared_cluster(390; CHECK: {391; CHECK-NEXT: .reg .b64 %rd<8>;392; CHECK-EMPTY:393; CHECK-NEXT: // %bb.0:394; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_shared_cluster_param_0];395; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_shared_cluster_param_1];396; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_shared_cluster_param_2];397; CHECK-NEXT: {398; CHECK-NEXT: .reg .b128 cmp, swap, dst;399; CHECK-NEXT: mov.b128 cmp, {%rd2, %rd3};400; CHECK-NEXT: mov.b128 swap, {%rd4, %rd5};401; CHECK-NEXT: atom.relaxed.sys.shared::cluster.cas.b128 dst, [%rd1], cmp, swap;402; CHECK-NEXT: mov.b128 {%rd6, %rd7}, dst;403; CHECK-NEXT: }404; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};405; CHECK-NEXT: ret;406 %pairold = cmpxchg ptr addrspace(7) %addr, i128 %cmp, i128 %new monotonic monotonic407 ret i128 %new408}409 410define i128 @test_cmpxchg_monotonic_monotonic(ptr %addr, i128 %cmp, i128 %new) {411; CHECK-LABEL: test_cmpxchg_monotonic_monotonic(412; CHECK: {413; CHECK-NEXT: .reg .b64 %rd<8>;414; CHECK-EMPTY:415; CHECK-NEXT: // %bb.0:416; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_monotonic_monotonic_param_0];417; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_monotonic_monotonic_param_1];418; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_monotonic_monotonic_param_2];419; CHECK-NEXT: {420; CHECK-NEXT: .reg .b128 cmp, swap, dst;421; CHECK-NEXT: mov.b128 cmp, {%rd2, %rd3};422; CHECK-NEXT: mov.b128 swap, {%rd4, %rd5};423; CHECK-NEXT: atom.relaxed.sys.cas.b128 dst, [%rd1], cmp, swap;424; CHECK-NEXT: mov.b128 {%rd6, %rd7}, dst;425; CHECK-NEXT: }426; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};427; CHECK-NEXT: ret;428 %pairold = cmpxchg ptr %addr, i128 %cmp, i128 %new monotonic monotonic429 ret i128 %new430}431 432define i128 @test_cmpxchg_monotonic_acquire(ptr %addr, i128 %cmp, i128 %new) {433; CHECK-LABEL: test_cmpxchg_monotonic_acquire(434; CHECK: {435; CHECK-NEXT: .reg .b64 %rd<8>;436; CHECK-EMPTY:437; CHECK-NEXT: // %bb.0:438; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_monotonic_acquire_param_0];439; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_monotonic_acquire_param_1];440; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_monotonic_acquire_param_2];441; CHECK-NEXT: {442; CHECK-NEXT: .reg .b128 cmp, swap, dst;443; CHECK-NEXT: mov.b128 cmp, {%rd2, %rd3};444; CHECK-NEXT: mov.b128 swap, {%rd4, %rd5};445; CHECK-NEXT: atom.acquire.sys.cas.b128 dst, [%rd1], cmp, swap;446; CHECK-NEXT: mov.b128 {%rd6, %rd7}, dst;447; CHECK-NEXT: }448; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};449; CHECK-NEXT: ret;450 %pairold = cmpxchg ptr %addr, i128 %cmp, i128 %new monotonic acquire451 ret i128 %new452}453 454define i128 @test_cmpxchg_monotonic_seq_cst(ptr %addr, i128 %cmp, i128 %new) {455; CHECK-LABEL: test_cmpxchg_monotonic_seq_cst(456; CHECK: {457; CHECK-NEXT: .reg .b64 %rd<8>;458; CHECK-EMPTY:459; CHECK-NEXT: // %bb.0:460; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_monotonic_seq_cst_param_0];461; CHECK-NEXT: fence.sc.sys;462; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_monotonic_seq_cst_param_1];463; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_monotonic_seq_cst_param_2];464; CHECK-NEXT: {465; CHECK-NEXT: .reg .b128 cmp, swap, dst;466; CHECK-NEXT: mov.b128 cmp, {%rd2, %rd3};467; CHECK-NEXT: mov.b128 swap, {%rd4, %rd5};468; CHECK-NEXT: atom.acquire.sys.cas.b128 dst, [%rd1], cmp, swap;469; CHECK-NEXT: mov.b128 {%rd6, %rd7}, dst;470; CHECK-NEXT: }471; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};472; CHECK-NEXT: ret;473 %pairold = cmpxchg ptr %addr, i128 %cmp, i128 %new monotonic seq_cst474 ret i128 %new475}476 477define i128 @test_cmpxchg_acquire_monotonic(ptr %addr, i128 %cmp, i128 %new) {478; CHECK-LABEL: test_cmpxchg_acquire_monotonic(479; CHECK: {480; CHECK-NEXT: .reg .b64 %rd<8>;481; CHECK-EMPTY:482; CHECK-NEXT: // %bb.0:483; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_acquire_monotonic_param_0];484; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_acquire_monotonic_param_1];485; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_acquire_monotonic_param_2];486; CHECK-NEXT: {487; CHECK-NEXT: .reg .b128 cmp, swap, dst;488; CHECK-NEXT: mov.b128 cmp, {%rd2, %rd3};489; CHECK-NEXT: mov.b128 swap, {%rd4, %rd5};490; CHECK-NEXT: atom.acquire.sys.cas.b128 dst, [%rd1], cmp, swap;491; CHECK-NEXT: mov.b128 {%rd6, %rd7}, dst;492; CHECK-NEXT: }493; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};494; CHECK-NEXT: ret;495 %pairold = cmpxchg ptr %addr, i128 %cmp, i128 %new acquire monotonic496 ret i128 %new497}498 499define i128 @test_cmpxchg_acquire_acquire(ptr %addr, i128 %cmp, i128 %new) {500; CHECK-LABEL: test_cmpxchg_acquire_acquire(501; CHECK: {502; CHECK-NEXT: .reg .b64 %rd<8>;503; CHECK-EMPTY:504; CHECK-NEXT: // %bb.0:505; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_acquire_acquire_param_0];506; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_acquire_acquire_param_1];507; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_acquire_acquire_param_2];508; CHECK-NEXT: {509; CHECK-NEXT: .reg .b128 cmp, swap, dst;510; CHECK-NEXT: mov.b128 cmp, {%rd2, %rd3};511; CHECK-NEXT: mov.b128 swap, {%rd4, %rd5};512; CHECK-NEXT: atom.acquire.sys.cas.b128 dst, [%rd1], cmp, swap;513; CHECK-NEXT: mov.b128 {%rd6, %rd7}, dst;514; CHECK-NEXT: }515; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};516; CHECK-NEXT: ret;517 %pairold = cmpxchg ptr %addr, i128 %cmp, i128 %new acquire acquire518 ret i128 %new519}520 521define i128 @test_cmpxchg_acquire_seq_cst(ptr %addr, i128 %cmp, i128 %new) {522; CHECK-LABEL: test_cmpxchg_acquire_seq_cst(523; CHECK: {524; CHECK-NEXT: .reg .b64 %rd<8>;525; CHECK-EMPTY:526; CHECK-NEXT: // %bb.0:527; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_acquire_seq_cst_param_0];528; CHECK-NEXT: fence.sc.sys;529; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_acquire_seq_cst_param_1];530; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_acquire_seq_cst_param_2];531; CHECK-NEXT: {532; CHECK-NEXT: .reg .b128 cmp, swap, dst;533; CHECK-NEXT: mov.b128 cmp, {%rd2, %rd3};534; CHECK-NEXT: mov.b128 swap, {%rd4, %rd5};535; CHECK-NEXT: atom.acquire.sys.cas.b128 dst, [%rd1], cmp, swap;536; CHECK-NEXT: mov.b128 {%rd6, %rd7}, dst;537; CHECK-NEXT: }538; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};539; CHECK-NEXT: ret;540 %pairold = cmpxchg ptr %addr, i128 %cmp, i128 %new acquire seq_cst541 ret i128 %new542}543 544define i128 @test_cmpxchg_release_monotonic(ptr %addr, i128 %cmp, i128 %new) {545; CHECK-LABEL: test_cmpxchg_release_monotonic(546; CHECK: {547; CHECK-NEXT: .reg .b64 %rd<8>;548; CHECK-EMPTY:549; CHECK-NEXT: // %bb.0:550; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_release_monotonic_param_0];551; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_release_monotonic_param_1];552; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_release_monotonic_param_2];553; CHECK-NEXT: {554; CHECK-NEXT: .reg .b128 cmp, swap, dst;555; CHECK-NEXT: mov.b128 cmp, {%rd2, %rd3};556; CHECK-NEXT: mov.b128 swap, {%rd4, %rd5};557; CHECK-NEXT: atom.release.sys.cas.b128 dst, [%rd1], cmp, swap;558; CHECK-NEXT: mov.b128 {%rd6, %rd7}, dst;559; CHECK-NEXT: }560; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};561; CHECK-NEXT: ret;562 %pairold = cmpxchg ptr %addr, i128 %cmp, i128 %new release monotonic563 ret i128 %new564}565 566define i128 @test_cmpxchg_release_acquire(ptr %addr, i128 %cmp, i128 %new) {567; CHECK-LABEL: test_cmpxchg_release_acquire(568; CHECK: {569; CHECK-NEXT: .reg .b64 %rd<8>;570; CHECK-EMPTY:571; CHECK-NEXT: // %bb.0:572; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_release_acquire_param_0];573; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_release_acquire_param_1];574; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_release_acquire_param_2];575; CHECK-NEXT: {576; CHECK-NEXT: .reg .b128 cmp, swap, dst;577; CHECK-NEXT: mov.b128 cmp, {%rd2, %rd3};578; CHECK-NEXT: mov.b128 swap, {%rd4, %rd5};579; CHECK-NEXT: atom.acq_rel.sys.cas.b128 dst, [%rd1], cmp, swap;580; CHECK-NEXT: mov.b128 {%rd6, %rd7}, dst;581; CHECK-NEXT: }582; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};583; CHECK-NEXT: ret;584 %pairold = cmpxchg ptr %addr, i128 %cmp, i128 %new release acquire585 ret i128 %new586}587 588define i128 @test_cmpxchg_release_seq_cst(ptr %addr, i128 %cmp, i128 %new) {589; CHECK-LABEL: test_cmpxchg_release_seq_cst(590; CHECK: {591; CHECK-NEXT: .reg .b64 %rd<8>;592; CHECK-EMPTY:593; CHECK-NEXT: // %bb.0:594; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_release_seq_cst_param_0];595; CHECK-NEXT: fence.sc.sys;596; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_release_seq_cst_param_1];597; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_release_seq_cst_param_2];598; CHECK-NEXT: {599; CHECK-NEXT: .reg .b128 cmp, swap, dst;600; CHECK-NEXT: mov.b128 cmp, {%rd2, %rd3};601; CHECK-NEXT: mov.b128 swap, {%rd4, %rd5};602; CHECK-NEXT: atom.acquire.sys.cas.b128 dst, [%rd1], cmp, swap;603; CHECK-NEXT: mov.b128 {%rd6, %rd7}, dst;604; CHECK-NEXT: }605; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};606; CHECK-NEXT: ret;607 %pairold = cmpxchg ptr %addr, i128 %cmp, i128 %new release seq_cst608 ret i128 %new609}610 611define i128 @test_cmpxchg_acq_rel_monotonic(ptr %addr, i128 %cmp, i128 %new) {612; CHECK-LABEL: test_cmpxchg_acq_rel_monotonic(613; CHECK: {614; CHECK-NEXT: .reg .b64 %rd<8>;615; CHECK-EMPTY:616; CHECK-NEXT: // %bb.0:617; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_acq_rel_monotonic_param_0];618; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_acq_rel_monotonic_param_1];619; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_acq_rel_monotonic_param_2];620; CHECK-NEXT: {621; CHECK-NEXT: .reg .b128 cmp, swap, dst;622; CHECK-NEXT: mov.b128 cmp, {%rd2, %rd3};623; CHECK-NEXT: mov.b128 swap, {%rd4, %rd5};624; CHECK-NEXT: atom.acq_rel.sys.cas.b128 dst, [%rd1], cmp, swap;625; CHECK-NEXT: mov.b128 {%rd6, %rd7}, dst;626; CHECK-NEXT: }627; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};628; CHECK-NEXT: ret;629 %pairold = cmpxchg ptr %addr, i128 %cmp, i128 %new acq_rel monotonic630 ret i128 %new631}632 633define i128 @test_cmpxchg_acq_rel_acquire(ptr %addr, i128 %cmp, i128 %new) {634; CHECK-LABEL: test_cmpxchg_acq_rel_acquire(635; CHECK: {636; CHECK-NEXT: .reg .b64 %rd<8>;637; CHECK-EMPTY:638; CHECK-NEXT: // %bb.0:639; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_acq_rel_acquire_param_0];640; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_acq_rel_acquire_param_1];641; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_acq_rel_acquire_param_2];642; CHECK-NEXT: {643; CHECK-NEXT: .reg .b128 cmp, swap, dst;644; CHECK-NEXT: mov.b128 cmp, {%rd2, %rd3};645; CHECK-NEXT: mov.b128 swap, {%rd4, %rd5};646; CHECK-NEXT: atom.acq_rel.sys.cas.b128 dst, [%rd1], cmp, swap;647; CHECK-NEXT: mov.b128 {%rd6, %rd7}, dst;648; CHECK-NEXT: }649; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};650; CHECK-NEXT: ret;651 %pairold = cmpxchg ptr %addr, i128 %cmp, i128 %new acq_rel acquire652 ret i128 %new653}654 655define i128 @test_cmpxchg_acq_rel_seq_cst(ptr %addr, i128 %cmp, i128 %new) {656; CHECK-LABEL: test_cmpxchg_acq_rel_seq_cst(657; CHECK: {658; CHECK-NEXT: .reg .b64 %rd<8>;659; CHECK-EMPTY:660; CHECK-NEXT: // %bb.0:661; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_acq_rel_seq_cst_param_0];662; CHECK-NEXT: fence.sc.sys;663; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_acq_rel_seq_cst_param_1];664; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_acq_rel_seq_cst_param_2];665; CHECK-NEXT: {666; CHECK-NEXT: .reg .b128 cmp, swap, dst;667; CHECK-NEXT: mov.b128 cmp, {%rd2, %rd3};668; CHECK-NEXT: mov.b128 swap, {%rd4, %rd5};669; CHECK-NEXT: atom.acquire.sys.cas.b128 dst, [%rd1], cmp, swap;670; CHECK-NEXT: mov.b128 {%rd6, %rd7}, dst;671; CHECK-NEXT: }672; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};673; CHECK-NEXT: ret;674 %pairold = cmpxchg ptr %addr, i128 %cmp, i128 %new acq_rel seq_cst675 ret i128 %new676}677 678define i128 @test_cmpxchg_seq_cst_monotonic(ptr %addr, i128 %cmp, i128 %new) {679; CHECK-LABEL: test_cmpxchg_seq_cst_monotonic(680; CHECK: {681; CHECK-NEXT: .reg .b64 %rd<8>;682; CHECK-EMPTY:683; CHECK-NEXT: // %bb.0:684; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_seq_cst_monotonic_param_0];685; CHECK-NEXT: fence.sc.sys;686; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_seq_cst_monotonic_param_1];687; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_seq_cst_monotonic_param_2];688; CHECK-NEXT: {689; CHECK-NEXT: .reg .b128 cmp, swap, dst;690; CHECK-NEXT: mov.b128 cmp, {%rd2, %rd3};691; CHECK-NEXT: mov.b128 swap, {%rd4, %rd5};692; CHECK-NEXT: atom.acquire.sys.cas.b128 dst, [%rd1], cmp, swap;693; CHECK-NEXT: mov.b128 {%rd6, %rd7}, dst;694; CHECK-NEXT: }695; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};696; CHECK-NEXT: ret;697 %pairold = cmpxchg ptr %addr, i128 %cmp, i128 %new seq_cst monotonic698 ret i128 %new699}700 701define i128 @test_cmpxchg_seq_cst_acquire(ptr %addr, i128 %cmp, i128 %new) {702; CHECK-LABEL: test_cmpxchg_seq_cst_acquire(703; CHECK: {704; CHECK-NEXT: .reg .b64 %rd<8>;705; CHECK-EMPTY:706; CHECK-NEXT: // %bb.0:707; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_seq_cst_acquire_param_0];708; CHECK-NEXT: fence.sc.sys;709; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_seq_cst_acquire_param_1];710; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_seq_cst_acquire_param_2];711; CHECK-NEXT: {712; CHECK-NEXT: .reg .b128 cmp, swap, dst;713; CHECK-NEXT: mov.b128 cmp, {%rd2, %rd3};714; CHECK-NEXT: mov.b128 swap, {%rd4, %rd5};715; CHECK-NEXT: atom.acquire.sys.cas.b128 dst, [%rd1], cmp, swap;716; CHECK-NEXT: mov.b128 {%rd6, %rd7}, dst;717; CHECK-NEXT: }718; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};719; CHECK-NEXT: ret;720 %pairold = cmpxchg ptr %addr, i128 %cmp, i128 %new seq_cst acquire721 ret i128 %new722}723 724define i128 @test_cmpxchg_seq_cst_seq_cst(ptr %addr, i128 %cmp, i128 %new) {725; CHECK-LABEL: test_cmpxchg_seq_cst_seq_cst(726; CHECK: {727; CHECK-NEXT: .reg .b64 %rd<8>;728; CHECK-EMPTY:729; CHECK-NEXT: // %bb.0:730; CHECK-NEXT: ld.param.b64 %rd1, [test_cmpxchg_seq_cst_seq_cst_param_0];731; CHECK-NEXT: fence.sc.sys;732; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_cmpxchg_seq_cst_seq_cst_param_1];733; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_cmpxchg_seq_cst_seq_cst_param_2];734; CHECK-NEXT: {735; CHECK-NEXT: .reg .b128 cmp, swap, dst;736; CHECK-NEXT: mov.b128 cmp, {%rd2, %rd3};737; CHECK-NEXT: mov.b128 swap, {%rd4, %rd5};738; CHECK-NEXT: atom.acquire.sys.cas.b128 dst, [%rd1], cmp, swap;739; CHECK-NEXT: mov.b128 {%rd6, %rd7}, dst;740; CHECK-NEXT: }741; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd5};742; CHECK-NEXT: ret;743 %pairold = cmpxchg ptr %addr, i128 %cmp, i128 %new seq_cst seq_cst744 ret i128 %new745}746 747define i128 @test_atomicrmw_and(ptr %ptr, i128 %val) {748; CHECK-LABEL: test_atomicrmw_and(749; CHECK: {750; CHECK-NEXT: .reg .pred %p<2>;751; CHECK-NEXT: .reg .b64 %rd<13>;752; CHECK-EMPTY:753; CHECK-NEXT: // %bb.0:754; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_atomicrmw_and_param_1];755; CHECK-NEXT: ld.param.b64 %rd3, [test_atomicrmw_and_param_0];756; CHECK-NEXT: ld.v2.b64 {%rd11, %rd12}, [%rd3];757; CHECK-NEXT: $L__BB34_1: // %atomicrmw.start758; CHECK-NEXT: // =>This Inner Loop Header: Depth=1759; CHECK-NEXT: and.b64 %rd6, %rd11, %rd4;760; CHECK-NEXT: and.b64 %rd7, %rd12, %rd5;761; CHECK-NEXT: {762; CHECK-NEXT: .reg .b128 cmp, swap, dst;763; CHECK-NEXT: mov.b128 cmp, {%rd11, %rd12};764; CHECK-NEXT: mov.b128 swap, {%rd6, %rd7};765; CHECK-NEXT: atom.relaxed.sys.cas.b128 dst, [%rd3], cmp, swap;766; CHECK-NEXT: mov.b128 {%rd1, %rd2}, dst;767; CHECK-NEXT: }768; CHECK-NEXT: xor.b64 %rd8, %rd2, %rd12;769; CHECK-NEXT: xor.b64 %rd9, %rd1, %rd11;770; CHECK-NEXT: or.b64 %rd10, %rd9, %rd8;771; CHECK-NEXT: setp.ne.b64 %p1, %rd10, 0;772; CHECK-NEXT: mov.b64 %rd11, %rd1;773; CHECK-NEXT: mov.b64 %rd12, %rd2;774; CHECK-NEXT: @%p1 bra $L__BB34_1;775; CHECK-NEXT: // %bb.2: // %atomicrmw.end776; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd1, %rd2};777; CHECK-NEXT: ret;778 %ret = atomicrmw and ptr %ptr, i128 %val monotonic779 ret i128 %ret780}781 782define i128 @test_atomicrmw_or(ptr %ptr, i128 %val) {783; CHECK-LABEL: test_atomicrmw_or(784; CHECK: {785; CHECK-NEXT: .reg .pred %p<2>;786; CHECK-NEXT: .reg .b64 %rd<13>;787; CHECK-EMPTY:788; CHECK-NEXT: // %bb.0:789; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_atomicrmw_or_param_1];790; CHECK-NEXT: ld.param.b64 %rd3, [test_atomicrmw_or_param_0];791; CHECK-NEXT: ld.v2.b64 {%rd11, %rd12}, [%rd3];792; CHECK-NEXT: $L__BB35_1: // %atomicrmw.start793; CHECK-NEXT: // =>This Inner Loop Header: Depth=1794; CHECK-NEXT: or.b64 %rd6, %rd11, %rd4;795; CHECK-NEXT: or.b64 %rd7, %rd12, %rd5;796; CHECK-NEXT: {797; CHECK-NEXT: .reg .b128 cmp, swap, dst;798; CHECK-NEXT: mov.b128 cmp, {%rd11, %rd12};799; CHECK-NEXT: mov.b128 swap, {%rd6, %rd7};800; CHECK-NEXT: atom.relaxed.sys.cas.b128 dst, [%rd3], cmp, swap;801; CHECK-NEXT: mov.b128 {%rd1, %rd2}, dst;802; CHECK-NEXT: }803; CHECK-NEXT: xor.b64 %rd8, %rd2, %rd12;804; CHECK-NEXT: xor.b64 %rd9, %rd1, %rd11;805; CHECK-NEXT: or.b64 %rd10, %rd9, %rd8;806; CHECK-NEXT: setp.ne.b64 %p1, %rd10, 0;807; CHECK-NEXT: mov.b64 %rd11, %rd1;808; CHECK-NEXT: mov.b64 %rd12, %rd2;809; CHECK-NEXT: @%p1 bra $L__BB35_1;810; CHECK-NEXT: // %bb.2: // %atomicrmw.end811; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd1, %rd2};812; CHECK-NEXT: ret;813 %ret = atomicrmw or ptr %ptr, i128 %val monotonic814 ret i128 %ret815}816 817define i128 @test_atomicrmw_xor(ptr %ptr, i128 %val) {818; CHECK-LABEL: test_atomicrmw_xor(819; CHECK: {820; CHECK-NEXT: .reg .pred %p<2>;821; CHECK-NEXT: .reg .b64 %rd<13>;822; CHECK-EMPTY:823; CHECK-NEXT: // %bb.0:824; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_atomicrmw_xor_param_1];825; CHECK-NEXT: ld.param.b64 %rd3, [test_atomicrmw_xor_param_0];826; CHECK-NEXT: ld.v2.b64 {%rd11, %rd12}, [%rd3];827; CHECK-NEXT: $L__BB36_1: // %atomicrmw.start828; CHECK-NEXT: // =>This Inner Loop Header: Depth=1829; CHECK-NEXT: xor.b64 %rd6, %rd11, %rd4;830; CHECK-NEXT: xor.b64 %rd7, %rd12, %rd5;831; CHECK-NEXT: {832; CHECK-NEXT: .reg .b128 cmp, swap, dst;833; CHECK-NEXT: mov.b128 cmp, {%rd11, %rd12};834; CHECK-NEXT: mov.b128 swap, {%rd6, %rd7};835; CHECK-NEXT: atom.relaxed.sys.cas.b128 dst, [%rd3], cmp, swap;836; CHECK-NEXT: mov.b128 {%rd1, %rd2}, dst;837; CHECK-NEXT: }838; CHECK-NEXT: xor.b64 %rd8, %rd2, %rd12;839; CHECK-NEXT: xor.b64 %rd9, %rd1, %rd11;840; CHECK-NEXT: or.b64 %rd10, %rd9, %rd8;841; CHECK-NEXT: setp.ne.b64 %p1, %rd10, 0;842; CHECK-NEXT: mov.b64 %rd11, %rd1;843; CHECK-NEXT: mov.b64 %rd12, %rd2;844; CHECK-NEXT: @%p1 bra $L__BB36_1;845; CHECK-NEXT: // %bb.2: // %atomicrmw.end846; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd1, %rd2};847; CHECK-NEXT: ret;848 %ret = atomicrmw xor ptr %ptr, i128 %val monotonic849 ret i128 %ret850}851 852define i128 @test_atomicrmw_min(ptr %ptr, i128 %val) {853; CHECK-LABEL: test_atomicrmw_min(854; CHECK: {855; CHECK-NEXT: .reg .pred %p<7>;856; CHECK-NEXT: .reg .b64 %rd<13>;857; CHECK-EMPTY:858; CHECK-NEXT: // %bb.0:859; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_atomicrmw_min_param_1];860; CHECK-NEXT: ld.param.b64 %rd3, [test_atomicrmw_min_param_0];861; CHECK-NEXT: ld.v2.b64 {%rd11, %rd12}, [%rd3];862; CHECK-NEXT: $L__BB37_1: // %atomicrmw.start863; CHECK-NEXT: // =>This Inner Loop Header: Depth=1864; CHECK-NEXT: setp.lt.u64 %p1, %rd11, %rd4;865; CHECK-NEXT: setp.eq.b64 %p2, %rd12, %rd5;866; CHECK-NEXT: and.pred %p3, %p2, %p1;867; CHECK-NEXT: setp.lt.s64 %p4, %rd12, %rd5;868; CHECK-NEXT: or.pred %p5, %p3, %p4;869; CHECK-NEXT: selp.b64 %rd6, %rd12, %rd5, %p5;870; CHECK-NEXT: selp.b64 %rd7, %rd11, %rd4, %p5;871; CHECK-NEXT: {872; CHECK-NEXT: .reg .b128 cmp, swap, dst;873; CHECK-NEXT: mov.b128 cmp, {%rd11, %rd12};874; CHECK-NEXT: mov.b128 swap, {%rd7, %rd6};875; CHECK-NEXT: atom.relaxed.sys.cas.b128 dst, [%rd3], cmp, swap;876; CHECK-NEXT: mov.b128 {%rd1, %rd2}, dst;877; CHECK-NEXT: }878; CHECK-NEXT: xor.b64 %rd8, %rd2, %rd12;879; CHECK-NEXT: xor.b64 %rd9, %rd1, %rd11;880; CHECK-NEXT: or.b64 %rd10, %rd9, %rd8;881; CHECK-NEXT: setp.ne.b64 %p6, %rd10, 0;882; CHECK-NEXT: mov.b64 %rd11, %rd1;883; CHECK-NEXT: mov.b64 %rd12, %rd2;884; CHECK-NEXT: @%p6 bra $L__BB37_1;885; CHECK-NEXT: // %bb.2: // %atomicrmw.end886; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd1, %rd2};887; CHECK-NEXT: ret;888 %ret = atomicrmw min ptr %ptr, i128 %val monotonic889 ret i128 %ret890}891 892define i128 @test_atomicrmw_max(ptr %ptr, i128 %val) {893; CHECK-LABEL: test_atomicrmw_max(894; CHECK: {895; CHECK-NEXT: .reg .pred %p<7>;896; CHECK-NEXT: .reg .b64 %rd<13>;897; CHECK-EMPTY:898; CHECK-NEXT: // %bb.0:899; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_atomicrmw_max_param_1];900; CHECK-NEXT: ld.param.b64 %rd3, [test_atomicrmw_max_param_0];901; CHECK-NEXT: ld.v2.b64 {%rd11, %rd12}, [%rd3];902; CHECK-NEXT: $L__BB38_1: // %atomicrmw.start903; CHECK-NEXT: // =>This Inner Loop Header: Depth=1904; CHECK-NEXT: setp.gt.u64 %p1, %rd11, %rd4;905; CHECK-NEXT: setp.eq.b64 %p2, %rd12, %rd5;906; CHECK-NEXT: and.pred %p3, %p2, %p1;907; CHECK-NEXT: setp.gt.s64 %p4, %rd12, %rd5;908; CHECK-NEXT: or.pred %p5, %p3, %p4;909; CHECK-NEXT: selp.b64 %rd6, %rd12, %rd5, %p5;910; CHECK-NEXT: selp.b64 %rd7, %rd11, %rd4, %p5;911; CHECK-NEXT: {912; CHECK-NEXT: .reg .b128 cmp, swap, dst;913; CHECK-NEXT: mov.b128 cmp, {%rd11, %rd12};914; CHECK-NEXT: mov.b128 swap, {%rd7, %rd6};915; CHECK-NEXT: atom.relaxed.sys.cas.b128 dst, [%rd3], cmp, swap;916; CHECK-NEXT: mov.b128 {%rd1, %rd2}, dst;917; CHECK-NEXT: }918; CHECK-NEXT: xor.b64 %rd8, %rd2, %rd12;919; CHECK-NEXT: xor.b64 %rd9, %rd1, %rd11;920; CHECK-NEXT: or.b64 %rd10, %rd9, %rd8;921; CHECK-NEXT: setp.ne.b64 %p6, %rd10, 0;922; CHECK-NEXT: mov.b64 %rd11, %rd1;923; CHECK-NEXT: mov.b64 %rd12, %rd2;924; CHECK-NEXT: @%p6 bra $L__BB38_1;925; CHECK-NEXT: // %bb.2: // %atomicrmw.end926; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd1, %rd2};927; CHECK-NEXT: ret;928 %ret = atomicrmw max ptr %ptr, i128 %val monotonic929 ret i128 %ret930}931 932define i128 @test_atomicrmw_umin(ptr %ptr, i128 %val) {933; CHECK-LABEL: test_atomicrmw_umin(934; CHECK: {935; CHECK-NEXT: .reg .pred %p<7>;936; CHECK-NEXT: .reg .b64 %rd<13>;937; CHECK-EMPTY:938; CHECK-NEXT: // %bb.0:939; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_atomicrmw_umin_param_1];940; CHECK-NEXT: ld.param.b64 %rd3, [test_atomicrmw_umin_param_0];941; CHECK-NEXT: ld.v2.b64 {%rd11, %rd12}, [%rd3];942; CHECK-NEXT: $L__BB39_1: // %atomicrmw.start943; CHECK-NEXT: // =>This Inner Loop Header: Depth=1944; CHECK-NEXT: setp.lt.u64 %p1, %rd11, %rd4;945; CHECK-NEXT: setp.eq.b64 %p2, %rd12, %rd5;946; CHECK-NEXT: and.pred %p3, %p2, %p1;947; CHECK-NEXT: setp.lt.u64 %p4, %rd12, %rd5;948; CHECK-NEXT: or.pred %p5, %p3, %p4;949; CHECK-NEXT: selp.b64 %rd6, %rd12, %rd5, %p5;950; CHECK-NEXT: selp.b64 %rd7, %rd11, %rd4, %p5;951; CHECK-NEXT: {952; CHECK-NEXT: .reg .b128 cmp, swap, dst;953; CHECK-NEXT: mov.b128 cmp, {%rd11, %rd12};954; CHECK-NEXT: mov.b128 swap, {%rd7, %rd6};955; CHECK-NEXT: atom.relaxed.sys.cas.b128 dst, [%rd3], cmp, swap;956; CHECK-NEXT: mov.b128 {%rd1, %rd2}, dst;957; CHECK-NEXT: }958; CHECK-NEXT: xor.b64 %rd8, %rd2, %rd12;959; CHECK-NEXT: xor.b64 %rd9, %rd1, %rd11;960; CHECK-NEXT: or.b64 %rd10, %rd9, %rd8;961; CHECK-NEXT: setp.ne.b64 %p6, %rd10, 0;962; CHECK-NEXT: mov.b64 %rd11, %rd1;963; CHECK-NEXT: mov.b64 %rd12, %rd2;964; CHECK-NEXT: @%p6 bra $L__BB39_1;965; CHECK-NEXT: // %bb.2: // %atomicrmw.end966; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd1, %rd2};967; CHECK-NEXT: ret;968 %ret = atomicrmw umin ptr %ptr, i128 %val monotonic969 ret i128 %ret970}971 972define i128 @test_atomicrmw_umax(ptr %ptr, i128 %val) {973; CHECK-LABEL: test_atomicrmw_umax(974; CHECK: {975; CHECK-NEXT: .reg .pred %p<7>;976; CHECK-NEXT: .reg .b64 %rd<13>;977; CHECK-EMPTY:978; CHECK-NEXT: // %bb.0:979; CHECK-NEXT: ld.param.v2.b64 {%rd4, %rd5}, [test_atomicrmw_umax_param_1];980; CHECK-NEXT: ld.param.b64 %rd3, [test_atomicrmw_umax_param_0];981; CHECK-NEXT: ld.v2.b64 {%rd11, %rd12}, [%rd3];982; CHECK-NEXT: $L__BB40_1: // %atomicrmw.start983; CHECK-NEXT: // =>This Inner Loop Header: Depth=1984; CHECK-NEXT: setp.gt.u64 %p1, %rd11, %rd4;985; CHECK-NEXT: setp.eq.b64 %p2, %rd12, %rd5;986; CHECK-NEXT: and.pred %p3, %p2, %p1;987; CHECK-NEXT: setp.gt.u64 %p4, %rd12, %rd5;988; CHECK-NEXT: or.pred %p5, %p3, %p4;989; CHECK-NEXT: selp.b64 %rd6, %rd12, %rd5, %p5;990; CHECK-NEXT: selp.b64 %rd7, %rd11, %rd4, %p5;991; CHECK-NEXT: {992; CHECK-NEXT: .reg .b128 cmp, swap, dst;993; CHECK-NEXT: mov.b128 cmp, {%rd11, %rd12};994; CHECK-NEXT: mov.b128 swap, {%rd7, %rd6};995; CHECK-NEXT: atom.relaxed.sys.cas.b128 dst, [%rd3], cmp, swap;996; CHECK-NEXT: mov.b128 {%rd1, %rd2}, dst;997; CHECK-NEXT: }998; CHECK-NEXT: xor.b64 %rd8, %rd2, %rd12;999; CHECK-NEXT: xor.b64 %rd9, %rd1, %rd11;1000; CHECK-NEXT: or.b64 %rd10, %rd9, %rd8;1001; CHECK-NEXT: setp.ne.b64 %p6, %rd10, 0;1002; CHECK-NEXT: mov.b64 %rd11, %rd1;1003; CHECK-NEXT: mov.b64 %rd12, %rd2;1004; CHECK-NEXT: @%p6 bra $L__BB40_1;1005; CHECK-NEXT: // %bb.2: // %atomicrmw.end1006; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd1, %rd2};1007; CHECK-NEXT: ret;1008 %ret = atomicrmw umax ptr %ptr, i128 %val monotonic1009 ret i128 %ret1010}1011 1012 1013@si128 = internal addrspace(3) global i128 0, align 161014 1015define void @test_atomicrmw_xchg_const() {1016; CHECK-LABEL: test_atomicrmw_xchg_const(1017; CHECK: {1018; CHECK-NEXT: .reg .b64 %rd<5>;1019; CHECK-NEXT: // demoted variable1020; CHECK-NEXT: .shared .align 16 .b8 si128[16];1021; CHECK-NEXT: // %bb.0:1022; CHECK-NEXT: mov.b64 %rd1, 0;1023; CHECK-NEXT: mov.b64 %rd2, 23;1024; CHECK-NEXT: {1025; CHECK-NEXT: .reg .b128 amt, dst;1026; CHECK-NEXT: mov.b128 amt, {%rd2, %rd1};1027; CHECK-NEXT: atom.relaxed.sys.shared.exch.b128 dst, [si128], amt;1028; CHECK-NEXT: mov.b128 {%rd3, %rd4}, dst;1029; CHECK-NEXT: }1030; CHECK-NEXT: ret;1031 %res = atomicrmw xchg ptr addrspace(3) @si128, i128 23 monotonic1032 ret void1033}1034