//
// 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<36>;
	.reg .b16 	%rs<97>;
	.reg .b32 	%r<365>;
	.reg .b64 	%rd<27>;
	.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 	%rd8, [quant_per_block_int8_kernel_param_0];
	ld.param.b64 	%rd9, [quant_per_block_int8_kernel_param_1];
$L__tmp0:
	.loc	1 12 15                         // quant_per_block.py:12:15
	mov.u32 	%r28, %ctaid.x;
	ld.param.b64 	%rd10, [quant_per_block_int8_kernel_param_2];
	.loc	1 13 13                         // quant_per_block.py:13:13
	mov.u32 	%r29, %ctaid.y;
	ld.param.b32 	%r30, [quant_per_block_int8_kernel_param_3];
	.loc	1 14 13                         // quant_per_block.py:14:13
	mov.u32 	%r31, %ctaid.z;
	.loc	1 16 14                         // quant_per_block.py:16:14
	shl.b32 	%r32, %r28, 6;
	ld.param.b32 	%r33, [quant_per_block_int8_kernel_param_4];
	ld.param.b32 	%r34, [quant_per_block_int8_kernel_param_5];
	.loc	1 16 30                         // quant_per_block.py:16:30
	mov.u32 	%r35, %tid.x;
	ld.param.b32 	%r36, [quant_per_block_int8_kernel_param_6];
	bfe.u32 	%r37, %r35, 2, 5;
	ld.param.b32 	%r38, [quant_per_block_int8_kernel_param_7];
	.loc	1 16 14                         // quant_per_block.py:16:14
	or.b32 	%r39, %r37, %r32;
	ld.param.b32 	%r40, [quant_per_block_int8_kernel_param_8];
	or.b32 	%r41, %r39, 32;
	ld.param.b32 	%r42, [quant_per_block_int8_kernel_param_9];
	.loc	1 19 26                         // quant_per_block.py:19:26
	mul.lo.s32 	%r43, %r33, %r31;
	ld.param.b32 	%r44, [quant_per_block_int8_kernel_param_10];
	ld.param.b32 	%r45, [quant_per_block_int8_kernel_param_11];
	.loc	1 19 18                         // quant_per_block.py:19:18
	mad.wide.s32 	%rd11, %r43, 2, %rd8;
	ld.param.b32 	%r46, [quant_per_block_int8_kernel_param_12];
	.loc	1 19 46                         // quant_per_block.py:19:46
	mul.lo.s32 	%r47, %r34, %r29;
	.loc	1 19 18                         // quant_per_block.py:19:18
	mad.wide.s32 	%rd12, %r47, 2, %rd11;
	.loc	1 19 66                         // quant_per_block.py:19:66
	mul.lo.s32 	%r48, %r36, %r39;
	mul.lo.s32 	%r49, %r36, %r41;
	.loc	1 19 18                         // quant_per_block.py:19:18
	mad.wide.s32 	%rd13, %r48, 2, %rd12;
	mad.wide.s32 	%rd14, %r49, 2, %rd12;
	.loc	1 19 96                         // quant_per_block.py:19:96
	and.b32 	%r50, %r35, 3;
	shl.b32 	%r51, %r50, 4;
	.loc	1 19 18                         // quant_per_block.py:19:18
	cvt.u64.u32 	%rd15, %r51;
	mul.wide.u32 	%rd16, %r51, 2;
	add.s64 	%rd1, %rd13, %rd16;
	add.s64 	%rd2, %rd1, 16;
	add.s64 	%rd3, %rd14, %rd16;
	add.s64 	%rd4, %rd3, 16;
	.loc	1 20 28                         // quant_per_block.py:20:28
	mul.lo.s32 	%r52, %r38, %r31;
	.loc	1 20 19                         // quant_per_block.py:20:19
	cvt.s64.s32 	%rd17, %r52;
	add.s64 	%rd18, %rd9, %rd17;
	.loc	1 20 48                         // quant_per_block.py:20:48
	mul.lo.s32 	%r53, %r40, %r29;
	.loc	1 20 19                         // quant_per_block.py:20:19
	cvt.s64.s32 	%rd19, %r53;
	add.s64 	%rd20, %rd18, %rd19;
	.loc	1 20 68                         // quant_per_block.py:20:68
	mul.lo.s32 	%r54, %r42, %r39;
	mul.lo.s32 	%r55, %r42, %r41;
	.loc	1 20 19                         // quant_per_block.py:20:19
	cvt.s64.s32 	%rd21, %r54;
	add.s64 	%rd22, %rd20, %rd21;
	cvt.s64.s32 	%rd23, %r55;
	add.s64 	%rd24, %rd20, %rd23;
	add.s64 	%rd5, %rd22, %rd15;
	add.s64 	%rd6, %rd24, %rd15;
	.loc	1 21 26                         // quant_per_block.py:21:26
	mul.lo.s32 	%r56, %r44, %r31;
	.loc	1 21 18                         // quant_per_block.py:21:18
	mad.wide.s32 	%rd25, %r56, 4, %rd10;
	.loc	1 21 46                         // quant_per_block.py:21:46
	mul.lo.s32 	%r57, %r45, %r29;
	.loc	1 21 18                         // quant_per_block.py:21:18
	mad.wide.s32 	%rd26, %r57, 4, %rd25;
	mad.wide.u32 	%rd7, %r28, 4, %rd26;
	.loc	1 23 34                         // quant_per_block.py:23:34
	setp.lt.s32 	%p1, %r39, %r30;
	setp.lt.s32 	%p2, %r41, %r30;
	.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
	.loc	1 24 9                          // quant_per_block.py:24:9
	mov.b32 	{%rs1, %rs2}, %r16;
	cvt.f32.bf16 	%r58, %rs2;
	cvt.f32.bf16 	%r59, %rs1;
	mov.b32 	{%rs3, %rs4}, %r15;
	cvt.f32.bf16 	%r60, %rs4;
	cvt.f32.bf16 	%r61, %rs3;
	mov.b32 	{%rs5, %rs6}, %r14;
	cvt.f32.bf16 	%r62, %rs6;
	cvt.f32.bf16 	%r63, %rs5;
	mov.b32 	{%rs7, %rs8}, %r13;
	cvt.f32.bf16 	%r64, %rs8;
	cvt.f32.bf16 	%r65, %rs7;
	mov.b32 	{%rs9, %rs10}, %r12;
	cvt.f32.bf16 	%r66, %rs10;
	cvt.f32.bf16 	%r67, %rs9;
	mov.b32 	{%rs11, %rs12}, %r11;
	cvt.f32.bf16 	%r68, %rs12;
	cvt.f32.bf16 	%r69, %rs11;
	mov.b32 	{%rs13, %rs14}, %r10;
	cvt.f32.bf16 	%r70, %rs14;
	cvt.f32.bf16 	%r71, %rs13;
	mov.b32 	{%rs15, %rs16}, %r9;
	cvt.f32.bf16 	%r72, %rs16;
	cvt.f32.bf16 	%r73, %rs15;
	mov.b32 	{%rs17, %rs18}, %r8;
	cvt.f32.bf16 	%r74, %rs18;
	cvt.f32.bf16 	%r75, %rs17;
	mov.b32 	{%rs19, %rs20}, %r7;
	cvt.f32.bf16 	%r76, %rs20;
	cvt.f32.bf16 	%r77, %rs19;
	mov.b32 	{%rs21, %rs22}, %r6;
	cvt.f32.bf16 	%r78, %rs22;
	cvt.f32.bf16 	%r79, %rs21;
	mov.b32 	{%rs23, %rs24}, %r5;
	cvt.f32.bf16 	%r80, %rs24;
	cvt.f32.bf16 	%r81, %rs23;
	mov.b32 	{%rs25, %rs26}, %r4;
	cvt.f32.bf16 	%r82, %rs26;
	cvt.f32.bf16 	%r83, %rs25;
	mov.b32 	{%rs27, %rs28}, %r3;
	cvt.f32.bf16 	%r84, %rs28;
	cvt.f32.bf16 	%r85, %rs27;
	mov.b32 	{%rs29, %rs30}, %r2;
	cvt.f32.bf16 	%r86, %rs30;
	cvt.f32.bf16 	%r87, %rs29;
	mov.b32 	{%rs31, %rs32}, %r1;
	cvt.f32.bf16 	%r88, %rs32;
	cvt.f32.bf16 	%r89, %rs31;
	.loc	1 25 5                          // quant_per_block.py:25:5
	mul.f32 	%r90, %r46, %r89;
	mul.f32 	%r91, %r46, %r88;
	mul.f32 	%r92, %r46, %r87;
	mul.f32 	%r93, %r46, %r86;
	mul.f32 	%r94, %r46, %r85;
	mul.f32 	%r95, %r46, %r84;
	mul.f32 	%r96, %r46, %r83;
	mul.f32 	%r97, %r46, %r82;
	mul.f32 	%r98, %r46, %r81;
	mul.f32 	%r99, %r46, %r80;
	mul.f32 	%r100, %r46, %r79;
	mul.f32 	%r101, %r46, %r78;
	mul.f32 	%r102, %r46, %r77;
	mul.f32 	%r103, %r46, %r76;
	mul.f32 	%r104, %r46, %r75;
	mul.f32 	%r105, %r46, %r74;
	mul.f32 	%r106, %r46, %r73;
	mul.f32 	%r107, %r46, %r72;
	mul.f32 	%r108, %r46, %r71;
	mul.f32 	%r109, %r46, %r70;
	mul.f32 	%r110, %r46, %r69;
	mul.f32 	%r111, %r46, %r68;
	mul.f32 	%r112, %r46, %r67;
	mul.f32 	%r113, %r46, %r66;
	mul.f32 	%r114, %r46, %r65;
	mul.f32 	%r115, %r46, %r64;
	mul.f32 	%r116, %r46, %r63;
	mul.f32 	%r117, %r46, %r62;
	mul.f32 	%r118, %r46, %r61;
	mul.f32 	%r119, %r46, %r60;
	mul.f32 	%r120, %r46, %r59;
	mul.f32 	%r121, %r46, %r58;
	.loc	1 26 20                         // quant_per_block.py:26:20
	abs.f32 	%r122, %r121;
	abs.f32 	%r123, %r120;
	abs.f32 	%r124, %r119;
	abs.f32 	%r125, %r118;
	abs.f32 	%r126, %r117;
	abs.f32 	%r127, %r116;
	abs.f32 	%r128, %r115;
	abs.f32 	%r129, %r114;
	abs.f32 	%r130, %r113;
	abs.f32 	%r131, %r112;
	abs.f32 	%r132, %r111;
	abs.f32 	%r133, %r110;
	abs.f32 	%r134, %r109;
	abs.f32 	%r135, %r108;
	abs.f32 	%r136, %r107;
	abs.f32 	%r137, %r106;
	abs.f32 	%r138, %r105;
	abs.f32 	%r139, %r104;
	abs.f32 	%r140, %r103;
	abs.f32 	%r141, %r102;
	abs.f32 	%r142, %r101;
	abs.f32 	%r143, %r100;
	abs.f32 	%r144, %r99;
	abs.f32 	%r145, %r98;
	abs.f32 	%r146, %r97;
	abs.f32 	%r147, %r96;
	abs.f32 	%r148, %r95;
	abs.f32 	%r149, %r94;
	abs.f32 	%r150, %r93;
	abs.f32 	%r151, %r92;
	abs.f32 	%r152, %r91;
	abs.f32 	%r153, %r90;
