.version 7.8
.target sm_60
.address_size 64

	// .globl	_Z27dequant_gemv_group32_batch623DequantGemvKernelParams
// _ZZ9gemv_int4ILi4ELi32ELi6EEvP6__halfPKS0_S3_PKjPKhS3_iiiiffbE12temp_storage has been demoted
.global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_cf3b3ebd6thrust6system6detail10sequential3seqE[1];
.global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_cf3b3ebd6thrust12placeholders2_1E[1];
.global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_cf3b3ebd6thrust12placeholders2_2E[1];
.global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_cf3b3ebd6thrust12placeholders2_3E[1];
.global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_cf3b3ebd6thrust12placeholders2_4E[1];
.global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_cf3b3ebd6thrust12placeholders2_5E[1];
.global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_cf3b3ebd6thrust12placeholders2_6E[1];
.global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_cf3b3ebd6thrust12placeholders2_7E[1];
.global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_cf3b3ebd6thrust12placeholders2_8E[1];
.global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_cf3b3ebd6thrust12placeholders2_9E[1];
.global .align 1 .b8 _ZN37_INTERNAL_36609967_7_gemv_cu_cf3b3ebd6thrust12placeholders3_10E[1];

.visible .entry _Z27dequant_gemv_group32_batch623DequantGemvKernelParams(
	.param .align 8 .b8 _Z27dequant_gemv_group32_batch623DequantGemvKernelParams_param_0[80]
)
{
	.reg .pred 	%p<22>;
	.reg .b16 	%rs<247>;
	.reg .f32 	%f<458>;
	.reg .b32 	%r<215>;
	.reg .b64 	%rd<60>;
	// demoted variable
	.shared .align 16 .b8 _ZZ9gemv_int4ILi4ELi32ELi6EEvP6__halfPKS0_S3_PKjPKhS3_iiiiffbE12temp_storage[3840];

	ld.param.v2.u32 	{%r24, %r25}, [_Z27dequant_gemv_group32_batch623DequantGemvKernelParams_param_0+48];
	ld.param.v2.u32 	{%r26, %r27}, [_Z27dequant_gemv_group32_batch623DequantGemvKernelParams_param_0+56];
	ld.param.v2.f32 	{%f59, %f60}, [_Z27dequant_gemv_group32_batch623DequantGemvKernelParams_param_0+64];
	ld.param.v4.u8 	{%rs76, %rs77, %rs78, %rs79}, [_Z27dequant_gemv_group32_batch623DequantGemvKernelParams_param_0+72];
	ld.param.u64 	%rd20, [_Z27dequant_gemv_group32_batch623DequantGemvKernelParams_param_0+40];
	ld.param.u64 	%rd19, [_Z27dequant_gemv_group32_batch623DequantGemvKernelParams_param_0+32];
	ld.param.u64 	%rd18, [_Z27dequant_gemv_group32_batch623DequantGemvKernelParams_param_0+24];
	ld.param.u64 	%rd17, [_Z27dequant_gemv_group32_batch623DequantGemvKernelParams_param_0+16];
	mov.u32 	%r1, %ctaid.x;
	mov.u32 	%r214, %tid.y;
	shl.b32 	%r28, %r214, 5;
	mov.u32 	%r4, %tid.x;
	add.s32 	%r213, %r28, %r4;
	setp.ge.u32 	%p1, %r213, %r26;
	mov.f32 	%f440, 0f00000000;
	mov.f32 	%f441, %f440;
	mov.f32 	%f442, %f440;
	mov.f32 	%f443, %f440;
	mov.f32 	%f444, %f440;
	mov.f32 	%f445, %f440;
	@%p1 bra 	$L__BB0_9;

	cvta.to.global.u64 	%rd1, %rd17;
	mul.lo.s32 	%r7, %r26, %r1;
	shr.u32 	%r8, %r4, 2;
	mul.lo.s32 	%r9, %r27, %r1;
	shl.b16 	%rs2, %rs76, 3;
	mul.wide.s32 	%rd3, %r24, 2;
	cvta.to.global.u64 	%rd4, %rd19;
	cvta.to.global.u64 	%rd5, %rd18;
	cvta.to.global.u64 	%rd6, %rd20;

$L__BB0_2:
	add.s32 	%r29, %r213, %r7;
	mul.wide.u32 	%rd21, %r29, 4;
	add.s64 	%rd22, %rd5, %rd21;
	ld.global.u32 	%r12, [%rd22];
	shl.b32 	%r30, %r214, 3;
	add.s32 	%r13, %r30, %r8;
	add.s32 	%r14, %r13, %r9;
	mul.wide.s32 	%rd23, %r14, 2;
	add.s64 	%rd24, %rd6, %rd23;
	ld.global.u16 	%rs84, [%rd24];
	// begin inline asm
	{  cvt.f32.f16 %f67, %rs84;}

	// end inline asm
	setp.eq.s64 	%p2, %rd19, 0;
	mov.u16 	%rs246, %rs2;
	@%p2 bra 	$L__BB0_4;

	shr.u32 	%r31, %r14, 31;
	add.s32 	%r32, %r14, %r31;
	shr.s32 	%r33, %r32, 1;
	cvt.s64.s32 	%rd25, %r33;
	add.s64 	%rd26, %rd4, %rd25;
	ld.global.u8 	%r34, [%rd26];
	shl.b32 	%r35, %r13, 2;
	and.b32  	%r36, %r35, 4;
	shr.u32 	%r37, %r34, %r36;
	cvt.u16.u32 	%rs85, %r37;
	and.b16  	%rs246, %rs85, 15;

$L__BB0_4:
	shl.b32 	%r15, %r213, 3;
	setp.ge.s32 	%p3, %r15, %r24;
	@%p3 bra 	$L__BB0_8;

	setp.eq.s16 	%p4, %rs76, 0;
	mul.wide.s32 	%rd27, %r15, 2;
	add.s64 	%rd28, %rd1, %rd27;
	ld.global.v4.u32 	{%r38, %r39, %r40, %r41}, [%rd28];
	add.s64 	%rd30, %rd28, %rd3;
	ld.global.v4.u32 	{%r46, %r47, %r48, %r49}, [%rd30];
	add.s32 	%r54, %r24, %r15;
	add.s32 	%r55, %r54, %r24;
	mul.wide.s32 	%rd31, %r55, 2;
	add.s64 	%rd32, %rd1, %rd31;
	ld.global.v4.u32 	{%r56, %r57, %r58, %r59}, [%rd32];
	add.s64 	%rd33, %rd32, %rd3;
	ld.global.v4.u32 	{%r64, %r65, %r66, %r67}, [%rd33];
	add.s64 	%rd34, %rd33, %rd3;
	ld.global.v4.u32 	{%r72, %r73, %r74, %r75}, [%rd34];
	add.s64 	%rd35, %rd34, %rd3;
	ld.global.v4.u32 	{%r80, %r81, %r82, %r83}, [%rd35];
	shr.u16 	%rs87, %rs246, 3;
	and.b16  	%rs88, %rs87, 1;
	setp.eq.b16 	%p5, %rs88, 1;
	and.pred  	%p6, %p4, %p5;
	selp.b16 	%rs89, -16, 0, %p6;
	or.b16  	%rs90, %rs89, %rs246;
	cvt.s16.s8 	%rs91, %rs90;
	cvt.rn.f32.s16 	%f8, %rs91;
	cvt.u16.u32 	%rs5, %r12;
	and.b16  	%rs6, %rs5, 15;
	mov.b32 	{%rs7, %rs15}, %r38;
	mov.b32 	{%rs8, %rs16}, %r46;
	mov.b32 	{%rs9, %rs17}, %r56;
	mov.b32 	{%rs10, %rs18}, %r64;
	mov.b32 	{%rs11, %rs19}, %r72;
	mov.b32 	{%rs12, %rs20}, %r80;
	shr.u32 	%r88, %r12, 4;
	cvt.u16.u32 	%rs13, %r88;
	and.b16  	%rs14, %rs13, 15;
	shr.u32 	%r89, %r12, 8;
	cvt.u16.u32 	%rs21, %r89;
	and.b16  	%rs22, %rs21, 15;
	mov.b32 	{%rs23, %rs31}, %r39;
	mov.b32 	{%rs24, %rs32}, %r47;
	mov.b32 	{%rs25, %rs33}, %r57;
	mov.b32 	{%rs26, %rs34}, %r65;
	mov.b32 	{%rs27, %rs35}, %r73;
	mov.b32 	{%rs28, %rs36}, %r81;
	shr.u32 	%r90, %r12, 12;
	cvt.u16.u32 	%rs29, %r90;
	and.b16  	%rs30, %rs29, 15;
	shr.u32 	%r91, %r12, 16;
	cvt.u16.u32 	%rs37, %r91;
	and.b16  	%rs38, %rs37, 15;
	mov.b32 	{%rs39, %rs47}, %r40;
	mov.b32 	{%rs40, %rs48}, %r48;
	mov.b32 	{%rs41, %rs49}, %r58;
	mov.b32 	{%rs42, %rs50}, %r66;
	mov.b32 	{%rs43, %rs51}, %r74;
	mov.b32 	{%rs44, %rs52}, %r82;
	shr.u32 	%r92, %r12, 20;
	cvt.u16.u32 	%rs45, %r92;
	and.b16  	%rs46, %rs45, 15;
	shr.u32 	%r93, %r12, 24;
	cvt.u16.u32 	%rs53, %r93;
	and.b16  	%rs54, %rs53, 15;
	mov.b32 	{%rs55, %rs62}, %r41;
	mov.b32 	{%rs56, %rs63}, %r49;
	mov.b32 	{%rs57, %rs64}, %r59;
	mov.b32 	{%rs58, %rs65}, %r67;
	mov.b32 	{%rs59, %rs66}, %r75;
	mov.b32 	{%rs60, %rs67}, %r83;
	shr.u32 	%r94, %r12, 28;
	cvt.u16.u32 	%rs61, %r94;
	@%p4 bra 	$L__BB0_7;

	cvt.rn.f32.s16 	%f116, %rs6;
	sub.ftz.f32 	%f117, %f116, %f8;
	mul.ftz.f32 	%f118, %f67, %f117;
	// begin inline asm
	{  cvt.f32.f16 %f68, %rs7;}

	// end inline asm
	fma.rn.ftz.f32 	%f119, %f118, %f68, %f440;
	// begin inline asm
	{  cvt.f32.f16 %f69, %rs8;}

	// end inline asm
	fma.rn.ftz.f32 	%f120, %f118, %f69, %f441;
	// begin inline asm
	{  cvt.f32.f16 %f70, %rs9;}

	// end inline asm
	fma.rn.ftz.f32 	%f121, %f118, %f70, %f442;
	// begin inline asm
	{  cvt.f32.f16 %f71, %rs10;}

	// end inline asm
	fma.rn.ftz.f32 	%f122, %f118, %f71, %f443;
	// begin inline asm
	{  cvt.f32.f16 %f72, %rs11;}

	// end inline asm
	fma.rn.ftz.f32 	%f123, %f118, %f72, %f444;
	// begin inline asm
	{  cvt.f32.f16 %f73, %rs12;}

	// end inline asm
	fma.rn.ftz.f32 	%f124, %f118, %f73, %f445;
	cvt.rn.f32.s16 	%f125, %rs14;
	sub.ftz.f32 	%f126, %f125, %f8;
	mul.ftz.f32 	%f127, %f67, %f126;
	// begin inline asm
	{  cvt.f32.f16 %f74, %rs15;}

	// end inline asm
	fma.rn.ftz.f32 	%f128, %f127, %f74, %f119;
	// begin inline asm
	{  cvt.f32.f16 %f75, %rs16;}

	// end inline asm
	fma.rn.ftz.f32 	%f129, %f127, %f75, %f120;
	// begin inline asm
	{  cvt.f32.f16 %f76, %rs17;}

	// end inline asm
	fma.rn.ftz.f32 	%f130, %f127, %f76, %f121;
	// begin inline asm
	{  cvt.f32.f16 %f77, %rs18;}

	// end inline asm
	fma.rn.ftz.f32 	%f131, %f127, %f77, %f122;
	// begin inline asm
	{  cvt.f32.f16 %f78, %rs19;}

	// end inline asm
	fma.rn.ftz.f32 	%f132, %f127, %f78, %f123;
	// begin inline asm
	{  cvt.f32.f16 %f79, %rs20;}

	// end inline asm
	fma.rn.ftz.f32 	%f133, %f127, %f79, %f124;
	cvt.rn.f32.s16 	%f134, %rs22;
	sub.ftz.f32 	%f135, %f134, %f8;
	mul.ftz.f32 	%f136, %f67, %f135;
	// begin inline asm
	{  cvt.f32.f16 %f80, %rs23;}

	// end inline asm
	fma.rn.ftz.f32 	%f137, %f136, %f80, %f128;
	// begin inline asm
	{  cvt.f32.f16 %f81, %rs24;}

	// end inline asm
	fma.rn.ftz.f32 	%f138, %f136, %f81, %f129;
	// begin inline asm
	{  cvt.f32.f16 %f82, %rs25;}

	// end inline asm
	fma.rn.ftz.f32 	%f139, %f136, %f82, %f130;
	// begin inline asm
	{  cvt.f32.f16 %f83, %rs26;}

	// end inline asm
	fma.rn.ftz.f32 	%f140, %f136, %f83, %f131;
	// begin inline asm
	{  cvt.f32.f16 %f84, %rs27;}

	// end inline asm
	fma.rn.ftz.f32 	%f141, %f136, %f84, %f132;
	// begin inline asm
	{  cvt.f32.f16 %f85, %rs28;}

	// end inline asm
	fma.rn.ftz.f32 	%f142, %f136, %f85, %f133;
	cvt.rn.f32.s16 	%f143, %rs30;
	sub.ftz.f32 	%f144, %f143, %f8;
	mul.ftz.f32 	%f145, %f67, %f144;
	// begin inline asm
	{  cvt.f32.f16 %f86, %rs31;}

	// end inline asm
	fma.rn.ftz.f32 	%f146, %f145, %f86, %f137;
	// begin inline asm
	{  cvt.f32.f16 %f87, %rs32;}

	// end inline asm
	fma.rn.ftz.f32 	%f147, %f145, %f87, %f138;
	// begin inline asm
	{  cvt.f32.f16 %f88, %rs33;}

	// end inline asm
	fma.rn.ftz.f32 	%f148, %f145, %f88, %f139;
	// begin inline asm
	{  cvt.f32.f16 %f89, %rs34;}

	// end inline asm
	fma.rn.ftz.f32 	%f149, %f145, %f89, %f140;
	// begin inline asm
	{  cvt.f32.f16 %f90, %rs35;}

	// end inline asm
	fma.rn.ftz.f32 	%f150, %f145, %f90, %f141;
	// begin inline asm
	{  cvt.f32.f16 %f91, %rs36;}

	// end inline asm
	fma.rn.ftz.f32 	%f151, %f145, %f91, %f142;
	cvt.rn.f32.s16 	%f152, %rs38;
	sub.ftz.f32 	%f153, %f152, %f8;
	mul.ftz.f32 	%f154, %f67, %f153;
	// begin inline asm
	{  cvt.f32.f16 %f92, %rs39;}

	// end inline asm
	fma.rn.ftz.f32 	%f155, %f154, %f92, %f146;
	// begin inline asm
	{  cvt.f32.f16 %f93, %rs40;}

	// end inline asm
	fma.rn.ftz.f32 	%f156, %f154, %f93, %f147;
	// begin inline asm
	{  cvt.f32.f16 %f94, %rs41;}

	// end inline asm
	fma.rn.ftz.f32 	%f157, %f154, %f94, %f148;
	// begin inline asm
	{  cvt.f32.f16 %f95, %rs42;}

	// end inline asm
	fma.rn.ftz.f32 	%f158, %f154, %f95, %f149;
	// begin inline asm
	{  cvt.f32.f16 %f96, %rs43;}

	// end inline asm
	fma.rn.ftz.f32 	%f159, %f154, %f96, %f150;
	// begin inline asm
	{  cvt.f32.f16 %f97, %rs44;}

	// end inline asm
	fma.rn.ftz.f32 	%f160, %f154, %f97, %f151;
	cvt.rn.f32.s16 	%f161, %rs46;
	sub.ftz.f32 	%f162, %f161, %f8;
	mul.ftz.f32 	%f163, %f67, %f162;
	// begin inline asm
	{  cvt.f32.f16 %f98, %rs47;}

	// end inline asm
	fma.rn.ftz.f32 	%f164, %f163, %f98, %f155;
	// begin inline asm
	{  cvt.f32.f16 %f99, %rs48;}

	// end inline asm
	fma.rn.ftz.f32 	%f165, %f163, %f99, %f156;
	// begin inline asm
	{  cvt.f32.f16 %f100, %rs49;}

	// end inline asm
	fma.rn.ftz.f32 	%f166, %f163, %f100, %f157;
	// begin inline asm
	{  cvt.f32.f16 %f101, %rs50;}

	// end inline asm
	fma.rn.ftz.f32 	%f167, %f163, %f101, %f158;
	// begin inline asm
	{  cvt.f32.f16 %f102, %rs51;}

	// end inline asm
	fma.rn.ftz.f32 	%f168, %f163, %f102, %f159;
	// begin inline asm
	{  cvt.f32.f16 %f103, %rs52;}

	// end inline asm
	fma.rn.ftz.f32 	%f169, %f163, %f103, %f160;
	cvt.rn.f32.s16 	%f170, %rs54;
	sub.ftz.f32 	%f171, %f170, %f8;
	mul.ftz.f32 	%f172, %f67, %f171;
	// begin inline asm
	{  cvt.f32.f16 %f104, %rs55;}

	// end inline asm
	fma.rn.ftz.f32 	%f173, %f172, %f104, %f164;
	// begin inline asm
	{  cvt.f32.f16 %f105, %rs56;}

	// end inline asm
	fma.rn.ftz.f32 	%f174, %f172, %f105, %f165;
	// begin inline asm
	{  cvt.f32.f16 %f106, %rs57;}

	// end inline asm
	fma.rn.ftz.f32 	%f175, %f172, %f106, %f166;
	// begin inline asm
	{  cvt.f32.f16 %f107, %rs58;}

	// end inline asm
	fma.rn.ftz.f32 	%f176, %f172, %f107, %f167;
	// begin inline asm
	{  cvt.f32.f16 %f108, %rs59;}

	// end inline asm
	fma.rn.ftz.f32 	%f177, %f172, %f108, %f168;
	// begin inline asm
	{  cvt.f32.f16 %f109, %rs60;}

	// end inline asm
	fma.rn.ftz.f32 	%f178, %f172, %f109, %f169;
	cvt.rn.f32.s16 	%f179, %rs61;
	sub.ftz.f32 	%f180, %f179, %f8;
	mul.ftz.f32 	%f181, %f67, %f180;
	// begin inline asm
	{  cvt.f32.f16 %f110, %rs62;}

	// end inline asm
	fma.rn.ftz.f32 	%f440, %f181, %f110, %f173;
	// begin inline asm
	{  cvt.f32.f16 %f111, %rs63;}

	// end inline asm
	fma.rn.ftz.f32 	%f441, %f181, %f111, %f174;
	// begin inline asm
	{  cvt.f32.f16 %f112, %rs64;}

	// end inline asm
	fma.rn.ftz.f32 	%f442, %f181, %f112, %f175;
	// begin inline asm
	{  cvt.f32.f16 %f113, %rs65;}

	// end inline asm
	fma.rn.ftz.f32 	%f443, %f181, %f113, %f176;
	// begin inline asm
	{  cvt.f32.f16 %f114, %rs66;}

	// end inline asm
	fma.rn.ftz.f32 	%f444, %f181, %f114, %f177;
	// begin inline asm
	{  cvt.f32.f16 %f115, %rs67;}

	// end inline asm
	fma.rn.ftz.f32 	%f445, %f181, %f115, %f178;
	bra.uni 	$L__BB0_8;

$L__BB0_7:
	shl.b16 	%rs188, %rs5, 4;
	cvt.s16.s8 	%rs189, %rs188;
	shr.s16 	%rs190, %rs189, 7;
	and.b16  	%rs191, %rs190, -16;
	or.b16  	%rs192, %rs191, %rs6;
	cvt.rn.f32.s16 	%f230, %rs192;
	sub.ftz.f32 	%f231, %f230, %f8;
	mul.ftz.f32 	%f232, %f67, %f231;
	// begin inline asm
	{  cvt.f32.f16 %f182, %rs7;}

	// end inline asm
	fma.rn.ftz.f32 	%f233, %f232, %f182, %f440;
	// begin inline asm
	{  cvt.f32.f16 %f183, %rs8;}

	// end inline asm
	fma.rn.ftz.f32 	%f234, %f232, %f183, %f441;
	// begin inline asm
	{  cvt.f32.f16 %f184, %rs9;}

	// end inline asm
	fma.rn.ftz.f32 	%f235, %f232, %f184, %f442;
	// begin inline asm
	{  cvt.f32.f16 %f185, %rs10;}

	// end inline asm
	fma.rn.ftz.f32 	%f236, %f232, %f185, %f443;
	// begin inline asm
	{  cvt.f32.f16 %f186, %rs11;}

	// end inline asm
	fma.rn.ftz.f32 	%f237, %f232, %f186, %f444;
	// begin inline asm
	{  cvt.f32.f16 %f187, %rs12;}

	// end inline asm
	fma.rn.ftz.f32 	%f238, %f232, %f187, %f445;
	shl.b16 	%rs193, %rs13, 4;
	cvt.s16.s8 	%rs194, %rs193;
	shr.s16 	%rs195, %rs194, 7;
	and.b16  	%rs196, %rs195, -16;
	or.b16  	%rs197, %rs196, %rs14;
	cvt.rn.f32.s16 	%f239, %rs197;
	sub.ftz.f32 	%f240, %f239, %f8;
	mul.ftz.f32 	%f241, %f67, %f240;
	// begin inline asm
	{  cvt.f32.f16 %f188, %rs15;}

	// end inline asm
	fma.rn.ftz.f32 	%f242, %f241, %f188, %f233;
	// begin inline asm
	{  cvt.f32.f16 %f189, %rs16;}

	// end inline asm
	fma.rn.ftz.f32 	%f243, %f241, %f189, %f234;
	// begin inline asm
	{  cvt.f32.f16 %f190, %rs17;}

	// end inline asm
	fma.rn.ftz.f32 	%f244, %f241, %f190, %f235;
	// begin inline asm
	{  cvt.f32.f16 %f191, %rs18;}

	// end inline asm
	fma.rn.ftz.f32 	%f245, %f241, %f191, %f236;
	// begin inline asm
	{  cvt.f32.f16 %f192, %rs19;}

	// end inline asm
	fma.rn.ftz.f32 	%f246, %f241, %f192, %f237;
	// begin inline asm
	{  cvt.f32.f16 %f193, %rs20;}

	// end inline asm
	fma.rn.ftz.f32 	%f247, %f241, %f193, %f238;
	shl.b16 	%rs198, %rs21, 4;
	cvt.s16.s8 	%rs199, %rs198;
	shr.s16 	%rs200, %rs199, 7;
	and.b16  	%rs201, %rs200, -16;
	or.b16  	%rs202, %rs201, %rs22;
	cvt.rn.f32.s16 	%f248, %rs202;
	sub.ftz.f32 	%f249, %f248, %f8;
	mul.ftz.f32 	%f250, %f67, %f249;
	// begin inline asm
	{  cvt.f32.f16 %f194, %rs23;}

	// end inline asm
	fma.rn.ftz.f32 	%f251, %f250, %f194, %f242;
	// begin inline asm
	{  cvt.f32.f16 %f195, %rs24;}

	// end inline asm
	fma.rn.ftz.f32 	%f252, %f250, %f195, %f243;
	// begin inline asm
	{  cvt.f32.f16 %f196, %rs25;}

	// end inline asm
	fma.rn.ftz.f32 	%f253, %f250, %f196, %f244;
	// begin inline asm
	{  cvt.f32.f16 %f197, %rs26;}

	// end inline asm
	fma.rn.ftz.f32 	%f254, %f250, %f197, %f245;
	// begin inline asm
	{  cvt.f32.f16 %f198, %rs27;}

	// end inline asm
	fma.rn.ftz.f32 	%f255, %f250, %f198, %f246;
	// begin inline asm
	{  cvt.f32.f16 %f199, %rs28;}

	// end inline asm
	fma.rn.ftz.f32 	%f256, %f250, %f199, %f247;
	shl.b16 	%rs203, %rs29, 4;
	cvt.s16.s8 	%rs204, %rs203;
	shr.s16 	%rs205, %rs204, 7;
	and.b16  	%rs206, %rs205, -16;
	or.b16  	%rs207, %rs206, %rs30;
	cvt.rn.f32.s16 	%f257, %rs207;
	sub.ftz.f32 	%f258, %f257, %f8;
	mul.ftz.f32 	%f259, %f67, %f258;
	// begin inline asm
	{  cvt.f32.f16 %f200, %rs31;}

	// end inline asm
	fma.rn.ftz.f32 	%f260, %f259, %f200, %f251;
	// begin inline asm
	{  cvt.f32.f16 %f201, %rs32;}

	// end inline asm
	fma.rn.ftz.f32 	%f261, %f259, %f201, %f252;
	// begin inline asm
	{  cvt.f32.f16 %f202, %rs33;}

	// end inline asm
	fma.rn.ftz.f32 	%f262, %f259, %f202, %f253;
	// begin inline asm
	{  cvt.f32.f16 %f203, %rs34;}

	// end inline asm
	fma.rn.ftz.f32 	%f263, %f259, %f203, %f254;
	// begin inline asm
	{  cvt.f32.f16 %f204, %rs35;}

	// end inline asm
	fma.rn.ftz.f32 	%f264, %f259, %f204, %f255;
	// begin inline asm
	{  cvt.f32.f16 %f205, %rs36;}

	// end inline asm
	fma.rn.ftz.f32 	%f265, %f259, %f205, %f256;
	shl.b16 	%rs208, %rs37, 4;
	cvt.s16.s8 	%rs209, %rs208;
	shr.s16 	%rs210, %rs209, 7;
	and.b16  	%rs211, %rs210, -16;
	or.b16  	%rs212, %rs211, %rs38;
	cvt.rn.f32.s16 	%f266, %rs212;
	sub.ftz.f32 	%f267, %f266, %f8;
	mul.ftz.f32 	%f268, %f67, %f267;
	// begin inline asm
	{  cvt.f32.f16 %f206, %rs39;}

	// end inline asm
	fma.rn.ftz.f32 	%f269, %f268, %f206, %f260;
	// begin inline asm
	{  cvt.f32.f16 %f207, %rs40;}

	// end inline asm
	fma.rn.ftz.f32 	%f270, %f268, %f207, %f261;
	// begin inline asm
	{  cvt.f32.f16 %f208, %rs41;}

	// end inline asm
	fma.rn.ftz.f32 	%f271, %f268, %f208, %f262;
	// begin inline asm
	{  cvt.f32.f16 %f209, %rs42;}

	// end inline asm
	fma.rn.ftz.f32 	%f272, %f268, %f209, %f263;
	// begin inline asm
	{  cvt.f32.f16 %f210, %rs43;}

	// end inline asm
	fma.rn.ftz.f32 	%f273, %f268, %f210, %f264;
	// begin inline asm
	{  cvt.f32.f16 %f211, %rs44;}

	// end inline asm
	fma.rn.ftz.f32 	%f274, %f268, %f211, %f265;
	shl.b16 	%rs213, %rs45, 4;
	cvt.s16.s8 	%rs214, %rs213;
	shr.s16 	%rs215, %rs214, 7;
	and.b16  	%rs216, %rs215, -16;
	or.b16  	%rs217, %rs216, %rs46;
	cvt.rn.f32.s16 	%f275, %rs217;
	sub.ftz.f32 	%f276, %f275, %f8;
	mul.ftz.f32 	%f277, %f67, %f276;
	// begin inline asm
	{  cvt.f32.f16 %f212, %rs47;}

	// end inline asm
	fma.rn.ftz.f32 	%f278, %f277, %f212, %f269;
	// begin inline asm
	{  cvt.f32.f16 %f213, %rs48;}

	// end inline asm
	fma.rn.ftz.f32 	%f279, %f277, %f213, %f270;
	// begin inline asm
	{  cvt.f32.f16 %f214, %rs49;}

	// end inline asm
	fma.rn.ftz.f32 	%f280, %f277, %f214, %f271;
	// begin inline asm
	{  cvt.f32.f16 %f215, %rs50;}

	// end inline asm
	fma.rn.ftz.f32 	%f281, %f277, %f215, %f272;
	// begin inline asm
	{  cvt.f32.f16 %f216, %rs51;}

	// end inline asm
	fma.rn.ftz.f32 	%f282, %f277, %f216, %f273;
	// begin inline asm
	{  cvt.f32.f16 %f217, %rs52;}

	// end inline asm
	fma.rn.ftz.f32 	%f283, %f277, %f217, %f274;
	shl.b16 	%rs218, %rs53, 4;
	cvt.s16.s8 	%rs219, %rs218;
	shr.s16 	%rs220, %rs219, 7;
	and.b16  	%rs221, %rs220, -16;
	or.b16  	%rs222, %rs221, %rs54;
	cvt.rn.f32.s16 	%f284, %rs222;
	sub.ftz.f32 	%f285, %f284, %f8;
	mul.ftz.f32 	%f286, %f67, %f285;
	// begin inline asm
	{  cvt.f32.f16 %f218, %rs55;}

	// end inline asm
	fma.rn.ftz.f32 	%f287, %f286, %f218, %f278;
	// begin inline asm
	{  cvt.f32.f16 %f219, %rs56;}

	// end inline asm
	fma.rn.ftz.f32 	%f288, %f286, %f219, %f279;
	// begin inline asm
	{  cvt.f32.f16 %f220, %rs57;}

	// end inline asm
	fma.rn.ftz.f32 	%f289, %f286, %f220, %f280;
	// begin inline asm
	{  cvt.f32.f16 %f221, %rs58;}

	// end inline asm
	fma.rn.ftz.f32 	%f290, %f286, %f221, %f281;
	// begin inline asm
	{  cvt.f32.f16 %f222, %rs59;}

	// end inline asm
	fma.rn.ftz.f32 	%f291, %f286, %f222, %f282;
	// begin inline asm
	{  cvt.f32.f16 %f223, %rs60;}

	// end inline asm
	fma.rn.ftz.f32 	%f292, %f286, %f223, %f283;
	shl.b16 	%rs223, %rs61, 4;
	cvt.s16.s8 	%rs224, %rs223;
	shr.s16 	%rs225, %rs224, 7;
	and.b16  	%rs226, %rs225, -16;
	or.b16  	%rs227, %rs226, %rs61;
	cvt.rn.f32.s16 	%f293, %rs227;
	sub.ftz.f32 	%f294, %f293, %f8;
	mul.ftz.f32 	%f295, %f67, %f294;
	// begin inline asm
	{  cvt.f32.f16 %f224, %rs62;}

	// end inline asm
	fma.rn.ftz.f32 	%f440, %f295, %f224, %f287;
	// begin inline asm
	{  cvt.f32.f16 %f225, %rs63;}

	// end inline asm
	fma.rn.ftz.f32 	%f441, %f295, %f225, %f288;
	// begin inline asm
	{  cvt.f32.f16 %f226, %rs64;}

	// end inline asm
	fma.rn.ftz.f32 	%f442, %f295, %f226, %f289;
	// begin inline asm
	{  cvt.f32.f16 %f227, %rs65;}

	// end inline asm
	fma.rn.ftz.f32 	%f443, %f295, %f227, %f290;
	// begin inline asm
	{  cvt.f32.f16 %f228, %rs66;}

	// end inline asm
	fma.rn.ftz.f32 	%f444, %f295, %f228, %f291;
	// begin inline asm
	{  cvt.f32.f16 %f229, %rs67;}

	// end inline asm
	fma.rn.ftz.f32 	%f445, %f295, %f229, %f292;

$L__BB0_8:
	add.s32 	%r214, %r214, 4;
	shl.b32 	%r95, %r214, 5;
	add.s32 	%r213, %r95, %r4;
	setp.lt.u32 	%p7, %r213, %r26;
	@%p7 bra 	$L__BB0_2;

$L__BB0_9:
	mov.u32 	%r209, %tid.y;
	shl.b32 	%r208, %r209, 5;
	add.s32 	%r207, %r208, %r4;
	shl.b32 	%r96, %r207, 2;
	mov.u32 	%r97, _ZZ9gemv_int4ILi4ELi32ELi6EEvP6__halfPKS0_S3_PKjPKhS3_iiiiffbE12temp_storage;
	add.s32 	%r98, %r97, %r96;
	setp.lt.u32 	%p8, %r207, 32;
	@%p8 bra 	$L__BB0_11;

	add.s32 	%r199, %r98, -112;
	st.shared.f32 	[%r199], %f440;

$L__BB0_11:
	mov.u32 	%r212, %tid.y;
	shl.b32 	%r211, %r212, 5;
	add.s32 	%r210, %r211, %r4;
	setp.gt.u32 	%p9, %r210, 31;
	bar.sync 	0;
	mad.lo.s32 	%r19, %r210, 12, %r97;
	@%p9 bra 	$L__BB0_13;

	mov.u32 	%r113, 16;
	ld.shared.f32 	%f311, [%r19+16];
	add.ftz.f32 	%f312, %f440, %f311;
	ld.shared.f32 	%f313, [%r19+20];
	add.ftz.f32 	%f314, %f312, %f313;
	ld.shared.f32 	%f315, [%r19+24];
	add.ftz.f32 	%f298, %f314, %f315;
	mov.u32 	%r101, 1;
	mov.u32 	%r114, 31;
	mov.u32 	%r115, -1;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f298, %r101, %r114, %r115;  @p add.f32 r0, r0, %f298;  mov.f32 %f296, r0;}
	// end inline asm
	mov.u32 	%r104, 2;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f296, %r104, %r114, %r115;  @p add.f32 r0, r0, %f296;  mov.f32 %f299, r0;}
	// end inline asm
	mov.u32 	%r107, 4;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f299, %r107, %r114, %r115;  @p add.f32 r0, r0, %f299;  mov.f32 %f302, r0;}
	// end inline asm
	mov.u32 	%r110, 8;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f302, %r110, %r114, %r115;  @p add.f32 r0, r0, %f302;  mov.f32 %f305, r0;}
	// end inline asm
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f305, %r113, %r114, %r115;  @p add.f32 r0, r0, %f305;  mov.f32 %f440, r0;}
	// end inline asm

