321 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -o - -mcpu=sm_90 -mattr=+ptx78 | FileCheck %s3; RUN: %if ptxas-sm_90 && ptxas-isa-7.8 %{ llc < %s -mcpu=sm_90 -mattr=+ptx78| %ptxas-verify -arch=sm_90 %}4 5target triple = "nvptx64-nvidia-cuda"6 7declare ptr addrspace(7) @llvm.nvvm.mapa.shared.cluster(ptr addrspace(3), i32)8declare i1 @llvm.nvvm.isspacep.shared.cluster(ptr)9declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()10declare ptr @llvm.nvvm.mapa(ptr, i32)11 12; Common setup for distributed shared memory cluster addressing13define i32 @test_distributed_shared_cluster_common(ptr %ptr, ptr addrspace(3) %smem_ptr) local_unnamed_addr {14; CHECK-LABEL: test_distributed_shared_cluster_common(15; CHECK: {16; CHECK-NEXT: .reg .pred %p<3>;17; CHECK-NEXT: .reg .b32 %r<8>;18; CHECK-NEXT: .reg .b64 %rd<5>;19; CHECK-EMPTY:20; CHECK-NEXT: // %bb.0: // %entry21; CHECK-NEXT: ld.param.b64 %rd1, [test_distributed_shared_cluster_common_param_0];22; CHECK-NEXT: ld.param.b64 %rd2, [test_distributed_shared_cluster_common_param_1];23; CHECK-NEXT: mov.u32 %r1, %ctaid.x;24; CHECK-NEXT: xor.b32 %r2, %r1, 1;25; CHECK-NEXT: isspacep.shared::cluster %p1, %rd1;26; CHECK-NEXT: mapa.u64 %rd3, %rd1, %r2;27; CHECK-NEXT: isspacep.shared::cluster %p2, %rd3;28; CHECK-NEXT: mapa.shared::cluster.u64 %rd4, %rd2, %r2;29; CHECK-NEXT: ld.shared::cluster.b32 %r3, [%rd4];30; CHECK-NEXT: add.s32 %r4, %r3, 42;31; CHECK-NEXT: st.shared::cluster.b32 [%rd4], %r4;32; CHECK-NEXT: selp.b32 %r5, 1, 0, %p1;33; CHECK-NEXT: selp.b32 %r6, 1, 0, %p2;34; CHECK-NEXT: add.s32 %r7, %r5, %r6;35; CHECK-NEXT: st.param.b32 [func_retval0], %r7;36; CHECK-NEXT: ret;37entry:38 %0 = tail call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()39 %1 = xor i32 %0, 140 %2 = tail call i1 @llvm.nvvm.isspacep.shared.cluster(ptr %ptr)41 %3 = tail call ptr @llvm.nvvm.mapa(ptr %ptr, i32 %1)42 %4 = tail call i1 @llvm.nvvm.isspacep.shared.cluster(ptr %3)43 %dsmem_ptr = call ptr addrspace(7) @llvm.nvvm.mapa.shared.cluster(ptr addrspace(3) %smem_ptr, i32 %1)44 45 ; Add load and store to the distributed shared memory cluster46 %loaded_val = load i32, ptr addrspace(7) %dsmem_ptr47 %updated_val = add i32 %loaded_val, 4248 store i32 %updated_val, ptr addrspace(7) %dsmem_ptr49 50 ; Return value preserves the isspacep test results plus the value operation51 %5 = zext i1 %2 to i3252 %6 = zext i1 %4 to i3253 %ret = add i32 %5, %654 ret i32 %ret55}56 57; Floating point atomic operations tests58define void @test_distributed_shared_cluster_float_atomic(ptr addrspace(7) %dsmem_ptr) local_unnamed_addr {59; CHECK-LABEL: test_distributed_shared_cluster_float_atomic(60; CHECK: {61; CHECK-NEXT: .reg .b16 %rs<5>;62; CHECK-NEXT: .reg .b32 %r<2>;63; CHECK-NEXT: .reg .b64 %rd<3>;64; CHECK-EMPTY:65; CHECK-NEXT: // %bb.0: // %entry66; CHECK-NEXT: ld.param.b64 %rd1, [test_distributed_shared_cluster_float_atomic_param_0];67; CHECK-NEXT: mov.b16 %rs1, 0x3C00;68; CHECK-NEXT: atom.shared::cluster.add.noftz.f16 %rs2, [%rd1], %rs1;69; CHECK-NEXT: mov.b16 %rs3, 0x3F80;70; CHECK-NEXT: atom.shared::cluster.add.noftz.bf16 %rs4, [%rd1], %rs3;71; CHECK-NEXT: atom.shared::cluster.add.f32 %r1, [%rd1], 0f3F800000;72; CHECK-NEXT: atom.shared::cluster.add.f64 %rd2, [%rd1], 0d3FF0000000000000;73; CHECK-NEXT: ret;74entry:75 ; Floating point atomic operations76 %0 = atomicrmw fadd ptr addrspace(7) %dsmem_ptr, half 1.000000e+00 seq_cst77 %1 = atomicrmw fadd ptr addrspace(7) %dsmem_ptr, bfloat 1.000000e+00 seq_cst78 %2 = atomicrmw fadd ptr addrspace(7) %dsmem_ptr, float 1.000000e+00 seq_cst79 %3 = atomicrmw fadd ptr addrspace(7) %dsmem_ptr, double 1.000000e+00 seq_cst80 81 ret void82}83 84; Integer atomic operations tests85define void @test_distributed_shared_cluster_int_atomic(ptr addrspace(7) %dsmem_ptr) local_unnamed_addr {86; CHECK-LABEL: test_distributed_shared_cluster_int_atomic(87; CHECK: {88; CHECK-NEXT: .reg .b32 %r<9>;89; CHECK-NEXT: .reg .b64 %rd<8>;90; CHECK-EMPTY:91; CHECK-NEXT: // %bb.0: // %entry92; CHECK-NEXT: ld.param.b64 %rd1, [test_distributed_shared_cluster_int_atomic_param_0];93; CHECK-NEXT: atom.shared::cluster.add.u32 %r1, [%rd1], 1;94; CHECK-NEXT: atom.shared::cluster.add.u64 %rd2, [%rd1], 1;95; CHECK-NEXT: atom.shared::cluster.exch.b32 %r2, [%rd1], 1;96; CHECK-NEXT: atom.shared::cluster.exch.b64 %rd3, [%rd1], 1;97; CHECK-NEXT: atom.shared::cluster.min.s32 %r3, [%rd1], 1;98; CHECK-NEXT: atom.shared::cluster.min.s64 %rd4, [%rd1], 1;99; CHECK-NEXT: atom.shared::cluster.min.u32 %r4, [%rd1], 1;100; CHECK-NEXT: atom.shared::cluster.min.u64 %rd5, [%rd1], 1;101; CHECK-NEXT: atom.shared::cluster.max.s32 %r5, [%rd1], 1;102; CHECK-NEXT: atom.shared::cluster.max.s64 %rd6, [%rd1], 1;103; CHECK-NEXT: atom.shared::cluster.max.u32 %r6, [%rd1], 1;104; CHECK-NEXT: atom.shared::cluster.max.u64 %rd7, [%rd1], 1;105; CHECK-NEXT: atom.shared::cluster.inc.u32 %r7, [%rd1], 1;106; CHECK-NEXT: atom.shared::cluster.dec.u32 %r8, [%rd1], 1;107; CHECK-NEXT: ret;108entry:109 ; Integer add operations110 %0 = atomicrmw add ptr addrspace(7) %dsmem_ptr, i32 1 monotonic111 %1 = atomicrmw add ptr addrspace(7) %dsmem_ptr, i64 1 monotonic112 113 ; Exchange operations114 %2 = atomicrmw xchg ptr addrspace(7) %dsmem_ptr, i32 1 monotonic115 %3 = atomicrmw xchg ptr addrspace(7) %dsmem_ptr, i64 1 monotonic116 117 ; Min operations (signed and unsigned)118 %4 = atomicrmw min ptr addrspace(7) %dsmem_ptr, i32 1 monotonic119 %5 = atomicrmw min ptr addrspace(7) %dsmem_ptr, i64 1 monotonic120 %6 = atomicrmw umin ptr addrspace(7) %dsmem_ptr, i32 1 monotonic121 %7 = atomicrmw umin ptr addrspace(7) %dsmem_ptr, i64 1 monotonic122 123 ; Max operations (signed and unsigned)124 %8 = atomicrmw max ptr addrspace(7) %dsmem_ptr, i32 1 monotonic125 %9 = atomicrmw max ptr addrspace(7) %dsmem_ptr, i64 1 monotonic126 %10 = atomicrmw umax ptr addrspace(7) %dsmem_ptr, i32 1 monotonic127 %11 = atomicrmw umax ptr addrspace(7) %dsmem_ptr, i64 1 monotonic128 129 ; Inc/Dec operations (32-bit only)130 %12 = atomicrmw uinc_wrap ptr addrspace(7) %dsmem_ptr, i32 1 monotonic131 %13 = atomicrmw udec_wrap ptr addrspace(7) %dsmem_ptr, i32 1 monotonic132 133 ret void134}135 136; Bitwise atomic operations tests137define void @test_distributed_shared_cluster_bitwise_atomic(ptr addrspace(7) %dsmem_ptr) local_unnamed_addr {138; CHECK-LABEL: test_distributed_shared_cluster_bitwise_atomic(139; CHECK: {140; CHECK-NEXT: .reg .b32 %r<4>;141; CHECK-NEXT: .reg .b64 %rd<5>;142; CHECK-EMPTY:143; CHECK-NEXT: // %bb.0: // %entry144; CHECK-NEXT: ld.param.b64 %rd1, [test_distributed_shared_cluster_bitwise_atomic_param_0];145; CHECK-NEXT: atom.shared::cluster.and.b32 %r1, [%rd1], 1;146; CHECK-NEXT: atom.shared::cluster.and.b64 %rd2, [%rd1], 1;147; CHECK-NEXT: atom.shared::cluster.or.b32 %r2, [%rd1], 1;148; CHECK-NEXT: atom.shared::cluster.or.b64 %rd3, [%rd1], 1;149; CHECK-NEXT: atom.shared::cluster.xor.b32 %r3, [%rd1], 1;150; CHECK-NEXT: atom.shared::cluster.xor.b64 %rd4, [%rd1], 1;151; CHECK-NEXT: ret;152entry:153 ; Bitwise operations154 %0 = atomicrmw and ptr addrspace(7) %dsmem_ptr, i32 1 monotonic155 %1 = atomicrmw and ptr addrspace(7) %dsmem_ptr, i64 1 monotonic156 %2 = atomicrmw or ptr addrspace(7) %dsmem_ptr, i32 1 monotonic157 %3 = atomicrmw or ptr addrspace(7) %dsmem_ptr, i64 1 monotonic158 %4 = atomicrmw xor ptr addrspace(7) %dsmem_ptr, i32 1 monotonic159 %5 = atomicrmw xor ptr addrspace(7) %dsmem_ptr, i64 1 monotonic160 161 ret void162}163 164; Compare-exchange operations tests165define void @test_distributed_shared_cluster_cmpxchg(ptr addrspace(7) %dsmem_ptr) local_unnamed_addr {166; CHECK-LABEL: test_distributed_shared_cluster_cmpxchg(167; CHECK: {168; CHECK-NEXT: .reg .pred %p<11>;169; CHECK-NEXT: .reg .b32 %r<43>;170; CHECK-NEXT: .reg .b64 %rd<12>;171; CHECK-EMPTY:172; CHECK-NEXT: // %bb.0: // %entry173; CHECK-NEXT: ld.param.b64 %rd2, [test_distributed_shared_cluster_cmpxchg_param_0];174; CHECK-NEXT: atom.relaxed.sys.shared::cluster.cas.b32 %r14, [%rd2], 1, 0;175; CHECK-NEXT: atom.acquire.sys.shared::cluster.cas.b32 %r15, [%rd2], 1, 0;176; CHECK-NEXT: atom.acquire.sys.shared::cluster.cas.b32 %r16, [%rd2], 1, 0;177; CHECK-NEXT: atom.release.sys.shared::cluster.cas.b32 %r17, [%rd2], 1, 0;178; CHECK-NEXT: atom.acq_rel.sys.shared::cluster.cas.b32 %r18, [%rd2], 1, 0;179; CHECK-NEXT: atom.acq_rel.sys.shared::cluster.cas.b32 %r19, [%rd2], 1, 0;180; CHECK-NEXT: fence.sc.sys;181; CHECK-NEXT: atom.acquire.sys.shared::cluster.cas.b32 %r20, [%rd2], 1, 0;182; CHECK-NEXT: fence.sc.sys;183; CHECK-NEXT: atom.acquire.sys.shared::cluster.cas.b32 %r21, [%rd2], 1, 0;184; CHECK-NEXT: fence.sc.sys;185; CHECK-NEXT: atom.acquire.sys.shared::cluster.cas.b32 %r22, [%rd2], 1, 0;186; CHECK-NEXT: atom.relaxed.sys.shared::cluster.cas.b64 %rd3, [%rd2], 1, 0;187; CHECK-NEXT: atom.acquire.sys.shared::cluster.cas.b64 %rd4, [%rd2], 1, 0;188; CHECK-NEXT: atom.acquire.sys.shared::cluster.cas.b64 %rd5, [%rd2], 1, 0;189; CHECK-NEXT: atom.release.sys.shared::cluster.cas.b64 %rd6, [%rd2], 1, 0;190; CHECK-NEXT: atom.acq_rel.sys.shared::cluster.cas.b64 %rd7, [%rd2], 1, 0;191; CHECK-NEXT: atom.acq_rel.sys.shared::cluster.cas.b64 %rd8, [%rd2], 1, 0;192; CHECK-NEXT: fence.sc.sys;193; CHECK-NEXT: atom.acquire.sys.shared::cluster.cas.b64 %rd9, [%rd2], 1, 0;194; CHECK-NEXT: fence.sc.sys;195; CHECK-NEXT: atom.acquire.sys.shared::cluster.cas.b64 %rd10, [%rd2], 1, 0;196; CHECK-NEXT: fence.sc.sys;197; CHECK-NEXT: atom.acquire.sys.shared::cluster.cas.b64 %rd11, [%rd2], 1, 0;198; CHECK-NEXT: and.b64 %rd1, %rd2, -4;199; CHECK-NEXT: cvt.u32.u64 %r23, %rd2;200; CHECK-NEXT: and.b32 %r24, %r23, 3;201; CHECK-NEXT: shl.b32 %r1, %r24, 3;202; CHECK-NEXT: mov.b32 %r25, 65535;203; CHECK-NEXT: shl.b32 %r26, %r25, %r1;204; CHECK-NEXT: not.b32 %r2, %r26;205; CHECK-NEXT: mov.b32 %r27, 1;206; CHECK-NEXT: shl.b32 %r3, %r27, %r1;207; CHECK-NEXT: ld.shared::cluster.b32 %r28, [%rd1];208; CHECK-NEXT: and.b32 %r38, %r28, %r2;209; CHECK-NEXT: $L__BB4_1: // %partword.cmpxchg.loop33210; CHECK-NEXT: // =>This Inner Loop Header: Depth=1211; CHECK-NEXT: or.b32 %r29, %r38, %r3;212; CHECK-NEXT: atom.relaxed.sys.shared::cluster.cas.b32 %r4, [%rd1], %r29, %r38;213; CHECK-NEXT: setp.eq.b32 %p1, %r4, %r29;214; CHECK-NEXT: @%p1 bra $L__BB4_3;215; CHECK-NEXT: // %bb.2: // %partword.cmpxchg.failure32216; CHECK-NEXT: // in Loop: Header=BB4_1 Depth=1217; CHECK-NEXT: and.b32 %r5, %r4, %r2;218; CHECK-NEXT: setp.ne.b32 %p2, %r38, %r5;219; CHECK-NEXT: mov.b32 %r38, %r5;220; CHECK-NEXT: @%p2 bra $L__BB4_1;221; CHECK-NEXT: $L__BB4_3: // %partword.cmpxchg.end31222; CHECK-NEXT: ld.shared::cluster.b32 %r30, [%rd1];223; CHECK-NEXT: and.b32 %r39, %r30, %r2;224; CHECK-NEXT: $L__BB4_4: // %partword.cmpxchg.loop23225; CHECK-NEXT: // =>This Inner Loop Header: Depth=1226; CHECK-NEXT: or.b32 %r31, %r39, %r3;227; CHECK-NEXT: atom.relaxed.sys.shared::cluster.cas.b32 %r6, [%rd1], %r31, %r39;228; CHECK-NEXT: setp.eq.b32 %p3, %r6, %r31;229; CHECK-NEXT: @%p3 bra $L__BB4_6;230; CHECK-NEXT: // %bb.5: // %partword.cmpxchg.failure22231; CHECK-NEXT: // in Loop: Header=BB4_4 Depth=1232; CHECK-NEXT: and.b32 %r7, %r6, %r2;233; CHECK-NEXT: setp.ne.b32 %p4, %r39, %r7;234; CHECK-NEXT: mov.b32 %r39, %r7;235; CHECK-NEXT: @%p4 bra $L__BB4_4;236; CHECK-NEXT: $L__BB4_6: // %partword.cmpxchg.end21237; CHECK-NEXT: fence.acq_rel.sys;238; CHECK-NEXT: fence.acq_rel.sys;239; CHECK-NEXT: ld.shared::cluster.b32 %r32, [%rd1];240; CHECK-NEXT: and.b32 %r40, %r32, %r2;241; CHECK-NEXT: $L__BB4_7: // %partword.cmpxchg.loop13242; CHECK-NEXT: // =>This Inner Loop Header: Depth=1243; CHECK-NEXT: or.b32 %r33, %r40, %r3;244; CHECK-NEXT: atom.relaxed.sys.shared::cluster.cas.b32 %r8, [%rd1], %r33, %r40;245; CHECK-NEXT: setp.eq.b32 %p5, %r8, %r33;246; CHECK-NEXT: @%p5 bra $L__BB4_9;247; CHECK-NEXT: // %bb.8: // %partword.cmpxchg.failure12248; CHECK-NEXT: // in Loop: Header=BB4_7 Depth=1249; CHECK-NEXT: and.b32 %r9, %r8, %r2;250; CHECK-NEXT: setp.ne.b32 %p6, %r40, %r9;251; CHECK-NEXT: mov.b32 %r40, %r9;252; CHECK-NEXT: @%p6 bra $L__BB4_7;253; CHECK-NEXT: $L__BB4_9: // %partword.cmpxchg.end11254; CHECK-NEXT: fence.acq_rel.sys;255; CHECK-NEXT: ld.shared::cluster.b32 %r34, [%rd1];256; CHECK-NEXT: and.b32 %r41, %r34, %r2;257; CHECK-NEXT: $L__BB4_10: // %partword.cmpxchg.loop3258; CHECK-NEXT: // =>This Inner Loop Header: Depth=1259; CHECK-NEXT: or.b32 %r35, %r41, %r3;260; CHECK-NEXT: atom.relaxed.sys.shared::cluster.cas.b32 %r10, [%rd1], %r35, %r41;261; CHECK-NEXT: setp.eq.b32 %p7, %r10, %r35;262; CHECK-NEXT: @%p7 bra $L__BB4_12;263; CHECK-NEXT: // %bb.11: // %partword.cmpxchg.failure2264; CHECK-NEXT: // in Loop: Header=BB4_10 Depth=1265; CHECK-NEXT: and.b32 %r11, %r10, %r2;266; CHECK-NEXT: setp.ne.b32 %p8, %r41, %r11;267; CHECK-NEXT: mov.b32 %r41, %r11;268; CHECK-NEXT: @%p8 bra $L__BB4_10;269; CHECK-NEXT: $L__BB4_12: // %partword.cmpxchg.end1270; CHECK-NEXT: fence.acq_rel.sys;271; CHECK-NEXT: fence.sc.sys;272; CHECK-NEXT: ld.shared::cluster.b32 %r36, [%rd1];273; CHECK-NEXT: and.b32 %r42, %r36, %r2;274; CHECK-NEXT: $L__BB4_13: // %partword.cmpxchg.loop275; CHECK-NEXT: // =>This Inner Loop Header: Depth=1276; CHECK-NEXT: or.b32 %r37, %r42, %r3;277; CHECK-NEXT: atom.relaxed.sys.shared::cluster.cas.b32 %r12, [%rd1], %r37, %r42;278; CHECK-NEXT: setp.eq.b32 %p9, %r12, %r37;279; CHECK-NEXT: @%p9 bra $L__BB4_15;280; CHECK-NEXT: // %bb.14: // %partword.cmpxchg.failure281; CHECK-NEXT: // in Loop: Header=BB4_13 Depth=1282; CHECK-NEXT: and.b32 %r13, %r12, %r2;283; CHECK-NEXT: setp.ne.b32 %p10, %r42, %r13;284; CHECK-NEXT: mov.b32 %r42, %r13;285; CHECK-NEXT: @%p10 bra $L__BB4_13;286; CHECK-NEXT: $L__BB4_15: // %partword.cmpxchg.end287; CHECK-NEXT: fence.acq_rel.sys;288; CHECK-NEXT: ret;289entry:290 ; Compare-exchange operation - all memory ordering combinations for 32-bit291 %0 = cmpxchg ptr addrspace(7) %dsmem_ptr, i32 1, i32 0 monotonic monotonic292 %1 = cmpxchg ptr addrspace(7) %dsmem_ptr, i32 1, i32 0 acquire monotonic293 %2 = cmpxchg ptr addrspace(7) %dsmem_ptr, i32 1, i32 0 acquire acquire294 %3 = cmpxchg ptr addrspace(7) %dsmem_ptr, i32 1, i32 0 release monotonic295 %4 = cmpxchg ptr addrspace(7) %dsmem_ptr, i32 1, i32 0 acq_rel monotonic296 %5 = cmpxchg ptr addrspace(7) %dsmem_ptr, i32 1, i32 0 acq_rel acquire297 %6 = cmpxchg ptr addrspace(7) %dsmem_ptr, i32 1, i32 0 seq_cst monotonic298 %7 = cmpxchg ptr addrspace(7) %dsmem_ptr, i32 1, i32 0 seq_cst acquire299 %8 = cmpxchg ptr addrspace(7) %dsmem_ptr, i32 1, i32 0 seq_cst seq_cst300 301 ; Compare-exchange operation - all memory ordering combinations for 64-bit302 %9 = cmpxchg ptr addrspace(7) %dsmem_ptr, i64 1, i64 0 monotonic monotonic303 %10 = cmpxchg ptr addrspace(7) %dsmem_ptr, i64 1, i64 0 acquire monotonic304 %11 = cmpxchg ptr addrspace(7) %dsmem_ptr, i64 1, i64 0 acquire acquire305 %12 = cmpxchg ptr addrspace(7) %dsmem_ptr, i64 1, i64 0 release monotonic306 %13 = cmpxchg ptr addrspace(7) %dsmem_ptr, i64 1, i64 0 acq_rel monotonic307 %14 = cmpxchg ptr addrspace(7) %dsmem_ptr, i64 1, i64 0 acq_rel acquire308 %15 = cmpxchg ptr addrspace(7) %dsmem_ptr, i64 1, i64 0 seq_cst monotonic309 %16 = cmpxchg ptr addrspace(7) %dsmem_ptr, i64 1, i64 0 seq_cst acquire310 %17 = cmpxchg ptr addrspace(7) %dsmem_ptr, i64 1, i64 0 seq_cst seq_cst311 312 ; Compare-exchange operation - 16-bit313 %18 = cmpxchg ptr addrspace(7) %dsmem_ptr, i16 1, i16 0 monotonic monotonic314 %19 = cmpxchg ptr addrspace(7) %dsmem_ptr, i16 1, i16 0 acquire acquire315 %20 = cmpxchg ptr addrspace(7) %dsmem_ptr, i16 1, i16 0 release monotonic316 %21 = cmpxchg ptr addrspace(7) %dsmem_ptr, i16 1, i16 0 acq_rel acquire317 %22 = cmpxchg ptr addrspace(7) %dsmem_ptr, i16 1, i16 0 seq_cst seq_cst318 319 ret void320}321