$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 	%r154, %r153, %r152;
	max.f32 	%r155, %r151, %r150;
	max.f32 	%r156, %r149, %r148;
	max.f32 	%r157, %r147, %r146;
	max.f32 	%r158, %r145, %r144;
	max.f32 	%r159, %r143, %r142;
	max.f32 	%r160, %r141, %r140;
	max.f32 	%r161, %r139, %r138;
	max.f32 	%r162, %r137, %r136;
	max.f32 	%r163, %r135, %r134;
	max.f32 	%r164, %r133, %r132;
	max.f32 	%r165, %r131, %r130;
	max.f32 	%r166, %r129, %r128;
	max.f32 	%r167, %r127, %r126;
	max.f32 	%r168, %r125, %r124;
	max.f32 	%r169, %r123, %r122;
	max.f32 	%r170, %r154, %r155;
	max.f32 	%r171, %r156, %r157;
	max.f32 	%r172, %r158, %r159;
	max.f32 	%r173, %r160, %r161;
	max.f32 	%r174, %r162, %r163;
	max.f32 	%r175, %r164, %r165;
	max.f32 	%r176, %r166, %r167;
	max.f32 	%r177, %r168, %r169;
	max.f32 	%r178, %r170, %r171;
	max.f32 	%r179, %r172, %r173;
	max.f32 	%r180, %r174, %r175;
	max.f32 	%r181, %r176, %r177;
	max.f32 	%r182, %r178, %r179;
	max.f32 	%r183, %r180, %r181;
	max.f32 	%r184, %r182, %r183;