$L__BB0_13:
	@%p8 bra 	$L__BB0_15;

	add.s32 	%r200, %r98, -112;
	st.shared.f32 	[%r200+640], %f441;

$L__BB0_15:
	bar.sync 	0;
	@%p9 bra 	$L__BB0_17;

	ld.shared.f32 	%f331, [%r19+656];
	add.ftz.f32 	%f332, %f441, %f331;
	ld.shared.f32 	%f333, [%r19+660];
	add.ftz.f32 	%f334, %f332, %f333;
	ld.shared.f32 	%f335, [%r19+664];
	add.ftz.f32 	%f318, %f334, %f335;
	mov.u32 	%r117, 1;
	mov.u32 	%r130, 31;
	mov.u32 	%r131, -1;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f318, %r117, %r130, %r131;  @p add.f32 r0, r0, %f318;  mov.f32 %f316, r0;}
	// end inline asm
	mov.u32 	%r120, 2;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f316, %r120, %r130, %r131;  @p add.f32 r0, r0, %f316;  mov.f32 %f319, r0;}
	// end inline asm
	mov.u32 	%r123, 4;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f319, %r123, %r130, %r131;  @p add.f32 r0, r0, %f319;  mov.f32 %f322, r0;}
	// end inline asm
	mov.u32 	%r126, 8;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f322, %r126, %r130, %r131;  @p add.f32 r0, r0, %f322;  mov.f32 %f325, r0;}
	// end inline asm
	mov.u32 	%r129, 16;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f325, %r129, %r130, %r131;  @p add.f32 r0, r0, %f325;  mov.f32 %f441, r0;}
	// end inline asm

