//
// Generated by LLVM NVPTX Back-End
//

.version 8.8
.target sm_89
.address_size 64

	// .globl	quant_per_block_int8_kernel // -- Begin function quant_per_block_int8_kernel
.extern .shared .align 16 .b8 global_smem[];
                                        // @quant_per_block_int8_kernel
.visible .entry quant_per_block_int8_kernel(
	.param .u64 .ptr .global .align 1 quant_per_block_int8_kernel_param_0,
	.param .u64 .ptr .global .align 1 quant_per_block_int8_kernel_param_1,
	.param .u64 .ptr .global .align 1 quant_per_block_int8_kernel_param_2,
	.param .u32 quant_per_block_int8_kernel_param_3,
	.param .u32 quant_per_block_int8_kernel_param_4,
	.param .u32 quant_per_block_int8_kernel_param_5,
	.param .u32 quant_per_block_int8_kernel_param_6,
	.param .u32 quant_per_block_int8_kernel_param_7,
	.param .u32 quant_per_block_int8_kernel_param_8,
	.param .u32 quant_per_block_int8_kernel_param_9,
	.param .u32 quant_per_block_int8_kernel_param_10,
	.param .f32 quant_per_block_int8_kernel_param_11,
	.param .u64 .ptr .global .align 1 quant_per_block_int8_kernel_param_12,
	.param .u64 .ptr .global .align 1 quant_per_block_int8_kernel_param_13
)
.reqntid 128
{
	.reg .pred 	%p<70>;
	.reg .b16 	%rs<193>;
	.reg .b32 	%r<681>;
	.reg .b64 	%rd<39>;
	.loc	1 6 0                           // quant_per_block.py:6:0
$L__func_begin0:
	.loc	1 6 0                           // quant_per_block.py:6:0

// %bb.0:
	ld.param.b64 	%rd14, [quant_per_block_int8_kernel_param_0];
	ld.param.b64 	%rd15, [quant_per_block_int8_kernel_param_1];
$L__tmp0:
	.loc	1 12 15                         // quant_per_block.py:12:15
	mov.u32 	%r52, %ctaid.x;
	ld.param.b64 	%rd16, [quant_per_block_int8_kernel_param_2];
	.loc	1 13 13                         // quant_per_block.py:13:13
	mov.u32 	%r53, %ctaid.y;
	ld.param.b32 	%r54, [quant_per_block_int8_kernel_param_3];
	.loc	1 14 13                         // quant_per_block.py:14:13
	mov.u32 	%r55, %ctaid.z;
	.loc	1 16 14                         // quant_per_block.py:16:14
	shl.b32 	%r56, %r52, 7;
	ld.param.b32 	%r57, [quant_per_block_int8_kernel_param_4];
	ld.param.b32 	%r58, [quant_per_block_int8_kernel_param_5];
	.loc	1 16 30                         // quant_per_block.py:16:30
	mov.u32 	%r59, %tid.x;
	ld.param.b32 	%r60, [quant_per_block_int8_kernel_param_6];
	bfe.u32 	%r61, %r59, 2, 5;
	ld.param.b32 	%r62, [quant_per_block_int8_kernel_param_7];
	.loc	1 16 14                         // quant_per_block.py:16:14
	or.b32 	%r63, %r61, %r56;
	ld.param.b32 	%r64, [quant_per_block_int8_kernel_param_8];
	or.b32 	%r65, %r63, 32;
	ld.param.b32 	%r66, [quant_per_block_int8_kernel_param_9];
	or.b32 	%r67, %r63, 64;
	ld.param.b32 	%r68, [quant_per_block_int8_kernel_param_10];
	or.b32 	%r69, %r63, 96;
	ld.param.b32 	%r70, [quant_per_block_int8_kernel_param_11];
	.loc	1 19 26                         // quant_per_block.py:19:26
	mul.lo.s32 	%r71, %r57, %r55;
	.loc	1 19 18                         // quant_per_block.py:19:18
	mad.wide.s32 	%rd17, %r71, 2, %rd14;
	.loc	1 19 46                         // quant_per_block.py:19:46
	mul.lo.s32 	%r72, %r58, %r53;
	.loc	1 19 18                         // quant_per_block.py:19:18
	mad.wide.s32 	%rd18, %r72, 2, %rd17;
	.loc	1 19 66                         // quant_per_block.py:19:66
	mul.lo.s32 	%r73, %r60, %r63;
	mul.lo.s32 	%r74, %r60, %r65;
	mul.lo.s32 	%r75, %r60, %r67;
	mul.lo.s32 	%r76, %r60, %r69;
	.loc	1 19 18                         // quant_per_block.py:19:18
	mad.wide.s32 	%rd19, %r73, 2, %rd18;
	mad.wide.s32 	%rd20, %r74, 2, %rd18;
	mad.wide.s32 	%rd21, %r75, 2, %rd18;
	mad.wide.s32 	%rd22, %r76, 2, %rd18;
	.loc	1 19 96                         // quant_per_block.py:19:96
	and.b32 	%r77, %r59, 3;
	shl.b32 	%r78, %r77, 4;
	.loc	1 19 18                         // quant_per_block.py:19:18
	cvt.u64.u32 	%rd23, %r78;
	mul.wide.u32 	%rd24, %r78, 2;
	add.s64 	%rd1, %rd19, %rd24;
	add.s64 	%rd2, %rd1, 16;
	add.s64 	%rd3, %rd20, %rd24;
	add.s64 	%rd4, %rd3, 16;
	add.s64 	%rd5, %rd21, %rd24;
	add.s64 	%rd6, %rd5, 16;
	add.s64 	%rd7, %rd22, %rd24;
	add.s64 	%rd8, %rd7, 16;
	.loc	1 20 28                         // quant_per_block.py:20:28
	mul.lo.s32 	%r79, %r62, %r55;
	.loc	1 20 19                         // quant_per_block.py:20:19
	cvt.s64.s32 	%rd25, %r79;
	add.s64 	%rd26, %rd15, %rd25;
	.loc	1 20 48                         // quant_per_block.py:20:48
	mul.lo.s32 	%r80, %r64, %r53;
	.loc	1 20 19                         // quant_per_block.py:20:19
	cvt.s64.s32 	%rd27, %r80;
	add.s64 	%rd28, %rd26, %rd27;
	.loc	1 20 68                         // quant_per_block.py:20:68
	mul.lo.s32 	%r81, %r66, %r63;
	mul.lo.s32 	%r82, %r66, %r65;
	mul.lo.s32 	%r83, %r66, %r67;
	mul.lo.s32 	%r84, %r66, %r69;
	.loc	1 20 19                         // quant_per_block.py:20:19
	cvt.s64.s32 	%rd29, %r81;
	add.s64 	%rd30, %rd28, %rd29;
	cvt.s64.s32 	%rd31, %r82;
	add.s64 	%rd32, %rd28, %rd31;
	cvt.s64.s32 	%rd33, %r83;
	add.s64 	%rd34, %rd28, %rd33;
	cvt.s64.s32 	%rd35, %r84;
	add.s64 	%rd36, %rd28, %rd35;
	add.s64 	%rd9, %rd30, %rd23;
	add.s64 	%rd10, %rd32, %rd23;
	add.s64 	%rd11, %rd34, %rd23;
	add.s64 	%rd12, %rd36, %rd23;
	.loc	1 21 26                         // quant_per_block.py:21:26
	mul.lo.s32 	%r85, %r68, %r55;
	.loc	1 21 18                         // quant_per_block.py:21:18
	mad.wide.s32 	%rd37, %r85, 4, %rd16;
	mad.wide.u32 	%rd38, %r53, 4, %rd37;
	mad.wide.u32 	%rd13, %r52, 4, %rd38;
	.loc	1 23 34                         // quant_per_block.py:23:34
	setp.lt.s32 	%p1, %r63, %r54;
	setp.lt.s32 	%p2, %r65, %r54;
	setp.lt.s32 	%p3, %r67, %r54;
	setp.lt.s32 	%p4, %r69, %r54;
	.loc	1 23 9                          // quant_per_block.py:23:9
	// begin inline asm
	.loc	1 23 9                          // quant_per_block.py:23:9
	mov.u32 %r1, 0x0;
	.loc	1 24 9                          // quant_per_block.py:24:9
	mov.u32 %r2, 0x0;
	.loc	1 25 9                          // quant_per_block.py:25:9
	mov.u32 %r3, 0x0;
	.loc	1 26 9                          // quant_per_block.py:26:9
	mov.u32 %r4, 0x0;
	@%p1 ld.global.v4.b32 { %r1, %r2, %r3, %r4 }, [ %rd1 + 0 ];

	// end inline asm
	// begin inline asm
	.loc	1 23 9                          // quant_per_block.py:23:9
	mov.u32 %r5, 0x0;
	.loc	1 24 9                          // quant_per_block.py:24:9
	mov.u32 %r6, 0x0;
	.loc	1 25 9                          // quant_per_block.py:25:9
	mov.u32 %r7, 0x0;
	.loc	1 26 9                          // quant_per_block.py:26:9
	mov.u32 %r8, 0x0;
	@%p1 ld.global.v4.b32 { %r5, %r6, %r7, %r8 }, [ %rd2 + 0 ];

	// end inline asm
	// begin inline asm
	.loc	1 23 9                          // quant_per_block.py:23:9
	mov.u32 %r9, 0x0;
	.loc	1 24 9                          // quant_per_block.py:24:9
	mov.u32 %r10, 0x0;
	.loc	1 25 9                          // quant_per_block.py:25:9
	mov.u32 %r11, 0x0;
	.loc	1 26 9                          // quant_per_block.py:26:9
	mov.u32 %r12, 0x0;
	@%p2 ld.global.v4.b32 { %r9, %r10, %r11, %r12 }, [ %rd3 + 0 ];

	// end inline asm
	// begin inline asm
	.loc	1 23 9                          // quant_per_block.py:23:9
	mov.u32 %r13, 0x0;
	.loc	1 24 9                          // quant_per_block.py:24:9
	mov.u32 %r14, 0x0;
	.loc	1 25 9                          // quant_per_block.py:25:9
	mov.u32 %r15, 0x0;
	.loc	1 26 9                          // quant_per_block.py:26:9
	mov.u32 %r16, 0x0;
	@%p2 ld.global.v4.b32 { %r13, %r14, %r15, %r16 }, [ %rd4 + 0 ];

	// end inline asm
	// begin inline asm
	.loc	1 23 9                          // quant_per_block.py:23:9
	mov.u32 %r17, 0x0;
	.loc	1 24 9                          // quant_per_block.py:24:9
	mov.u32 %r18, 0x0;
	.loc	1 25 9                          // quant_per_block.py:25:9
	mov.u32 %r19, 0x0;
	.loc	1 26 9                          // quant_per_block.py:26:9
	mov.u32 %r20, 0x0;
	@%p3 ld.global.v4.b32 { %r17, %r18, %r19, %r20 }, [ %rd5 + 0 ];

	// end inline asm
	// begin inline asm
	.loc	1 23 9                          // quant_per_block.py:23:9
	mov.u32 %r21, 0x0;
	.loc	1 24 9                          // quant_per_block.py:24:9
	mov.u32 %r22, 0x0;
	.loc	1 25 9                          // quant_per_block.py:25:9
	mov.u32 %r23, 0x0;
	.loc	1 26 9                          // quant_per_block.py:26:9
	mov.u32 %r24, 0x0;
	@%p3 ld.global.v4.b32 { %r21, %r22, %r23, %r24 }, [ %rd6 + 0 ];

	// end inline asm
	// begin inline asm
	.loc	1 23 9                          // quant_per_block.py:23:9
	mov.u32 %r25, 0x0;
	.loc	1 24 9                          // quant_per_block.py:24:9
	mov.u32 %r26, 0x0;
	.loc	1 25 9                          // quant_per_block.py:25:9
	mov.u32 %r27, 0x0;
	.loc	1 26 9                          // quant_per_block.py:26:9
	mov.u32 %r28, 0x0;
	@%p4 ld.global.v4.b32 { %r25, %r26, %r27, %r28 }, [ %rd7 + 0 ];

	// end inline asm
	// begin inline asm
	.loc	1 23 9                          // quant_per_block.py:23:9
	mov.u32 %r29, 0x0;
	.loc	1 24 9                          // quant_per_block.py:24:9
	mov.u32 %r30, 0x0;
	.loc	1 25 9                          // quant_per_block.py:25:9
	mov.u32 %r31, 0x0;
	.loc	1 26 9                          // quant_per_block.py:26:9
	mov.u32 %r32, 0x0;
	@%p4 ld.global.v4.b32 { %r29, %r30, %r31, %r32 }, [ %rd8 + 0 ];

	// end inline asm
	.loc	1 24 9                          // quant_per_block.py:24:9
	mov.b32 	{%rs1, %rs2}, %r32;
	cvt.f32.bf16 	%r86, %rs2;
	cvt.f32.bf16 	%r87, %rs1;
	mov.b32 	{%rs3, %rs4}, %r31;
	cvt.f32.bf16 	%r88, %rs4;
	cvt.f32.bf16 	%r89, %rs3;
	mov.b32 	{%rs5, %rs6}, %r30;
	cvt.f32.bf16 	%r90, %rs6;
	cvt.f32.bf16 	%r91, %rs5;
	mov.b32 	{%rs7, %rs8}, %r29;
	cvt.f32.bf16 	%r92, %rs8;
	cvt.f32.bf16 	%r93, %rs7;
	mov.b32 	{%rs9, %rs10}, %r28;
	cvt.f32.bf16 	%r94, %rs10;
	cvt.f32.bf16 	%r95, %rs9;
	mov.b32 	{%rs11, %rs12}, %r27;
	cvt.f32.bf16 	%r96, %rs12;
	cvt.f32.bf16 	%r97, %rs11;
	mov.b32 	{%rs13, %rs14}, %r26;
	cvt.f32.bf16 	%r98, %rs14;
	cvt.f32.bf16 	%r99, %rs13;
	mov.b32 	{%rs15, %rs16}, %r25;
	cvt.f32.bf16 	%r100, %rs16;
	cvt.f32.bf16 	%r101, %rs15;
	mov.b32 	{%rs17, %rs18}, %r24;
	cvt.f32.bf16 	%r102, %rs18;
	cvt.f32.bf16 	%r103, %rs17;
	mov.b32 	{%rs19, %rs20}, %r23;
	cvt.f32.bf16 	%r104, %rs20;
	cvt.f32.bf16 	%r105, %rs19;
	mov.b32 	{%rs21, %rs22}, %r22;
	cvt.f32.bf16 	%r106, %rs22;
	cvt.f32.bf16 	%r107, %rs21;
	mov.b32 	{%rs23, %rs24}, %r21;
	cvt.f32.bf16 	%r108, %rs24;
	cvt.f32.bf16 	%r109, %rs23;
	mov.b32 	{%rs25, %rs26}, %r20;
	cvt.f32.bf16 	%r110, %rs26;
	cvt.f32.bf16 	%r111, %rs25;
	mov.b32 	{%rs27, %rs28}, %r19;
	cvt.f32.bf16 	%r112, %rs28;
	cvt.f32.bf16 	%r113, %rs27;
	mov.b32 	{%rs29, %rs30}, %r18;
	cvt.f32.bf16 	%r114, %rs30;
	cvt.f32.bf16 	%r115, %rs29;
	mov.b32 	{%rs31, %rs32}, %r17;
	cvt.f32.bf16 	%r116, %rs32;
	cvt.f32.bf16 	%r117, %rs31;
	mov.b32 	{%rs33, %rs34}, %r16;
	cvt.f32.bf16 	%r118, %rs34;
	cvt.f32.bf16 	%r119, %rs33;
	mov.b32 	{%rs35, %rs36}, %r15;
	cvt.f32.bf16 	%r120, %rs36;
	cvt.f32.bf16 	%r121, %rs35;
	mov.b32 	{%rs37, %rs38}, %r14;
	cvt.f32.bf16 	%r122, %rs38;
	cvt.f32.bf16 	%r123, %rs37;
	mov.b32 	{%rs39, %rs40}, %r13;
	cvt.f32.bf16 	%r124, %rs40;
	cvt.f32.bf16 	%r125, %rs39;
	mov.b32 	{%rs41, %rs42}, %r12;
	cvt.f32.bf16 	%r126, %rs42;
	cvt.f32.bf16 	%r127, %rs41;
	mov.b32 	{%rs43, %rs44}, %r11;
	cvt.f32.bf16 	%r128, %rs44;
	cvt.f32.bf16 	%r129, %rs43;
	mov.b32 	{%rs45, %rs46}, %r10;
	cvt.f32.bf16 	%r130, %rs46;
	cvt.f32.bf16 	%r131, %rs45;
	mov.b32 	{%rs47, %rs48}, %r9;
	cvt.f32.bf16 	%r132, %rs48;
	cvt.f32.bf16 	%r133, %rs47;
	mov.b32 	{%rs49, %rs50}, %r8;
	cvt.f32.bf16 	%r134, %rs50;
	cvt.f32.bf16 	%r135, %rs49;
	mov.b32 	{%rs51, %rs52}, %r7;
	cvt.f32.bf16 	%r136, %rs52;
	cvt.f32.bf16 	%r137, %rs51;
	mov.b32 	{%rs53, %rs54}, %r6;
	cvt.f32.bf16 	%r138, %rs54;
	cvt.f32.bf16 	%r139, %rs53;
	mov.b32 	{%rs55, %rs56}, %r5;
	cvt.f32.bf16 	%r140, %rs56;
	cvt.f32.bf16 	%r141, %rs55;
	mov.b32 	{%rs57, %rs58}, %r4;
	cvt.f32.bf16 	%r142, %rs58;
	cvt.f32.bf16 	%r143, %rs57;
	mov.b32 	{%rs59, %rs60}, %r3;
	cvt.f32.bf16 	%r144, %rs60;
	cvt.f32.bf16 	%r145, %rs59;
	mov.b32 	{%rs61, %rs62}, %r2;
	cvt.f32.bf16 	%r146, %rs62;
	cvt.f32.bf16 	%r147, %rs61;
	mov.b32 	{%rs63, %rs64}, %r1;
	cvt.f32.bf16 	%r148, %rs64;
	cvt.f32.bf16 	%r149, %rs63;
	.loc	1 25 5                          // quant_per_block.py:25:5
	mul.f32 	%r150, %r70, %r149;
	mul.f32 	%r151, %r70, %r148;
	mul.f32 	%r152, %r70, %r147;
	mul.f32 	%r153, %r70, %r146;
	mul.f32 	%r154, %r70, %r145;
	mul.f32 	%r155, %r70, %r144;
	mul.f32 	%r156, %r70, %r143;
	mul.f32 	%r157, %r70, %r142;
	mul.f32 	%r158, %r70, %r141;
	mul.f32 	%r159, %r70, %r140;
	mul.f32 	%r160, %r70, %r139;
	mul.f32 	%r161, %r70, %r138;
	mul.f32 	%r162, %r70, %r137;
	mul.f32 	%r163, %r70, %r136;
	mul.f32 	%r164, %r70, %r135;
	mul.f32 	%r165, %r70, %r134;
	mul.f32 	%r166, %r70, %r133;
	mul.f32 	%r167, %r70, %r132;
	mul.f32 	%r168, %r70, %r131;
	mul.f32 	%r169, %r70, %r130;
	mul.f32 	%r170, %r70, %r129;
	mul.f32 	%r171, %r70, %r128;
	mul.f32 	%r172, %r70, %r127;
	mul.f32 	%r173, %r70, %r126;
	mul.f32 	%r174, %r70, %r125;
	mul.f32 	%r175, %r70, %r124;
	mul.f32 	%r176, %r70, %r123;
	mul.f32 	%r177, %r70, %r122;
	mul.f32 	%r178, %r70, %r121;
	mul.f32 	%r179, %r70, %r120;
	mul.f32 	%r180, %r70, %r119;
	mul.f32 	%r181, %r70, %r118;
	mul.f32 	%r182, %r70, %r117;
	mul.f32 	%r183, %r70, %r116;
	mul.f32 	%r184, %r70, %r115;
	mul.f32 	%r185, %r70, %r114;
	mul.f32 	%r186, %r70, %r113;
	mul.f32 	%r187, %r70, %r112;
	mul.f32 	%r188, %r70, %r111;
	mul.f32 	%r189, %r70, %r110;
	mul.f32 	%r190, %r70, %r109;
	mul.f32 	%r191, %r70, %r108;
	mul.f32 	%r192, %r70, %r107;
	mul.f32 	%r193, %r70, %r106;
	mul.f32 	%r194, %r70, %r105;
	mul.f32 	%r195, %r70, %r104;
	mul.f32 	%r196, %r70, %r103;
	mul.f32 	%r197, %r70, %r102;
	mul.f32 	%r198, %r70, %r101;
	mul.f32 	%r199, %r70, %r100;
	mul.f32 	%r200, %r70, %r99;
	mul.f32 	%r201, %r70, %r98;
	mul.f32 	%r202, %r70, %r97;
	mul.f32 	%r203, %r70, %r96;
	mul.f32 	%r204, %r70, %r95;
	mul.f32 	%r205, %r70, %r94;
	mul.f32 	%r206, %r70, %r93;
	mul.f32 	%r207, %r70, %r92;
	mul.f32 	%r208, %r70, %r91;
	mul.f32 	%r209, %r70, %r90;
	mul.f32 	%r210, %r70, %r89;
	mul.f32 	%r211, %r70, %r88;
	mul.f32 	%r212, %r70, %r87;
	mul.f32 	%r213, %r70, %r86;
	.loc	1 26 20                         // quant_per_block.py:26:20
	abs.f32 	%r214, %r213;
	abs.f32 	%r215, %r212;
	abs.f32 	%r216, %r211;
	abs.f32 	%r217, %r210;
	abs.f32 	%r218, %r209;
	abs.f32 	%r219, %r208;
	abs.f32 	%r220, %r207;
	abs.f32 	%r221, %r206;
	abs.f32 	%r222, %r205;
	abs.f32 	%r223, %r204;
	abs.f32 	%r224, %r203;
	abs.f32 	%r225, %r202;
	abs.f32 	%r226, %r201;
	abs.f32 	%r227, %r200;
	abs.f32 	%r228, %r199;
	abs.f32 	%r229, %r198;
	abs.f32 	%r230, %r197;
	abs.f32 	%r231, %r196;
	abs.f32 	%r232, %r195;
	abs.f32 	%r233, %r194;
	abs.f32 	%r234, %r193;
	abs.f32 	%r235, %r192;
	abs.f32 	%r236, %r191;
	abs.f32 	%r237, %r190;
	abs.f32 	%r238, %r189;
	abs.f32 	%r239, %r188;
	abs.f32 	%r240, %r187;
	abs.f32 	%r241, %r186;
	abs.f32 	%r242, %r185;
	abs.f32 	%r243, %r184;
	abs.f32 	%r244, %r183;
	abs.f32 	%r245, %r182;
	abs.f32 	%r246, %r181;
	abs.f32 	%r247, %r180;
	abs.f32 	%r248, %r179;
	abs.f32 	%r249, %r178;
	abs.f32 	%r250, %r177;
	abs.f32 	%r251, %r176;
	abs.f32 	%r252, %r175;
	abs.f32 	%r253, %r174;
	abs.f32 	%r254, %r173;
	abs.f32 	%r255, %r172;
	abs.f32 	%r256, %r171;
	abs.f32 	%r257, %r170;
	abs.f32 	%r258, %r169;
	abs.f32 	%r259, %r168;
	abs.f32 	%r260, %r167;
	abs.f32 	%r261, %r166;
	abs.f32 	%r262, %r165;
	abs.f32 	%r263, %r164;
	abs.f32 	%r264, %r163;
	abs.f32 	%r265, %r162;
	abs.f32 	%r266, %r161;
	abs.f32 	%r267, %r160;
	abs.f32 	%r268, %r159;
	abs.f32 	%r269, %r158;
	abs.f32 	%r270, %r157;
	abs.f32 	%r271, %r156;
	abs.f32 	%r272, %r155;
	abs.f32 	%r273, %r154;
	abs.f32 	%r274, %r153;
	abs.f32 	%r275, %r152;
	abs.f32 	%r276, %r151;
	abs.f32 	%r277, %r150;
$L__tmp1:
	.loc	1 26 13                         // quant_per_block.py:26:13
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 26 13 // standard.py:191:16
	.loc	2 170 12, function_name $L__info_string0, inlined_at 2 191 16 // standard.py:170:12
	max.f32 	%r278, %r277, %r276;
	max.f32 	%r279, %r275, %r274;
	max.f32 	%r280, %r273, %r272;
	max.f32 	%r281, %r271, %r270;
	max.f32 	%r282, %r269, %r268;
	max.f32 	%r283, %r267, %r266;
	max.f32 	%r284, %r265, %r264;
	max.f32 	%r285, %r263, %r262;
	max.f32 	%r286, %r261, %r260;
	max.f32 	%r287, %r259, %r258;
	max.f32 	%r288, %r257, %r256;
	max.f32 	%r289, %r255, %r254;
	max.f32 	%r290, %r253, %r252;
	max.f32 	%r291, %r251, %r250;
	max.f32 	%r292, %r249, %r248;
	max.f32 	%r293, %r247, %r246;
	max.f32 	%r294, %r245, %r244;
	max.f32 	%r295, %r243, %r242;
	max.f32 	%r296, %r241, %r240;
	max.f32 	%r297, %r239, %r238;
	max.f32 	%r298, %r237, %r236;
	max.f32 	%r299, %r235, %r234;
	max.f32 	%r300, %r233, %r232;
	max.f32 	%r301, %r231, %r230;
	max.f32 	%r302, %r229, %r228;
	max.f32 	%r303, %r227, %r226;
	max.f32 	%r304, %r225, %r224;
	max.f32 	%r305, %r223, %r222;
	max.f32 	%r306, %r221, %r220;
	max.f32 	%r307, %r219, %r218;
	max.f32 	%r308, %r217, %r216;
	max.f32 	%r309, %r215, %r214;
	max.f32 	%r310, %r278, %r279;
	max.f32 	%r311, %r280, %r281;
	max.f32 	%r312, %r282, %r283;
	max.f32 	%r313, %r284, %r285;
	max.f32 	%r314, %r286, %r287;
	max.f32 	%r315, %r288, %r289;
	max.f32 	%r316, %r290, %r291;
	max.f32 	%r317, %r292, %r293;
	max.f32 	%r318, %r294, %r295;
	max.f32 	%r319, %r296, %r297;
	max.f32 	%r320, %r298, %r299;
	max.f32 	%r321, %r300, %r301;
	max.f32 	%r322, %r302, %r303;
	max.f32 	%r323, %r304, %r305;
	max.f32 	%r324, %r306, %r307;
	max.f32 	%r325, %r308, %r309;
	max.f32 	%r326, %r310, %r311;
	max.f32 	%r327, %r312, %r313;
	max.f32 	%r328, %r314, %r315;
	max.f32 	%r329, %r316, %r317;
	max.f32 	%r330, %r318, %r319;
	max.f32 	%r331, %r320, %r321;
	max.f32 	%r332, %r322, %r323;
	max.f32 	%r333, %r324, %r325;
	max.f32 	%r334, %r326, %r327;
	max.f32 	%r335, %r328, %r329;
	max.f32 	%r336, %r330, %r331;
	max.f32 	%r337, %r332, %r333;
	max.f32 	%r338, %r334, %r335;
	max.f32 	%r339, %r336, %r337;
	max.f32 	%r340, %r338, %r339;
$L__tmp2:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 26 13 // standard.py:191:16
	shfl.sync.bfly.b32 	%r341, %r340, 16, 31, -1;
$L__tmp3:
	.loc	2 170 12, function_name $L__info_string0, inlined_at 2 191 16 // standard.py:170:12
	max.f32 	%r342, %r340, %r341;
$L__tmp4:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 26 13 // standard.py:191:16
	shfl.sync.bfly.b32 	%r343, %r342, 8, 31, -1;
$L__tmp5:
	.loc	2 170 12, function_name $L__info_string0, inlined_at 2 191 16 // standard.py:170:12
	max.f32 	%r344, %r342, %r343;
$L__tmp6:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 26 13 // standard.py:191:16
	shfl.sync.bfly.b32 	%r345, %r344, 4, 31, -1;
$L__tmp7:
	.loc	2 170 12, function_name $L__info_string0, inlined_at 2 191 16 // standard.py:170:12
	max.f32 	%r346, %r344, %r345;
$L__tmp8:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 26 13 // standard.py:191:16
	shfl.sync.bfly.b32 	%r347, %r346, 2, 31, -1;
$L__tmp9:
	.loc	2 170 12, function_name $L__info_string0, inlined_at 2 191 16 // standard.py:170:12
	max.f32 	%r348, %r346, %r347;
$L__tmp10:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 26 13 // standard.py:191:16
	shfl.sync.bfly.b32 	%r349, %r348, 1, 31, -1;
$L__tmp11:
	.loc	2 170 12, function_name $L__info_string0, inlined_at 2 191 16 // standard.py:170:12
	max.f32 	%r34, %r348, %r349;
$L__tmp12:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 26 13 // standard.py:191:16
	shr.u32 	%r350, %r59, 3;
	and.b32 	%r351, %r350, 12;
	mov.b32 	%r352, global_smem;
	add.s32 	%r33, %r352, %r351;
	// begin inline asm
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 26 13 // standard.py:191:16
	st.shared::cta.b32 [ %r33 + 0 ], %r34;

	// end inline asm
	bar.sync 	0;
	shl.b32 	%r353, %r77, 2;
	add.s32 	%r354, %r352, %r353;
	ld.shared.b32 	%r355, [%r354];
	shfl.sync.bfly.b32 	%r356, %r355, 2, 31, -1;
$L__tmp13:
	.loc	2 170 12, function_name $L__info_string0, inlined_at 2 191 16 // standard.py:170:12
	max.f32 	%r357, %r355, %r356;
$L__tmp14:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 26 13 // standard.py:191:16
	shfl.sync.bfly.b32 	%r358, %r357, 1, 31, -1;
$L__tmp15:
	.loc	2 170 12, function_name $L__info_string0, inlined_at 2 191 16 // standard.py:170:12
	max.f32 	%r359, %r357, %r358;
	mov.b32 	%r360, 0f42FE0000;
$L__tmp16:
	.loc	1 26 13                         // quant_per_block.py:26:13
	div.full.f32 	%r51, %r359, %r360;
	.loc	1 27 14                         // quant_per_block.py:27:14
	div.full.f32 	%r361, %r150, %r51;
	div.full.f32 	%r362, %r151, %r51;
	div.full.f32 	%r363, %r152, %r51;
	div.full.f32 	%r364, %r153, %r51;
	div.full.f32 	%r365, %r154, %r51;
	div.full.f32 	%r366, %r155, %r51;
	div.full.f32 	%r367, %r156, %r51;
	div.full.f32 	%r368, %r157, %r51;
	div.full.f32 	%r369, %r158, %r51;
	div.full.f32 	%r370, %r159, %r51;
	div.full.f32 	%r371, %r160, %r51;
	div.full.f32 	%r372, %r161, %r51;
	div.full.f32 	%r373, %r162, %r51;
	div.full.f32 	%r374, %r163, %r51;
	div.full.f32 	%r375, %r164, %r51;
	div.full.f32 	%r376, %r165, %r51;
	div.full.f32 	%r377, %r166, %r51;
	div.full.f32 	%r378, %r167, %r51;
	div.full.f32 	%r379, %r168, %r51;
	div.full.f32 	%r380, %r169, %r51;
	div.full.f32 	%r381, %r170, %r51;
	div.full.f32 	%r382, %r171, %r51;
	div.full.f32 	%r383, %r172, %r51;
	div.full.f32 	%r384, %r173, %r51;
	div.full.f32 	%r385, %r174, %r51;
	div.full.f32 	%r386, %r175, %r51;
	div.full.f32 	%r387, %r176, %r51;
	div.full.f32 	%r388, %r177, %r51;
	div.full.f32 	%r389, %r178, %r51;
	div.full.f32 	%r390, %r179, %r51;
	div.full.f32 	%r391, %r180, %r51;
	div.full.f32 	%r392, %r181, %r51;
	div.full.f32 	%r393, %r182, %r51;
	div.full.f32 	%r394, %r183, %r51;
	div.full.f32 	%r395, %r184, %r51;
	div.full.f32 	%r396, %r185, %r51;
	div.full.f32 	%r397, %r186, %r51;
	div.full.f32 	%r398, %r187, %r51;
	div.full.f32 	%r399, %r188, %r51;
	div.full.f32 	%r400, %r189, %r51;
	div.full.f32 	%r401, %r190, %r51;
	div.full.f32 	%r402, %r191, %r51;
	div.full.f32 	%r403, %r192, %r51;
	div.full.f32 	%r404, %r193, %r51;
	div.full.f32 	%r405, %r194, %r51;
	div.full.f32 	%r406, %r195, %r51;
	div.full.f32 	%r407, %r196, %r51;
	div.full.f32 	%r408, %r197, %r51;
	div.full.f32 	%r409, %r198, %r51;
	div.full.f32 	%r410, %r199, %r51;
	div.full.f32 	%r411, %r200, %r51;
	div.full.f32 	%r412, %r201, %r51;
	div.full.f32 	%r413, %r202, %r51;
	div.full.f32 	%r414, %r203, %r51;
	div.full.f32 	%r415, %r204, %r51;
	div.full.f32 	%r416, %r205, %r51;
	div.full.f32 	%r417, %r206, %r51;
	div.full.f32 	%r418, %r207, %r51;
	div.full.f32 	%r419, %r208, %r51;
	div.full.f32 	%r420, %r209, %r51;
	div.full.f32 	%r421, %r210, %r51;
	div.full.f32 	%r422, %r211, %r51;
	div.full.f32 	%r423, %r212, %r51;
	div.full.f32 	%r424, %r213, %r51;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p6, %r361, 0f00000000;
	setp.ge.f32 	%p7, %r362, 0f00000000;
	setp.ge.f32 	%p8, %r363, 0f00000000;
	setp.ge.f32 	%p9, %r364, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r425, 0f3F000000, 0fBF000000, %p9;
	selp.f32 	%r426, 0f3F000000, 0fBF000000, %p8;
	selp.f32 	%r427, 0f3F000000, 0fBF000000, %p7;
	selp.f32 	%r428, 0f3F000000, 0fBF000000, %p6;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r429, %r361, %r428;
	add.f32 	%r430, %r362, %r427;
	add.f32 	%r431, %r363, %r426;
	add.f32 	%r432, %r364, %r425;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs65, %r432;
	cvt.rzi.s16.f32 	%rs66, %r431;
	mov.b32 	%r433, {%rs66, %rs65};
	mov.b32 	{%rs67, %rs68}, %r433;
	cvt.u32.u16 	%r434, %rs68;
	cvt.u32.u16 	%r435, %rs67;
	prmt.b32 	%r436, %r435, %r434, 0x3340U;
	cvt.rzi.s16.f32 	%rs69, %r430;
	cvt.rzi.s16.f32 	%rs70, %r429;
	mov.b32 	%r437, {%rs70, %rs69};
	mov.b32 	{%rs71, %rs72}, %r437;
	cvt.u32.u16 	%r438, %rs72;
	cvt.u32.u16 	%r439, %rs71;
	prmt.b32 	%r440, %r439, %r438, 0x3340U;
	prmt.b32 	%r35, %r440, %r436, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p10, %r365, 0f00000000;
	setp.ge.f32 	%p11, %r366, 0f00000000;
	setp.ge.f32 	%p12, %r367, 0f00000000;
	setp.ge.f32 	%p13, %r368, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r441, 0f3F000000, 0fBF000000, %p13;
	selp.f32 	%r442, 0f3F000000, 0fBF000000, %p12;
	selp.f32 	%r443, 0f3F000000, 0fBF000000, %p11;
	selp.f32 	%r444, 0f3F000000, 0fBF000000, %p10;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r445, %r365, %r444;
	add.f32 	%r446, %r366, %r443;
	add.f32 	%r447, %r367, %r442;
	add.f32 	%r448, %r368, %r441;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs73, %r448;
	cvt.rzi.s16.f32 	%rs74, %r447;
	mov.b32 	%r449, {%rs74, %rs73};
	mov.b32 	{%rs75, %rs76}, %r449;
	cvt.u32.u16 	%r450, %rs76;
	cvt.u32.u16 	%r451, %rs75;
	prmt.b32 	%r452, %r451, %r450, 0x3340U;
	cvt.rzi.s16.f32 	%rs77, %r446;
	cvt.rzi.s16.f32 	%rs78, %r445;
	mov.b32 	%r453, {%rs78, %rs77};
	mov.b32 	{%rs79, %rs80}, %r453;
	cvt.u32.u16 	%r454, %rs80;
	cvt.u32.u16 	%r455, %rs79;
	prmt.b32 	%r456, %r455, %r454, 0x3340U;
	prmt.b32 	%r36, %r456, %r452, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p14, %r369, 0f00000000;
	setp.ge.f32 	%p15, %r370, 0f00000000;
	setp.ge.f32 	%p16, %r371, 0f00000000;
	setp.ge.f32 	%p17, %r372, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r457, 0f3F000000, 0fBF000000, %p17;
	selp.f32 	%r458, 0f3F000000, 0fBF000000, %p16;
	selp.f32 	%r459, 0f3F000000, 0fBF000000, %p15;
	selp.f32 	%r460, 0f3F000000, 0fBF000000, %p14;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r461, %r369, %r460;
	add.f32 	%r462, %r370, %r459;
	add.f32 	%r463, %r371, %r458;
	add.f32 	%r464, %r372, %r457;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs81, %r464;
	cvt.rzi.s16.f32 	%rs82, %r463;
	mov.b32 	%r465, {%rs82, %rs81};
	mov.b32 	{%rs83, %rs84}, %r465;
	cvt.u32.u16 	%r466, %rs84;
	cvt.u32.u16 	%r467, %rs83;
	prmt.b32 	%r468, %r467, %r466, 0x3340U;
	cvt.rzi.s16.f32 	%rs85, %r462;
	cvt.rzi.s16.f32 	%rs86, %r461;
	mov.b32 	%r469, {%rs86, %rs85};
	mov.b32 	{%rs87, %rs88}, %r469;
	cvt.u32.u16 	%r470, %rs88;
	cvt.u32.u16 	%r471, %rs87;
	prmt.b32 	%r472, %r471, %r470, 0x3340U;
	prmt.b32 	%r37, %r472, %r468, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p18, %r373, 0f00000000;
	setp.ge.f32 	%p19, %r374, 0f00000000;
	setp.ge.f32 	%p20, %r375, 0f00000000;
	setp.ge.f32 	%p21, %r376, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r473, 0f3F000000, 0fBF000000, %p21;
	selp.f32 	%r474, 0f3F000000, 0fBF000000, %p20;
	selp.f32 	%r475, 0f3F000000, 0fBF000000, %p19;
	selp.f32 	%r476, 0f3F000000, 0fBF000000, %p18;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r477, %r373, %r476;
	add.f32 	%r478, %r374, %r475;
	add.f32 	%r479, %r375, %r474;
	add.f32 	%r480, %r376, %r473;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs89, %r480;
	cvt.rzi.s16.f32 	%rs90, %r479;
	mov.b32 	%r481, {%rs90, %rs89};
	mov.b32 	{%rs91, %rs92}, %r481;
	cvt.u32.u16 	%r482, %rs92;
	cvt.u32.u16 	%r483, %rs91;
	prmt.b32 	%r484, %r483, %r482, 0x3340U;
	cvt.rzi.s16.f32 	%rs93, %r478;
	cvt.rzi.s16.f32 	%rs94, %r477;
	mov.b32 	%r485, {%rs94, %rs93};
	mov.b32 	{%rs95, %rs96}, %r485;
	cvt.u32.u16 	%r486, %rs96;
	cvt.u32.u16 	%r487, %rs95;
	prmt.b32 	%r488, %r487, %r486, 0x3340U;
	prmt.b32 	%r38, %r488, %r484, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p22, %r377, 0f00000000;
	setp.ge.f32 	%p23, %r378, 0f00000000;
	setp.ge.f32 	%p24, %r379, 0f00000000;
	setp.ge.f32 	%p25, %r380, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r489, 0f3F000000, 0fBF000000, %p25;
	selp.f32 	%r490, 0f3F000000, 0fBF000000, %p24;
	selp.f32 	%r491, 0f3F000000, 0fBF000000, %p23;
	selp.f32 	%r492, 0f3F000000, 0fBF000000, %p22;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r493, %r377, %r492;
	add.f32 	%r494, %r378, %r491;
	add.f32 	%r495, %r379, %r490;
	add.f32 	%r496, %r380, %r489;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs97, %r496;
	cvt.rzi.s16.f32 	%rs98, %r495;
	mov.b32 	%r497, {%rs98, %rs97};
	mov.b32 	{%rs99, %rs100}, %r497;
	cvt.u32.u16 	%r498, %rs100;
	cvt.u32.u16 	%r499, %rs99;
	prmt.b32 	%r500, %r499, %r498, 0x3340U;
	cvt.rzi.s16.f32 	%rs101, %r494;
	cvt.rzi.s16.f32 	%rs102, %r493;
	mov.b32 	%r501, {%rs102, %rs101};
	mov.b32 	{%rs103, %rs104}, %r501;
	cvt.u32.u16 	%r502, %rs104;
	cvt.u32.u16 	%r503, %rs103;
	prmt.b32 	%r504, %r503, %r502, 0x3340U;
	prmt.b32 	%r39, %r504, %r500, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p26, %r381, 0f00000000;
	setp.ge.f32 	%p27, %r382, 0f00000000;
	setp.ge.f32 	%p28, %r383, 0f00000000;
	setp.ge.f32 	%p29, %r384, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r505, 0f3F000000, 0fBF000000, %p29;
	selp.f32 	%r506, 0f3F000000, 0fBF000000, %p28;
	selp.f32 	%r507, 0f3F000000, 0fBF000000, %p27;
	selp.f32 	%r508, 0f3F000000, 0fBF000000, %p26;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r509, %r381, %r508;
	add.f32 	%r510, %r382, %r507;
	add.f32 	%r511, %r383, %r506;
	add.f32 	%r512, %r384, %r505;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs105, %r512;
	cvt.rzi.s16.f32 	%rs106, %r511;
	mov.b32 	%r513, {%rs106, %rs105};
	mov.b32 	{%rs107, %rs108}, %r513;
	cvt.u32.u16 	%r514, %rs108;
	cvt.u32.u16 	%r515, %rs107;
	prmt.b32 	%r516, %r515, %r514, 0x3340U;
	cvt.rzi.s16.f32 	%rs109, %r510;
	cvt.rzi.s16.f32 	%rs110, %r509;
	mov.b32 	%r517, {%rs110, %rs109};
	mov.b32 	{%rs111, %rs112}, %r517;
	cvt.u32.u16 	%r518, %rs112;
	cvt.u32.u16 	%r519, %rs111;
	prmt.b32 	%r520, %r519, %r518, 0x3340U;
	prmt.b32 	%r40, %r520, %r516, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p30, %r385, 0f00000000;
	setp.ge.f32 	%p31, %r386, 0f00000000;
	setp.ge.f32 	%p32, %r387, 0f00000000;
	setp.ge.f32 	%p33, %r388, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r521, 0f3F000000, 0fBF000000, %p33;
	selp.f32 	%r522, 0f3F000000, 0fBF000000, %p32;
	selp.f32 	%r523, 0f3F000000, 0fBF000000, %p31;
	selp.f32 	%r524, 0f3F000000, 0fBF000000, %p30;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r525, %r385, %r524;
	add.f32 	%r526, %r386, %r523;
	add.f32 	%r527, %r387, %r522;
	add.f32 	%r528, %r388, %r521;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs113, %r528;
	cvt.rzi.s16.f32 	%rs114, %r527;
	mov.b32 	%r529, {%rs114, %rs113};
	mov.b32 	{%rs115, %rs116}, %r529;
	cvt.u32.u16 	%r530, %rs116;
	cvt.u32.u16 	%r531, %rs115;
	prmt.b32 	%r532, %r531, %r530, 0x3340U;
	cvt.rzi.s16.f32 	%rs117, %r526;
	cvt.rzi.s16.f32 	%rs118, %r525;
	mov.b32 	%r533, {%rs118, %rs117};
	mov.b32 	{%rs119, %rs120}, %r533;
	cvt.u32.u16 	%r534, %rs120;
	cvt.u32.u16 	%r535, %rs119;
	prmt.b32 	%r536, %r535, %r534, 0x3340U;
	prmt.b32 	%r41, %r536, %r532, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p34, %r389, 0f00000000;
	setp.ge.f32 	%p35, %r390, 0f00000000;
	setp.ge.f32 	%p36, %r391, 0f00000000;
	setp.ge.f32 	%p37, %r392, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r537, 0f3F000000, 0fBF000000, %p37;
	selp.f32 	%r538, 0f3F000000, 0fBF000000, %p36;
	selp.f32 	%r539, 0f3F000000, 0fBF000000, %p35;
	selp.f32 	%r540, 0f3F000000, 0fBF000000, %p34;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r541, %r389, %r540;
	add.f32 	%r542, %r390, %r539;
	add.f32 	%r543, %r391, %r538;
	add.f32 	%r544, %r392, %r537;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs121, %r544;
	cvt.rzi.s16.f32 	%rs122, %r543;
	mov.b32 	%r545, {%rs122, %rs121};
	mov.b32 	{%rs123, %rs124}, %r545;
	cvt.u32.u16 	%r546, %rs124;
	cvt.u32.u16 	%r547, %rs123;
	prmt.b32 	%r548, %r547, %r546, 0x3340U;
	cvt.rzi.s16.f32 	%rs125, %r542;
	cvt.rzi.s16.f32 	%rs126, %r541;
	mov.b32 	%r549, {%rs126, %rs125};
	mov.b32 	{%rs127, %rs128}, %r549;
	cvt.u32.u16 	%r550, %rs128;
	cvt.u32.u16 	%r551, %rs127;
	prmt.b32 	%r552, %r551, %r550, 0x3340U;
	prmt.b32 	%r42, %r552, %r548, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p38, %r393, 0f00000000;
	setp.ge.f32 	%p39, %r394, 0f00000000;
	setp.ge.f32 	%p40, %r395, 0f00000000;
	setp.ge.f32 	%p41, %r396, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r553, 0f3F000000, 0fBF000000, %p41;
	selp.f32 	%r554, 0f3F000000, 0fBF000000, %p40;
	selp.f32 	%r555, 0f3F000000, 0fBF000000, %p39;
	selp.f32 	%r556, 0f3F000000, 0fBF000000, %p38;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r557, %r393, %r556;
	add.f32 	%r558, %r394, %r555;
	add.f32 	%r559, %r395, %r554;
	add.f32 	%r560, %r396, %r553;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs129, %r560;
	cvt.rzi.s16.f32 	%rs130, %r559;
	mov.b32 	%r561, {%rs130, %rs129};
	mov.b32 	{%rs131, %rs132}, %r561;
	cvt.u32.u16 	%r562, %rs132;
	cvt.u32.u16 	%r563, %rs131;
	prmt.b32 	%r564, %r563, %r562, 0x3340U;
	cvt.rzi.s16.f32 	%rs133, %r558;
	cvt.rzi.s16.f32 	%rs134, %r557;
	mov.b32 	%r565, {%rs134, %rs133};
	mov.b32 	{%rs135, %rs136}, %r565;
	cvt.u32.u16 	%r566, %rs136;
	cvt.u32.u16 	%r567, %rs135;
	prmt.b32 	%r568, %r567, %r566, 0x3340U;
	prmt.b32 	%r43, %r568, %r564, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p42, %r397, 0f00000000;
	setp.ge.f32 	%p43, %r398, 0f00000000;
	setp.ge.f32 	%p44, %r399, 0f00000000;
	setp.ge.f32 	%p45, %r400, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r569, 0f3F000000, 0fBF000000, %p45;
	selp.f32 	%r570, 0f3F000000, 0fBF000000, %p44;
	selp.f32 	%r571, 0f3F000000, 0fBF000000, %p43;
	selp.f32 	%r572, 0f3F000000, 0fBF000000, %p42;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r573, %r397, %r572;
	add.f32 	%r574, %r398, %r571;
	add.f32 	%r575, %r399, %r570;
	add.f32 	%r576, %r400, %r569;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs137, %r576;
	cvt.rzi.s16.f32 	%rs138, %r575;
	mov.b32 	%r577, {%rs138, %rs137};
	mov.b32 	{%rs139, %rs140}, %r577;
	cvt.u32.u16 	%r578, %rs140;
	cvt.u32.u16 	%r579, %rs139;
	prmt.b32 	%r580, %r579, %r578, 0x3340U;
	cvt.rzi.s16.f32 	%rs141, %r574;
	cvt.rzi.s16.f32 	%rs142, %r573;
	mov.b32 	%r581, {%rs142, %rs141};
	mov.b32 	{%rs143, %rs144}, %r581;
	cvt.u32.u16 	%r582, %rs144;
	cvt.u32.u16 	%r583, %rs143;
	prmt.b32 	%r584, %r583, %r582, 0x3340U;
	prmt.b32 	%r44, %r584, %r580, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p46, %r401, 0f00000000;
	setp.ge.f32 	%p47, %r402, 0f00000000;
	setp.ge.f32 	%p48, %r403, 0f00000000;
	setp.ge.f32 	%p49, %r404, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r585, 0f3F000000, 0fBF000000, %p49;
	selp.f32 	%r586, 0f3F000000, 0fBF000000, %p48;
	selp.f32 	%r587, 0f3F000000, 0fBF000000, %p47;
	selp.f32 	%r588, 0f3F000000, 0fBF000000, %p46;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r589, %r401, %r588;
	add.f32 	%r590, %r402, %r587;
	add.f32 	%r591, %r403, %r586;
	add.f32 	%r592, %r404, %r585;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs145, %r592;
	cvt.rzi.s16.f32 	%rs146, %r591;
	mov.b32 	%r593, {%rs146, %rs145};
	mov.b32 	{%rs147, %rs148}, %r593;
	cvt.u32.u16 	%r594, %rs148;
	cvt.u32.u16 	%r595, %rs147;
	prmt.b32 	%r596, %r595, %r594, 0x3340U;
	cvt.rzi.s16.f32 	%rs149, %r590;
	cvt.rzi.s16.f32 	%rs150, %r589;
	mov.b32 	%r597, {%rs150, %rs149};
	mov.b32 	{%rs151, %rs152}, %r597;
	cvt.u32.u16 	%r598, %rs152;
	cvt.u32.u16 	%r599, %rs151;
	prmt.b32 	%r600, %r599, %r598, 0x3340U;
	prmt.b32 	%r45, %r600, %r596, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p50, %r405, 0f00000000;
	setp.ge.f32 	%p51, %r406, 0f00000000;
	setp.ge.f32 	%p52, %r407, 0f00000000;
	setp.ge.f32 	%p53, %r408, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r601, 0f3F000000, 0fBF000000, %p53;
	selp.f32 	%r602, 0f3F000000, 0fBF000000, %p52;
	selp.f32 	%r603, 0f3F000000, 0fBF000000, %p51;
	selp.f32 	%r604, 0f3F000000, 0fBF000000, %p50;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r605, %r405, %r604;
	add.f32 	%r606, %r406, %r603;
	add.f32 	%r607, %r407, %r602;
	add.f32 	%r608, %r408, %r601;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs153, %r608;
	cvt.rzi.s16.f32 	%rs154, %r607;
	mov.b32 	%r609, {%rs154, %rs153};
	mov.b32 	{%rs155, %rs156}, %r609;
	cvt.u32.u16 	%r610, %rs156;
	cvt.u32.u16 	%r611, %rs155;
	prmt.b32 	%r612, %r611, %r610, 0x3340U;
	cvt.rzi.s16.f32 	%rs157, %r606;
	cvt.rzi.s16.f32 	%rs158, %r605;
	mov.b32 	%r613, {%rs158, %rs157};
	mov.b32 	{%rs159, %rs160}, %r613;
	cvt.u32.u16 	%r614, %rs160;
	cvt.u32.u16 	%r615, %rs159;
	prmt.b32 	%r616, %r615, %r614, 0x3340U;
	prmt.b32 	%r46, %r616, %r612, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p54, %r409, 0f00000000;
	setp.ge.f32 	%p55, %r410, 0f00000000;
	setp.ge.f32 	%p56, %r411, 0f00000000;
	setp.ge.f32 	%p57, %r412, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r617, 0f3F000000, 0fBF000000, %p57;
	selp.f32 	%r618, 0f3F000000, 0fBF000000, %p56;
	selp.f32 	%r619, 0f3F000000, 0fBF000000, %p55;
	selp.f32 	%r620, 0f3F000000, 0fBF000000, %p54;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r621, %r409, %r620;
	add.f32 	%r622, %r410, %r619;
	add.f32 	%r623, %r411, %r618;
	add.f32 	%r624, %r412, %r617;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs161, %r624;
	cvt.rzi.s16.f32 	%rs162, %r623;
	mov.b32 	%r625, {%rs162, %rs161};
	mov.b32 	{%rs163, %rs164}, %r625;
	cvt.u32.u16 	%r626, %rs164;
	cvt.u32.u16 	%r627, %rs163;
	prmt.b32 	%r628, %r627, %r626, 0x3340U;
	cvt.rzi.s16.f32 	%rs165, %r622;
	cvt.rzi.s16.f32 	%rs166, %r621;
	mov.b32 	%r629, {%rs166, %rs165};
	mov.b32 	{%rs167, %rs168}, %r629;
	cvt.u32.u16 	%r630, %rs168;
	cvt.u32.u16 	%r631, %rs167;
	prmt.b32 	%r632, %r631, %r630, 0x3340U;
	prmt.b32 	%r47, %r632, %r628, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p58, %r413, 0f00000000;
	setp.ge.f32 	%p59, %r414, 0f00000000;
	setp.ge.f32 	%p60, %r415, 0f00000000;
	setp.ge.f32 	%p61, %r416, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r633, 0f3F000000, 0fBF000000, %p61;
	selp.f32 	%r634, 0f3F000000, 0fBF000000, %p60;
	selp.f32 	%r635, 0f3F000000, 0fBF000000, %p59;
	selp.f32 	%r636, 0f3F000000, 0fBF000000, %p58;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r637, %r413, %r636;
	add.f32 	%r638, %r414, %r635;
	add.f32 	%r639, %r415, %r634;
	add.f32 	%r640, %r416, %r633;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs169, %r640;
	cvt.rzi.s16.f32 	%rs170, %r639;
	mov.b32 	%r641, {%rs170, %rs169};
	mov.b32 	{%rs171, %rs172}, %r641;
	cvt.u32.u16 	%r642, %rs172;
	cvt.u32.u16 	%r643, %rs171;
	prmt.b32 	%r644, %r643, %r642, 0x3340U;
	cvt.rzi.s16.f32 	%rs173, %r638;
	cvt.rzi.s16.f32 	%rs174, %r637;
	mov.b32 	%r645, {%rs174, %rs173};
	mov.b32 	{%rs175, %rs176}, %r645;
	cvt.u32.u16 	%r646, %rs176;
	cvt.u32.u16 	%r647, %rs175;
	prmt.b32 	%r648, %r647, %r646, 0x3340U;
	prmt.b32 	%r48, %r648, %r644, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p62, %r417, 0f00000000;
	setp.ge.f32 	%p63, %r418, 0f00000000;
	setp.ge.f32 	%p64, %r419, 0f00000000;
	setp.ge.f32 	%p65, %r420, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r649, 0f3F000000, 0fBF000000, %p65;
	selp.f32 	%r650, 0f3F000000, 0fBF000000, %p64;
	selp.f32 	%r651, 0f3F000000, 0fBF000000, %p63;
	selp.f32 	%r652, 0f3F000000, 0fBF000000, %p62;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r653, %r417, %r652;
	add.f32 	%r654, %r418, %r651;
	add.f32 	%r655, %r419, %r650;
	add.f32 	%r656, %r420, %r649;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs177, %r656;
	cvt.rzi.s16.f32 	%rs178, %r655;
	mov.b32 	%r657, {%rs178, %rs177};
	mov.b32 	{%rs179, %rs180}, %r657;
	cvt.u32.u16 	%r658, %rs180;
	cvt.u32.u16 	%r659, %rs179;
	prmt.b32 	%r660, %r659, %r658, 0x3340U;
	cvt.rzi.s16.f32 	%rs181, %r654;
	cvt.rzi.s16.f32 	%rs182, %r653;
	mov.b32 	%r661, {%rs182, %rs181};
	mov.b32 	{%rs183, %rs184}, %r661;
	cvt.u32.u16 	%r662, %rs184;
	cvt.u32.u16 	%r663, %rs183;
	prmt.b32 	%r664, %r663, %r662, 0x3340U;
	prmt.b32 	%r49, %r664, %r660, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p66, %r421, 0f00000000;
	setp.ge.f32 	%p67, %r422, 0f00000000;
	setp.ge.f32 	%p68, %r423, 0f00000000;
	setp.ge.f32 	%p69, %r424, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r665, 0f3F000000, 0fBF000000, %p69;
	selp.f32 	%r666, 0f3F000000, 0fBF000000, %p68;
	selp.f32 	%r667, 0f3F000000, 0fBF000000, %p67;
	selp.f32 	%r668, 0f3F000000, 0fBF000000, %p66;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r669, %r421, %r668;
	add.f32 	%r670, %r422, %r667;
	add.f32 	%r671, %r423, %r666;
	add.f32 	%r672, %r424, %r665;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs185, %r672;
	cvt.rzi.s16.f32 	%rs186, %r671;
	mov.b32 	%r673, {%rs186, %rs185};
	mov.b32 	{%rs187, %rs188}, %r673;
	cvt.u32.u16 	%r674, %rs188;
	cvt.u32.u16 	%r675, %rs187;
	prmt.b32 	%r676, %r675, %r674, 0x3340U;
	cvt.rzi.s16.f32 	%rs189, %r670;
	cvt.rzi.s16.f32 	%rs190, %r669;
	mov.b32 	%r677, {%rs190, %rs189};
	mov.b32 	{%rs191, %rs192}, %r677;
	cvt.u32.u16 	%r678, %rs192;
	cvt.u32.u16 	%r679, %rs191;
	prmt.b32 	%r680, %r679, %r678, 0x3340U;
	prmt.b32 	%r50, %r680, %r676, 0x5410U;
	.loc	1 30 5                          // quant_per_block.py:30:5
	// begin inline asm
	@%p1 st.global.v4.b32 [ %rd9 + 0 ], { %r35, %r36, %r37, %r38 };

	// end inline asm
	// begin inline asm
	@%p2 st.global.v4.b32 [ %rd10 + 0 ], { %r39, %r40, %r41, %r42 };

	// end inline asm
	// begin inline asm
	@%p3 st.global.v4.b32 [ %rd11 + 0 ], { %r43, %r44, %r45, %r46 };

	// end inline asm
	// begin inline asm
	@%p4 st.global.v4.b32 [ %rd12 + 0 ], { %r47, %r48, %r49, %r50 };

	// end inline asm
	.loc	1 31 5                          // quant_per_block.py:31:5
	setp.eq.b32 	%p5, %r59, 0;
	// begin inline asm
	@%p5 st.global.b32 [ %rd13 + 0 ], { %r51 };

	// end inline asm
	.loc	1 6 1                           // quant_per_block.py:6:1
	ret;
$L__tmp17:
$L__func_end0:
                                        // -- End function
}
	.file	1 "D:\\Comfy\\App\\ComfyUI_windows_portable\\python_embeded\\Lib\\site-packages\\sageattention\\quant_per_block.py"
	.file	2 "D:\\Comfy\\App\\ComfyUI_windows_portable\\python_embeded\\Lib\\site-packages\\triton\\language\\standard.py"
	.section	.debug_abbrev
	{
.b8 1                                   // Abbreviation Code
.b8 17                                  // DW_TAG_compile_unit
.b8 1                                   // DW_CHILDREN_yes
.b8 37                                  // DW_AT_producer
.b8 8                                   // DW_FORM_string
.b8 19                                  // DW_AT_language
.b8 5                                   // DW_FORM_data2
.b8 3                                   // DW_AT_name
.b8 8                                   // DW_FORM_string
.b8 16                                  // DW_AT_stmt_list
.b8 6                                   // DW_FORM_data4
.b8 27                                  // DW_AT_comp_dir
.b8 8                                   // DW_FORM_string
.b8 0                                   // EOM(1)
.b8 0                                   // EOM(2)
.b8 2                                   // Abbreviation Code
.b8 46                                  // DW_TAG_subprogram
.b8 0                                   // DW_CHILDREN_no
.b8 3                                   // DW_AT_name
.b8 8                                   // DW_FORM_string
.b8 32                                  // DW_AT_inline
.b8 11                                  // DW_FORM_data1
.b8 0                                   // EOM(1)
.b8 0                                   // EOM(2)
.b8 3                                   // Abbreviation Code
.b8 46                                  // DW_TAG_subprogram
.b8 1                                   // DW_CHILDREN_yes
.b8 17                                  // DW_AT_low_pc
.b8 1                                   // DW_FORM_addr
.b8 18                                  // DW_AT_high_pc
.b8 1                                   // DW_FORM_addr
.b8 49                                  // DW_AT_abstract_origin
.b8 19                                  // DW_FORM_ref4
.b8 0                                   // EOM(1)
.b8 0                                   // EOM(2)
.b8 4                                   // Abbreviation Code
.b8 29                                  // DW_TAG_inlined_subroutine
.b8 1                                   // DW_CHILDREN_yes
.b8 49                                  // DW_AT_abstract_origin
.b8 19                                  // DW_FORM_ref4
.b8 17                                  // DW_AT_low_pc
.b8 1                                   // DW_FORM_addr
.b8 18                                  // DW_AT_high_pc
.b8 1                                   // DW_FORM_addr
.b8 88                                  // DW_AT_call_file
.b8 11                                  // DW_FORM_data1
.b8 89                                  // DW_AT_call_line
.b8 11                                  // DW_FORM_data1
.b8 87                                  // DW_AT_call_column
.b8 11                                  // DW_FORM_data1
.b8 0                                   // EOM(1)
.b8 0                                   // EOM(2)
.b8 5                                   // Abbreviation Code
.b8 29                                  // DW_TAG_inlined_subroutine
.b8 0                                   // DW_CHILDREN_no
.b8 49                                  // DW_AT_abstract_origin
.b8 19                                  // DW_FORM_ref4
.b8 17                                  // DW_AT_low_pc
.b8 1                                   // DW_FORM_addr
.b8 18                                  // DW_AT_high_pc
.b8 1                                   // DW_FORM_addr
.b8 88                                  // DW_AT_call_file
.b8 11                                  // DW_FORM_data1
.b8 89                                  // DW_AT_call_line
.b8 11                                  // DW_FORM_data1
.b8 87                                  // DW_AT_call_column
.b8 11                                  // DW_FORM_data1
.b8 0                                   // EOM(1)
.b8 0                                   // EOM(2)
.b8 0                                   // EOM(3)
	}
	.section	.debug_info
	{
.b32 227                                // Length of Unit
.b8 2                                   // DWARF version number
.b8 0
.b32 .debug_abbrev                      // Offset Into Abbrev. Section
.b8 8                                   // Address Size (in bytes)
.b8 1                                   // Abbrev [1] 0xb:0xdc DW_TAG_compile_unit
.b8 116                                 // DW_AT_producer
.b8 114
.b8 105
.b8 116
.b8 111
.b8 110
.b8 0
.b8 2                                   // DW_AT_language
.b8 0
.b8 113                                 // DW_AT_name
.b8 117
.b8 97
.b8 110
.b8 116
.b8 95
.b8 112
.b8 101
.b8 114
.b8 95
.b8 98
.b8 108
.b8 111
.b8 99
.b8 107
.b8 46
.b8 112
.b8 121
.b8 0
.b32 .debug_line                        // DW_AT_stmt_list
.b8 68                                  // DW_AT_comp_dir
.b8 58
.b8 92
.b8 67
.b8 111
.b8 109
.b8 102
.b8 121
.b8 92
.b8 65
.b8 112
.b8 112
.b8 92
.b8 67
.b8 111
.b8 109
.b8 102
.b8 121
.b8 85
.b8 73
.b8 95
.b8 119
.b8 105
.b8 110
.b8 100
.b8 111
.b8 119
.b8 115
.b8 95
.b8 112
.b8 111
.b8 114
.b8 116
.b8 97
.b8 98
.b8 108
.b8 101
.b8 92
.b8 112
.b8 121
.b8 116
.b8 104
.b8 111
.b8 110
.b8 95
.b8 101
.b8 109
.b8 98
.b8 101
.b8 100
.b8 101
.b8 100
.b8 92
.b8 76
.b8 105
.b8 98
.b8 92
.b8 115
.b8 105
.b8 116
.b8 101
.b8 45
.b8 112
.b8 97
.b8 99
.b8 107
.b8 97
.b8 103
.b8 101
.b8 115
.b8 92
.b8 115
.b8 97
.b8 103
.b8 101
.b8 97
.b8 116
.b8 116
.b8 101
.b8 110
.b8 116
.b8 105
.b8 111
.b8 110
.b8 0
.b8 2                                   // Abbrev [2] 0x81:0x1e DW_TAG_subprogram
.b8 113                                 // DW_AT_name
.b8 117
.b8 97
.b8 110
.b8 116
.b8 95
.b8 112
.b8 101
.b8 114
.b8 95
.b8 98
.b8 108
.b8 111
.b8 99
.b8 107
.b8 95
.b8 105
.b8 110
.b8 116
.b8 56
.b8 95
.b8 107
.b8 101
.b8 114
.b8 110
.b8 101
.b8 108
.b8 0
.b8 1                                   // DW_AT_inline
.b8 3                                   // Abbrev [3] 0x9f:0x47 DW_TAG_subprogram
.b64 $L__func_begin0                    // DW_AT_low_pc
.b64 $L__func_end0                      // DW_AT_high_pc
.b32 129                                // DW_AT_abstract_origin
.b8 4                                   // Abbrev [4] 0xb4:0x31 DW_TAG_inlined_subroutine
.b32 129                                // DW_AT_abstract_origin
.b64 $L__tmp1                           // DW_AT_low_pc
.b64 $L__tmp16                          // DW_AT_high_pc
.b8 1                                   // DW_AT_call_file
.b8 26                                  // DW_AT_call_line
.b8 13                                  // DW_AT_call_column
.b8 5                                   // Abbrev [5] 0xcc:0x18 DW_TAG_inlined_subroutine
.b32 129                                // DW_AT_abstract_origin
.b64 $L__tmp1                           // DW_AT_low_pc
.b64 $L__tmp16                          // DW_AT_high_pc
.b8 2                                   // DW_AT_call_file
.b8 191                                 // DW_AT_call_line
.b8 16                                  // DW_AT_call_column
.b8 0                                   // End Of Children Mark
.b8 0                                   // End Of Children Mark
.b8 0                                   // End Of Children Mark
	}
	.section	.debug_str
	{
$L__info_string0:
.b8 113                                 // string offset=0 ; quant_per_block_int8_kernel
.b8 117
.b8 97
.b8 110
.b8 116
.b8 95
.b8 112
.b8 101
.b8 114
.b8 95
.b8 98
.b8 108
.b8 111
.b8 99
.b8 107
.b8 95
.b8 105
.b8 110
.b8 116
.b8 56
.b8 95
.b8 107
.b8 101
.b8 114
.b8 110
.b8 101
.b8 108
.b8 0
	}
	.section	.debug_macinfo	{	}