Added many other

2026-08-06 18:56:32 +03:00
parent 9f2fd08ee5
commit d9d7ee49d2
2144 changed files with 1115424 additions and 1284 deletions
@@ -0,0 +1,140 @@
+//
+// Generated by NVIDIA NVVM Compiler
+//
+// Compiler Build ID: CL-23083092
+// Cuda compilation tools, release 9.1, V9.1.85
+// Based on LLVM 3.4svn
+//
+
+.version 6.1
+.target sm_52
+.address_size 64
+
+	// .globl	oxMain
+.const .align 16 .b8 params[1184];
+
+.visible .entry oxMain(
+
+)
+{
+	.reg .b16 	%rs<31>;
+	.reg .f32 	%f<20>;
+	.reg .b32 	%r<15>;
+	.reg .b64 	%rd<13>;
+
+
+	mov.u32 	%r1, %ctaid.x;
+	mov.u32 	%r2, %ntid.x;
+	mov.u32 	%r3, %tid.x;
+	mad.lo.s32 	%r4, %r2, %r1, %r3;
+	mov.u32 	%r5, %ctaid.y;
+	mov.u32 	%r6, %ntid.y;
+	mov.u32 	%r7, %tid.y;
+	mad.lo.s32 	%r8, %r6, %r5, %r7;
+	ld.const.u64 	%rd1, [params+144];
+	cvta.to.global.u64 	%rd2, %rd1;
+	ld.const.u32 	%r9, [params+136];
+	mad.lo.s32 	%r10, %r9, %r8, %r4;
+	mul.wide.u32 	%rd3, %r10, 8;
+	add.s64 	%rd4, %rd2, %rd3;
+	ld.global.v4.u16 	{%rs19, %rs20, %rs21, %rs22}, [%rd4];
+	// inline asm
+	{  cvt.f32.f16 %f1, %rs22;}
+
+	// inline asm
+	// inline asm
+	{  cvt.f32.f16 %f2, %rs19;}
+
+	// inline asm
+	// inline asm
+	{  cvt.f32.f16 %f3, %rs20;}
+
+	// inline asm
+	// inline asm
+	{  cvt.f32.f16 %f4, %rs21;}
+
+	// inline asm
+	ld.const.u64 	%rd5, [params+160];
+	cvta.to.global.u64 	%rd6, %rd5;
+	ld.const.u32 	%r11, [params+152];
+	mad.lo.s32 	%r12, %r11, %r8, %r4;
+	mul.wide.u32 	%rd7, %r12, 8;
+	add.s64 	%rd8, %rd6, %rd7;
+	ld.global.v4.u16 	{%rs23, %rs24, %rs25, %rs26}, [%rd8];
+	// inline asm
+	{  cvt.f32.f16 %f5, %rs26;}
+
+	// inline asm
+	// inline asm
+	{  cvt.f32.f16 %f6, %rs23;}
+
+	// inline asm
+	// inline asm
+	{  cvt.f32.f16 %f7, %rs24;}
+
+	// inline asm
+	// inline asm
+	{  cvt.f32.f16 %f8, %rs25;}
+
+	// inline asm
+	add.f32 	%f9, %f2, %f6;
+	add.f32 	%f10, %f3, %f7;
+	add.f32 	%f11, %f4, %f8;
+	min.f32 	%f12, %f1, %f5;
+	// inline asm
+	{  cvt.rn.f16.f32 %rs12, %f12;}
+
+	// inline asm
+	// inline asm
+	{  cvt.rn.f16.f32 %rs11, %f11;}
+
+	// inline asm
+	// inline asm
+	{  cvt.rn.f16.f32 %rs10, %f10;}
+
+	// inline asm
+	// inline asm
+	{  cvt.rn.f16.f32 %rs9, %f9;}
+
+	// inline asm
+	st.global.v4.u16 	[%rd8], {%rs9, %rs10, %rs11, %rs12};
+	ld.const.u64 	%rd9, [params+744];
+	cvta.to.global.u64 	%rd10, %rd9;
+	ld.const.u32 	%r13, [params+736];
+	mad.lo.s32 	%r14, %r13, %r8, %r4;
+	mul.wide.u32 	%rd11, %r14, 8;
+	add.s64 	%rd12, %rd10, %rd11;
+	ld.global.v4.u16 	{%rs27, %rs28, %rs29, %rs30}, [%rd12];
+	// inline asm
+	{  cvt.f32.f16 %f13, %rs27;}
+
+	// inline asm
+	// inline asm
+	{  cvt.f32.f16 %f14, %rs28;}
+
+	// inline asm
+	// inline asm
+	{  cvt.f32.f16 %f15, %rs29;}
+
+	// inline asm
+	ld.const.f32 	%f19, [params+752];
+	fma.rn.f32 	%f16, %f2, %f19, %f13;
+	fma.rn.f32 	%f17, %f3, %f19, %f14;
+	fma.rn.f32 	%f18, %f4, %f19, %f15;
+	// inline asm
+	{  cvt.rn.f16.f32 %rs18, %f18;}
+
+	// inline asm
+	// inline asm
+	{  cvt.rn.f16.f32 %rs17, %f17;}
+
+	// inline asm
+	// inline asm
+	{  cvt.rn.f16.f32 %rs16, %f16;}
+
+	// inline asm
+	st.global.v4.u16 	[%rd12], {%rs16, %rs17, %rs18, %rs12};
+	ret;
+}
+
+