$L__BB0_17:
	@%p8 bra 	$L__BB0_19;

	add.s32 	%r201, %r98, -112;
	st.shared.f32 	[%r201+1280], %f442;

$L__BB0_19:
	bar.sync 	0;
	@%p9 bra 	$L__BB0_21;

	ld.shared.f32 	%f351, [%r19+1296];
	add.ftz.f32 	%f352, %f442, %f351;
	ld.shared.f32 	%f353, [%r19+1300];
	add.ftz.f32 	%f354, %f352, %f353;
	ld.shared.f32 	%f355, [%r19+1304];
	add.ftz.f32 	%f338, %f354, %f355;
	mov.u32 	%r133, 1;
	mov.u32 	%r146, 31;
	mov.u32 	%r147, -1;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f338, %r133, %r146, %r147;  @p add.f32 r0, r0, %f338;  mov.f32 %f336, r0;}
	// end inline asm
	mov.u32 	%r136, 2;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f336, %r136, %r146, %r147;  @p add.f32 r0, r0, %f336;  mov.f32 %f339, r0;}
	// end inline asm
	mov.u32 	%r139, 4;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f339, %r139, %r146, %r147;  @p add.f32 r0, r0, %f339;  mov.f32 %f342, r0;}
	// end inline asm
	mov.u32 	%r142, 8;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f342, %r142, %r146, %r147;  @p add.f32 r0, r0, %f342;  mov.f32 %f345, r0;}
	// end inline asm
	mov.u32 	%r145, 16;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f345, %r145, %r146, %r147;  @p add.f32 r0, r0, %f345;  mov.f32 %f442, r0;}
	// end inline asm

