//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-22781540
// Cuda compilation tools, release 9.0, V9.0.176
// Based on LLVM 3.4svn
//

.version 6.0
.target sm_70
.address_size 64

	// .globl	_Z12wmma_exampleP6__halfS0_Pfiiiff
.extern .func  (.param .b32 func_retval0) vprintf
(
	.param .b64 vprintf_param_0,
	.param .b64 vprintf_param_1
)
;
.global .align 16 .b8 $str[9] = {99, 108, 111, 99, 107, 61, 37, 100, 0};

.visible .entry _Z12wmma_exampleP6__halfS0_Pfiiiff(
	.param .u64 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_0,
	.param .u64 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_1,
	.param .u64 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_2,
	.param .u32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_3,
	.param .u32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_4,
	.param .u32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_5,
	.param .f32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_6,
	.param .f32 _Z12wmma_exampleP6__halfS0_Pfiiiff_param_7
)
{
	.local .align 8 .b8 	__local_depot0[8];
	.reg .b64 	%SP;
	.reg .b64 	%SPL;
	.reg .pred 	%p<6>;
	.reg .f32 	%f<34>;
	.reg .b32 	%r<38>;
	.reg .b64 	%rd<18>;


	mov.u64 	%rd17, __local_depot0;
	cvta.local.u64 	%SP, %rd17;
	ld.param.u64 	%rd1, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_0];
	ld.param.u64 	%rd2, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_1];
	ld.param.u64 	%rd3, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_2];
	ld.param.u32 	%r4, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_3];
	ld.param.u32 	%r7, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_4];
	ld.param.u32 	%r5, [_Z12wmma_exampleP6__halfS0_Pfiiiff_param_5];
	// inline asm
	mov.u32 	%r6, %clock;
	// inline asm
	mov.u32 	%r8, %ntid.x;
	mov.u32 	%r9, %ctaid.x;
	mov.u32 	%r10, %tid.x;
	mad.lo.s32 	%r11, %r8, %r9, %r10;
	mov.u32 	%r12, WARP_SZ;
	div.u32 	%r13, %r11, %r12;
	mov.u32 	%r14, %ntid.y;
	mov.u32 	%r15, %ctaid.y;
	mov.u32 	%r16, %tid.y;
	mad.lo.s32 	%r17, %r14, %r15, %r16;
	shl.b32 	%r2, %r13, 4;
	shl.b32 	%r3, %r17, 4;
	setp.lt.s32	%p1, %r2, %r4;
	setp.gt.s32	%p2, %r5, 0;
	and.pred  	%p3, %p1, %p2;
	setp.lt.s32	%p4, %r3, %r7;
	and.pred  	%p5, %p3, %p4;
	mov.f32 	%f26, 0f00000000;
	mov.f32 	%f27, %f26;
	mov.f32 	%f28, %f26;
	mov.f32 	%f29, %f26;
	mov.f32 	%f30, %f26;
	mov.f32 	%f31, %f26;
	mov.f32 	%f32, %f26;
	mov.f32 	%f33, %f26;
	@!%p5 bra 	BB0_2;
	bra.uni 	BB0_1;

BB0_1:
	mul.wide.s32 	%rd4, %r2, 2;
	add.s64 	%rd5, %rd1, %rd4;
	wmma.load.a.sync.row.m16n16k16.f16 	{%r18, %r19, %r20, %r21, %r22, %r23, %r24, %r25}, [%rd5], %r4;
	mul.wide.s32 	%rd6, %r3, 2;
	add.s64 	%rd7, %rd2, %rd6;
	wmma.load.b.sync.col.m16n16k16.f16 	{%r26, %r27, %r28, %r29, %r30, %r31, %r32, %r33}, [%rd7], %r5;
	mov.f32 	%f25, 0f00000000;
	wmma.mma.sync.row.col.m16n16k16.f32.f32 {%f33, %f32, %f31, %f30, %f29, %f28, %f27, %f26}, {%r18, %r19, %r20, %r21, %r22, %r23, %r24, %r25}, {%r26, %r27, %r28, %r29, %r30, %r31, %r32, %r33}, {%f25, %f25, %f25, %f25, %f25, %f25, %f25, %f25};

BB0_2:
	add.u64 	%rd8, %SP, 0;
	cvta.to.local.u64 	%rd9, %rd8;
	mul.lo.s32 	%r35, %r3, %r4;
	cvt.s64.s32	%rd10, %r35;
	cvt.s64.s32	%rd11, %r2;
	add.s64 	%rd12, %rd10, %rd11;
	shl.b64 	%rd13, %rd12, 2;
	add.s64 	%rd14, %rd3, %rd13;
	wmma.store.d.sync.col.m16n16k16.f32 	[%rd14], {%f33, %f32, %f31, %f30, %f29, %f28, %f27, %f26}, %r4;
	// inline asm
	mov.u32 	%r34, %clock;
	// inline asm
	sub.s32 	%r36, %r34, %r6;
	st.local.u32 	[%rd9], %r36;
	mov.u64 	%rd15, $str;
	cvta.global.u64 	%rd16, %rd15;
	// Callseq Start 0
	{
	.reg .b32 temp_param_reg;
	// <end>}
	.param .b64 param0;
	st.param.b64	[param0+0], %rd16;
	.param .b64 param1;
	st.param.b64	[param1+0], %rd8;
	.param .b32 retval0;
	call.uni (retval0), 
	vprintf, 
	(
	param0, 
	param1
	);
	ld.param.b32	%r37, [retval0+0];
	
	//{
	}// Callseq End 0
	ret;
}

	// .globl	_Z17convertFp32ToFp16P6__halfPfi
.visible .entry _Z17convertFp32ToFp16P6__halfPfi(
	.param .u64 _Z17convertFp32ToFp16P6__halfPfi_param_0,
	.param .u64 _Z17convertFp32ToFp16P6__halfPfi_param_1,
	.param .u32 _Z17convertFp32ToFp16P6__halfPfi_param_2
)
{
	.reg .pred 	%p<2>;
	.reg .b16 	%rs<2>;
	.reg .f32 	%f<2>;
	.reg .b32 	%r<6>;
	.reg .b64 	%rd<9>;


	ld.param.u64 	%rd1, [_Z17convertFp32ToFp16P6__halfPfi_param_0];
	ld.param.u64 	%rd2, [_Z17convertFp32ToFp16P6__halfPfi_param_1];
	ld.param.u32 	%r2, [_Z17convertFp32ToFp16P6__halfPfi_param_2];
	mov.u32 	%r3, %ntid.x;
	mov.u32 	%r4, %ctaid.x;
	mov.u32 	%r5, %tid.x;
	mad.lo.s32 	%r1, %r4, %r3, %r5;
	setp.ge.s32	%p1, %r1, %r2;
	@%p1 bra 	BB1_2;

	cvta.to.global.u64 	%rd3, %rd2;
	mul.wide.s32 	%rd4, %r1, 4;
	add.s64 	%rd5, %rd3, %rd4;
	ld.global.f32 	%f1, [%rd5];
	// inline asm
	{  cvt.rn.f16.f32 %rs1, %f1;}

	// inline asm
	cvta.to.global.u64 	%rd6, %rd1;
	mul.wide.s32 	%rd7, %r1, 2;
	add.s64 	%rd8, %rd6, %rd7;
	st.global.u16 	[%rd8], %rs1;

BB1_2:
	ret;
}