632 lines
14 KiB
Plaintext
632 lines
14 KiB
Plaintext
//
|
|
// Generated by NVIDIA NVVM Compiler
|
|
//
|
|
// Compiler Build ID: CL-31968024
|
|
// Cuda compilation tools, release 12.0, V12.0.76
|
|
// Based on NVVM 7.0.1
|
|
//
|
|
|
|
.version 8.0
|
|
.target sm_52
|
|
.address_size 64
|
|
|
|
// .globl oxMain
|
|
.const .align 16 .b8 params[1184];
|
|
|
|
.visible .entry oxMain()
|
|
{
|
|
.reg .b16 %rs<93>;
|
|
.reg .f32 %f<120>;
|
|
.reg .b32 %r<45>;
|
|
.reg .b64 %rd<73>;
|
|
|
|
|
|
mov.u32 %r1, %ctaid.x;
|
|
mov.u32 %r2, %ntid.x;
|
|
mov.u32 %r3, %tid.x;
|
|
mad.lo.s32 %r4, %r1, %r2, %r3;
|
|
mov.u32 %r5, %ctaid.y;
|
|
mov.u32 %r6, %ntid.y;
|
|
mov.u32 %r7, %tid.y;
|
|
mad.lo.s32 %r8, %r5, %r6, %r7;
|
|
ld.const.u64 %rd1, [params+784];
|
|
cvta.to.global.u64 %rd2, %rd1;
|
|
ld.const.u32 %r9, [params+776];
|
|
mad.lo.s32 %r10, %r9, %r8, %r4;
|
|
mul.wide.u32 %rd3, %r10, 8;
|
|
add.s64 %rd4, %rd2, %rd3;
|
|
ld.global.u16 %rs1, [%rd4];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f1, %rs1;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs2, [%rd4+2];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f2, %rs2;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs3, [%rd4+4];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f3, %rs3;}
|
|
|
|
// end inline asm
|
|
ld.const.u64 %rd5, [params+800];
|
|
cvta.to.global.u64 %rd6, %rd5;
|
|
ld.const.u32 %r11, [params+792];
|
|
mad.lo.s32 %r12, %r11, %r8, %r4;
|
|
mul.wide.u32 %rd7, %r12, 8;
|
|
add.s64 %rd8, %rd6, %rd7;
|
|
ld.global.u16 %rs4, [%rd8];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f4, %rs4;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs5, [%rd8+2];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f5, %rs5;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs6, [%rd8+4];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f6, %rs6;}
|
|
|
|
// end inline asm
|
|
ld.const.u64 %rd9, [params+816];
|
|
cvta.to.global.u64 %rd10, %rd9;
|
|
ld.const.u32 %r13, [params+808];
|
|
mad.lo.s32 %r14, %r13, %r8, %r4;
|
|
mul.wide.u32 %rd11, %r14, 8;
|
|
add.s64 %rd12, %rd10, %rd11;
|
|
ld.global.u16 %rs7, [%rd12];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f7, %rs7;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs8, [%rd12+2];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f8, %rs8;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs9, [%rd12+4];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f9, %rs9;}
|
|
|
|
// end inline asm
|
|
ld.const.u64 %rd13, [params+832];
|
|
cvta.to.global.u64 %rd14, %rd13;
|
|
ld.const.u32 %r15, [params+824];
|
|
mad.lo.s32 %r16, %r15, %r8, %r4;
|
|
mul.wide.u32 %rd15, %r16, 8;
|
|
add.s64 %rd16, %rd14, %rd15;
|
|
ld.global.u16 %rs10, [%rd16];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f10, %rs10;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs11, [%rd16+2];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f11, %rs11;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs12, [%rd16+4];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f12, %rs12;}
|
|
|
|
// end inline asm
|
|
ld.const.u64 %rd17, [params+848];
|
|
cvta.to.global.u64 %rd18, %rd17;
|
|
ld.const.u32 %r17, [params+840];
|
|
mad.lo.s32 %r18, %r17, %r8, %r4;
|
|
mul.wide.u32 %rd19, %r18, 8;
|
|
add.s64 %rd20, %rd18, %rd19;
|
|
ld.global.u16 %rs13, [%rd20];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f13, %rs13;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs14, [%rd20+2];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f14, %rs14;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs15, [%rd20+4];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f15, %rs15;}
|
|
|
|
// end inline asm
|
|
ld.const.u64 %rd21, [params+864];
|
|
cvta.to.global.u64 %rd22, %rd21;
|
|
ld.const.u32 %r19, [params+856];
|
|
mad.lo.s32 %r20, %r19, %r8, %r4;
|
|
mul.wide.u32 %rd23, %r20, 8;
|
|
add.s64 %rd24, %rd22, %rd23;
|
|
ld.global.u16 %rs16, [%rd24];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f16, %rs16;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs17, [%rd24+2];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f17, %rs17;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs18, [%rd24+4];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f18, %rs18;}
|
|
|
|
// end inline asm
|
|
ld.const.u64 %rd25, [params+880];
|
|
cvta.to.global.u64 %rd26, %rd25;
|
|
ld.const.u32 %r21, [params+872];
|
|
mad.lo.s32 %r22, %r21, %r8, %r4;
|
|
mul.wide.u32 %rd27, %r22, 8;
|
|
add.s64 %rd28, %rd26, %rd27;
|
|
ld.global.u16 %rs19, [%rd28];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f19, %rs19;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs20, [%rd28+2];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f20, %rs20;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs21, [%rd28+4];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f21, %rs21;}
|
|
|
|
// end inline asm
|
|
ld.const.u64 %rd29, [params+896];
|
|
cvta.to.global.u64 %rd30, %rd29;
|
|
ld.const.u32 %r23, [params+888];
|
|
mad.lo.s32 %r24, %r23, %r8, %r4;
|
|
mul.wide.u32 %rd31, %r24, 8;
|
|
add.s64 %rd32, %rd30, %rd31;
|
|
ld.global.u16 %rs22, [%rd32];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f22, %rs22;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs23, [%rd32+2];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f23, %rs23;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs24, [%rd32+4];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f24, %rs24;}
|
|
|
|
// end inline asm
|
|
ld.const.u64 %rd33, [params+912];
|
|
cvta.to.global.u64 %rd34, %rd33;
|
|
ld.const.u32 %r25, [params+904];
|
|
mad.lo.s32 %r26, %r25, %r8, %r4;
|
|
mul.wide.u32 %rd35, %r26, 8;
|
|
add.s64 %rd36, %rd34, %rd35;
|
|
ld.global.u16 %rs25, [%rd36];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f25, %rs25;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs26, [%rd36+2];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f26, %rs26;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs27, [%rd36+4];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f27, %rs27;}
|
|
|
|
// end inline asm
|
|
ld.const.u64 %rd37, [params+928];
|
|
cvta.to.global.u64 %rd38, %rd37;
|
|
ld.const.u32 %r27, [params+920];
|
|
mad.lo.s32 %r28, %r27, %r8, %r4;
|
|
mul.wide.u32 %rd39, %r28, 8;
|
|
add.s64 %rd40, %rd38, %rd39;
|
|
ld.global.u16 %rs28, [%rd40];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f28, %rs28;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs29, [%rd40+2];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f29, %rs29;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs30, [%rd40+4];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f30, %rs30;}
|
|
|
|
// end inline asm
|
|
ld.const.u64 %rd41, [params+944];
|
|
cvta.to.global.u64 %rd42, %rd41;
|
|
ld.const.u32 %r29, [params+936];
|
|
mad.lo.s32 %r30, %r29, %r8, %r4;
|
|
mul.wide.u32 %rd43, %r30, 8;
|
|
add.s64 %rd44, %rd42, %rd43;
|
|
ld.global.u16 %rs31, [%rd44];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f31, %rs31;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs32, [%rd44+2];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f32, %rs32;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs33, [%rd44+4];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f33, %rs33;}
|
|
|
|
// end inline asm
|
|
ld.const.u64 %rd45, [params+960];
|
|
cvta.to.global.u64 %rd46, %rd45;
|
|
ld.const.u32 %r31, [params+952];
|
|
mad.lo.s32 %r32, %r31, %r8, %r4;
|
|
mul.wide.u32 %rd47, %r32, 8;
|
|
add.s64 %rd48, %rd46, %rd47;
|
|
ld.global.u16 %rs34, [%rd48];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f34, %rs34;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs35, [%rd48+2];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f35, %rs35;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs36, [%rd48+4];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f36, %rs36;}
|
|
|
|
// end inline asm
|
|
ld.const.u64 %rd49, [params+976];
|
|
cvta.to.global.u64 %rd50, %rd49;
|
|
ld.const.u32 %r33, [params+968];
|
|
mad.lo.s32 %r34, %r33, %r8, %r4;
|
|
mul.wide.u32 %rd51, %r34, 8;
|
|
add.s64 %rd52, %rd50, %rd51;
|
|
ld.global.u16 %rs37, [%rd52];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f37, %rs37;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs38, [%rd52+2];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f38, %rs38;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs39, [%rd52+4];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f39, %rs39;}
|
|
|
|
// end inline asm
|
|
ld.const.u64 %rd53, [params+992];
|
|
cvta.to.global.u64 %rd54, %rd53;
|
|
ld.const.u32 %r35, [params+984];
|
|
mad.lo.s32 %r36, %r35, %r8, %r4;
|
|
mul.wide.u32 %rd55, %r36, 8;
|
|
add.s64 %rd56, %rd54, %rd55;
|
|
ld.global.u16 %rs40, [%rd56];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f40, %rs40;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs41, [%rd56+2];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f41, %rs41;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs42, [%rd56+4];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f42, %rs42;}
|
|
|
|
// end inline asm
|
|
ld.const.u64 %rd57, [params+1008];
|
|
cvta.to.global.u64 %rd58, %rd57;
|
|
ld.const.u32 %r37, [params+1000];
|
|
mad.lo.s32 %r38, %r37, %r8, %r4;
|
|
mul.wide.u32 %rd59, %r38, 8;
|
|
add.s64 %rd60, %rd58, %rd59;
|
|
ld.global.u16 %rs43, [%rd60];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f43, %rs43;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs44, [%rd60+2];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f44, %rs44;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs45, [%rd60+4];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f45, %rs45;}
|
|
|
|
// end inline asm
|
|
ld.const.u64 %rd61, [params+1024];
|
|
cvta.to.global.u64 %rd62, %rd61;
|
|
ld.const.u32 %r39, [params+1016];
|
|
mad.lo.s32 %r40, %r39, %r8, %r4;
|
|
mul.wide.u32 %rd63, %r40, 8;
|
|
add.s64 %rd64, %rd62, %rd63;
|
|
ld.global.u16 %rs46, [%rd64];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f46, %rs46;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs47, [%rd64+2];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f47, %rs47;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs48, [%rd64+4];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f48, %rs48;}
|
|
|
|
// end inline asm
|
|
ld.const.u64 %rd65, [params+1040];
|
|
cvta.to.global.u64 %rd66, %rd65;
|
|
ld.const.u32 %r41, [params+1032];
|
|
mad.lo.s32 %r42, %r41, %r8, %r4;
|
|
mul.wide.u32 %rd67, %r42, 8;
|
|
add.s64 %rd68, %rd66, %rd67;
|
|
ld.global.u16 %rs49, [%rd68];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f49, %rs49;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs50, [%rd68+2];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f50, %rs50;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs51, [%rd68+4];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f51, %rs51;}
|
|
|
|
// end inline asm
|
|
ld.const.u64 %rd69, [params+1056];
|
|
cvta.to.global.u64 %rd70, %rd69;
|
|
ld.const.u32 %r43, [params+1048];
|
|
mad.lo.s32 %r44, %r43, %r8, %r4;
|
|
mul.wide.u32 %rd71, %r44, 8;
|
|
add.s64 %rd72, %rd70, %rd71;
|
|
ld.global.u16 %rs52, [%rd72];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f52, %rs52;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs53, [%rd72+2];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f53, %rs53;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs54, [%rd72+4];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f54, %rs54;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs55, [%rd4+6];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f55, %rs55;}
|
|
|
|
// end inline asm
|
|
ld.global.u16 %rs56, [%rd40+6];
|
|
// begin inline asm
|
|
{ cvt.f32.f16 %f56, %rs56;}
|
|
|
|
// end inline asm
|
|
min.f32 %f92, %f55, %f56;
|
|
add.f32 %f93, %f1, %f28;
|
|
add.f32 %f94, %f2, %f29;
|
|
add.f32 %f95, %f3, %f30;
|
|
add.f32 %f96, %f4, %f31;
|
|
add.f32 %f97, %f5, %f32;
|
|
add.f32 %f98, %f6, %f33;
|
|
add.f32 %f99, %f7, %f34;
|
|
add.f32 %f100, %f8, %f35;
|
|
add.f32 %f101, %f9, %f36;
|
|
add.f32 %f102, %f10, %f37;
|
|
add.f32 %f103, %f11, %f38;
|
|
add.f32 %f104, %f12, %f39;
|
|
add.f32 %f105, %f13, %f40;
|
|
add.f32 %f106, %f14, %f41;
|
|
add.f32 %f107, %f15, %f42;
|
|
add.f32 %f108, %f16, %f43;
|
|
add.f32 %f109, %f17, %f44;
|
|
add.f32 %f110, %f18, %f45;
|
|
add.f32 %f111, %f19, %f46;
|
|
add.f32 %f112, %f20, %f47;
|
|
add.f32 %f113, %f21, %f48;
|
|
add.f32 %f114, %f22, %f49;
|
|
add.f32 %f115, %f23, %f50;
|
|
add.f32 %f116, %f24, %f51;
|
|
add.f32 %f117, %f25, %f52;
|
|
add.f32 %f118, %f26, %f53;
|
|
add.f32 %f119, %f27, %f54;
|
|
mul.f32 %f57, %f93, %f92;
|
|
mul.f32 %f58, %f94, %f92;
|
|
mul.f32 %f59, %f95, %f92;
|
|
mul.f32 %f61, %f96, %f92;
|
|
mul.f32 %f62, %f97, %f92;
|
|
mul.f32 %f63, %f98, %f92;
|
|
mul.f32 %f65, %f99, %f92;
|
|
mul.f32 %f66, %f100, %f92;
|
|
mul.f32 %f67, %f101, %f92;
|
|
mul.f32 %f69, %f102, %f92;
|
|
mul.f32 %f70, %f103, %f92;
|
|
mul.f32 %f71, %f104, %f92;
|
|
mul.f32 %f73, %f105, %f92;
|
|
mul.f32 %f74, %f106, %f92;
|
|
mul.f32 %f75, %f107, %f92;
|
|
mul.f32 %f77, %f108, %f92;
|
|
mul.f32 %f78, %f109, %f92;
|
|
mul.f32 %f79, %f110, %f92;
|
|
mul.f32 %f81, %f111, %f92;
|
|
mul.f32 %f82, %f112, %f92;
|
|
mul.f32 %f83, %f113, %f92;
|
|
mul.f32 %f85, %f114, %f92;
|
|
mul.f32 %f86, %f115, %f92;
|
|
mul.f32 %f87, %f116, %f92;
|
|
mul.f32 %f89, %f117, %f92;
|
|
mul.f32 %f90, %f118, %f92;
|
|
mul.f32 %f91, %f119, %f92;
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs59, %f59;}
|
|
|
|
// end inline asm
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs58, %f58;}
|
|
|
|
// end inline asm
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs57, %f57;}
|
|
|
|
// end inline asm
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs60, %f92;}
|
|
|
|
// end inline asm
|
|
st.global.v4.u16 [%rd40], {%rs57, %rs58, %rs59, %rs60};
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs63, %f63;}
|
|
|
|
// end inline asm
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs62, %f62;}
|
|
|
|
// end inline asm
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs61, %f61;}
|
|
|
|
// end inline asm
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs64, %f92;}
|
|
|
|
// end inline asm
|
|
st.global.v4.u16 [%rd44], {%rs61, %rs62, %rs63, %rs64};
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs67, %f67;}
|
|
|
|
// end inline asm
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs66, %f66;}
|
|
|
|
// end inline asm
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs65, %f65;}
|
|
|
|
// end inline asm
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs68, %f92;}
|
|
|
|
// end inline asm
|
|
st.global.v4.u16 [%rd48], {%rs65, %rs66, %rs67, %rs68};
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs71, %f71;}
|
|
|
|
// end inline asm
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs70, %f70;}
|
|
|
|
// end inline asm
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs69, %f69;}
|
|
|
|
// end inline asm
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs72, %f92;}
|
|
|
|
// end inline asm
|
|
st.global.v4.u16 [%rd52], {%rs69, %rs70, %rs71, %rs72};
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs75, %f75;}
|
|
|
|
// end inline asm
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs74, %f74;}
|
|
|
|
// end inline asm
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs73, %f73;}
|
|
|
|
// end inline asm
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs76, %f92;}
|
|
|
|
// end inline asm
|
|
st.global.v4.u16 [%rd56], {%rs73, %rs74, %rs75, %rs76};
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs79, %f79;}
|
|
|
|
// end inline asm
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs78, %f78;}
|
|
|
|
// end inline asm
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs77, %f77;}
|
|
|
|
// end inline asm
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs80, %f92;}
|
|
|
|
// end inline asm
|
|
st.global.v4.u16 [%rd60], {%rs77, %rs78, %rs79, %rs80};
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs83, %f83;}
|
|
|
|
// end inline asm
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs82, %f82;}
|
|
|
|
// end inline asm
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs81, %f81;}
|
|
|
|
// end inline asm
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs84, %f92;}
|
|
|
|
// end inline asm
|
|
st.global.v4.u16 [%rd64], {%rs81, %rs82, %rs83, %rs84};
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs87, %f87;}
|
|
|
|
// end inline asm
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs86, %f86;}
|
|
|
|
// end inline asm
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs85, %f85;}
|
|
|
|
// end inline asm
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs88, %f92;}
|
|
|
|
// end inline asm
|
|
st.global.v4.u16 [%rd68], {%rs85, %rs86, %rs87, %rs88};
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs91, %f91;}
|
|
|
|
// end inline asm
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs90, %f90;}
|
|
|
|
// end inline asm
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs89, %f89;}
|
|
|
|
// end inline asm
|
|
// begin inline asm
|
|
{ cvt.rn.f16.f32 %rs92, %f92;}
|
|
|
|
// end inline asm
|
|
st.global.v4.u16 [%rd72], {%rs89, %rs90, %rs91, %rs92};
|
|
ret;
|
|
|
|
}
|
|
|