$L__BB0_21:
	@%p8 bra 	$L__BB0_23;

	add.s32 	%r202, %r98, -112;
	st.shared.f32 	[%r202+1920], %f443;

$L__BB0_23:
	bar.sync 	0;
	@%p9 bra 	$L__BB0_25;

	ld.shared.f32 	%f371, [%r19+1936];
	add.ftz.f32 	%f372, %f443, %f371;
	ld.shared.f32 	%f373, [%r19+1940];
	add.ftz.f32 	%f374, %f372, %f373;
	ld.shared.f32 	%f375, [%r19+1944];
	add.ftz.f32 	%f358, %f374, %f375;
	mov.u32 	%r149, 1;
	mov.u32 	%r162, 31;
	mov.u32 	%r163, -1;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f358, %r149, %r162, %r163;  @p add.f32 r0, r0, %f358;  mov.f32 %f356, r0;}
	// end inline asm
	mov.u32 	%r152, 2;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f356, %r152, %r162, %r163;  @p add.f32 r0, r0, %f356;  mov.f32 %f359, r0;}
	// end inline asm
	mov.u32 	%r155, 4;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f359, %r155, %r162, %r163;  @p add.f32 r0, r0, %f359;  mov.f32 %f362, r0;}
	// end inline asm
	mov.u32 	%r158, 8;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f362, %r158, %r162, %r163;  @p add.f32 r0, r0, %f362;  mov.f32 %f365, r0;}
	// end inline asm
	mov.u32 	%r161, 16;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f365, %r161, %r162, %r163;  @p add.f32 r0, r0, %f365;  mov.f32 %f443, r0;}
	// end inline asm

