//
// 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 .u32 quant_per_block_int8_kernel_param_11,
	.param .f32 quant_per_block_int8_kernel_param_12,
	.param .u64 .ptr .global .align 1 quant_per_block_int8_kernel_param_13,
	.param .u64 .ptr .global .align 1 quant_per_block_int8_kernel_param_14
)
.reqntid 128
{
	.reg .pred 	%p<138>;
	.reg .b16 	%rs<385>;
	.reg .b32 	%r<1321>;
	.reg .b64 	%rd<63>;
	.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 	%rd26, [quant_per_block_int8_kernel_param_0];
	ld.param.b64 	%rd27, [quant_per_block_int8_kernel_param_1];
$L__tmp0:
	.loc	1 12 15                         // quant_per_block.py:12:15
	mov.u32 	%r100, %ctaid.x;
	ld.param.b64 	%rd28, [quant_per_block_int8_kernel_param_2];
	.loc	1 13 13                         // quant_per_block.py:13:13
	mov.u32 	%r101, %ctaid.y;
	ld.param.b32 	%r102, [quant_per_block_int8_kernel_param_3];
	.loc	1 14 13                         // quant_per_block.py:14:13
	mov.u32 	%r103, %ctaid.z;
	.loc	1 16 14                         // quant_per_block.py:16:14
	shl.b32 	%r104, %r100, 7;
	ld.param.b32 	%r105, [quant_per_block_int8_kernel_param_4];
	ld.param.b32 	%r106, [quant_per_block_int8_kernel_param_5];
	.loc	1 16 30                         // quant_per_block.py:16:30
	mov.u32 	%r107, %tid.x;
	shr.u32 	%r108, %r107, 3;
	ld.param.b32 	%r109, [quant_per_block_int8_kernel_param_6];
	bfe.u32 	%r110, %r107, 3, 4;
	ld.param.b32 	%r111, [quant_per_block_int8_kernel_param_7];
	.loc	1 16 14                         // quant_per_block.py:16:14
	or.b32 	%r112, %r110, %r104;
	ld.param.b32 	%r113, [quant_per_block_int8_kernel_param_8];
	or.b32 	%r114, %r112, 16;
	ld.param.b32 	%r115, [quant_per_block_int8_kernel_param_9];
	or.b32 	%r116, %r112, 32;
	ld.param.b32 	%r117, [quant_per_block_int8_kernel_param_10];
	or.b32 	%r118, %r112, 48;
	ld.param.b32 	%r119, [quant_per_block_int8_kernel_param_11];
	or.b32 	%r120, %r112, 64;
	ld.param.b32 	%r121, [quant_per_block_int8_kernel_param_12];
	or.b32 	%r122, %r112, 80;
	or.b32 	%r123, %r112, 96;
	or.b32 	%r124, %r108, %r104;
	or.b32 	%r125, %r124, 112;
	.loc	1 19 26                         // quant_per_block.py:19:26
	mul.lo.s32 	%r126, %r105, %r103;
	.loc	1 19 18                         // quant_per_block.py:19:18
	mad.wide.s32 	%rd29, %r126, 2, %rd26;
	.loc	1 19 46                         // quant_per_block.py:19:46
	mul.lo.s32 	%r127, %r106, %r101;
	.loc	1 19 18                         // quant_per_block.py:19:18
	mad.wide.s32 	%rd30, %r127, 2, %rd29;
	.loc	1 19 66                         // quant_per_block.py:19:66
	mul.lo.s32 	%r128, %r109, %r112;
	mul.lo.s32 	%r129, %r109, %r114;
	mul.lo.s32 	%r130, %r109, %r116;
	mul.lo.s32 	%r131, %r109, %r118;
	mul.lo.s32 	%r132, %r109, %r120;
	mul.lo.s32 	%r133, %r109, %r122;
	mul.lo.s32 	%r134, %r109, %r123;
	mul.lo.s32 	%r135, %r109, %r125;
	.loc	1 19 18                         // quant_per_block.py:19:18
	mad.wide.s32 	%rd31, %r128, 2, %rd30;
	mad.wide.s32 	%rd32, %r129, 2, %rd30;
	mad.wide.s32 	%rd33, %r130, 2, %rd30;
	mad.wide.s32 	%rd34, %r131, 2, %rd30;
	mad.wide.s32 	%rd35, %r132, 2, %rd30;
	mad.wide.s32 	%rd36, %r133, 2, %rd30;
	mad.wide.s32 	%rd37, %r134, 2, %rd30;
	mad.wide.s32 	%rd38, %r135, 2, %rd30;
	.loc	1 19 96                         // quant_per_block.py:19:96
	shl.b32 	%r136, %r107, 4;
	and.b32 	%r137, %r136, 112;
	.loc	1 19 18                         // quant_per_block.py:19:18
	cvt.u64.u32 	%rd39, %r137;
	mul.wide.u32 	%rd40, %r137, 2;
	add.s64 	%rd1, %rd31, %rd40;
	add.s64 	%rd2, %rd1, 16;
	add.s64 	%rd3, %rd32, %rd40;
	add.s64 	%rd4, %rd3, 16;
	add.s64 	%rd5, %rd33, %rd40;
	add.s64 	%rd6, %rd5, 16;
	add.s64 	%rd7, %rd34, %rd40;
	add.s64 	%rd8, %rd7, 16;
	add.s64 	%rd9, %rd35, %rd40;
	add.s64 	%rd10, %rd9, 16;
	add.s64 	%rd11, %rd36, %rd40;
	add.s64 	%rd12, %rd11, 16;
	add.s64 	%rd13, %rd37, %rd40;
	add.s64 	%rd14, %rd13, 16;
	add.s64 	%rd15, %rd38, %rd40;
	add.s64 	%rd16, %rd15, 16;
	.loc	1 20 28                         // quant_per_block.py:20:28
	mul.lo.s32 	%r138, %r111, %r103;
	.loc	1 20 19                         // quant_per_block.py:20:19
	cvt.s64.s32 	%rd41, %r138;
	add.s64 	%rd42, %rd27, %rd41;
	.loc	1 20 48                         // quant_per_block.py:20:48
	mul.lo.s32 	%r139, %r113, %r101;
	.loc	1 20 19                         // quant_per_block.py:20:19
	cvt.s64.s32 	%rd43, %r139;
	add.s64 	%rd44, %rd42, %rd43;
	.loc	1 20 68                         // quant_per_block.py:20:68
	mul.lo.s32 	%r140, %r115, %r112;
	mul.lo.s32 	%r141, %r115, %r114;
	mul.lo.s32 	%r142, %r115, %r116;
	mul.lo.s32 	%r143, %r115, %r118;
	mul.lo.s32 	%r144, %r115, %r120;
	mul.lo.s32 	%r145, %r115, %r122;
	mul.lo.s32 	%r146, %r115, %r123;
	mul.lo.s32 	%r147, %r115, %r125;
	.loc	1 20 19                         // quant_per_block.py:20:19
	cvt.s64.s32 	%rd45, %r140;
	add.s64 	%rd46, %rd44, %rd45;
	cvt.s64.s32 	%rd47, %r141;
	add.s64 	%rd48, %rd44, %rd47;
	cvt.s64.s32 	%rd49, %r142;
	add.s64 	%rd50, %rd44, %rd49;
	cvt.s64.s32 	%rd51, %r143;
	add.s64 	%rd52, %rd44, %rd51;
	cvt.s64.s32 	%rd53, %r144;
	add.s64 	%rd54, %rd44, %rd53;
	cvt.s64.s32 	%rd55, %r145;
	add.s64 	%rd56, %rd44, %rd55;
	cvt.s64.s32 	%rd57, %r146;
	add.s64 	%rd58, %rd44, %rd57;
	cvt.s64.s32 	%rd59, %r147;
	add.s64 	%rd60, %rd44, %rd59;
	add.s64 	%rd17, %rd46, %rd39;
	add.s64 	%rd18, %rd48, %rd39;
	add.s64 	%rd19, %rd50, %rd39;
	add.s64 	%rd20, %rd52, %rd39;
	add.s64 	%rd21, %rd54, %rd39;
	add.s64 	%rd22, %rd56, %rd39;
	add.s64 	%rd23, %rd58, %rd39;
	add.s64 	%rd24, %rd60, %rd39;
	.loc	1 21 26                         // quant_per_block.py:21:26
	mul.lo.s32 	%r148, %r117, %r103;
	.loc	1 21 18                         // quant_per_block.py:21:18
	mad.wide.s32 	%rd61, %r148, 4, %rd28;
	.loc	1 21 46                         // quant_per_block.py:21:46
	mul.lo.s32 	%r149, %r119, %r101;
	.loc	1 21 18                         // quant_per_block.py:21:18
	mad.wide.s32 	%rd62, %r149, 4, %rd61;
	mad.wide.u32 	%rd25, %r100, 4, %rd62;
	.loc	1 23 34                         // quant_per_block.py:23:34
	setp.lt.s32 	%p1, %r112, %r102;
	setp.lt.s32 	%p2, %r114, %r102;
	setp.lt.s32 	%p3, %r116, %r102;
	setp.lt.s32 	%p4, %r118, %r102;
	setp.lt.s32 	%p5, %r120, %r102;
	setp.lt.s32 	%p6, %r122, %r102;
	setp.lt.s32 	%p7, %r123, %r102;
	setp.lt.s32 	%p8, %r125, %r102;
	.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
	// begin inline asm
	.loc	1 23 9                          // quant_per_block.py:23:9
	mov.u32 %r33, 0x0;
	.loc	1 24 9                          // quant_per_block.py:24:9
	mov.u32 %r34, 0x0;
	.loc	1 25 9                          // quant_per_block.py:25:9
	mov.u32 %r35, 0x0;
	.loc	1 26 9                          // quant_per_block.py:26:9
	mov.u32 %r36, 0x0;
	@%p5 ld.global.v4.b32 { %r33, %r34, %r35, %r36 }, [ %rd9 + 0 ];

	// end inline asm
	// begin inline asm
	.loc	1 23 9                          // quant_per_block.py:23:9
	mov.u32 %r37, 0x0;
	.loc	1 24 9                          // quant_per_block.py:24:9
	mov.u32 %r38, 0x0;
	.loc	1 25 9                          // quant_per_block.py:25:9
	mov.u32 %r39, 0x0;
	.loc	1 26 9                          // quant_per_block.py:26:9
	mov.u32 %r40, 0x0;
	@%p5 ld.global.v4.b32 { %r37, %r38, %r39, %r40 }, [ %rd10 + 0 ];

	// end inline asm
	// begin inline asm
	.loc	1 23 9                          // quant_per_block.py:23:9
	mov.u32 %r41, 0x0;
	.loc	1 24 9                          // quant_per_block.py:24:9
	mov.u32 %r42, 0x0;
	.loc	1 25 9                          // quant_per_block.py:25:9
	mov.u32 %r43, 0x0;
	.loc	1 26 9                          // quant_per_block.py:26:9
	mov.u32 %r44, 0x0;
	@%p6 ld.global.v4.b32 { %r41, %r42, %r43, %r44 }, [ %rd11 + 0 ];

	// end inline asm
	// begin inline asm
	.loc	1 23 9                          // quant_per_block.py:23:9
	mov.u32 %r45, 0x0;
	.loc	1 24 9                          // quant_per_block.py:24:9
	mov.u32 %r46, 0x0;
	.loc	1 25 9                          // quant_per_block.py:25:9
	mov.u32 %r47, 0x0;
	.loc	1 26 9                          // quant_per_block.py:26:9
	mov.u32 %r48, 0x0;
	@%p6 ld.global.v4.b32 { %r45, %r46, %r47, %r48 }, [ %rd12 + 0 ];

	// end inline asm
	// begin inline asm
	.loc	1 23 9                          // quant_per_block.py:23:9
	mov.u32 %r49, 0x0;
	.loc	1 24 9                          // quant_per_block.py:24:9
	mov.u32 %r50, 0x0;
	.loc	1 25 9                          // quant_per_block.py:25:9
	mov.u32 %r51, 0x0;
	.loc	1 26 9                          // quant_per_block.py:26:9
	mov.u32 %r52, 0x0;
	@%p7 ld.global.v4.b32 { %r49, %r50, %r51, %r52 }, [ %rd13 + 0 ];

	// end inline asm
	// begin inline asm
	.loc	1 23 9                          // quant_per_block.py:23:9
	mov.u32 %r53, 0x0;
	.loc	1 24 9                          // quant_per_block.py:24:9
	mov.u32 %r54, 0x0;
	.loc	1 25 9                          // quant_per_block.py:25:9
	mov.u32 %r55, 0x0;
	.loc	1 26 9                          // quant_per_block.py:26:9
	mov.u32 %r56, 0x0;
	@%p7 ld.global.v4.b32 { %r53, %r54, %r55, %r56 }, [ %rd14 + 0 ];

	// end inline asm
	// begin inline asm
	.loc	1 23 9                          // quant_per_block.py:23:9
	mov.u32 %r57, 0x0;
	.loc	1 24 9                          // quant_per_block.py:24:9
	mov.u32 %r58, 0x0;
	.loc	1 25 9                          // quant_per_block.py:25:9
	mov.u32 %r59, 0x0;
	.loc	1 26 9                          // quant_per_block.py:26:9
	mov.u32 %r60, 0x0;
	@%p8 ld.global.v4.b32 { %r57, %r58, %r59, %r60 }, [ %rd15 + 0 ];

	// end inline asm
	// begin inline asm
	.loc	1 23 9                          // quant_per_block.py:23:9
	mov.u32 %r61, 0x0;
	.loc	1 24 9                          // quant_per_block.py:24:9
	mov.u32 %r62, 0x0;
	.loc	1 25 9                          // quant_per_block.py:25:9
	mov.u32 %r63, 0x0;
	.loc	1 26 9                          // quant_per_block.py:26:9
	mov.u32 %r64, 0x0;
	@%p8 ld.global.v4.b32 { %r61, %r62, %r63, %r64 }, [ %rd16 + 0 ];

	// end inline asm
	.loc	1 24 9                          // quant_per_block.py:24:9
	mov.b32 	{%rs1, %rs2}, %r64;
	cvt.f32.bf16 	%r150, %rs2;
	cvt.f32.bf16 	%r151, %rs1;
	mov.b32 	{%rs3, %rs4}, %r63;
	cvt.f32.bf16 	%r152, %rs4;
	cvt.f32.bf16 	%r153, %rs3;
	mov.b32 	{%rs5, %rs6}, %r62;
	cvt.f32.bf16 	%r154, %rs6;
	cvt.f32.bf16 	%r155, %rs5;
	mov.b32 	{%rs7, %rs8}, %r61;
	cvt.f32.bf16 	%r156, %rs8;
	cvt.f32.bf16 	%r157, %rs7;
	mov.b32 	{%rs9, %rs10}, %r60;
	cvt.f32.bf16 	%r158, %rs10;
	cvt.f32.bf16 	%r159, %rs9;
	mov.b32 	{%rs11, %rs12}, %r59;
	cvt.f32.bf16 	%r160, %rs12;
	cvt.f32.bf16 	%r161, %rs11;
	mov.b32 	{%rs13, %rs14}, %r58;
	cvt.f32.bf16 	%r162, %rs14;
	cvt.f32.bf16 	%r163, %rs13;
	mov.b32 	{%rs15, %rs16}, %r57;
	cvt.f32.bf16 	%r164, %rs16;
	cvt.f32.bf16 	%r165, %rs15;
	mov.b32 	{%rs17, %rs18}, %r56;
	cvt.f32.bf16 	%r166, %rs18;
	cvt.f32.bf16 	%r167, %rs17;
	mov.b32 	{%rs19, %rs20}, %r55;
	cvt.f32.bf16 	%r168, %rs20;
	cvt.f32.bf16 	%r169, %rs19;
	mov.b32 	{%rs21, %rs22}, %r54;
	cvt.f32.bf16 	%r170, %rs22;
	cvt.f32.bf16 	%r171, %rs21;
	mov.b32 	{%rs23, %rs24}, %r53;
	cvt.f32.bf16 	%r172, %rs24;
	cvt.f32.bf16 	%r173, %rs23;
	mov.b32 	{%rs25, %rs26}, %r52;
	cvt.f32.bf16 	%r174, %rs26;
	cvt.f32.bf16 	%r175, %rs25;
	mov.b32 	{%rs27, %rs28}, %r51;
	cvt.f32.bf16 	%r176, %rs28;
	cvt.f32.bf16 	%r177, %rs27;
	mov.b32 	{%rs29, %rs30}, %r50;
	cvt.f32.bf16 	%r178, %rs30;
	cvt.f32.bf16 	%r179, %rs29;
	mov.b32 	{%rs31, %rs32}, %r49;
	cvt.f32.bf16 	%r180, %rs32;
	cvt.f32.bf16 	%r181, %rs31;
	mov.b32 	{%rs33, %rs34}, %r48;
	cvt.f32.bf16 	%r182, %rs34;
	cvt.f32.bf16 	%r183, %rs33;
	mov.b32 	{%rs35, %rs36}, %r47;
	cvt.f32.bf16 	%r184, %rs36;
	cvt.f32.bf16 	%r185, %rs35;
	mov.b32 	{%rs37, %rs38}, %r46;
	cvt.f32.bf16 	%r186, %rs38;
	cvt.f32.bf16 	%r187, %rs37;
	mov.b32 	{%rs39, %rs40}, %r45;
	cvt.f32.bf16 	%r188, %rs40;
	cvt.f32.bf16 	%r189, %rs39;
	mov.b32 	{%rs41, %rs42}, %r44;
	cvt.f32.bf16 	%r190, %rs42;
	cvt.f32.bf16 	%r191, %rs41;
	mov.b32 	{%rs43, %rs44}, %r43;
	cvt.f32.bf16 	%r192, %rs44;
	cvt.f32.bf16 	%r193, %rs43;
	mov.b32 	{%rs45, %rs46}, %r42;
	cvt.f32.bf16 	%r194, %rs46;
	cvt.f32.bf16 	%r195, %rs45;
	mov.b32 	{%rs47, %rs48}, %r41;
	cvt.f32.bf16 	%r196, %rs48;
	cvt.f32.bf16 	%r197, %rs47;
	mov.b32 	{%rs49, %rs50}, %r40;
	cvt.f32.bf16 	%r198, %rs50;
	cvt.f32.bf16 	%r199, %rs49;
	mov.b32 	{%rs51, %rs52}, %r39;
	cvt.f32.bf16 	%r200, %rs52;
	cvt.f32.bf16 	%r201, %rs51;
	mov.b32 	{%rs53, %rs54}, %r38;
	cvt.f32.bf16 	%r202, %rs54;
	cvt.f32.bf16 	%r203, %rs53;
	mov.b32 	{%rs55, %rs56}, %r37;
	cvt.f32.bf16 	%r204, %rs56;
	cvt.f32.bf16 	%r205, %rs55;
	mov.b32 	{%rs57, %rs58}, %r36;
	cvt.f32.bf16 	%r206, %rs58;
	cvt.f32.bf16 	%r207, %rs57;
	mov.b32 	{%rs59, %rs60}, %r35;
	cvt.f32.bf16 	%r208, %rs60;
	cvt.f32.bf16 	%r209, %rs59;
	mov.b32 	{%rs61, %rs62}, %r34;
	cvt.f32.bf16 	%r210, %rs62;
	cvt.f32.bf16 	%r211, %rs61;
	mov.b32 	{%rs63, %rs64}, %r33;
	cvt.f32.bf16 	%r212, %rs64;
	cvt.f32.bf16 	%r213, %rs63;
	mov.b32 	{%rs65, %rs66}, %r32;
	cvt.f32.bf16 	%r214, %rs66;
	cvt.f32.bf16 	%r215, %rs65;
	mov.b32 	{%rs67, %rs68}, %r31;
	cvt.f32.bf16 	%r216, %rs68;
	cvt.f32.bf16 	%r217, %rs67;
	mov.b32 	{%rs69, %rs70}, %r30;
	cvt.f32.bf16 	%r218, %rs70;
	cvt.f32.bf16 	%r219, %rs69;
	mov.b32 	{%rs71, %rs72}, %r29;
	cvt.f32.bf16 	%r220, %rs72;
	cvt.f32.bf16 	%r221, %rs71;
	mov.b32 	{%rs73, %rs74}, %r28;
	cvt.f32.bf16 	%r222, %rs74;
	cvt.f32.bf16 	%r223, %rs73;
	mov.b32 	{%rs75, %rs76}, %r27;
	cvt.f32.bf16 	%r224, %rs76;
	cvt.f32.bf16 	%r225, %rs75;
	mov.b32 	{%rs77, %rs78}, %r26;
	cvt.f32.bf16 	%r226, %rs78;
	cvt.f32.bf16 	%r227, %rs77;
	mov.b32 	{%rs79, %rs80}, %r25;
	cvt.f32.bf16 	%r228, %rs80;
	cvt.f32.bf16 	%r229, %rs79;
	mov.b32 	{%rs81, %rs82}, %r24;
	cvt.f32.bf16 	%r230, %rs82;
	cvt.f32.bf16 	%r231, %rs81;
	mov.b32 	{%rs83, %rs84}, %r23;
	cvt.f32.bf16 	%r232, %rs84;
	cvt.f32.bf16 	%r233, %rs83;
	mov.b32 	{%rs85, %rs86}, %r22;
	cvt.f32.bf16 	%r234, %rs86;
	cvt.f32.bf16 	%r235, %rs85;
	mov.b32 	{%rs87, %rs88}, %r21;
	cvt.f32.bf16 	%r236, %rs88;
	cvt.f32.bf16 	%r237, %rs87;
	mov.b32 	{%rs89, %rs90}, %r20;
	cvt.f32.bf16 	%r238, %rs90;
	cvt.f32.bf16 	%r239, %rs89;
	mov.b32 	{%rs91, %rs92}, %r19;
	cvt.f32.bf16 	%r240, %rs92;
	cvt.f32.bf16 	%r241, %rs91;
	mov.b32 	{%rs93, %rs94}, %r18;
	cvt.f32.bf16 	%r242, %rs94;
	cvt.f32.bf16 	%r243, %rs93;
	mov.b32 	{%rs95, %rs96}, %r17;
	cvt.f32.bf16 	%r244, %rs96;
	cvt.f32.bf16 	%r245, %rs95;
	mov.b32 	{%rs97, %rs98}, %r16;
	cvt.f32.bf16 	%r246, %rs98;
	cvt.f32.bf16 	%r247, %rs97;
	mov.b32 	{%rs99, %rs100}, %r15;
	cvt.f32.bf16 	%r248, %rs100;
	cvt.f32.bf16 	%r249, %rs99;
	mov.b32 	{%rs101, %rs102}, %r14;
	cvt.f32.bf16 	%r250, %rs102;
	cvt.f32.bf16 	%r251, %rs101;
	mov.b32 	{%rs103, %rs104}, %r13;
	cvt.f32.bf16 	%r252, %rs104;
	cvt.f32.bf16 	%r253, %rs103;
	mov.b32 	{%rs105, %rs106}, %r12;
	cvt.f32.bf16 	%r254, %rs106;
	cvt.f32.bf16 	%r255, %rs105;
	mov.b32 	{%rs107, %rs108}, %r11;
	cvt.f32.bf16 	%r256, %rs108;
	cvt.f32.bf16 	%r257, %rs107;
	mov.b32 	{%rs109, %rs110}, %r10;
	cvt.f32.bf16 	%r258, %rs110;
	cvt.f32.bf16 	%r259, %rs109;
	mov.b32 	{%rs111, %rs112}, %r9;
	cvt.f32.bf16 	%r260, %rs112;
	cvt.f32.bf16 	%r261, %rs111;
	mov.b32 	{%rs113, %rs114}, %r8;
	cvt.f32.bf16 	%r262, %rs114;
	cvt.f32.bf16 	%r263, %rs113;
	mov.b32 	{%rs115, %rs116}, %r7;
	cvt.f32.bf16 	%r264, %rs116;
	cvt.f32.bf16 	%r265, %rs115;
	mov.b32 	{%rs117, %rs118}, %r6;
	cvt.f32.bf16 	%r266, %rs118;
	cvt.f32.bf16 	%r267, %rs117;
	mov.b32 	{%rs119, %rs120}, %r5;
	cvt.f32.bf16 	%r268, %rs120;
	cvt.f32.bf16 	%r269, %rs119;
	mov.b32 	{%rs121, %rs122}, %r4;
	cvt.f32.bf16 	%r270, %rs122;
	cvt.f32.bf16 	%r271, %rs121;
	mov.b32 	{%rs123, %rs124}, %r3;
	cvt.f32.bf16 	%r272, %rs124;
	cvt.f32.bf16 	%r273, %rs123;
	mov.b32 	{%rs125, %rs126}, %r2;
	cvt.f32.bf16 	%r274, %rs126;
	cvt.f32.bf16 	%r275, %rs125;
	mov.b32 	{%rs127, %rs128}, %r1;
	cvt.f32.bf16 	%r276, %rs128;
	cvt.f32.bf16 	%r277, %rs127;
	.loc	1 25 5                          // quant_per_block.py:25:5
	mul.f32 	%r278, %r121, %r277;
	mul.f32 	%r279, %r121, %r276;
	mul.f32 	%r280, %r121, %r275;
	mul.f32 	%r281, %r121, %r274;
	mul.f32 	%r282, %r121, %r273;
	mul.f32 	%r283, %r121, %r272;
	mul.f32 	%r284, %r121, %r271;
	mul.f32 	%r285, %r121, %r270;
	mul.f32 	%r286, %r121, %r269;
	mul.f32 	%r287, %r121, %r268;
	mul.f32 	%r288, %r121, %r267;
	mul.f32 	%r289, %r121, %r266;
	mul.f32 	%r290, %r121, %r265;
	mul.f32 	%r291, %r121, %r264;
	mul.f32 	%r292, %r121, %r263;
	mul.f32 	%r293, %r121, %r262;
	mul.f32 	%r294, %r121, %r261;
	mul.f32 	%r295, %r121, %r260;
	mul.f32 	%r296, %r121, %r259;
	mul.f32 	%r297, %r121, %r258;
	mul.f32 	%r298, %r121, %r257;
	mul.f32 	%r299, %r121, %r256;
	mul.f32 	%r300, %r121, %r255;
	mul.f32 	%r301, %r121, %r254;
	mul.f32 	%r302, %r121, %r253;
	mul.f32 	%r303, %r121, %r252;
	mul.f32 	%r304, %r121, %r251;
	mul.f32 	%r305, %r121, %r250;
	mul.f32 	%r306, %r121, %r249;
	mul.f32 	%r307, %r121, %r248;
	mul.f32 	%r308, %r121, %r247;
	mul.f32 	%r309, %r121, %r246;
	mul.f32 	%r310, %r121, %r245;
	mul.f32 	%r311, %r121, %r244;
	mul.f32 	%r312, %r121, %r243;
	mul.f32 	%r313, %r121, %r242;
	mul.f32 	%r314, %r121, %r241;
	mul.f32 	%r315, %r121, %r240;
	mul.f32 	%r316, %r121, %r239;
	mul.f32 	%r317, %r121, %r238;
	mul.f32 	%r318, %r121, %r237;
	mul.f32 	%r319, %r121, %r236;
	mul.f32 	%r320, %r121, %r235;
	mul.f32 	%r321, %r121, %r234;
	mul.f32 	%r322, %r121, %r233;
	mul.f32 	%r323, %r121, %r232;
	mul.f32 	%r324, %r121, %r231;
	mul.f32 	%r325, %r121, %r230;
	mul.f32 	%r326, %r121, %r229;
	mul.f32 	%r327, %r121, %r228;
	mul.f32 	%r328, %r121, %r227;
	mul.f32 	%r329, %r121, %r226;
	mul.f32 	%r330, %r121, %r225;
	mul.f32 	%r331, %r121, %r224;
	mul.f32 	%r332, %r121, %r223;
	mul.f32 	%r333, %r121, %r222;
	mul.f32 	%r334, %r121, %r221;
	mul.f32 	%r335, %r121, %r220;
	mul.f32 	%r336, %r121, %r219;
	mul.f32 	%r337, %r121, %r218;
	mul.f32 	%r338, %r121, %r217;
	mul.f32 	%r339, %r121, %r216;
	mul.f32 	%r340, %r121, %r215;
	mul.f32 	%r341, %r121, %r214;
	mul.f32 	%r342, %r121, %r213;
	mul.f32 	%r343, %r121, %r212;
	mul.f32 	%r344, %r121, %r211;
	mul.f32 	%r345, %r121, %r210;
	mul.f32 	%r346, %r121, %r209;
	mul.f32 	%r347, %r121, %r208;
	mul.f32 	%r348, %r121, %r207;
	mul.f32 	%r349, %r121, %r206;
	mul.f32 	%r350, %r121, %r205;
	mul.f32 	%r351, %r121, %r204;
	mul.f32 	%r352, %r121, %r203;
	mul.f32 	%r353, %r121, %r202;
	mul.f32 	%r354, %r121, %r201;
	mul.f32 	%r355, %r121, %r200;
	mul.f32 	%r356, %r121, %r199;
	mul.f32 	%r357, %r121, %r198;
	mul.f32 	%r358, %r121, %r197;
	mul.f32 	%r359, %r121, %r196;
	mul.f32 	%r360, %r121, %r195;
	mul.f32 	%r361, %r121, %r194;
	mul.f32 	%r362, %r121, %r193;
	mul.f32 	%r363, %r121, %r192;
	mul.f32 	%r364, %r121, %r191;
	mul.f32 	%r365, %r121, %r190;
	mul.f32 	%r366, %r121, %r189;
	mul.f32 	%r367, %r121, %r188;
	mul.f32 	%r368, %r121, %r187;
	mul.f32 	%r369, %r121, %r186;
	mul.f32 	%r370, %r121, %r185;
	mul.f32 	%r371, %r121, %r184;
	mul.f32 	%r372, %r121, %r183;
	mul.f32 	%r373, %r121, %r182;
	mul.f32 	%r374, %r121, %r181;
	mul.f32 	%r375, %r121, %r180;
	mul.f32 	%r376, %r121, %r179;
	mul.f32 	%r377, %r121, %r178;
	mul.f32 	%r378, %r121, %r177;
	mul.f32 	%r379, %r121, %r176;
	mul.f32 	%r380, %r121, %r175;
	mul.f32 	%r381, %r121, %r174;
	mul.f32 	%r382, %r121, %r173;
	mul.f32 	%r383, %r121, %r172;
	mul.f32 	%r384, %r121, %r171;
	mul.f32 	%r385, %r121, %r170;
	mul.f32 	%r386, %r121, %r169;
	mul.f32 	%r387, %r121, %r168;
	mul.f32 	%r388, %r121, %r167;
	mul.f32 	%r389, %r121, %r166;
	mul.f32 	%r390, %r121, %r165;
	mul.f32 	%r391, %r121, %r164;
	mul.f32 	%r392, %r121, %r163;
	mul.f32 	%r393, %r121, %r162;
	mul.f32 	%r394, %r121, %r161;
	mul.f32 	%r395, %r121, %r160;
	mul.f32 	%r396, %r121, %r159;
	mul.f32 	%r397, %r121, %r158;
	mul.f32 	%r398, %r121, %r157;
	mul.f32 	%r399, %r121, %r156;
	mul.f32 	%r400, %r121, %r155;
	mul.f32 	%r401, %r121, %r154;
	mul.f32 	%r402, %r121, %r153;
	mul.f32 	%r403, %r121, %r152;
	mul.f32 	%r404, %r121, %r151;
	mul.f32 	%r405, %r121, %r150;
	.loc	1 26 20                         // quant_per_block.py:26:20
	abs.f32 	%r406, %r405;
	abs.f32 	%r407, %r404;
	abs.f32 	%r408, %r403;
	abs.f32 	%r409, %r402;
	abs.f32 	%r410, %r401;
	abs.f32 	%r411, %r400;
	abs.f32 	%r412, %r399;
	abs.f32 	%r413, %r398;
	abs.f32 	%r414, %r397;
	abs.f32 	%r415, %r396;
	abs.f32 	%r416, %r395;
	abs.f32 	%r417, %r394;
	abs.f32 	%r418, %r393;
	abs.f32 	%r419, %r392;
	abs.f32 	%r420, %r391;
	abs.f32 	%r421, %r390;
	abs.f32 	%r422, %r389;
	abs.f32 	%r423, %r388;
	abs.f32 	%r424, %r387;
	abs.f32 	%r425, %r386;
	abs.f32 	%r426, %r385;
	abs.f32 	%r427, %r384;
	abs.f32 	%r428, %r383;
	abs.f32 	%r429, %r382;
	abs.f32 	%r430, %r381;
	abs.f32 	%r431, %r380;
	abs.f32 	%r432, %r379;
	abs.f32 	%r433, %r378;
	abs.f32 	%r434, %r377;
	abs.f32 	%r435, %r376;
	abs.f32 	%r436, %r375;
	abs.f32 	%r437, %r374;
	abs.f32 	%r438, %r373;
	abs.f32 	%r439, %r372;
	abs.f32 	%r440, %r371;
	abs.f32 	%r441, %r370;
	abs.f32 	%r442, %r369;
	abs.f32 	%r443, %r368;
	abs.f32 	%r444, %r367;
	abs.f32 	%r445, %r366;
	abs.f32 	%r446, %r365;
	abs.f32 	%r447, %r364;
	abs.f32 	%r448, %r363;
	abs.f32 	%r449, %r362;
	abs.f32 	%r450, %r361;
	abs.f32 	%r451, %r360;
	abs.f32 	%r452, %r359;
	abs.f32 	%r453, %r358;
	abs.f32 	%r454, %r357;
	abs.f32 	%r455, %r356;
	abs.f32 	%r456, %r355;
	abs.f32 	%r457, %r354;
	abs.f32 	%r458, %r353;
	abs.f32 	%r459, %r352;
	abs.f32 	%r460, %r351;
	abs.f32 	%r461, %r350;
	abs.f32 	%r462, %r349;
	abs.f32 	%r463, %r348;
	abs.f32 	%r464, %r347;
	abs.f32 	%r465, %r346;
	abs.f32 	%r466, %r345;
	abs.f32 	%r467, %r344;
	abs.f32 	%r468, %r343;
	abs.f32 	%r469, %r342;
	abs.f32 	%r470, %r341;
	abs.f32 	%r471, %r340;
	abs.f32 	%r472, %r339;
	abs.f32 	%r473, %r338;
	abs.f32 	%r474, %r337;
	abs.f32 	%r475, %r336;
	abs.f32 	%r476, %r335;
	abs.f32 	%r477, %r334;
	abs.f32 	%r478, %r333;
	abs.f32 	%r479, %r332;
	abs.f32 	%r480, %r331;
	abs.f32 	%r481, %r330;
	abs.f32 	%r482, %r329;
	abs.f32 	%r483, %r328;
	abs.f32 	%r484, %r327;
	abs.f32 	%r485, %r326;
	abs.f32 	%r486, %r325;
	abs.f32 	%r487, %r324;
	abs.f32 	%r488, %r323;
	abs.f32 	%r489, %r322;
	abs.f32 	%r490, %r321;
	abs.f32 	%r491, %r320;
	abs.f32 	%r492, %r319;
	abs.f32 	%r493, %r318;
	abs.f32 	%r494, %r317;
	abs.f32 	%r495, %r316;
	abs.f32 	%r496, %r315;
	abs.f32 	%r497, %r314;
	abs.f32 	%r498, %r313;
	abs.f32 	%r499, %r312;
	abs.f32 	%r500, %r311;
	abs.f32 	%r501, %r310;
	abs.f32 	%r502, %r309;
	abs.f32 	%r503, %r308;
	abs.f32 	%r504, %r307;
	abs.f32 	%r505, %r306;
	abs.f32 	%r506, %r305;
	abs.f32 	%r507, %r304;
	abs.f32 	%r508, %r303;
	abs.f32 	%r509, %r302;
	abs.f32 	%r510, %r301;
	abs.f32 	%r511, %r300;
	abs.f32 	%r512, %r299;
	abs.f32 	%r513, %r298;
	abs.f32 	%r514, %r297;
	abs.f32 	%r515, %r296;
	abs.f32 	%r516, %r295;
	abs.f32 	%r517, %r294;
	abs.f32 	%r518, %r293;
	abs.f32 	%r519, %r292;
	abs.f32 	%r520, %r291;
	abs.f32 	%r521, %r290;
	abs.f32 	%r522, %r289;
	abs.f32 	%r523, %r288;
	abs.f32 	%r524, %r287;
	abs.f32 	%r525, %r286;
	abs.f32 	%r526, %r285;
	abs.f32 	%r527, %r284;
	abs.f32 	%r528, %r283;
	abs.f32 	%r529, %r282;
	abs.f32 	%r530, %r281;
	abs.f32 	%r531, %r280;
	abs.f32 	%r532, %r279;
	abs.f32 	%r533, %r278;
$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 	%r534, %r533, %r532;
	max.f32 	%r535, %r531, %r530;
	max.f32 	%r536, %r529, %r528;
	max.f32 	%r537, %r527, %r526;
	max.f32 	%r538, %r525, %r524;
	max.f32 	%r539, %r523, %r522;
	max.f32 	%r540, %r521, %r520;
	max.f32 	%r541, %r519, %r518;
	max.f32 	%r542, %r517, %r516;
	max.f32 	%r543, %r515, %r514;
	max.f32 	%r544, %r513, %r512;
	max.f32 	%r545, %r511, %r510;
	max.f32 	%r546, %r509, %r508;
	max.f32 	%r547, %r507, %r506;
	max.f32 	%r548, %r505, %r504;
	max.f32 	%r549, %r503, %r502;
	max.f32 	%r550, %r501, %r500;
	max.f32 	%r551, %r499, %r498;
	max.f32 	%r552, %r497, %r496;
	max.f32 	%r553, %r495, %r494;
	max.f32 	%r554, %r493, %r492;
	max.f32 	%r555, %r491, %r490;
	max.f32 	%r556, %r489, %r488;
	max.f32 	%r557, %r487, %r486;
	max.f32 	%r558, %r485, %r484;
	max.f32 	%r559, %r483, %r482;
	max.f32 	%r560, %r481, %r480;
	max.f32 	%r561, %r479, %r478;
	max.f32 	%r562, %r477, %r476;
	max.f32 	%r563, %r475, %r474;
	max.f32 	%r564, %r473, %r472;
	max.f32 	%r565, %r471, %r470;
	max.f32 	%r566, %r469, %r468;
	max.f32 	%r567, %r467, %r466;
	max.f32 	%r568, %r465, %r464;
	max.f32 	%r569, %r463, %r462;
	max.f32 	%r570, %r461, %r460;
	max.f32 	%r571, %r459, %r458;
	max.f32 	%r572, %r457, %r456;
	max.f32 	%r573, %r455, %r454;
	max.f32 	%r574, %r453, %r452;
	max.f32 	%r575, %r451, %r450;
	max.f32 	%r576, %r449, %r448;
	max.f32 	%r577, %r447, %r446;
	max.f32 	%r578, %r445, %r444;
	max.f32 	%r579, %r443, %r442;
	max.f32 	%r580, %r441, %r440;
	max.f32 	%r581, %r439, %r438;
	max.f32 	%r582, %r437, %r436;
	max.f32 	%r583, %r435, %r434;
	max.f32 	%r584, %r433, %r432;
	max.f32 	%r585, %r431, %r430;
	max.f32 	%r586, %r429, %r428;
	max.f32 	%r587, %r427, %r426;
	max.f32 	%r588, %r425, %r424;
	max.f32 	%r589, %r423, %r422;
	max.f32 	%r590, %r421, %r420;
	max.f32 	%r591, %r419, %r418;
	max.f32 	%r592, %r417, %r416;
	max.f32 	%r593, %r415, %r414;
	max.f32 	%r594, %r413, %r412;
	max.f32 	%r595, %r411, %r410;
	max.f32 	%r596, %r409, %r408;
	max.f32 	%r597, %r407, %r406;
	max.f32 	%r598, %r534, %r535;
	max.f32 	%r599, %r536, %r537;
	max.f32 	%r600, %r538, %r539;
	max.f32 	%r601, %r540, %r541;
	max.f32 	%r602, %r542, %r543;
	max.f32 	%r603, %r544, %r545;
	max.f32 	%r604, %r546, %r547;
	max.f32 	%r605, %r548, %r549;
	max.f32 	%r606, %r550, %r551;
	max.f32 	%r607, %r552, %r553;
	max.f32 	%r608, %r554, %r555;
	max.f32 	%r609, %r556, %r557;
	max.f32 	%r610, %r558, %r559;
	max.f32 	%r611, %r560, %r561;
	max.f32 	%r612, %r562, %r563;
	max.f32 	%r613, %r564, %r565;
	max.f32 	%r614, %r566, %r567;
	max.f32 	%r615, %r568, %r569;
	max.f32 	%r616, %r570, %r571;
	max.f32 	%r617, %r572, %r573;
	max.f32 	%r618, %r574, %r575;
	max.f32 	%r619, %r576, %r577;
	max.f32 	%r620, %r578, %r579;
	max.f32 	%r621, %r580, %r581;
	max.f32 	%r622, %r582, %r583;
	max.f32 	%r623, %r584, %r585;
	max.f32 	%r624, %r586, %r587;
	max.f32 	%r625, %r588, %r589;
	max.f32 	%r626, %r590, %r591;
	max.f32 	%r627, %r592, %r593;
	max.f32 	%r628, %r594, %r595;
	max.f32 	%r629, %r596, %r597;
	max.f32 	%r630, %r598, %r599;
	max.f32 	%r631, %r600, %r601;
	max.f32 	%r632, %r602, %r603;
	max.f32 	%r633, %r604, %r605;
	max.f32 	%r634, %r606, %r607;
	max.f32 	%r635, %r608, %r609;
	max.f32 	%r636, %r610, %r611;
	max.f32 	%r637, %r612, %r613;
	max.f32 	%r638, %r614, %r615;
	max.f32 	%r639, %r616, %r617;
	max.f32 	%r640, %r618, %r619;
	max.f32 	%r641, %r620, %r621;
	max.f32 	%r642, %r622, %r623;
	max.f32 	%r643, %r624, %r625;
	max.f32 	%r644, %r626, %r627;
	max.f32 	%r645, %r628, %r629;
	max.f32 	%r646, %r630, %r631;
	max.f32 	%r647, %r632, %r633;
	max.f32 	%r648, %r634, %r635;
	max.f32 	%r649, %r636, %r637;
	max.f32 	%r650, %r638, %r639;
	max.f32 	%r651, %r640, %r641;
	max.f32 	%r652, %r642, %r643;
	max.f32 	%r653, %r644, %r645;
	max.f32 	%r654, %r646, %r647;
	max.f32 	%r655, %r648, %r649;
	max.f32 	%r656, %r650, %r651;
	max.f32 	%r657, %r652, %r653;
	max.f32 	%r658, %r654, %r655;
	max.f32 	%r659, %r656, %r657;
	max.f32 	%r660, %r658, %r659;
$L__tmp2:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 26 13 // standard.py:191:16
	shfl.sync.bfly.b32 	%r661, %r660, 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 	%r662, %r660, %r661;
$L__tmp4:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 26 13 // standard.py:191:16
	shfl.sync.bfly.b32 	%r663, %r662, 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 	%r664, %r662, %r663;
$L__tmp6:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 26 13 // standard.py:191:16
	shfl.sync.bfly.b32 	%r665, %r664, 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 	%r666, %r664, %r665;
$L__tmp8:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 26 13 // standard.py:191:16
	shfl.sync.bfly.b32 	%r667, %r666, 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 	%r668, %r666, %r667;
$L__tmp10:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 26 13 // standard.py:191:16
	shfl.sync.bfly.b32 	%r669, %r668, 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 	%r66, %r668, %r669;
$L__tmp12:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 26 13 // standard.py:191:16
	and.b32 	%r670, %r108, 12;
	mov.b32 	%r671, global_smem;
	add.s32 	%r65, %r671, %r670;
	// 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 [ %r65 + 0 ], %r66;

	// end inline asm
	bar.sync 	0;
	shl.b32 	%r672, %r107, 2;
	and.b32 	%r673, %r672, 12;
	add.s32 	%r674, %r671, %r673;
	ld.shared.b32 	%r675, [%r674];
	shfl.sync.bfly.b32 	%r676, %r675, 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 	%r677, %r675, %r676;
$L__tmp14:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 26 13 // standard.py:191:16
	shfl.sync.bfly.b32 	%r678, %r677, 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 	%r679, %r677, %r678;
	mov.b32 	%r680, 0f42FE0000;
$L__tmp16:
	.loc	1 26 13                         // quant_per_block.py:26:13
	div.full.f32 	%r99, %r679, %r680;
	.loc	1 27 14                         // quant_per_block.py:27:14
	div.full.f32 	%r681, %r278, %r99;
	div.full.f32 	%r682, %r279, %r99;
	div.full.f32 	%r683, %r280, %r99;
	div.full.f32 	%r684, %r281, %r99;
	div.full.f32 	%r685, %r282, %r99;
	div.full.f32 	%r686, %r283, %r99;
	div.full.f32 	%r687, %r284, %r99;
	div.full.f32 	%r688, %r285, %r99;
	div.full.f32 	%r689, %r286, %r99;
	div.full.f32 	%r690, %r287, %r99;
	div.full.f32 	%r691, %r288, %r99;
	div.full.f32 	%r692, %r289, %r99;
	div.full.f32 	%r693, %r290, %r99;
	div.full.f32 	%r694, %r291, %r99;
	div.full.f32 	%r695, %r292, %r99;
	div.full.f32 	%r696, %r293, %r99;
	div.full.f32 	%r697, %r294, %r99;
	div.full.f32 	%r698, %r295, %r99;
	div.full.f32 	%r699, %r296, %r99;
	div.full.f32 	%r700, %r297, %r99;
	div.full.f32 	%r701, %r298, %r99;
	div.full.f32 	%r702, %r299, %r99;
	div.full.f32 	%r703, %r300, %r99;
	div.full.f32 	%r704, %r301, %r99;
	div.full.f32 	%r705, %r302, %r99;
	div.full.f32 	%r706, %r303, %r99;
	div.full.f32 	%r707, %r304, %r99;
	div.full.f32 	%r708, %r305, %r99;
	div.full.f32 	%r709, %r306, %r99;
	div.full.f32 	%r710, %r307, %r99;
	div.full.f32 	%r711, %r308, %r99;
	div.full.f32 	%r712, %r309, %r99;
	div.full.f32 	%r713, %r310, %r99;
	div.full.f32 	%r714, %r311, %r99;
	div.full.f32 	%r715, %r312, %r99;
	div.full.f32 	%r716, %r313, %r99;
	div.full.f32 	%r717, %r314, %r99;
	div.full.f32 	%r718, %r315, %r99;
	div.full.f32 	%r719, %r316, %r99;
	div.full.f32 	%r720, %r317, %r99;
	div.full.f32 	%r721, %r318, %r99;
	div.full.f32 	%r722, %r319, %r99;
	div.full.f32 	%r723, %r320, %r99;
	div.full.f32 	%r724, %r321, %r99;
	div.full.f32 	%r725, %r322, %r99;
	div.full.f32 	%r726, %r323, %r99;
	div.full.f32 	%r727, %r324, %r99;
	div.full.f32 	%r728, %r325, %r99;
	div.full.f32 	%r729, %r326, %r99;
	div.full.f32 	%r730, %r327, %r99;
	div.full.f32 	%r731, %r328, %r99;
	div.full.f32 	%r732, %r329, %r99;
	div.full.f32 	%r733, %r330, %r99;
	div.full.f32 	%r734, %r331, %r99;
	div.full.f32 	%r735, %r332, %r99;
	div.full.f32 	%r736, %r333, %r99;
	div.full.f32 	%r737, %r334, %r99;
	div.full.f32 	%r738, %r335, %r99;
	div.full.f32 	%r739, %r336, %r99;
	div.full.f32 	%r740, %r337, %r99;
	div.full.f32 	%r741, %r338, %r99;
	div.full.f32 	%r742, %r339, %r99;
	div.full.f32 	%r743, %r340, %r99;
	div.full.f32 	%r744, %r341, %r99;
	div.full.f32 	%r745, %r342, %r99;
	div.full.f32 	%r746, %r343, %r99;
	div.full.f32 	%r747, %r344, %r99;
	div.full.f32 	%r748, %r345, %r99;
	div.full.f32 	%r749, %r346, %r99;
	div.full.f32 	%r750, %r347, %r99;
	div.full.f32 	%r751, %r348, %r99;
	div.full.f32 	%r752, %r349, %r99;
	div.full.f32 	%r753, %r350, %r99;
	div.full.f32 	%r754, %r351, %r99;
	div.full.f32 	%r755, %r352, %r99;
	div.full.f32 	%r756, %r353, %r99;
	div.full.f32 	%r757, %r354, %r99;
	div.full.f32 	%r758, %r355, %r99;
	div.full.f32 	%r759, %r356, %r99;
	div.full.f32 	%r760, %r357, %r99;
	div.full.f32 	%r761, %r358, %r99;
	div.full.f32 	%r762, %r359, %r99;
	div.full.f32 	%r763, %r360, %r99;
	div.full.f32 	%r764, %r361, %r99;
	div.full.f32 	%r765, %r362, %r99;
	div.full.f32 	%r766, %r363, %r99;
	div.full.f32 	%r767, %r364, %r99;
	div.full.f32 	%r768, %r365, %r99;
	div.full.f32 	%r769, %r366, %r99;
	div.full.f32 	%r770, %r367, %r99;
	div.full.f32 	%r771, %r368, %r99;
	div.full.f32 	%r772, %r369, %r99;
	div.full.f32 	%r773, %r370, %r99;
	div.full.f32 	%r774, %r371, %r99;
	div.full.f32 	%r775, %r372, %r99;
	div.full.f32 	%r776, %r373, %r99;
	div.full.f32 	%r777, %r374, %r99;
	div.full.f32 	%r778, %r375, %r99;
	div.full.f32 	%r779, %r376, %r99;
	div.full.f32 	%r780, %r377, %r99;
	div.full.f32 	%r781, %r378, %r99;
	div.full.f32 	%r782, %r379, %r99;
	div.full.f32 	%r783, %r380, %r99;
	div.full.f32 	%r784, %r381, %r99;
	div.full.f32 	%r785, %r382, %r99;
	div.full.f32 	%r786, %r383, %r99;
	div.full.f32 	%r787, %r384, %r99;
	div.full.f32 	%r788, %r385, %r99;
	div.full.f32 	%r789, %r386, %r99;
	div.full.f32 	%r790, %r387, %r99;
	div.full.f32 	%r791, %r388, %r99;
	div.full.f32 	%r792, %r389, %r99;
	div.full.f32 	%r793, %r390, %r99;
	div.full.f32 	%r794, %r391, %r99;
	div.full.f32 	%r795, %r392, %r99;
	div.full.f32 	%r796, %r393, %r99;
	div.full.f32 	%r797, %r394, %r99;
	div.full.f32 	%r798, %r395, %r99;
	div.full.f32 	%r799, %r396, %r99;
	div.full.f32 	%r800, %r397, %r99;
	div.full.f32 	%r801, %r398, %r99;
	div.full.f32 	%r802, %r399, %r99;
	div.full.f32 	%r803, %r400, %r99;
	div.full.f32 	%r804, %r401, %r99;
	div.full.f32 	%r805, %r402, %r99;
	div.full.f32 	%r806, %r403, %r99;
	div.full.f32 	%r807, %r404, %r99;
	div.full.f32 	%r808, %r405, %r99;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p10, %r681, 0f00000000;
	setp.ge.f32 	%p11, %r682, 0f00000000;
	setp.ge.f32 	%p12, %r683, 0f00000000;
	setp.ge.f32 	%p13, %r684, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r809, 0f3F000000, 0fBF000000, %p13;
	selp.f32 	%r810, 0f3F000000, 0fBF000000, %p12;
	selp.f32 	%r811, 0f3F000000, 0fBF000000, %p11;
	selp.f32 	%r812, 0f3F000000, 0fBF000000, %p10;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r813, %r681, %r812;
	add.f32 	%r814, %r682, %r811;
	add.f32 	%r815, %r683, %r810;
	add.f32 	%r816, %r684, %r809;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs129, %r816;
	cvt.rzi.s16.f32 	%rs130, %r815;
	mov.b32 	%r817, {%rs130, %rs129};
	mov.b32 	{%rs131, %rs132}, %r817;
	cvt.u32.u16 	%r818, %rs132;
	cvt.u32.u16 	%r819, %rs131;
	prmt.b32 	%r820, %r819, %r818, 0x3340U;
	cvt.rzi.s16.f32 	%rs133, %r814;
	cvt.rzi.s16.f32 	%rs134, %r813;
	mov.b32 	%r821, {%rs134, %rs133};
	mov.b32 	{%rs135, %rs136}, %r821;
	cvt.u32.u16 	%r822, %rs136;
	cvt.u32.u16 	%r823, %rs135;
	prmt.b32 	%r824, %r823, %r822, 0x3340U;
	prmt.b32 	%r67, %r824, %r820, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p14, %r685, 0f00000000;
	setp.ge.f32 	%p15, %r686, 0f00000000;
	setp.ge.f32 	%p16, %r687, 0f00000000;
	setp.ge.f32 	%p17, %r688, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r825, 0f3F000000, 0fBF000000, %p17;
	selp.f32 	%r826, 0f3F000000, 0fBF000000, %p16;
	selp.f32 	%r827, 0f3F000000, 0fBF000000, %p15;
	selp.f32 	%r828, 0f3F000000, 0fBF000000, %p14;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r829, %r685, %r828;
	add.f32 	%r830, %r686, %r827;
	add.f32 	%r831, %r687, %r826;
	add.f32 	%r832, %r688, %r825;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs137, %r832;
	cvt.rzi.s16.f32 	%rs138, %r831;
	mov.b32 	%r833, {%rs138, %rs137};
	mov.b32 	{%rs139, %rs140}, %r833;
	cvt.u32.u16 	%r834, %rs140;
	cvt.u32.u16 	%r835, %rs139;
	prmt.b32 	%r836, %r835, %r834, 0x3340U;
	cvt.rzi.s16.f32 	%rs141, %r830;
	cvt.rzi.s16.f32 	%rs142, %r829;
	mov.b32 	%r837, {%rs142, %rs141};
	mov.b32 	{%rs143, %rs144}, %r837;
	cvt.u32.u16 	%r838, %rs144;
	cvt.u32.u16 	%r839, %rs143;
	prmt.b32 	%r840, %r839, %r838, 0x3340U;
	prmt.b32 	%r68, %r840, %r836, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p18, %r689, 0f00000000;
	setp.ge.f32 	%p19, %r690, 0f00000000;
	setp.ge.f32 	%p20, %r691, 0f00000000;
	setp.ge.f32 	%p21, %r692, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r841, 0f3F000000, 0fBF000000, %p21;
	selp.f32 	%r842, 0f3F000000, 0fBF000000, %p20;
	selp.f32 	%r843, 0f3F000000, 0fBF000000, %p19;
	selp.f32 	%r844, 0f3F000000, 0fBF000000, %p18;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r845, %r689, %r844;
	add.f32 	%r846, %r690, %r843;
	add.f32 	%r847, %r691, %r842;
	add.f32 	%r848, %r692, %r841;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs145, %r848;
	cvt.rzi.s16.f32 	%rs146, %r847;
	mov.b32 	%r849, {%rs146, %rs145};
	mov.b32 	{%rs147, %rs148}, %r849;
	cvt.u32.u16 	%r850, %rs148;
	cvt.u32.u16 	%r851, %rs147;
	prmt.b32 	%r852, %r851, %r850, 0x3340U;
	cvt.rzi.s16.f32 	%rs149, %r846;
	cvt.rzi.s16.f32 	%rs150, %r845;
	mov.b32 	%r853, {%rs150, %rs149};
	mov.b32 	{%rs151, %rs152}, %r853;
	cvt.u32.u16 	%r854, %rs152;
	cvt.u32.u16 	%r855, %rs151;
	prmt.b32 	%r856, %r855, %r854, 0x3340U;
	prmt.b32 	%r69, %r856, %r852, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p22, %r693, 0f00000000;
	setp.ge.f32 	%p23, %r694, 0f00000000;
	setp.ge.f32 	%p24, %r695, 0f00000000;
	setp.ge.f32 	%p25, %r696, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r857, 0f3F000000, 0fBF000000, %p25;
	selp.f32 	%r858, 0f3F000000, 0fBF000000, %p24;
	selp.f32 	%r859, 0f3F000000, 0fBF000000, %p23;
	selp.f32 	%r860, 0f3F000000, 0fBF000000, %p22;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r861, %r693, %r860;
	add.f32 	%r862, %r694, %r859;
	add.f32 	%r863, %r695, %r858;
	add.f32 	%r864, %r696, %r857;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs153, %r864;
	cvt.rzi.s16.f32 	%rs154, %r863;
	mov.b32 	%r865, {%rs154, %rs153};
	mov.b32 	{%rs155, %rs156}, %r865;
	cvt.u32.u16 	%r866, %rs156;
	cvt.u32.u16 	%r867, %rs155;
	prmt.b32 	%r868, %r867, %r866, 0x3340U;
	cvt.rzi.s16.f32 	%rs157, %r862;
	cvt.rzi.s16.f32 	%rs158, %r861;
	mov.b32 	%r869, {%rs158, %rs157};
	mov.b32 	{%rs159, %rs160}, %r869;
	cvt.u32.u16 	%r870, %rs160;
	cvt.u32.u16 	%r871, %rs159;
	prmt.b32 	%r872, %r871, %r870, 0x3340U;
	prmt.b32 	%r70, %r872, %r868, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p26, %r697, 0f00000000;
	setp.ge.f32 	%p27, %r698, 0f00000000;
	setp.ge.f32 	%p28, %r699, 0f00000000;
	setp.ge.f32 	%p29, %r700, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r873, 0f3F000000, 0fBF000000, %p29;
	selp.f32 	%r874, 0f3F000000, 0fBF000000, %p28;
	selp.f32 	%r875, 0f3F000000, 0fBF000000, %p27;
	selp.f32 	%r876, 0f3F000000, 0fBF000000, %p26;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r877, %r697, %r876;
	add.f32 	%r878, %r698, %r875;
	add.f32 	%r879, %r699, %r874;
	add.f32 	%r880, %r700, %r873;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs161, %r880;
	cvt.rzi.s16.f32 	%rs162, %r879;
	mov.b32 	%r881, {%rs162, %rs161};
	mov.b32 	{%rs163, %rs164}, %r881;
	cvt.u32.u16 	%r882, %rs164;
	cvt.u32.u16 	%r883, %rs163;
	prmt.b32 	%r884, %r883, %r882, 0x3340U;
	cvt.rzi.s16.f32 	%rs165, %r878;
	cvt.rzi.s16.f32 	%rs166, %r877;
	mov.b32 	%r885, {%rs166, %rs165};
	mov.b32 	{%rs167, %rs168}, %r885;
	cvt.u32.u16 	%r886, %rs168;
	cvt.u32.u16 	%r887, %rs167;
	prmt.b32 	%r888, %r887, %r886, 0x3340U;
	prmt.b32 	%r71, %r888, %r884, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p30, %r701, 0f00000000;
	setp.ge.f32 	%p31, %r702, 0f00000000;
	setp.ge.f32 	%p32, %r703, 0f00000000;
	setp.ge.f32 	%p33, %r704, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r889, 0f3F000000, 0fBF000000, %p33;
	selp.f32 	%r890, 0f3F000000, 0fBF000000, %p32;
	selp.f32 	%r891, 0f3F000000, 0fBF000000, %p31;
	selp.f32 	%r892, 0f3F000000, 0fBF000000, %p30;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r893, %r701, %r892;
	add.f32 	%r894, %r702, %r891;
	add.f32 	%r895, %r703, %r890;
	add.f32 	%r896, %r704, %r889;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs169, %r896;
	cvt.rzi.s16.f32 	%rs170, %r895;
	mov.b32 	%r897, {%rs170, %rs169};
	mov.b32 	{%rs171, %rs172}, %r897;
	cvt.u32.u16 	%r898, %rs172;
	cvt.u32.u16 	%r899, %rs171;
	prmt.b32 	%r900, %r899, %r898, 0x3340U;
	cvt.rzi.s16.f32 	%rs173, %r894;
	cvt.rzi.s16.f32 	%rs174, %r893;
	mov.b32 	%r901, {%rs174, %rs173};
	mov.b32 	{%rs175, %rs176}, %r901;
	cvt.u32.u16 	%r902, %rs176;
	cvt.u32.u16 	%r903, %rs175;
	prmt.b32 	%r904, %r903, %r902, 0x3340U;
	prmt.b32 	%r72, %r904, %r900, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p34, %r705, 0f00000000;
	setp.ge.f32 	%p35, %r706, 0f00000000;
	setp.ge.f32 	%p36, %r707, 0f00000000;
	setp.ge.f32 	%p37, %r708, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r905, 0f3F000000, 0fBF000000, %p37;
	selp.f32 	%r906, 0f3F000000, 0fBF000000, %p36;
	selp.f32 	%r907, 0f3F000000, 0fBF000000, %p35;
	selp.f32 	%r908, 0f3F000000, 0fBF000000, %p34;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r909, %r705, %r908;
	add.f32 	%r910, %r706, %r907;
	add.f32 	%r911, %r707, %r906;
	add.f32 	%r912, %r708, %r905;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs177, %r912;
	cvt.rzi.s16.f32 	%rs178, %r911;
	mov.b32 	%r913, {%rs178, %rs177};
	mov.b32 	{%rs179, %rs180}, %r913;
	cvt.u32.u16 	%r914, %rs180;
	cvt.u32.u16 	%r915, %rs179;
	prmt.b32 	%r916, %r915, %r914, 0x3340U;
	cvt.rzi.s16.f32 	%rs181, %r910;
	cvt.rzi.s16.f32 	%rs182, %r909;
	mov.b32 	%r917, {%rs182, %rs181};
	mov.b32 	{%rs183, %rs184}, %r917;
	cvt.u32.u16 	%r918, %rs184;
	cvt.u32.u16 	%r919, %rs183;
	prmt.b32 	%r920, %r919, %r918, 0x3340U;
	prmt.b32 	%r73, %r920, %r916, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p38, %r709, 0f00000000;
	setp.ge.f32 	%p39, %r710, 0f00000000;
	setp.ge.f32 	%p40, %r711, 0f00000000;
	setp.ge.f32 	%p41, %r712, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r921, 0f3F000000, 0fBF000000, %p41;
	selp.f32 	%r922, 0f3F000000, 0fBF000000, %p40;
	selp.f32 	%r923, 0f3F000000, 0fBF000000, %p39;
	selp.f32 	%r924, 0f3F000000, 0fBF000000, %p38;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r925, %r709, %r924;
	add.f32 	%r926, %r710, %r923;
	add.f32 	%r927, %r711, %r922;
	add.f32 	%r928, %r712, %r921;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs185, %r928;
	cvt.rzi.s16.f32 	%rs186, %r927;
	mov.b32 	%r929, {%rs186, %rs185};
	mov.b32 	{%rs187, %rs188}, %r929;
	cvt.u32.u16 	%r930, %rs188;
	cvt.u32.u16 	%r931, %rs187;
	prmt.b32 	%r932, %r931, %r930, 0x3340U;
	cvt.rzi.s16.f32 	%rs189, %r926;
	cvt.rzi.s16.f32 	%rs190, %r925;
	mov.b32 	%r933, {%rs190, %rs189};
	mov.b32 	{%rs191, %rs192}, %r933;
	cvt.u32.u16 	%r934, %rs192;
	cvt.u32.u16 	%r935, %rs191;
	prmt.b32 	%r936, %r935, %r934, 0x3340U;
	prmt.b32 	%r74, %r936, %r932, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p42, %r713, 0f00000000;
	setp.ge.f32 	%p43, %r714, 0f00000000;
	setp.ge.f32 	%p44, %r715, 0f00000000;
	setp.ge.f32 	%p45, %r716, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r937, 0f3F000000, 0fBF000000, %p45;
	selp.f32 	%r938, 0f3F000000, 0fBF000000, %p44;
	selp.f32 	%r939, 0f3F000000, 0fBF000000, %p43;
	selp.f32 	%r940, 0f3F000000, 0fBF000000, %p42;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r941, %r713, %r940;
	add.f32 	%r942, %r714, %r939;
	add.f32 	%r943, %r715, %r938;
	add.f32 	%r944, %r716, %r937;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs193, %r944;
	cvt.rzi.s16.f32 	%rs194, %r943;
	mov.b32 	%r945, {%rs194, %rs193};
	mov.b32 	{%rs195, %rs196}, %r945;
	cvt.u32.u16 	%r946, %rs196;
	cvt.u32.u16 	%r947, %rs195;
	prmt.b32 	%r948, %r947, %r946, 0x3340U;
	cvt.rzi.s16.f32 	%rs197, %r942;
	cvt.rzi.s16.f32 	%rs198, %r941;
	mov.b32 	%r949, {%rs198, %rs197};
	mov.b32 	{%rs199, %rs200}, %r949;
	cvt.u32.u16 	%r950, %rs200;
	cvt.u32.u16 	%r951, %rs199;
	prmt.b32 	%r952, %r951, %r950, 0x3340U;
	prmt.b32 	%r75, %r952, %r948, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p46, %r717, 0f00000000;
	setp.ge.f32 	%p47, %r718, 0f00000000;
	setp.ge.f32 	%p48, %r719, 0f00000000;
	setp.ge.f32 	%p49, %r720, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r953, 0f3F000000, 0fBF000000, %p49;
	selp.f32 	%r954, 0f3F000000, 0fBF000000, %p48;
	selp.f32 	%r955, 0f3F000000, 0fBF000000, %p47;
	selp.f32 	%r956, 0f3F000000, 0fBF000000, %p46;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r957, %r717, %r956;
	add.f32 	%r958, %r718, %r955;
	add.f32 	%r959, %r719, %r954;
	add.f32 	%r960, %r720, %r953;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs201, %r960;
	cvt.rzi.s16.f32 	%rs202, %r959;
	mov.b32 	%r961, {%rs202, %rs201};
	mov.b32 	{%rs203, %rs204}, %r961;
	cvt.u32.u16 	%r962, %rs204;
	cvt.u32.u16 	%r963, %rs203;
	prmt.b32 	%r964, %r963, %r962, 0x3340U;
	cvt.rzi.s16.f32 	%rs205, %r958;
	cvt.rzi.s16.f32 	%rs206, %r957;
	mov.b32 	%r965, {%rs206, %rs205};
	mov.b32 	{%rs207, %rs208}, %r965;
	cvt.u32.u16 	%r966, %rs208;
	cvt.u32.u16 	%r967, %rs207;
	prmt.b32 	%r968, %r967, %r966, 0x3340U;
	prmt.b32 	%r76, %r968, %r964, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p50, %r721, 0f00000000;
	setp.ge.f32 	%p51, %r722, 0f00000000;
	setp.ge.f32 	%p52, %r723, 0f00000000;
	setp.ge.f32 	%p53, %r724, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r969, 0f3F000000, 0fBF000000, %p53;
	selp.f32 	%r970, 0f3F000000, 0fBF000000, %p52;
	selp.f32 	%r971, 0f3F000000, 0fBF000000, %p51;
	selp.f32 	%r972, 0f3F000000, 0fBF000000, %p50;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r973, %r721, %r972;
	add.f32 	%r974, %r722, %r971;
	add.f32 	%r975, %r723, %r970;
	add.f32 	%r976, %r724, %r969;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs209, %r976;
	cvt.rzi.s16.f32 	%rs210, %r975;
	mov.b32 	%r977, {%rs210, %rs209};
	mov.b32 	{%rs211, %rs212}, %r977;
	cvt.u32.u16 	%r978, %rs212;
	cvt.u32.u16 	%r979, %rs211;
	prmt.b32 	%r980, %r979, %r978, 0x3340U;
	cvt.rzi.s16.f32 	%rs213, %r974;
	cvt.rzi.s16.f32 	%rs214, %r973;
	mov.b32 	%r981, {%rs214, %rs213};
	mov.b32 	{%rs215, %rs216}, %r981;
	cvt.u32.u16 	%r982, %rs216;
	cvt.u32.u16 	%r983, %rs215;
	prmt.b32 	%r984, %r983, %r982, 0x3340U;
	prmt.b32 	%r77, %r984, %r980, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p54, %r725, 0f00000000;
	setp.ge.f32 	%p55, %r726, 0f00000000;
	setp.ge.f32 	%p56, %r727, 0f00000000;
	setp.ge.f32 	%p57, %r728, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r985, 0f3F000000, 0fBF000000, %p57;
	selp.f32 	%r986, 0f3F000000, 0fBF000000, %p56;
	selp.f32 	%r987, 0f3F000000, 0fBF000000, %p55;
	selp.f32 	%r988, 0f3F000000, 0fBF000000, %p54;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r989, %r725, %r988;
	add.f32 	%r990, %r726, %r987;
	add.f32 	%r991, %r727, %r986;
	add.f32 	%r992, %r728, %r985;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs217, %r992;
	cvt.rzi.s16.f32 	%rs218, %r991;
	mov.b32 	%r993, {%rs218, %rs217};
	mov.b32 	{%rs219, %rs220}, %r993;
	cvt.u32.u16 	%r994, %rs220;
	cvt.u32.u16 	%r995, %rs219;
	prmt.b32 	%r996, %r995, %r994, 0x3340U;
	cvt.rzi.s16.f32 	%rs221, %r990;
	cvt.rzi.s16.f32 	%rs222, %r989;
	mov.b32 	%r997, {%rs222, %rs221};
	mov.b32 	{%rs223, %rs224}, %r997;
	cvt.u32.u16 	%r998, %rs224;
	cvt.u32.u16 	%r999, %rs223;
	prmt.b32 	%r1000, %r999, %r998, 0x3340U;
	prmt.b32 	%r78, %r1000, %r996, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p58, %r729, 0f00000000;
	setp.ge.f32 	%p59, %r730, 0f00000000;
	setp.ge.f32 	%p60, %r731, 0f00000000;
	setp.ge.f32 	%p61, %r732, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r1001, 0f3F000000, 0fBF000000, %p61;
	selp.f32 	%r1002, 0f3F000000, 0fBF000000, %p60;
	selp.f32 	%r1003, 0f3F000000, 0fBF000000, %p59;
	selp.f32 	%r1004, 0f3F000000, 0fBF000000, %p58;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r1005, %r729, %r1004;
	add.f32 	%r1006, %r730, %r1003;
	add.f32 	%r1007, %r731, %r1002;
	add.f32 	%r1008, %r732, %r1001;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs225, %r1008;
	cvt.rzi.s16.f32 	%rs226, %r1007;
	mov.b32 	%r1009, {%rs226, %rs225};
	mov.b32 	{%rs227, %rs228}, %r1009;
	cvt.u32.u16 	%r1010, %rs228;
	cvt.u32.u16 	%r1011, %rs227;
	prmt.b32 	%r1012, %r1011, %r1010, 0x3340U;
	cvt.rzi.s16.f32 	%rs229, %r1006;
	cvt.rzi.s16.f32 	%rs230, %r1005;
	mov.b32 	%r1013, {%rs230, %rs229};
	mov.b32 	{%rs231, %rs232}, %r1013;
	cvt.u32.u16 	%r1014, %rs232;
	cvt.u32.u16 	%r1015, %rs231;
	prmt.b32 	%r1016, %r1015, %r1014, 0x3340U;
	prmt.b32 	%r79, %r1016, %r1012, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p62, %r733, 0f00000000;
	setp.ge.f32 	%p63, %r734, 0f00000000;
	setp.ge.f32 	%p64, %r735, 0f00000000;
	setp.ge.f32 	%p65, %r736, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r1017, 0f3F000000, 0fBF000000, %p65;
	selp.f32 	%r1018, 0f3F000000, 0fBF000000, %p64;
	selp.f32 	%r1019, 0f3F000000, 0fBF000000, %p63;
	selp.f32 	%r1020, 0f3F000000, 0fBF000000, %p62;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r1021, %r733, %r1020;
	add.f32 	%r1022, %r734, %r1019;
	add.f32 	%r1023, %r735, %r1018;
	add.f32 	%r1024, %r736, %r1017;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs233, %r1024;
	cvt.rzi.s16.f32 	%rs234, %r1023;
	mov.b32 	%r1025, {%rs234, %rs233};
	mov.b32 	{%rs235, %rs236}, %r1025;
	cvt.u32.u16 	%r1026, %rs236;
	cvt.u32.u16 	%r1027, %rs235;
	prmt.b32 	%r1028, %r1027, %r1026, 0x3340U;
	cvt.rzi.s16.f32 	%rs237, %r1022;
	cvt.rzi.s16.f32 	%rs238, %r1021;
	mov.b32 	%r1029, {%rs238, %rs237};
	mov.b32 	{%rs239, %rs240}, %r1029;
	cvt.u32.u16 	%r1030, %rs240;
	cvt.u32.u16 	%r1031, %rs239;
	prmt.b32 	%r1032, %r1031, %r1030, 0x3340U;
	prmt.b32 	%r80, %r1032, %r1028, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p66, %r737, 0f00000000;
	setp.ge.f32 	%p67, %r738, 0f00000000;
	setp.ge.f32 	%p68, %r739, 0f00000000;
	setp.ge.f32 	%p69, %r740, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r1033, 0f3F000000, 0fBF000000, %p69;
	selp.f32 	%r1034, 0f3F000000, 0fBF000000, %p68;
	selp.f32 	%r1035, 0f3F000000, 0fBF000000, %p67;
	selp.f32 	%r1036, 0f3F000000, 0fBF000000, %p66;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r1037, %r737, %r1036;
	add.f32 	%r1038, %r738, %r1035;
	add.f32 	%r1039, %r739, %r1034;
	add.f32 	%r1040, %r740, %r1033;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs241, %r1040;
	cvt.rzi.s16.f32 	%rs242, %r1039;
	mov.b32 	%r1041, {%rs242, %rs241};
	mov.b32 	{%rs243, %rs244}, %r1041;
	cvt.u32.u16 	%r1042, %rs244;
	cvt.u32.u16 	%r1043, %rs243;
	prmt.b32 	%r1044, %r1043, %r1042, 0x3340U;
	cvt.rzi.s16.f32 	%rs245, %r1038;
	cvt.rzi.s16.f32 	%rs246, %r1037;
	mov.b32 	%r1045, {%rs246, %rs245};
	mov.b32 	{%rs247, %rs248}, %r1045;
	cvt.u32.u16 	%r1046, %rs248;
	cvt.u32.u16 	%r1047, %rs247;
	prmt.b32 	%r1048, %r1047, %r1046, 0x3340U;
	prmt.b32 	%r81, %r1048, %r1044, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p70, %r741, 0f00000000;
	setp.ge.f32 	%p71, %r742, 0f00000000;
	setp.ge.f32 	%p72, %r743, 0f00000000;
	setp.ge.f32 	%p73, %r744, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r1049, 0f3F000000, 0fBF000000, %p73;
	selp.f32 	%r1050, 0f3F000000, 0fBF000000, %p72;
	selp.f32 	%r1051, 0f3F000000, 0fBF000000, %p71;
	selp.f32 	%r1052, 0f3F000000, 0fBF000000, %p70;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r1053, %r741, %r1052;
	add.f32 	%r1054, %r742, %r1051;
	add.f32 	%r1055, %r743, %r1050;
	add.f32 	%r1056, %r744, %r1049;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs249, %r1056;
	cvt.rzi.s16.f32 	%rs250, %r1055;
	mov.b32 	%r1057, {%rs250, %rs249};
	mov.b32 	{%rs251, %rs252}, %r1057;
	cvt.u32.u16 	%r1058, %rs252;
	cvt.u32.u16 	%r1059, %rs251;
	prmt.b32 	%r1060, %r1059, %r1058, 0x3340U;
	cvt.rzi.s16.f32 	%rs253, %r1054;
	cvt.rzi.s16.f32 	%rs254, %r1053;
	mov.b32 	%r1061, {%rs254, %rs253};
	mov.b32 	{%rs255, %rs256}, %r1061;
	cvt.u32.u16 	%r1062, %rs256;
	cvt.u32.u16 	%r1063, %rs255;
	prmt.b32 	%r1064, %r1063, %r1062, 0x3340U;
	prmt.b32 	%r82, %r1064, %r1060, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p74, %r745, 0f00000000;
	setp.ge.f32 	%p75, %r746, 0f00000000;
	setp.ge.f32 	%p76, %r747, 0f00000000;
	setp.ge.f32 	%p77, %r748, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r1065, 0f3F000000, 0fBF000000, %p77;
	selp.f32 	%r1066, 0f3F000000, 0fBF000000, %p76;
	selp.f32 	%r1067, 0f3F000000, 0fBF000000, %p75;
	selp.f32 	%r1068, 0f3F000000, 0fBF000000, %p74;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r1069, %r745, %r1068;
	add.f32 	%r1070, %r746, %r1067;
	add.f32 	%r1071, %r747, %r1066;
	add.f32 	%r1072, %r748, %r1065;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs257, %r1072;
	cvt.rzi.s16.f32 	%rs258, %r1071;
	mov.b32 	%r1073, {%rs258, %rs257};
	mov.b32 	{%rs259, %rs260}, %r1073;
	cvt.u32.u16 	%r1074, %rs260;
	cvt.u32.u16 	%r1075, %rs259;
	prmt.b32 	%r1076, %r1075, %r1074, 0x3340U;
	cvt.rzi.s16.f32 	%rs261, %r1070;
	cvt.rzi.s16.f32 	%rs262, %r1069;
	mov.b32 	%r1077, {%rs262, %rs261};
	mov.b32 	{%rs263, %rs264}, %r1077;
	cvt.u32.u16 	%r1078, %rs264;
	cvt.u32.u16 	%r1079, %rs263;
	prmt.b32 	%r1080, %r1079, %r1078, 0x3340U;
	prmt.b32 	%r83, %r1080, %r1076, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p78, %r749, 0f00000000;
	setp.ge.f32 	%p79, %r750, 0f00000000;
	setp.ge.f32 	%p80, %r751, 0f00000000;
	setp.ge.f32 	%p81, %r752, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r1081, 0f3F000000, 0fBF000000, %p81;
	selp.f32 	%r1082, 0f3F000000, 0fBF000000, %p80;
	selp.f32 	%r1083, 0f3F000000, 0fBF000000, %p79;
	selp.f32 	%r1084, 0f3F000000, 0fBF000000, %p78;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r1085, %r749, %r1084;
	add.f32 	%r1086, %r750, %r1083;
	add.f32 	%r1087, %r751, %r1082;
	add.f32 	%r1088, %r752, %r1081;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs265, %r1088;
	cvt.rzi.s16.f32 	%rs266, %r1087;
	mov.b32 	%r1089, {%rs266, %rs265};
	mov.b32 	{%rs267, %rs268}, %r1089;
	cvt.u32.u16 	%r1090, %rs268;
	cvt.u32.u16 	%r1091, %rs267;
	prmt.b32 	%r1092, %r1091, %r1090, 0x3340U;
	cvt.rzi.s16.f32 	%rs269, %r1086;
	cvt.rzi.s16.f32 	%rs270, %r1085;
	mov.b32 	%r1093, {%rs270, %rs269};
	mov.b32 	{%rs271, %rs272}, %r1093;
	cvt.u32.u16 	%r1094, %rs272;
	cvt.u32.u16 	%r1095, %rs271;
	prmt.b32 	%r1096, %r1095, %r1094, 0x3340U;
	prmt.b32 	%r84, %r1096, %r1092, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p82, %r753, 0f00000000;
	setp.ge.f32 	%p83, %r754, 0f00000000;
	setp.ge.f32 	%p84, %r755, 0f00000000;
	setp.ge.f32 	%p85, %r756, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r1097, 0f3F000000, 0fBF000000, %p85;
	selp.f32 	%r1098, 0f3F000000, 0fBF000000, %p84;
	selp.f32 	%r1099, 0f3F000000, 0fBF000000, %p83;
	selp.f32 	%r1100, 0f3F000000, 0fBF000000, %p82;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r1101, %r753, %r1100;
	add.f32 	%r1102, %r754, %r1099;
	add.f32 	%r1103, %r755, %r1098;
	add.f32 	%r1104, %r756, %r1097;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs273, %r1104;
	cvt.rzi.s16.f32 	%rs274, %r1103;
	mov.b32 	%r1105, {%rs274, %rs273};
	mov.b32 	{%rs275, %rs276}, %r1105;
	cvt.u32.u16 	%r1106, %rs276;
	cvt.u32.u16 	%r1107, %rs275;
	prmt.b32 	%r1108, %r1107, %r1106, 0x3340U;
	cvt.rzi.s16.f32 	%rs277, %r1102;
	cvt.rzi.s16.f32 	%rs278, %r1101;
	mov.b32 	%r1109, {%rs278, %rs277};
	mov.b32 	{%rs279, %rs280}, %r1109;
	cvt.u32.u16 	%r1110, %rs280;
	cvt.u32.u16 	%r1111, %rs279;
	prmt.b32 	%r1112, %r1111, %r1110, 0x3340U;
	prmt.b32 	%r85, %r1112, %r1108, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p86, %r757, 0f00000000;
	setp.ge.f32 	%p87, %r758, 0f00000000;
	setp.ge.f32 	%p88, %r759, 0f00000000;
	setp.ge.f32 	%p89, %r760, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r1113, 0f3F000000, 0fBF000000, %p89;
	selp.f32 	%r1114, 0f3F000000, 0fBF000000, %p88;
	selp.f32 	%r1115, 0f3F000000, 0fBF000000, %p87;
	selp.f32 	%r1116, 0f3F000000, 0fBF000000, %p86;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r1117, %r757, %r1116;
	add.f32 	%r1118, %r758, %r1115;
	add.f32 	%r1119, %r759, %r1114;
	add.f32 	%r1120, %r760, %r1113;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs281, %r1120;
	cvt.rzi.s16.f32 	%rs282, %r1119;
	mov.b32 	%r1121, {%rs282, %rs281};
	mov.b32 	{%rs283, %rs284}, %r1121;
	cvt.u32.u16 	%r1122, %rs284;
	cvt.u32.u16 	%r1123, %rs283;
	prmt.b32 	%r1124, %r1123, %r1122, 0x3340U;
	cvt.rzi.s16.f32 	%rs285, %r1118;
	cvt.rzi.s16.f32 	%rs286, %r1117;
	mov.b32 	%r1125, {%rs286, %rs285};
	mov.b32 	{%rs287, %rs288}, %r1125;
	cvt.u32.u16 	%r1126, %rs288;
	cvt.u32.u16 	%r1127, %rs287;
	prmt.b32 	%r1128, %r1127, %r1126, 0x3340U;
	prmt.b32 	%r86, %r1128, %r1124, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p90, %r761, 0f00000000;
	setp.ge.f32 	%p91, %r762, 0f00000000;
	setp.ge.f32 	%p92, %r763, 0f00000000;
	setp.ge.f32 	%p93, %r764, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r1129, 0f3F000000, 0fBF000000, %p93;
	selp.f32 	%r1130, 0f3F000000, 0fBF000000, %p92;
	selp.f32 	%r1131, 0f3F000000, 0fBF000000, %p91;
	selp.f32 	%r1132, 0f3F000000, 0fBF000000, %p90;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r1133, %r761, %r1132;
	add.f32 	%r1134, %r762, %r1131;
	add.f32 	%r1135, %r763, %r1130;
	add.f32 	%r1136, %r764, %r1129;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs289, %r1136;
	cvt.rzi.s16.f32 	%rs290, %r1135;
	mov.b32 	%r1137, {%rs290, %rs289};
	mov.b32 	{%rs291, %rs292}, %r1137;
	cvt.u32.u16 	%r1138, %rs292;
	cvt.u32.u16 	%r1139, %rs291;
	prmt.b32 	%r1140, %r1139, %r1138, 0x3340U;
	cvt.rzi.s16.f32 	%rs293, %r1134;
	cvt.rzi.s16.f32 	%rs294, %r1133;
	mov.b32 	%r1141, {%rs294, %rs293};
	mov.b32 	{%rs295, %rs296}, %r1141;
	cvt.u32.u16 	%r1142, %rs296;
	cvt.u32.u16 	%r1143, %rs295;
	prmt.b32 	%r1144, %r1143, %r1142, 0x3340U;
	prmt.b32 	%r87, %r1144, %r1140, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p94, %r765, 0f00000000;
	setp.ge.f32 	%p95, %r766, 0f00000000;
	setp.ge.f32 	%p96, %r767, 0f00000000;
	setp.ge.f32 	%p97, %r768, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r1145, 0f3F000000, 0fBF000000, %p97;
	selp.f32 	%r1146, 0f3F000000, 0fBF000000, %p96;
	selp.f32 	%r1147, 0f3F000000, 0fBF000000, %p95;
	selp.f32 	%r1148, 0f3F000000, 0fBF000000, %p94;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r1149, %r765, %r1148;
	add.f32 	%r1150, %r766, %r1147;
	add.f32 	%r1151, %r767, %r1146;
	add.f32 	%r1152, %r768, %r1145;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs297, %r1152;
	cvt.rzi.s16.f32 	%rs298, %r1151;
	mov.b32 	%r1153, {%rs298, %rs297};
	mov.b32 	{%rs299, %rs300}, %r1153;
	cvt.u32.u16 	%r1154, %rs300;
	cvt.u32.u16 	%r1155, %rs299;
	prmt.b32 	%r1156, %r1155, %r1154, 0x3340U;
	cvt.rzi.s16.f32 	%rs301, %r1150;
	cvt.rzi.s16.f32 	%rs302, %r1149;
	mov.b32 	%r1157, {%rs302, %rs301};
	mov.b32 	{%rs303, %rs304}, %r1157;
	cvt.u32.u16 	%r1158, %rs304;
	cvt.u32.u16 	%r1159, %rs303;
	prmt.b32 	%r1160, %r1159, %r1158, 0x3340U;
	prmt.b32 	%r88, %r1160, %r1156, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p98, %r769, 0f00000000;
	setp.ge.f32 	%p99, %r770, 0f00000000;
	setp.ge.f32 	%p100, %r771, 0f00000000;
	setp.ge.f32 	%p101, %r772, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r1161, 0f3F000000, 0fBF000000, %p101;
	selp.f32 	%r1162, 0f3F000000, 0fBF000000, %p100;
	selp.f32 	%r1163, 0f3F000000, 0fBF000000, %p99;
	selp.f32 	%r1164, 0f3F000000, 0fBF000000, %p98;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r1165, %r769, %r1164;
	add.f32 	%r1166, %r770, %r1163;
	add.f32 	%r1167, %r771, %r1162;
	add.f32 	%r1168, %r772, %r1161;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs305, %r1168;
	cvt.rzi.s16.f32 	%rs306, %r1167;
	mov.b32 	%r1169, {%rs306, %rs305};
	mov.b32 	{%rs307, %rs308}, %r1169;
	cvt.u32.u16 	%r1170, %rs308;
	cvt.u32.u16 	%r1171, %rs307;
	prmt.b32 	%r1172, %r1171, %r1170, 0x3340U;
	cvt.rzi.s16.f32 	%rs309, %r1166;
	cvt.rzi.s16.f32 	%rs310, %r1165;
	mov.b32 	%r1173, {%rs310, %rs309};
	mov.b32 	{%rs311, %rs312}, %r1173;
	cvt.u32.u16 	%r1174, %rs312;
	cvt.u32.u16 	%r1175, %rs311;
	prmt.b32 	%r1176, %r1175, %r1174, 0x3340U;
	prmt.b32 	%r89, %r1176, %r1172, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p102, %r773, 0f00000000;
	setp.ge.f32 	%p103, %r774, 0f00000000;
	setp.ge.f32 	%p104, %r775, 0f00000000;
	setp.ge.f32 	%p105, %r776, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r1177, 0f3F000000, 0fBF000000, %p105;
	selp.f32 	%r1178, 0f3F000000, 0fBF000000, %p104;
	selp.f32 	%r1179, 0f3F000000, 0fBF000000, %p103;
	selp.f32 	%r1180, 0f3F000000, 0fBF000000, %p102;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r1181, %r773, %r1180;
	add.f32 	%r1182, %r774, %r1179;
	add.f32 	%r1183, %r775, %r1178;
	add.f32 	%r1184, %r776, %r1177;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs313, %r1184;
	cvt.rzi.s16.f32 	%rs314, %r1183;
	mov.b32 	%r1185, {%rs314, %rs313};
	mov.b32 	{%rs315, %rs316}, %r1185;
	cvt.u32.u16 	%r1186, %rs316;
	cvt.u32.u16 	%r1187, %rs315;
	prmt.b32 	%r1188, %r1187, %r1186, 0x3340U;
	cvt.rzi.s16.f32 	%rs317, %r1182;
	cvt.rzi.s16.f32 	%rs318, %r1181;
	mov.b32 	%r1189, {%rs318, %rs317};
	mov.b32 	{%rs319, %rs320}, %r1189;
	cvt.u32.u16 	%r1190, %rs320;
	cvt.u32.u16 	%r1191, %rs319;
	prmt.b32 	%r1192, %r1191, %r1190, 0x3340U;
	prmt.b32 	%r90, %r1192, %r1188, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p106, %r777, 0f00000000;
	setp.ge.f32 	%p107, %r778, 0f00000000;
	setp.ge.f32 	%p108, %r779, 0f00000000;
	setp.ge.f32 	%p109, %r780, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r1193, 0f3F000000, 0fBF000000, %p109;
	selp.f32 	%r1194, 0f3F000000, 0fBF000000, %p108;
	selp.f32 	%r1195, 0f3F000000, 0fBF000000, %p107;
	selp.f32 	%r1196, 0f3F000000, 0fBF000000, %p106;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r1197, %r777, %r1196;
	add.f32 	%r1198, %r778, %r1195;
	add.f32 	%r1199, %r779, %r1194;
	add.f32 	%r1200, %r780, %r1193;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs321, %r1200;
	cvt.rzi.s16.f32 	%rs322, %r1199;
	mov.b32 	%r1201, {%rs322, %rs321};
	mov.b32 	{%rs323, %rs324}, %r1201;
	cvt.u32.u16 	%r1202, %rs324;
	cvt.u32.u16 	%r1203, %rs323;
	prmt.b32 	%r1204, %r1203, %r1202, 0x3340U;
	cvt.rzi.s16.f32 	%rs325, %r1198;
	cvt.rzi.s16.f32 	%rs326, %r1197;
	mov.b32 	%r1205, {%rs326, %rs325};
	mov.b32 	{%rs327, %rs328}, %r1205;
	cvt.u32.u16 	%r1206, %rs328;
	cvt.u32.u16 	%r1207, %rs327;
	prmt.b32 	%r1208, %r1207, %r1206, 0x3340U;
	prmt.b32 	%r91, %r1208, %r1204, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p110, %r781, 0f00000000;
	setp.ge.f32 	%p111, %r782, 0f00000000;
	setp.ge.f32 	%p112, %r783, 0f00000000;
	setp.ge.f32 	%p113, %r784, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r1209, 0f3F000000, 0fBF000000, %p113;
	selp.f32 	%r1210, 0f3F000000, 0fBF000000, %p112;
	selp.f32 	%r1211, 0f3F000000, 0fBF000000, %p111;
	selp.f32 	%r1212, 0f3F000000, 0fBF000000, %p110;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r1213, %r781, %r1212;
	add.f32 	%r1214, %r782, %r1211;
	add.f32 	%r1215, %r783, %r1210;
	add.f32 	%r1216, %r784, %r1209;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs329, %r1216;
	cvt.rzi.s16.f32 	%rs330, %r1215;
	mov.b32 	%r1217, {%rs330, %rs329};
	mov.b32 	{%rs331, %rs332}, %r1217;
	cvt.u32.u16 	%r1218, %rs332;
	cvt.u32.u16 	%r1219, %rs331;
	prmt.b32 	%r1220, %r1219, %r1218, 0x3340U;
	cvt.rzi.s16.f32 	%rs333, %r1214;
	cvt.rzi.s16.f32 	%rs334, %r1213;
	mov.b32 	%r1221, {%rs334, %rs333};
	mov.b32 	{%rs335, %rs336}, %r1221;
	cvt.u32.u16 	%r1222, %rs336;
	cvt.u32.u16 	%r1223, %rs335;
	prmt.b32 	%r1224, %r1223, %r1222, 0x3340U;
	prmt.b32 	%r92, %r1224, %r1220, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p114, %r785, 0f00000000;
	setp.ge.f32 	%p115, %r786, 0f00000000;
	setp.ge.f32 	%p116, %r787, 0f00000000;
	setp.ge.f32 	%p117, %r788, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r1225, 0f3F000000, 0fBF000000, %p117;
	selp.f32 	%r1226, 0f3F000000, 0fBF000000, %p116;
	selp.f32 	%r1227, 0f3F000000, 0fBF000000, %p115;
	selp.f32 	%r1228, 0f3F000000, 0fBF000000, %p114;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r1229, %r785, %r1228;
	add.f32 	%r1230, %r786, %r1227;
	add.f32 	%r1231, %r787, %r1226;
	add.f32 	%r1232, %r788, %r1225;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs337, %r1232;
	cvt.rzi.s16.f32 	%rs338, %r1231;
	mov.b32 	%r1233, {%rs338, %rs337};
	mov.b32 	{%rs339, %rs340}, %r1233;
	cvt.u32.u16 	%r1234, %rs340;
	cvt.u32.u16 	%r1235, %rs339;
	prmt.b32 	%r1236, %r1235, %r1234, 0x3340U;
	cvt.rzi.s16.f32 	%rs341, %r1230;
	cvt.rzi.s16.f32 	%rs342, %r1229;
	mov.b32 	%r1237, {%rs342, %rs341};
	mov.b32 	{%rs343, %rs344}, %r1237;
	cvt.u32.u16 	%r1238, %rs344;
	cvt.u32.u16 	%r1239, %rs343;
	prmt.b32 	%r1240, %r1239, %r1238, 0x3340U;
	prmt.b32 	%r93, %r1240, %r1236, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p118, %r789, 0f00000000;
	setp.ge.f32 	%p119, %r790, 0f00000000;
	setp.ge.f32 	%p120, %r791, 0f00000000;
	setp.ge.f32 	%p121, %r792, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r1241, 0f3F000000, 0fBF000000, %p121;
	selp.f32 	%r1242, 0f3F000000, 0fBF000000, %p120;
	selp.f32 	%r1243, 0f3F000000, 0fBF000000, %p119;
	selp.f32 	%r1244, 0f3F000000, 0fBF000000, %p118;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r1245, %r789, %r1244;
	add.f32 	%r1246, %r790, %r1243;
	add.f32 	%r1247, %r791, %r1242;
	add.f32 	%r1248, %r792, %r1241;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs345, %r1248;
	cvt.rzi.s16.f32 	%rs346, %r1247;
	mov.b32 	%r1249, {%rs346, %rs345};
	mov.b32 	{%rs347, %rs348}, %r1249;
	cvt.u32.u16 	%r1250, %rs348;
	cvt.u32.u16 	%r1251, %rs347;
	prmt.b32 	%r1252, %r1251, %r1250, 0x3340U;
	cvt.rzi.s16.f32 	%rs349, %r1246;
	cvt.rzi.s16.f32 	%rs350, %r1245;
	mov.b32 	%r1253, {%rs350, %rs349};
	mov.b32 	{%rs351, %rs352}, %r1253;
	cvt.u32.u16 	%r1254, %rs352;
	cvt.u32.u16 	%r1255, %rs351;
	prmt.b32 	%r1256, %r1255, %r1254, 0x3340U;
	prmt.b32 	%r94, %r1256, %r1252, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p122, %r793, 0f00000000;
	setp.ge.f32 	%p123, %r794, 0f00000000;
	setp.ge.f32 	%p124, %r795, 0f00000000;
	setp.ge.f32 	%p125, %r796, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r1257, 0f3F000000, 0fBF000000, %p125;
	selp.f32 	%r1258, 0f3F000000, 0fBF000000, %p124;
	selp.f32 	%r1259, 0f3F000000, 0fBF000000, %p123;
	selp.f32 	%r1260, 0f3F000000, 0fBF000000, %p122;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r1261, %r793, %r1260;
	add.f32 	%r1262, %r794, %r1259;
	add.f32 	%r1263, %r795, %r1258;
	add.f32 	%r1264, %r796, %r1257;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs353, %r1264;
	cvt.rzi.s16.f32 	%rs354, %r1263;
	mov.b32 	%r1265, {%rs354, %rs353};
	mov.b32 	{%rs355, %rs356}, %r1265;
	cvt.u32.u16 	%r1266, %rs356;
	cvt.u32.u16 	%r1267, %rs355;
	prmt.b32 	%r1268, %r1267, %r1266, 0x3340U;
	cvt.rzi.s16.f32 	%rs357, %r1262;
	cvt.rzi.s16.f32 	%rs358, %r1261;
	mov.b32 	%r1269, {%rs358, %rs357};
	mov.b32 	{%rs359, %rs360}, %r1269;
	cvt.u32.u16 	%r1270, %rs360;
	cvt.u32.u16 	%r1271, %rs359;
	prmt.b32 	%r1272, %r1271, %r1270, 0x3340U;
	prmt.b32 	%r95, %r1272, %r1268, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p126, %r797, 0f00000000;
	setp.ge.f32 	%p127, %r798, 0f00000000;
	setp.ge.f32 	%p128, %r799, 0f00000000;
	setp.ge.f32 	%p129, %r800, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r1273, 0f3F000000, 0fBF000000, %p129;
	selp.f32 	%r1274, 0f3F000000, 0fBF000000, %p128;
	selp.f32 	%r1275, 0f3F000000, 0fBF000000, %p127;
	selp.f32 	%r1276, 0f3F000000, 0fBF000000, %p126;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r1277, %r797, %r1276;
	add.f32 	%r1278, %r798, %r1275;
	add.f32 	%r1279, %r799, %r1274;
	add.f32 	%r1280, %r800, %r1273;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs361, %r1280;
	cvt.rzi.s16.f32 	%rs362, %r1279;
	mov.b32 	%r1281, {%rs362, %rs361};
	mov.b32 	{%rs363, %rs364}, %r1281;
	cvt.u32.u16 	%r1282, %rs364;
	cvt.u32.u16 	%r1283, %rs363;
	prmt.b32 	%r1284, %r1283, %r1282, 0x3340U;
	cvt.rzi.s16.f32 	%rs365, %r1278;
	cvt.rzi.s16.f32 	%rs366, %r1277;
	mov.b32 	%r1285, {%rs366, %rs365};
	mov.b32 	{%rs367, %rs368}, %r1285;
	cvt.u32.u16 	%r1286, %rs368;
	cvt.u32.u16 	%r1287, %rs367;
	prmt.b32 	%r1288, %r1287, %r1286, 0x3340U;
	prmt.b32 	%r96, %r1288, %r1284, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p130, %r801, 0f00000000;
	setp.ge.f32 	%p131, %r802, 0f00000000;
	setp.ge.f32 	%p132, %r803, 0f00000000;
	setp.ge.f32 	%p133, %r804, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r1289, 0f3F000000, 0fBF000000, %p133;
	selp.f32 	%r1290, 0f3F000000, 0fBF000000, %p132;
	selp.f32 	%r1291, 0f3F000000, 0fBF000000, %p131;
	selp.f32 	%r1292, 0f3F000000, 0fBF000000, %p130;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r1293, %r801, %r1292;
	add.f32 	%r1294, %r802, %r1291;
	add.f32 	%r1295, %r803, %r1290;
	add.f32 	%r1296, %r804, %r1289;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs369, %r1296;
	cvt.rzi.s16.f32 	%rs370, %r1295;
	mov.b32 	%r1297, {%rs370, %rs369};
	mov.b32 	{%rs371, %rs372}, %r1297;
	cvt.u32.u16 	%r1298, %rs372;
	cvt.u32.u16 	%r1299, %rs371;
	prmt.b32 	%r1300, %r1299, %r1298, 0x3340U;
	cvt.rzi.s16.f32 	%rs373, %r1294;
	cvt.rzi.s16.f32 	%rs374, %r1293;
	mov.b32 	%r1301, {%rs374, %rs373};
	mov.b32 	{%rs375, %rs376}, %r1301;
	cvt.u32.u16 	%r1302, %rs376;
	cvt.u32.u16 	%r1303, %rs375;
	prmt.b32 	%r1304, %r1303, %r1302, 0x3340U;
	prmt.b32 	%r97, %r1304, %r1300, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p134, %r805, 0f00000000;
	setp.ge.f32 	%p135, %r806, 0f00000000;
	setp.ge.f32 	%p136, %r807, 0f00000000;
	setp.ge.f32 	%p137, %r808, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r1305, 0f3F000000, 0fBF000000, %p137;
	selp.f32 	%r1306, 0f3F000000, 0fBF000000, %p136;
	selp.f32 	%r1307, 0f3F000000, 0fBF000000, %p135;
	selp.f32 	%r1308, 0f3F000000, 0fBF000000, %p134;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r1309, %r805, %r1308;
	add.f32 	%r1310, %r806, %r1307;
	add.f32 	%r1311, %r807, %r1306;
	add.f32 	%r1312, %r808, %r1305;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs377, %r1312;
	cvt.rzi.s16.f32 	%rs378, %r1311;
	mov.b32 	%r1313, {%rs378, %rs377};
	mov.b32 	{%rs379, %rs380}, %r1313;
	cvt.u32.u16 	%r1314, %rs380;
	cvt.u32.u16 	%r1315, %rs379;
	prmt.b32 	%r1316, %r1315, %r1314, 0x3340U;
	cvt.rzi.s16.f32 	%rs381, %r1310;
	cvt.rzi.s16.f32 	%rs382, %r1309;
	mov.b32 	%r1317, {%rs382, %rs381};
	mov.b32 	{%rs383, %rs384}, %r1317;
	cvt.u32.u16 	%r1318, %rs384;
	cvt.u32.u16 	%r1319, %rs383;
	prmt.b32 	%r1320, %r1319, %r1318, 0x3340U;
	prmt.b32 	%r98, %r1320, %r1316, 0x5410U;
	.loc	1 30 5                          // quant_per_block.py:30:5
	// begin inline asm
	@%p1 st.global.v4.b32 [ %rd17 + 0 ], { %r67, %r68, %r69, %r70 };

	// end inline asm
	// begin inline asm
	@%p2 st.global.v4.b32 [ %rd18 + 0 ], { %r71, %r72, %r73, %r74 };

	// end inline asm
	// begin inline asm
	@%p3 st.global.v4.b32 [ %rd19 + 0 ], { %r75, %r76, %r77, %r78 };

	// end inline asm
	// begin inline asm
	@%p4 st.global.v4.b32 [ %rd20 + 0 ], { %r79, %r80, %r81, %r82 };

	// end inline asm
	// begin inline asm
	@%p5 st.global.v4.b32 [ %rd21 + 0 ], { %r83, %r84, %r85, %r86 };

	// end inline asm
	// begin inline asm
	@%p6 st.global.v4.b32 [ %rd22 + 0 ], { %r87, %r88, %r89, %r90 };

	// end inline asm
	// begin inline asm
	@%p7 st.global.v4.b32 [ %rd23 + 0 ], { %r91, %r92, %r93, %r94 };

	// end inline asm
	// begin inline asm
	@%p8 st.global.v4.b32 [ %rd24 + 0 ], { %r95, %r96, %r97, %r98 };

	// end inline asm
	.loc	1 31 5                          // quant_per_block.py:31:5
	setp.eq.b32 	%p9, %r107, 0;
	// begin inline asm
	@%p9 st.global.b32 [ %rd25 + 0 ], { %r99 };

	// 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	{	}