$L__tmp2:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 26 13 // standard.py:191:16
	shfl.sync.bfly.b32 	%r185, %r184, 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 	%r186, %r184, %r185;
$L__tmp4:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 26 13 // standard.py:191:16
	shfl.sync.bfly.b32 	%r187, %r186, 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 	%r188, %r186, %r187;
$L__tmp6:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 26 13 // standard.py:191:16
	shfl.sync.bfly.b32 	%r189, %r188, 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 	%r190, %r188, %r189;
$L__tmp8:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 26 13 // standard.py:191:16
	shfl.sync.bfly.b32 	%r191, %r190, 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 	%r192, %r190, %r191;
$L__tmp10:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 26 13 // standard.py:191:16
	shfl.sync.bfly.b32 	%r193, %r192, 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 	%r18, %r192, %r193;
$L__tmp12:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 26 13 // standard.py:191:16
	shr.u32 	%r194, %r35, 3;
	and.b32 	%r195, %r194, 12;
	mov.b32 	%r196, global_smem;
	add.s32 	%r17, %r196, %r195;
	// 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 [ %r17 + 0 ], %r18;

	// end inline asm
	bar.sync 	0;
	shl.b32 	%r197, %r50, 2;
	add.s32 	%r198, %r196, %r197;
	ld.shared.b32 	%r199, [%r198];
	shfl.sync.bfly.b32 	%r200, %r199, 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 	%r201, %r199, %r200;