$L__BB0_25:
	@%p8 bra 	$L__BB0_27;

	add.s32 	%r203, %r98, -112;
	st.shared.f32 	[%r203+2560], %f444;

$L__BB0_27:
	bar.sync 	0;
	@%p9 bra 	$L__BB0_29;

	ld.shared.f32 	%f391, [%r19+2576];
	add.ftz.f32 	%f392, %f444, %f391;
	ld.shared.f32 	%f393, [%r19+2580];
	add.ftz.f32 	%f394, %f392, %f393;
	ld.shared.f32 	%f395, [%r19+2584];
	add.ftz.f32 	%f378, %f394, %f395;
	mov.u32 	%r165, 1;
	mov.u32 	%r178, 31;
	mov.u32 	%r179, -1;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f378, %r165, %r178, %r179;  @p add.f32 r0, r0, %f378;  mov.f32 %f376, r0;}
	// end inline asm
	mov.u32 	%r168, 2;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f376, %r168, %r178, %r179;  @p add.f32 r0, r0, %f376;  mov.f32 %f379, r0;}
	// end inline asm
	mov.u32 	%r171, 4;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f379, %r171, %r178, %r179;  @p add.f32 r0, r0, %f379;  mov.f32 %f382, r0;}
	// end inline asm
	mov.u32 	%r174, 8;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f382, %r174, %r178, %r179;  @p add.f32 r0, r0, %f382;  mov.f32 %f385, r0;}
	// end inline asm
	mov.u32 	%r177, 16;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f385, %r177, %r178, %r179;  @p add.f32 r0, r0, %f385;  mov.f32 %f444, r0;}
	// end inline asm

