148 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 42; RUN: llc < %s -mtriple=nvptx -mcpu=sm_90 -mattr=+ptx78 | FileCheck %s --check-prefixes=CHECK3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx78 | FileCheck %s --check-prefixes=CHECK644; RUN: llc < %s -mtriple=nvptx -mcpu=sm_86 -mattr=+ptx71 | FileCheck %s --check-prefixes=CHECKPTX715; RUN: %if ptxas-sm_90 && ptxas-isa-7.8 && ptxas-ptr32 %{ llc < %s -mtriple=nvptx -mcpu=sm_90 -mattr=+ptx78 | %ptxas-verify -arch=sm_90 %}6; RUN: %if ptxas-sm_90 && ptxas-isa-7.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx78 | %ptxas-verify -arch=sm_90 %}7; RUN: %if ptxas-sm_86 && ptxas-isa-7.1 && ptxas-ptr32 %{ llc < %s -mtriple=nvptx -mcpu=sm_86 -mattr=+ptx71 | %ptxas-verify -arch=sm_86 %}8 9target triple = "nvptx64-nvidia-cuda"10 11define void @test(ptr %dp0, ptr addrspace(1) %dp1, ptr addrspace(3) %dp3, bfloat %val) {12; CHECK-LABEL: test(13; CHECK: {14; CHECK-NEXT: .reg .b16 %rs<7>;15; CHECK-NEXT: .reg .b32 %r<4>;16; CHECK-EMPTY:17; CHECK-NEXT: // %bb.0:18; CHECK-NEXT: ld.param.b32 %r1, [test_param_0];19; CHECK-NEXT: ld.param.b16 %rs1, [test_param_3];20; CHECK-NEXT: atom.add.noftz.bf16 %rs2, [%r1], %rs1;21; CHECK-NEXT: ld.param.b32 %r2, [test_param_1];22; CHECK-NEXT: mov.b16 %rs3, 0x3F80;23; CHECK-NEXT: atom.add.noftz.bf16 %rs4, [%r1], %rs3;24; CHECK-NEXT: ld.param.b32 %r3, [test_param_2];25; CHECK-NEXT: atom.global.add.noftz.bf16 %rs5, [%r2], %rs1;26; CHECK-NEXT: atom.shared.add.noftz.bf16 %rs6, [%r3], %rs1;27; CHECK-NEXT: ret;28;29; CHECK64-LABEL: test(30; CHECK64: {31; CHECK64-NEXT: .reg .b16 %rs<7>;32; CHECK64-NEXT: .reg .b64 %rd<4>;33; CHECK64-EMPTY:34; CHECK64-NEXT: // %bb.0:35; CHECK64-NEXT: ld.param.b64 %rd1, [test_param_0];36; CHECK64-NEXT: ld.param.b16 %rs1, [test_param_3];37; CHECK64-NEXT: atom.add.noftz.bf16 %rs2, [%rd1], %rs1;38; CHECK64-NEXT: ld.param.b64 %rd2, [test_param_1];39; CHECK64-NEXT: mov.b16 %rs3, 0x3F80;40; CHECK64-NEXT: atom.add.noftz.bf16 %rs4, [%rd1], %rs3;41; CHECK64-NEXT: ld.param.b64 %rd3, [test_param_2];42; CHECK64-NEXT: atom.global.add.noftz.bf16 %rs5, [%rd2], %rs1;43; CHECK64-NEXT: atom.shared.add.noftz.bf16 %rs6, [%rd3], %rs1;44; CHECK64-NEXT: ret;45;46; CHECKPTX71-LABEL: test(47; CHECKPTX71: {48; CHECKPTX71-NEXT: .reg .pred %p<5>;49; CHECKPTX71-NEXT: .reg .b16 %rs<14>;50; CHECKPTX71-NEXT: .reg .b32 %r<50>;51; CHECKPTX71-EMPTY:52; CHECKPTX71-NEXT: // %bb.0:53; CHECKPTX71-NEXT: ld.param.b16 %rs1, [test_param_3];54; CHECKPTX71-NEXT: ld.param.b32 %r15, [test_param_2];55; CHECKPTX71-NEXT: ld.param.b32 %r14, [test_param_1];56; CHECKPTX71-NEXT: ld.param.b32 %r16, [test_param_0];57; CHECKPTX71-NEXT: and.b32 %r1, %r16, -4;58; CHECKPTX71-NEXT: and.b32 %r17, %r16, 3;59; CHECKPTX71-NEXT: shl.b32 %r2, %r17, 3;60; CHECKPTX71-NEXT: mov.b32 %r18, 65535;61; CHECKPTX71-NEXT: shl.b32 %r19, %r18, %r2;62; CHECKPTX71-NEXT: not.b32 %r3, %r19;63; CHECKPTX71-NEXT: ld.b32 %r46, [%r1];64; CHECKPTX71-NEXT: $L__BB0_1: // %atomicrmw.start4565; CHECKPTX71-NEXT: // =>This Inner Loop Header: Depth=166; CHECKPTX71-NEXT: shr.u32 %r20, %r46, %r2;67; CHECKPTX71-NEXT: cvt.u16.u32 %rs2, %r20;68; CHECKPTX71-NEXT: mov.b16 %rs3, 0x3F80;69; CHECKPTX71-NEXT: fma.rn.bf16 %rs4, %rs2, %rs3, %rs1;70; CHECKPTX71-NEXT: cvt.u32.u16 %r21, %rs4;71; CHECKPTX71-NEXT: shl.b32 %r22, %r21, %r2;72; CHECKPTX71-NEXT: and.b32 %r23, %r46, %r3;73; CHECKPTX71-NEXT: or.b32 %r24, %r23, %r22;74; CHECKPTX71-NEXT: atom.relaxed.sys.cas.b32 %r4, [%r1], %r46, %r24;75; CHECKPTX71-NEXT: setp.ne.b32 %p1, %r4, %r46;76; CHECKPTX71-NEXT: mov.b32 %r46, %r4;77; CHECKPTX71-NEXT: @%p1 bra $L__BB0_1;78; CHECKPTX71-NEXT: // %bb.2: // %atomicrmw.end4479; CHECKPTX71-NEXT: ld.b32 %r47, [%r1];80; CHECKPTX71-NEXT: $L__BB0_3: // %atomicrmw.start2781; CHECKPTX71-NEXT: // =>This Inner Loop Header: Depth=182; CHECKPTX71-NEXT: shr.u32 %r25, %r47, %r2;83; CHECKPTX71-NEXT: cvt.u16.u32 %rs5, %r25;84; CHECKPTX71-NEXT: mov.b16 %rs6, 0x3F80;85; CHECKPTX71-NEXT: fma.rn.bf16 %rs7, %rs5, %rs6, %rs6;86; CHECKPTX71-NEXT: cvt.u32.u16 %r26, %rs7;87; CHECKPTX71-NEXT: shl.b32 %r27, %r26, %r2;88; CHECKPTX71-NEXT: and.b32 %r28, %r47, %r3;89; CHECKPTX71-NEXT: or.b32 %r29, %r28, %r27;90; CHECKPTX71-NEXT: atom.relaxed.sys.cas.b32 %r5, [%r1], %r47, %r29;91; CHECKPTX71-NEXT: setp.ne.b32 %p2, %r5, %r47;92; CHECKPTX71-NEXT: mov.b32 %r47, %r5;93; CHECKPTX71-NEXT: @%p2 bra $L__BB0_3;94; CHECKPTX71-NEXT: // %bb.4: // %atomicrmw.end2695; CHECKPTX71-NEXT: and.b32 %r6, %r14, -4;96; CHECKPTX71-NEXT: shl.b32 %r30, %r14, 3;97; CHECKPTX71-NEXT: and.b32 %r7, %r30, 24;98; CHECKPTX71-NEXT: mov.b32 %r31, 65535;99; CHECKPTX71-NEXT: shl.b32 %r32, %r31, %r7;100; CHECKPTX71-NEXT: not.b32 %r8, %r32;101; CHECKPTX71-NEXT: ld.global.b32 %r48, [%r6];102; CHECKPTX71-NEXT: $L__BB0_5: // %atomicrmw.start9103; CHECKPTX71-NEXT: // =>This Inner Loop Header: Depth=1104; CHECKPTX71-NEXT: shr.u32 %r33, %r48, %r7;105; CHECKPTX71-NEXT: cvt.u16.u32 %rs8, %r33;106; CHECKPTX71-NEXT: mov.b16 %rs9, 0x3F80;107; CHECKPTX71-NEXT: fma.rn.bf16 %rs10, %rs8, %rs9, %rs1;108; CHECKPTX71-NEXT: cvt.u32.u16 %r34, %rs10;109; CHECKPTX71-NEXT: shl.b32 %r35, %r34, %r7;110; CHECKPTX71-NEXT: and.b32 %r36, %r48, %r8;111; CHECKPTX71-NEXT: or.b32 %r37, %r36, %r35;112; CHECKPTX71-NEXT: atom.relaxed.sys.global.cas.b32 %r9, [%r6], %r48, %r37;113; CHECKPTX71-NEXT: setp.ne.b32 %p3, %r9, %r48;114; CHECKPTX71-NEXT: mov.b32 %r48, %r9;115; CHECKPTX71-NEXT: @%p3 bra $L__BB0_5;116; CHECKPTX71-NEXT: // %bb.6: // %atomicrmw.end8117; CHECKPTX71-NEXT: and.b32 %r10, %r15, -4;118; CHECKPTX71-NEXT: shl.b32 %r38, %r15, 3;119; CHECKPTX71-NEXT: and.b32 %r11, %r38, 24;120; CHECKPTX71-NEXT: mov.b32 %r39, 65535;121; CHECKPTX71-NEXT: shl.b32 %r40, %r39, %r11;122; CHECKPTX71-NEXT: not.b32 %r12, %r40;123; CHECKPTX71-NEXT: ld.shared.b32 %r49, [%r10];124; CHECKPTX71-NEXT: $L__BB0_7: // %atomicrmw.start125; CHECKPTX71-NEXT: // =>This Inner Loop Header: Depth=1126; CHECKPTX71-NEXT: shr.u32 %r41, %r49, %r11;127; CHECKPTX71-NEXT: cvt.u16.u32 %rs11, %r41;128; CHECKPTX71-NEXT: mov.b16 %rs12, 0x3F80;129; CHECKPTX71-NEXT: fma.rn.bf16 %rs13, %rs11, %rs12, %rs1;130; CHECKPTX71-NEXT: cvt.u32.u16 %r42, %rs13;131; CHECKPTX71-NEXT: shl.b32 %r43, %r42, %r11;132; CHECKPTX71-NEXT: and.b32 %r44, %r49, %r12;133; CHECKPTX71-NEXT: or.b32 %r45, %r44, %r43;134; CHECKPTX71-NEXT: atom.relaxed.sys.shared.cas.b32 %r13, [%r10], %r49, %r45;135; CHECKPTX71-NEXT: setp.ne.b32 %p4, %r13, %r49;136; CHECKPTX71-NEXT: mov.b32 %r49, %r13;137; CHECKPTX71-NEXT: @%p4 bra $L__BB0_7;138; CHECKPTX71-NEXT: // %bb.8: // %atomicrmw.end139; CHECKPTX71-NEXT: ret;140 %r1 = atomicrmw fadd ptr %dp0, bfloat %val monotonic141 %r2 = atomicrmw fadd ptr %dp0, bfloat 1.0 monotonic142 %r3 = atomicrmw fadd ptr addrspace(1) %dp1, bfloat %val monotonic143 %r4 = atomicrmw fadd ptr addrspace(3) %dp3, bfloat %val monotonic144 ret void145}146 147attributes #1 = { argmemonly nounwind }148