$L__tmp14:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 26 13 // standard.py:191:16
	shfl.sync.bfly.b32 	%r202, %r201, 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 	%r203, %r201, %r202;
	mov.b32 	%r204, 0f42FE0000;
$L__tmp16:
	.loc	1 26 13                         // quant_per_block.py:26:13
	div.full.f32 	%r27, %r203, %r204;
	.loc	1 27 14                         // quant_per_block.py:27:14
	div.full.f32 	%r205, %r90, %r27;
	div.full.f32 	%r206, %r91, %r27;
	div.full.f32 	%r207, %r92, %r27;
	div.full.f32 	%r208, %r93, %r27;
	div.full.f32 	%r209, %r94, %r27;
	div.full.f32 	%r210, %r95, %r27;
	div.full.f32 	%r211, %r96, %r27;
	div.full.f32 	%r212, %r97, %r27;
	div.full.f32 	%r213, %r98, %r27;
	div.full.f32 	%r214, %r99, %r27;
	div.full.f32 	%r215, %r100, %r27;
	div.full.f32 	%r216, %r101, %r27;
	div.full.f32 	%r217, %r102, %r27;
	div.full.f32 	%r218, %r103, %r27;
	div.full.f32 	%r219, %r104, %r27;
	div.full.f32 	%r220, %r105, %r27;
	div.full.f32 	%r221, %r106, %r27;
	div.full.f32 	%r222, %r107, %r27;
	div.full.f32 	%r223, %r108, %r27;
	div.full.f32 	%r224, %r109, %r27;
	div.full.f32 	%r225, %r110, %r27;
	div.full.f32 	%r226, %r111, %r27;
	div.full.f32 	%r227, %r112, %r27;
	div.full.f32 	%r228, %r113, %r27;
	div.full.f32 	%r229, %r114, %r27;
	div.full.f32 	%r230, %r115, %r27;
	div.full.f32 	%r231, %r116, %r27;
	div.full.f32 	%r232, %r117, %r27;
	div.full.f32 	%r233, %r118, %r27;
	div.full.f32 	%r234, %r119, %r27;
	div.full.f32 	%r235, %r120, %r27;
	div.full.f32 	%r236, %r121, %r27;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p4, %r205, 0f00000000;
	setp.ge.f32 	%p5, %r206, 0f00000000;
	setp.ge.f32 	%p6, %r207, 0f00000000;
	setp.ge.f32 	%p7, %r208, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r237, 0f3F000000, 0fBF000000, %p7;
	selp.f32 	%r238, 0f3F000000, 0fBF000000, %p6;
	selp.f32 	%r239, 0f3F000000, 0fBF000000, %p5;
	selp.f32 	%r240, 0f3F000000, 0fBF000000, %p4;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r241, %r205, %r240;
	add.f32 	%r242, %r206, %r239;
	add.f32 	%r243, %r207, %r238;
	add.f32 	%r244, %r208, %r237;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs33, %r244;
	cvt.rzi.s16.f32 	%rs34, %r243;
	mov.b32 	%r245, {%rs34, %rs33};
	mov.b32 	{%rs35, %rs36}, %r245;
	cvt.u32.u16 	%r246, %rs36;
	cvt.u32.u16 	%r247, %rs35;
	prmt.b32 	%r248, %r247, %r246, 0x3340U;
	cvt.rzi.s16.f32 	%rs37, %r242;
	cvt.rzi.s16.f32 	%rs38, %r241;
	mov.b32 	%r249, {%rs38, %rs37};
	mov.b32 	{%rs39, %rs40}, %r249;
	cvt.u32.u16 	%r250, %rs40;
	cvt.u32.u16 	%r251, %rs39;
	prmt.b32 	%r252, %r251, %r250, 0x3340U;
	prmt.b32 	%r19, %r252, %r248, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p8, %r209, 0f00000000;
	setp.ge.f32 	%p9, %r210, 0f00000000;
	setp.ge.f32 	%p10, %r211, 0f00000000;
	setp.ge.f32 	%p11, %r212, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r253, 0f3F000000, 0fBF000000, %p11;
	selp.f32 	%r254, 0f3F000000, 0fBF000000, %p10;
	selp.f32 	%r255, 0f3F000000, 0fBF000000, %p9;
	selp.f32 	%r256, 0f3F000000, 0fBF000000, %p8;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r257, %r209, %r256;
	add.f32 	%r258, %r210, %r255;
	add.f32 	%r259, %r211, %r254;
	add.f32 	%r260, %r212, %r253;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs41, %r260;
	cvt.rzi.s16.f32 	%rs42, %r259;
	mov.b32 	%r261, {%rs42, %rs41};
	mov.b32 	{%rs43, %rs44}, %r261;
	cvt.u32.u16 	%r262, %rs44;
	cvt.u32.u16 	%r263, %rs43;
	prmt.b32 	%r264, %r263, %r262, 0x3340U;
	cvt.rzi.s16.f32 	%rs45, %r258;
	cvt.rzi.s16.f32 	%rs46, %r257;
	mov.b32 	%r265, {%rs46, %rs45};
	mov.b32 	{%rs47, %rs48}, %r265;
	cvt.u32.u16 	%r266, %rs48;
	cvt.u32.u16 	%r267, %rs47;
	prmt.b32 	%r268, %r267, %r266, 0x3340U;
	prmt.b32 	%r20, %r268, %r264, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p12, %r213, 0f00000000;
	setp.ge.f32 	%p13, %r214, 0f00000000;
	setp.ge.f32 	%p14, %r215, 0f00000000;
	setp.ge.f32 	%p15, %r216, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r269, 0f3F000000, 0fBF000000, %p15;
	selp.f32 	%r270, 0f3F000000, 0fBF000000, %p14;
	selp.f32 	%r271, 0f3F000000, 0fBF000000, %p13;
	selp.f32 	%r272, 0f3F000000, 0fBF000000, %p12;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r273, %r213, %r272;
	add.f32 	%r274, %r214, %r271;
	add.f32 	%r275, %r215, %r270;
	add.f32 	%r276, %r216, %r269;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs49, %r276;
	cvt.rzi.s16.f32 	%rs50, %r275;
	mov.b32 	%r277, {%rs50, %rs49};
	mov.b32 	{%rs51, %rs52}, %r277;
	cvt.u32.u16 	%r278, %rs52;
	cvt.u32.u16 	%r279, %rs51;
	prmt.b32 	%r280, %r279, %r278, 0x3340U;
	cvt.rzi.s16.f32 	%rs53, %r274;
	cvt.rzi.s16.f32 	%rs54, %r273;
	mov.b32 	%r281, {%rs54, %rs53};
	mov.b32 	{%rs55, %rs56}, %r281;
	cvt.u32.u16 	%r282, %rs56;
	cvt.u32.u16 	%r283, %rs55;
	prmt.b32 	%r284, %r283, %r282, 0x3340U;
	prmt.b32 	%r21, %r284, %r280, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p16, %r217, 0f00000000;
	setp.ge.f32 	%p17, %r218, 0f00000000;
	setp.ge.f32 	%p18, %r219, 0f00000000;
	setp.ge.f32 	%p19, %r220, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r285, 0f3F000000, 0fBF000000, %p19;
	selp.f32 	%r286, 0f3F000000, 0fBF000000, %p18;
	selp.f32 	%r287, 0f3F000000, 0fBF000000, %p17;
	selp.f32 	%r288, 0f3F000000, 0fBF000000, %p16;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r289, %r217, %r288;
	add.f32 	%r290, %r218, %r287;
	add.f32 	%r291, %r219, %r286;
	add.f32 	%r292, %r220, %r285;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs57, %r292;
	cvt.rzi.s16.f32 	%rs58, %r291;
	mov.b32 	%r293, {%rs58, %rs57};
	mov.b32 	{%rs59, %rs60}, %r293;
	cvt.u32.u16 	%r294, %rs60;
	cvt.u32.u16 	%r295, %rs59;
	prmt.b32 	%r296, %r295, %r294, 0x3340U;
	cvt.rzi.s16.f32 	%rs61, %r290;
	cvt.rzi.s16.f32 	%rs62, %r289;
	mov.b32 	%r297, {%rs62, %rs61};
	mov.b32 	{%rs63, %rs64}, %r297;
	cvt.u32.u16 	%r298, %rs64;
	cvt.u32.u16 	%r299, %rs63;
	prmt.b32 	%r300, %r299, %r298, 0x3340U;
	prmt.b32 	%r22, %r300, %r296, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p20, %r221, 0f00000000;
	setp.ge.f32 	%p21, %r222, 0f00000000;
	setp.ge.f32 	%p22, %r223, 0f00000000;
	setp.ge.f32 	%p23, %r224, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r301, 0f3F000000, 0fBF000000, %p23;
	selp.f32 	%r302, 0f3F000000, 0fBF000000, %p22;
	selp.f32 	%r303, 0f3F000000, 0fBF000000, %p21;
	selp.f32 	%r304, 0f3F000000, 0fBF000000, %p20;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r305, %r221, %r304;
	add.f32 	%r306, %r222, %r303;
	add.f32 	%r307, %r223, %r302;
	add.f32 	%r308, %r224, %r301;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs65, %r308;
	cvt.rzi.s16.f32 	%rs66, %r307;
	mov.b32 	%r309, {%rs66, %rs65};
	mov.b32 	{%rs67, %rs68}, %r309;
	cvt.u32.u16 	%r310, %rs68;
	cvt.u32.u16 	%r311, %rs67;
	prmt.b32 	%r312, %r311, %r310, 0x3340U;
	cvt.rzi.s16.f32 	%rs69, %r306;
	cvt.rzi.s16.f32 	%rs70, %r305;
	mov.b32 	%r313, {%rs70, %rs69};
	mov.b32 	{%rs71, %rs72}, %r313;
	cvt.u32.u16 	%r314, %rs72;
	cvt.u32.u16 	%r315, %rs71;
	prmt.b32 	%r316, %r315, %r314, 0x3340U;
	prmt.b32 	%r23, %r316, %r312, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p24, %r225, 0f00000000;
	setp.ge.f32 	%p25, %r226, 0f00000000;
	setp.ge.f32 	%p26, %r227, 0f00000000;
	setp.ge.f32 	%p27, %r228, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r317, 0f3F000000, 0fBF000000, %p27;
	selp.f32 	%r318, 0f3F000000, 0fBF000000, %p26;
	selp.f32 	%r319, 0f3F000000, 0fBF000000, %p25;
	selp.f32 	%r320, 0f3F000000, 0fBF000000, %p24;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r321, %r225, %r320;
	add.f32 	%r322, %r226, %r319;
	add.f32 	%r323, %r227, %r318;
	add.f32 	%r324, %r228, %r317;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs73, %r324;
	cvt.rzi.s16.f32 	%rs74, %r323;
	mov.b32 	%r325, {%rs74, %rs73};
	mov.b32 	{%rs75, %rs76}, %r325;
	cvt.u32.u16 	%r326, %rs76;
	cvt.u32.u16 	%r327, %rs75;
	prmt.b32 	%r328, %r327, %r326, 0x3340U;
	cvt.rzi.s16.f32 	%rs77, %r322;
	cvt.rzi.s16.f32 	%rs78, %r321;
	mov.b32 	%r329, {%rs78, %rs77};
	mov.b32 	{%rs79, %rs80}, %r329;
	cvt.u32.u16 	%r330, %rs80;
	cvt.u32.u16 	%r331, %rs79;
	prmt.b32 	%r332, %r331, %r330, 0x3340U;
	prmt.b32 	%r24, %r332, %r328, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p28, %r229, 0f00000000;
	setp.ge.f32 	%p29, %r230, 0f00000000;
	setp.ge.f32 	%p30, %r231, 0f00000000;
	setp.ge.f32 	%p31, %r232, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r333, 0f3F000000, 0fBF000000, %p31;
	selp.f32 	%r334, 0f3F000000, 0fBF000000, %p30;
	selp.f32 	%r335, 0f3F000000, 0fBF000000, %p29;
	selp.f32 	%r336, 0f3F000000, 0fBF000000, %p28;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r337, %r229, %r336;
	add.f32 	%r338, %r230, %r335;
	add.f32 	%r339, %r231, %r334;
	add.f32 	%r340, %r232, %r333;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs81, %r340;
	cvt.rzi.s16.f32 	%rs82, %r339;
	mov.b32 	%r341, {%rs82, %rs81};
	mov.b32 	{%rs83, %rs84}, %r341;
	cvt.u32.u16 	%r342, %rs84;
	cvt.u32.u16 	%r343, %rs83;
	prmt.b32 	%r344, %r343, %r342, 0x3340U;
	cvt.rzi.s16.f32 	%rs85, %r338;
	cvt.rzi.s16.f32 	%rs86, %r337;
	mov.b32 	%r345, {%rs86, %rs85};
	mov.b32 	{%rs87, %rs88}, %r345;
	cvt.u32.u16 	%r346, %rs88;
	cvt.u32.u16 	%r347, %rs87;
	prmt.b32 	%r348, %r347, %r346, 0x3340U;
	prmt.b32 	%r25, %r348, %r344, 0x5410U;
	.loc	1 28 30                         // quant_per_block.py:28:30
	setp.ge.f32 	%p32, %r233, 0f00000000;
	setp.ge.f32 	%p33, %r234, 0f00000000;
	setp.ge.f32 	%p34, %r235, 0f00000000;
	setp.ge.f32 	%p35, %r236, 0f00000000;
	.loc	1 28 21                         // quant_per_block.py:28:21
	selp.f32 	%r349, 0f3F000000, 0fBF000000, %p35;
	selp.f32 	%r350, 0f3F000000, 0fBF000000, %p34;
	selp.f32 	%r351, 0f3F000000, 0fBF000000, %p33;
	selp.f32 	%r352, 0f3F000000, 0fBF000000, %p32;
	.loc	1 28 5                          // quant_per_block.py:28:5
	add.f32 	%r353, %r233, %r352;
	add.f32 	%r354, %r234, %r351;
	add.f32 	%r355, %r235, %r350;
	add.f32 	%r356, %r236, %r349;
	.loc	1 29 14                         // quant_per_block.py:29:14
	cvt.rzi.s16.f32 	%rs89, %r356;
	cvt.rzi.s16.f32 	%rs90, %r355;
	mov.b32 	%r357, {%rs90, %rs89};
	mov.b32 	{%rs91, %rs92}, %r357;
	cvt.u32.u16 	%r358, %rs92;
	cvt.u32.u16 	%r359, %rs91;
	prmt.b32 	%r360, %r359, %r358, 0x3340U;
	cvt.rzi.s16.f32 	%rs93, %r354;
	cvt.rzi.s16.f32 	%rs94, %r353;
	mov.b32 	%r361, {%rs94, %rs93};
	mov.b32 	{%rs95, %rs96}, %r361;
	cvt.u32.u16 	%r362, %rs96;
	cvt.u32.u16 	%r363, %rs95;
	prmt.b32 	%r364, %r363, %r362, 0x3340U;
	prmt.b32 	%r26, %r364, %r360, 0x5410U;
	.loc	1 30 5                          // quant_per_block.py:30:5
	// begin inline asm
	@%p1 st.global.v4.b32 [ %rd5 + 0 ], { %r19, %r20, %r21, %r22 };

	// end inline asm
	// begin inline asm
	@%p2 st.global.v4.b32 [ %rd6 + 0 ], { %r23, %r24, %r25, %r26 };

	// end inline asm
	.loc	1 31 5                          // quant_per_block.py:31:5
	setp.eq.b32 	%p3, %r35, 0;
	// begin inline asm
	@%p3 st.global.b32 [ %rd7 + 0 ], { %r27 };

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