$L__BB0_29:
	@%p8 bra 	$L__BB0_31;

	add.s32 	%r204, %r98, -112;
	st.shared.f32 	[%r204+3200], %f445;

$L__BB0_31:
	bar.sync 	0;
	@%p9 bra 	$L__BB0_33;

	ld.shared.f32 	%f411, [%r19+3216];
	add.ftz.f32 	%f412, %f445, %f411;
	ld.shared.f32 	%f413, [%r19+3220];
	add.ftz.f32 	%f414, %f412, %f413;
	ld.shared.f32 	%f415, [%r19+3224];
	add.ftz.f32 	%f398, %f414, %f415;
	mov.u32 	%r181, 1;
	mov.u32 	%r194, 31;
	mov.u32 	%r195, -1;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f398, %r181, %r194, %r195;  @p add.f32 r0, r0, %f398;  mov.f32 %f396, r0;}
	// end inline asm
	mov.u32 	%r184, 2;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f396, %r184, %r194, %r195;  @p add.f32 r0, r0, %f396;  mov.f32 %f399, r0;}
	// end inline asm
	mov.u32 	%r187, 4;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f399, %r187, %r194, %r195;  @p add.f32 r0, r0, %f399;  mov.f32 %f402, r0;}
	// end inline asm
	mov.u32 	%r190, 8;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f402, %r190, %r194, %r195;  @p add.f32 r0, r0, %f402;  mov.f32 %f405, r0;}
	// end inline asm
	mov.u32 	%r193, 16;
	// begin inline asm
	{  .reg .f32 r0;  .reg .pred p;  shfl.sync.down.b32 r0|p, %f405, %r193, %r194, %r195;  @p add.f32 r0, r0, %f405;  mov.f32 %f445, r0;}
	// end inline asm

