145 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_70 -mattr=+ptx63 | FileCheck %s --check-prefixes=CHECK3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_70 -mattr=+ptx63 | FileCheck %s --check-prefixes=CHECK644; RUN: llc < %s -mtriple=nvptx -mcpu=sm_70 -mattr=+ptx62 | FileCheck %s --check-prefixes=CHECKPTX625; RUN: %if ptxas-sm_70 && ptxas-isa-6.3 && ptxas-ptr32 %{ llc < %s -mtriple=nvptx -mcpu=sm_70 -mattr=+ptx63 | %ptxas-verify -arch=sm_70 %}6; RUN: %if ptxas-sm_70 && ptxas-isa-6.3 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_70 -mattr=+ptx63 | %ptxas-verify -arch=sm_70 %}7; RUN: %if ptxas-sm_70 && ptxas-isa-6.2 && ptxas-ptr32 %{ llc < %s -mtriple=nvptx -mcpu=sm_70 -mattr=+ptx62 | %ptxas-verify -arch=sm_70 %}8 9target triple = "nvptx64-nvidia-cuda"10 11define void @test(ptr %dp0, ptr addrspace(1) %dp1, ptr addrspace(3) %dp3, half %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.f16 %rs2, [%r1], %rs1;21; CHECK-NEXT: ld.param.b32 %r2, [test_param_1];22; CHECK-NEXT: mov.b16 %rs3, 0x3C00;23; CHECK-NEXT: atom.add.noftz.f16 %rs4, [%r1], %rs3;24; CHECK-NEXT: ld.param.b32 %r3, [test_param_2];25; CHECK-NEXT: atom.global.add.noftz.f16 %rs5, [%r2], %rs1;26; CHECK-NEXT: atom.shared.add.noftz.f16 %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.f16 %rs2, [%rd1], %rs1;38; CHECK64-NEXT: ld.param.b64 %rd2, [test_param_1];39; CHECK64-NEXT: mov.b16 %rs3, 0x3C00;40; CHECK64-NEXT: atom.add.noftz.f16 %rs4, [%rd1], %rs3;41; CHECK64-NEXT: ld.param.b64 %rd3, [test_param_2];42; CHECK64-NEXT: atom.global.add.noftz.f16 %rs5, [%rd2], %rs1;43; CHECK64-NEXT: atom.shared.add.noftz.f16 %rs6, [%rd3], %rs1;44; CHECK64-NEXT: ret;45;46; CHECKPTX62-LABEL: test(47; CHECKPTX62: {48; CHECKPTX62-NEXT: .reg .pred %p<5>;49; CHECKPTX62-NEXT: .reg .b16 %rs<11>;50; CHECKPTX62-NEXT: .reg .b32 %r<50>;51; CHECKPTX62-EMPTY:52; CHECKPTX62-NEXT: // %bb.0:53; CHECKPTX62-NEXT: ld.param.b16 %rs1, [test_param_3];54; CHECKPTX62-NEXT: ld.param.b32 %r15, [test_param_2];55; CHECKPTX62-NEXT: ld.param.b32 %r14, [test_param_1];56; CHECKPTX62-NEXT: ld.param.b32 %r16, [test_param_0];57; CHECKPTX62-NEXT: and.b32 %r1, %r16, -4;58; CHECKPTX62-NEXT: and.b32 %r17, %r16, 3;59; CHECKPTX62-NEXT: shl.b32 %r2, %r17, 3;60; CHECKPTX62-NEXT: mov.b32 %r18, 65535;61; CHECKPTX62-NEXT: shl.b32 %r19, %r18, %r2;62; CHECKPTX62-NEXT: not.b32 %r3, %r19;63; CHECKPTX62-NEXT: ld.b32 %r46, [%r1];64; CHECKPTX62-NEXT: $L__BB0_1: // %atomicrmw.start4565; CHECKPTX62-NEXT: // =>This Inner Loop Header: Depth=166; CHECKPTX62-NEXT: shr.u32 %r20, %r46, %r2;67; CHECKPTX62-NEXT: cvt.u16.u32 %rs2, %r20;68; CHECKPTX62-NEXT: add.rn.f16 %rs3, %rs2, %rs1;69; CHECKPTX62-NEXT: cvt.u32.u16 %r21, %rs3;70; CHECKPTX62-NEXT: shl.b32 %r22, %r21, %r2;71; CHECKPTX62-NEXT: and.b32 %r23, %r46, %r3;72; CHECKPTX62-NEXT: or.b32 %r24, %r23, %r22;73; CHECKPTX62-NEXT: atom.relaxed.sys.cas.b32 %r4, [%r1], %r46, %r24;74; CHECKPTX62-NEXT: setp.ne.b32 %p1, %r4, %r46;75; CHECKPTX62-NEXT: mov.b32 %r46, %r4;76; CHECKPTX62-NEXT: @%p1 bra $L__BB0_1;77; CHECKPTX62-NEXT: // %bb.2: // %atomicrmw.end4478; CHECKPTX62-NEXT: ld.b32 %r47, [%r1];79; CHECKPTX62-NEXT: $L__BB0_3: // %atomicrmw.start2780; CHECKPTX62-NEXT: // =>This Inner Loop Header: Depth=181; CHECKPTX62-NEXT: shr.u32 %r25, %r47, %r2;82; CHECKPTX62-NEXT: cvt.u16.u32 %rs4, %r25;83; CHECKPTX62-NEXT: mov.b16 %rs5, 0x3C00;84; CHECKPTX62-NEXT: add.rn.f16 %rs6, %rs4, %rs5;85; CHECKPTX62-NEXT: cvt.u32.u16 %r26, %rs6;86; CHECKPTX62-NEXT: shl.b32 %r27, %r26, %r2;87; CHECKPTX62-NEXT: and.b32 %r28, %r47, %r3;88; CHECKPTX62-NEXT: or.b32 %r29, %r28, %r27;89; CHECKPTX62-NEXT: atom.relaxed.sys.cas.b32 %r5, [%r1], %r47, %r29;90; CHECKPTX62-NEXT: setp.ne.b32 %p2, %r5, %r47;91; CHECKPTX62-NEXT: mov.b32 %r47, %r5;92; CHECKPTX62-NEXT: @%p2 bra $L__BB0_3;93; CHECKPTX62-NEXT: // %bb.4: // %atomicrmw.end2694; CHECKPTX62-NEXT: and.b32 %r6, %r14, -4;95; CHECKPTX62-NEXT: shl.b32 %r30, %r14, 3;96; CHECKPTX62-NEXT: and.b32 %r7, %r30, 24;97; CHECKPTX62-NEXT: mov.b32 %r31, 65535;98; CHECKPTX62-NEXT: shl.b32 %r32, %r31, %r7;99; CHECKPTX62-NEXT: not.b32 %r8, %r32;100; CHECKPTX62-NEXT: ld.global.b32 %r48, [%r6];101; CHECKPTX62-NEXT: $L__BB0_5: // %atomicrmw.start9102; CHECKPTX62-NEXT: // =>This Inner Loop Header: Depth=1103; CHECKPTX62-NEXT: shr.u32 %r33, %r48, %r7;104; CHECKPTX62-NEXT: cvt.u16.u32 %rs7, %r33;105; CHECKPTX62-NEXT: add.rn.f16 %rs8, %rs7, %rs1;106; CHECKPTX62-NEXT: cvt.u32.u16 %r34, %rs8;107; CHECKPTX62-NEXT: shl.b32 %r35, %r34, %r7;108; CHECKPTX62-NEXT: and.b32 %r36, %r48, %r8;109; CHECKPTX62-NEXT: or.b32 %r37, %r36, %r35;110; CHECKPTX62-NEXT: atom.relaxed.sys.global.cas.b32 %r9, [%r6], %r48, %r37;111; CHECKPTX62-NEXT: setp.ne.b32 %p3, %r9, %r48;112; CHECKPTX62-NEXT: mov.b32 %r48, %r9;113; CHECKPTX62-NEXT: @%p3 bra $L__BB0_5;114; CHECKPTX62-NEXT: // %bb.6: // %atomicrmw.end8115; CHECKPTX62-NEXT: and.b32 %r10, %r15, -4;116; CHECKPTX62-NEXT: shl.b32 %r38, %r15, 3;117; CHECKPTX62-NEXT: and.b32 %r11, %r38, 24;118; CHECKPTX62-NEXT: mov.b32 %r39, 65535;119; CHECKPTX62-NEXT: shl.b32 %r40, %r39, %r11;120; CHECKPTX62-NEXT: not.b32 %r12, %r40;121; CHECKPTX62-NEXT: ld.shared.b32 %r49, [%r10];122; CHECKPTX62-NEXT: $L__BB0_7: // %atomicrmw.start123; CHECKPTX62-NEXT: // =>This Inner Loop Header: Depth=1124; CHECKPTX62-NEXT: shr.u32 %r41, %r49, %r11;125; CHECKPTX62-NEXT: cvt.u16.u32 %rs9, %r41;126; CHECKPTX62-NEXT: add.rn.f16 %rs10, %rs9, %rs1;127; CHECKPTX62-NEXT: cvt.u32.u16 %r42, %rs10;128; CHECKPTX62-NEXT: shl.b32 %r43, %r42, %r11;129; CHECKPTX62-NEXT: and.b32 %r44, %r49, %r12;130; CHECKPTX62-NEXT: or.b32 %r45, %r44, %r43;131; CHECKPTX62-NEXT: atom.relaxed.sys.shared.cas.b32 %r13, [%r10], %r49, %r45;132; CHECKPTX62-NEXT: setp.ne.b32 %p4, %r13, %r49;133; CHECKPTX62-NEXT: mov.b32 %r49, %r13;134; CHECKPTX62-NEXT: @%p4 bra $L__BB0_7;135; CHECKPTX62-NEXT: // %bb.8: // %atomicrmw.end136; CHECKPTX62-NEXT: ret;137 %r1 = atomicrmw fadd ptr %dp0, half %val monotonic138 %r2 = atomicrmw fadd ptr %dp0, half 1.0 monotonic139 %r3 = atomicrmw fadd ptr addrspace(1) %dp1, half %val monotonic140 %r4 = atomicrmw fadd ptr addrspace(3) %dp3, half %val monotonic141 ret void142}143 144attributes #1 = { argmemonly nounwind }145