$L__BB0_33:
	mov.u32 	%r205, %tid.y;
	or.b32  	%r196, %r4, %r205;
	setp.ne.s32 	%p20, %r196, 0;
	@%p20 bra 	$L__BB0_37;

	ld.param.u64 	%rd58, [_Z27dequant_gemv_group32_batch623DequantGemvKernelParams_param_0+8];
	ld.param.u64 	%rd57, [_Z27dequant_gemv_group32_batch623DequantGemvKernelParams_param_0];
	mov.u32 	%r206, %ctaid.x;
	cvta.to.global.u64 	%rd36, %rd57;
	setp.eq.s64 	%p21, %rd58, 0;
	mul.ftz.f32 	%f45, %f59, %f440;
	cvt.s64.s32 	%rd8, %r206;
	mul.wide.s32 	%rd37, %r206, 2;
	add.s64 	%rd9, %rd36, %rd37;
	mul.ftz.f32 	%f46, %f59, %f441;
	cvt.s64.s32 	%rd10, %r25;
	mul.wide.s32 	%rd38, %r25, 2;
	add.s64 	%rd11, %rd9, %rd38;
	mul.ftz.f32 	%f47, %f59, %f442;
	add.s32 	%r197, %r25, %r206;
	add.s32 	%r198, %r197, %r25;
	cvt.s64.s32 	%rd12, %r198;
	mul.wide.s32 	%rd39, %r198, 2;
	add.s64 	%rd14, %rd36, %rd39;
	mul.ftz.f32 	%f48, %f59, %f443;
	mul.ftz.f32 	%f49, %f59, %f444;
	mul.ftz.f32 	%f50, %f59, %f445;
	@%p21 bra 	$L__BB0_36;

	ld.param.u64 	%rd59, [_Z27dequant_gemv_group32_batch623DequantGemvKernelParams_param_0+8];
	cvta.to.global.u64 	%rd40, %rd59;
	shl.b64 	%rd41, %rd8, 1;
	add.s64 	%rd42, %rd40, %rd41;
	ld.global.u16 	%rs228, [%rd42];
	// begin inline asm
	{  cvt.f32.f16 %f416, %rs228;}

	// end inline asm
	fma.rn.ftz.f32 	%f417, %f60, %f416, %f45;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs229, %f417;}

	// end inline asm
	st.global.u16 	[%rd9], %rs229;
	shl.b64 	%rd43, %rd10, 1;
	add.s64 	%rd44, %rd42, %rd43;
	ld.global.u16 	%rs230, [%rd44];
	// begin inline asm
	{  cvt.f32.f16 %f418, %rs230;}

	// end inline asm
	fma.rn.ftz.f32 	%f419, %f60, %f418, %f46;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs231, %f419;}

	// end inline asm
	st.global.u16 	[%rd11], %rs231;
	shl.b64 	%rd45, %rd12, 1;
	add.s64 	%rd46, %rd40, %rd45;
	ld.global.u16 	%rs232, [%rd46];
	// begin inline asm
	{  cvt.f32.f16 %f420, %rs232;}

	// end inline asm
	fma.rn.ftz.f32 	%f421, %f60, %f420, %f47;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs233, %f421;}

	// end inline asm
	st.global.u16 	[%rd14], %rs233;
	add.s64 	%rd47, %rd46, %rd43;
	ld.global.u16 	%rs234, [%rd47];
	// begin inline asm
	{  cvt.f32.f16 %f422, %rs234;}

	// end inline asm
	fma.rn.ftz.f32 	%f423, %f60, %f422, %f48;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs235, %f423;}

	// end inline asm
	add.s64 	%rd48, %rd14, %rd43;
	st.global.u16 	[%rd48], %rs235;
	add.s64 	%rd49, %rd47, %rd43;
	ld.global.u16 	%rs236, [%rd49];
	// begin inline asm
	{  cvt.f32.f16 %f424, %rs236;}

	// end inline asm
	fma.rn.ftz.f32 	%f425, %f60, %f424, %f49;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs237, %f425;}

	// end inline asm
	add.s64 	%rd50, %rd48, %rd43;
	st.global.u16 	[%rd50], %rs237;
	add.s64 	%rd51, %rd49, %rd43;
	ld.global.u16 	%rs238, [%rd51];
	// begin inline asm
	{  cvt.f32.f16 %f426, %rs238;}

	// end inline asm
	fma.rn.ftz.f32 	%f427, %f60, %f426, %f50;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs239, %f427;}

	// end inline asm
	add.s64 	%rd52, %rd50, %rd43;
	st.global.u16 	[%rd52], %rs239;
	bra.uni 	$L__BB0_37;

$L__BB0_36:
	// begin inline asm
	{  cvt.rn.f16.f32 %rs240, %f45;}

	// end inline asm
	st.global.u16 	[%rd9], %rs240;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs241, %f46;}

	// end inline asm
	st.global.u16 	[%rd11], %rs241;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs242, %f47;}

	// end inline asm
	st.global.u16 	[%rd14], %rs242;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs243, %f48;}

	// end inline asm
	shl.b64 	%rd53, %rd10, 1;
	add.s64 	%rd54, %rd14, %rd53;
	st.global.u16 	[%rd54], %rs243;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs244, %f49;}

	// end inline asm
	add.s64 	%rd55, %rd54, %rd53;
	st.global.u16 	[%rd55], %rs244;
	// begin inline asm
	{  cvt.rn.f16.f32 %rs245, %f50;}

	// end inline asm
	add.s64 	%rd56, %rd55, %rd53;
	st.global.u16 	[%rd56], %rs245;

$L__BB0_37:
	ret;

}
	// .globl	_ZN3cub11EmptyKernelIvEEvv
.visible .entry _ZN3cub11EmptyKernelIvEEvv()
{



	ret;

}