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

.version 8.8
.target sm_89
.address_size 64

	// .globl	_quant_key_per_thread_int8_i64_kernel // -- Begin function _quant_key_per_thread_int8_i64_kernel
.extern .shared .align 16 .b8 global_smem[];
                                        // @_quant_key_per_thread_int8_i64_kernel
.visible .entry _quant_key_per_thread_int8_i64_kernel(
	.param .u64 .ptr .global .align 1 _quant_key_per_thread_int8_i64_kernel_param_0,
	.param .u64 .ptr .global .align 1 _quant_key_per_thread_int8_i64_kernel_param_1,
	.param .u64 .ptr .global .align 1 _quant_key_per_thread_int8_i64_kernel_param_2,
	.param .u32 _quant_key_per_thread_int8_i64_kernel_param_3,
	.param .u32 _quant_key_per_thread_int8_i64_kernel_param_4,
	.param .u32 _quant_key_per_thread_int8_i64_kernel_param_5,
	.param .u32 _quant_key_per_thread_int8_i64_kernel_param_6,
	.param .u32 _quant_key_per_thread_int8_i64_kernel_param_7,
	.param .u32 _quant_key_per_thread_int8_i64_kernel_param_8,
	.param .u32 _quant_key_per_thread_int8_i64_kernel_param_9,
	.param .u32 _quant_key_per_thread_int8_i64_kernel_param_10,
	.param .u32 _quant_key_per_thread_int8_i64_kernel_param_11,
	.param .u64 .ptr .global .align 1 _quant_key_per_thread_int8_i64_kernel_param_12,
	.param .u64 .ptr .global .align 1 _quant_key_per_thread_int8_i64_kernel_param_13
)
.reqntid 128
{
	.reg .pred 	%p<20>;
	.reg .b16 	%rs<49>;
	.reg .b32 	%r<211>;
	.reg .b64 	%rd<28>;
	.loc	1 1785 0                        // ltxv_nodes.py:1785:0
$L__func_begin0:
	.loc	1 1785 0                        // ltxv_nodes.py:1785:0

// %bb.0:
	ld.param.b64 	%rd6, [_quant_key_per_thread_int8_i64_kernel_param_0];
	ld.param.b64 	%rd7, [_quant_key_per_thread_int8_i64_kernel_param_1];
$L__tmp0:
	.loc	1 1790 19                       // ltxv_nodes.py:1790:19
	mov.u32 	%r17, %ctaid.x;
	.loc	1 1791 19                       // ltxv_nodes.py:1791:19
	and.b32 	%r18, %r17, 3;
	ld.param.b64 	%rd8, [_quant_key_per_thread_int8_i64_kernel_param_2];
	ld.param.b32 	%r19, [_quant_key_per_thread_int8_i64_kernel_param_3];
	.loc	1 1792 17                       // ltxv_nodes.py:1792:17
	mov.u32 	%r20, %ctaid.y;
	.loc	1 1793 17                       // ltxv_nodes.py:1793:17
	mov.u32 	%r21, %ctaid.z;
	ld.param.b32 	%r22, [_quant_key_per_thread_int8_i64_kernel_param_4];
	.loc	1 1795 19                       // ltxv_nodes.py:1795:19
	shl.b32 	%r23, %r17, 4;
	ld.param.b32 	%r24, [_quant_key_per_thread_int8_i64_kernel_param_5];
	and.b32 	%r25, %r23, -64;
	ld.param.b32 	%r26, [_quant_key_per_thread_int8_i64_kernel_param_6];
	ld.param.b32 	%r27, [_quant_key_per_thread_int8_i64_kernel_param_7];
	.loc	1 1795 35                       // ltxv_nodes.py:1795:35
	mov.u32 	%r28, %tid.x;
	shr.u32 	%r29, %r28, 1;
	ld.param.b32 	%r30, [_quant_key_per_thread_int8_i64_kernel_param_8];
	and.b32 	%r31, %r29, 56;
	ld.param.b32 	%r32, [_quant_key_per_thread_int8_i64_kernel_param_9];
	.loc	1 1795 19                       // ltxv_nodes.py:1795:19
	or.b32 	%r33, %r25, %r31;
	ld.param.b32 	%r34, [_quant_key_per_thread_int8_i64_kernel_param_10];
	.loc	1 1795 64                       // ltxv_nodes.py:1795:64
	shl.b32 	%r35, %r18, 1;
	ld.param.b32 	%r36, [_quant_key_per_thread_int8_i64_kernel_param_11];
	.loc	1 1795 19                       // ltxv_nodes.py:1795:19
	or.b32 	%r37, %r33, %r35;
	.loc	1 1796 19                       // ltxv_nodes.py:1796:19
	or.b32 	%r38, %r37, 1;
	.loc	1 1799 31                       // ltxv_nodes.py:1799:31
	mul.lo.s32 	%r39, %r22, %r21;
	.loc	1 1799 23                       // ltxv_nodes.py:1799:23
	mad.wide.s32 	%rd9, %r39, 2, %rd6;
	.loc	1 1799 51                       // ltxv_nodes.py:1799:51
	mul.lo.s32 	%r40, %r24, %r20;
	.loc	1 1799 23                       // ltxv_nodes.py:1799:23
	mad.wide.s32 	%rd10, %r40, 2, %rd9;
	.loc	1 1799 71                       // ltxv_nodes.py:1799:71
	mul.wide.s32 	%rd11, %r26, %r37;
	.loc	1 1799 23                       // ltxv_nodes.py:1799:23
	shl.b64 	%rd12, %rd11, 1;
	add.s64 	%rd13, %rd10, %rd12;
	.loc	1 1799 115                      // ltxv_nodes.py:1799:115
	shl.b32 	%r41, %r28, 3;
	and.b32 	%r42, %r41, 120;
	.loc	1 1799 23                       // ltxv_nodes.py:1799:23
	cvt.u64.u32 	%rd14, %r42;
	mul.wide.u32 	%rd15, %r42, 2;
	add.s64 	%rd1, %rd13, %rd15;
	.loc	1 1800 71                       // ltxv_nodes.py:1800:71
	mul.wide.s32 	%rd16, %r26, %r38;
	.loc	1 1800 23                       // ltxv_nodes.py:1800:23
	shl.b64 	%rd17, %rd16, 1;
	add.s64 	%rd18, %rd10, %rd17;
	add.s64 	%rd2, %rd18, %rd15;
	.loc	1 1801 33                       // ltxv_nodes.py:1801:33
	mul.lo.s32 	%r43, %r27, %r21;
	.loc	1 1801 24                       // ltxv_nodes.py:1801:24
	cvt.s64.s32 	%rd19, %r43;
	add.s64 	%rd20, %rd7, %rd19;
	.loc	1 1801 53                       // ltxv_nodes.py:1801:53
	mul.lo.s32 	%r44, %r30, %r20;
	.loc	1 1801 24                       // ltxv_nodes.py:1801:24
	cvt.s64.s32 	%rd21, %r44;
	add.s64 	%rd22, %rd20, %rd21;
	mad.wide.s32 	%rd23, %r32, %r37, %rd22;
	add.s64 	%rd3, %rd23, %rd14;
	.loc	1 1802 24                       // ltxv_nodes.py:1802:24
	mad.wide.s32 	%rd24, %r32, %r38, %rd22;
	add.s64 	%rd4, %rd24, %rd14;
	.loc	1 1803 30                       // ltxv_nodes.py:1803:30
	mul.lo.s32 	%r45, %r34, %r21;
	.loc	1 1803 22                       // ltxv_nodes.py:1803:22
	mad.wide.s32 	%rd25, %r45, 4, %rd8;
	.loc	1 1803 50                       // ltxv_nodes.py:1803:50
	mul.lo.s32 	%r46, %r36, %r20;
	.loc	1 1803 22                       // ltxv_nodes.py:1803:22
	mad.wide.s32 	%rd26, %r46, 4, %rd25;
	.loc	1 1803 70                       // ltxv_nodes.py:1803:70
	and.b32 	%r47, %r17, 2147483644;
	.loc	1 1803 22                       // ltxv_nodes.py:1803:22
	mad.wide.u32 	%rd27, %r47, 4, %rd26;
	mad.wide.u32 	%rd5, %r18, 4, %rd27;
	.loc	1 1805 40                       // ltxv_nodes.py:1805:40
	setp.lt.s32 	%p1, %r37, %r19;
	.loc	1 1805 14                       // ltxv_nodes.py:1805:14
	// begin inline asm
	.loc	1 1805 14                       // ltxv_nodes.py:1805:14
	mov.u32 %r1, 0x0;
	.loc	1 1806 14                       // ltxv_nodes.py:1806:14
	mov.u32 %r2, 0x0;
	.loc	1 1807 14                       // ltxv_nodes.py:1807:14
	mov.u32 %r3, 0x0;
	.loc	1 1808 14                       // ltxv_nodes.py:1808:14
	mov.u32 %r4, 0x0;
	@%p1 ld.global.v4.b32 { %r1, %r2, %r3, %r4 }, [ %rd1 + 0 ];

	// end inline asm
	mov.b32 	{%rs1, %rs2}, %r1;
	mov.b32 	{%rs3, %rs4}, %r2;
	mov.b32 	{%rs5, %rs6}, %r3;
	mov.b32 	{%rs7, %rs8}, %r4;
	.loc	1 1806 40                       // ltxv_nodes.py:1806:40
	setp.lt.s32 	%p2, %r38, %r19;
	.loc	1 1806 14                       // ltxv_nodes.py:1806:14
	// begin inline asm
	.loc	1 1806 14                       // ltxv_nodes.py:1806:14
	mov.u32 %r5, 0x0;
	.loc	1 1807 14                       // ltxv_nodes.py:1807:14
	mov.u32 %r6, 0x0;
	.loc	1 1808 14                       // ltxv_nodes.py:1808:14
	mov.u32 %r7, 0x0;
	.loc	1 1809 14                       // ltxv_nodes.py:1809:14
	mov.u32 %r8, 0x0;
	@%p2 ld.global.v4.b32 { %r5, %r6, %r7, %r8 }, [ %rd2 + 0 ];

	// end inline asm
	mov.b32 	{%rs9, %rs10}, %r5;
	mov.b32 	{%rs11, %rs12}, %r6;
	mov.b32 	{%rs13, %rs14}, %r7;
	mov.b32 	{%rs15, %rs16}, %r8;
	.loc	1 1807 14                       // ltxv_nodes.py:1807:14
	cvt.f32.bf16 	%r48, %rs1;
	cvt.f32.bf16 	%r49, %rs2;
	cvt.f32.bf16 	%r50, %rs3;
	cvt.f32.bf16 	%r51, %rs4;
	cvt.f32.bf16 	%r52, %rs5;
	cvt.f32.bf16 	%r53, %rs6;
	cvt.f32.bf16 	%r54, %rs7;
	cvt.f32.bf16 	%r55, %rs8;
	.loc	1 1808 14                       // ltxv_nodes.py:1808:14
	cvt.f32.bf16 	%r56, %rs9;
	cvt.f32.bf16 	%r57, %rs10;
	cvt.f32.bf16 	%r58, %rs11;
	cvt.f32.bf16 	%r59, %rs12;
	cvt.f32.bf16 	%r60, %rs13;
	cvt.f32.bf16 	%r61, %rs14;
	cvt.f32.bf16 	%r62, %rs15;
	cvt.f32.bf16 	%r63, %rs16;
	.loc	1 1809 28                       // ltxv_nodes.py:1809:28
	abs.f32 	%r64, %r48;
	abs.f32 	%r65, %r49;
	abs.f32 	%r66, %r50;
	abs.f32 	%r67, %r51;
	abs.f32 	%r68, %r52;
	abs.f32 	%r69, %r53;
	abs.f32 	%r70, %r54;
	abs.f32 	%r71, %r55;
$L__tmp1:
	.loc	1 1809 21                       // ltxv_nodes.py:1809:21
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 1809 21 // standard.py:191:16
	.loc	2 170 12, function_name $L__info_string0, inlined_at 2 191 16 // standard.py:170:12
	max.f32 	%r72, %r64, %r65;
	max.f32 	%r73, %r66, %r67;
	max.f32 	%r74, %r68, %r69;
	max.f32 	%r75, %r70, %r71;
	max.f32 	%r76, %r72, %r73;
	max.f32 	%r77, %r74, %r75;
	max.f32 	%r78, %r76, %r77;
$L__tmp2:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 1809 21 // standard.py:191:16
	shfl.sync.bfly.b32 	%r79, %r78, 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 	%r80, %r78, %r79;
$L__tmp4:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 1809 21 // standard.py:191:16
	shfl.sync.bfly.b32 	%r81, %r80, 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 	%r82, %r80, %r81;
$L__tmp6:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 1809 21 // standard.py:191:16
	shfl.sync.bfly.b32 	%r83, %r82, 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 	%r84, %r82, %r83;
$L__tmp8:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 1809 21 // standard.py:191:16
	shfl.sync.bfly.b32 	%r85, %r84, 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 	%r86, %r84, %r85;
$L__tmp10:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 1809 21 // standard.py:191:16
	shfl.sync.bfly.b32 	%r87, %r86, 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 	%r10, %r86, %r87;
$L__tmp12:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 1809 21 // standard.py:191:16
	shr.u32 	%r88, %r28, 3;
	and.b32 	%r89, %r88, 12;
	mov.b32 	%r90, global_smem;
	add.s32 	%r9, %r90, %r89;
	// begin inline asm
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 1809 21 // standard.py:191:16
	st.shared::cta.b32 [ %r9 + 0 ], %r10;

	// end inline asm
	bar.sync 	0;
	shl.b32 	%r91, %r28, 2;
	and.b32 	%r92, %r91, 12;
	add.s32 	%r93, %r90, %r92;
	ld.shared.b32 	%r94, [%r93];
	shfl.sync.bfly.b32 	%r95, %r94, 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 	%r96, %r94, %r95;
$L__tmp14:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 1809 21 // standard.py:191:16
	shfl.sync.bfly.b32 	%r97, %r96, 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 	%r98, %r96, %r97;
$L__tmp16:
	.loc	1 1809 48                       // ltxv_nodes.py:1809:48
	abs.f32 	%r99, %r56;
	abs.f32 	%r100, %r57;
	abs.f32 	%r101, %r58;
	abs.f32 	%r102, %r59;
	abs.f32 	%r103, %r60;
	abs.f32 	%r104, %r61;
	abs.f32 	%r105, %r62;
	abs.f32 	%r106, %r63;
$L__tmp17:
	.loc	1 1809 41                       // ltxv_nodes.py:1809:41
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 1809 41 // standard.py:191:16
	bar.sync 	0;
$L__tmp18:
	.loc	2 170 12, function_name $L__info_string0, inlined_at 2 191 16 // standard.py:170:12
	max.f32 	%r107, %r99, %r100;
	max.f32 	%r108, %r101, %r102;
	max.f32 	%r109, %r103, %r104;
	max.f32 	%r110, %r105, %r106;
	max.f32 	%r111, %r107, %r108;
	max.f32 	%r112, %r109, %r110;
	max.f32 	%r113, %r111, %r112;
$L__tmp19:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 1809 41 // standard.py:191:16
	shfl.sync.bfly.b32 	%r114, %r113, 16, 31, -1;
$L__tmp20:
	.loc	2 170 12, function_name $L__info_string0, inlined_at 2 191 16 // standard.py:170:12
	max.f32 	%r115, %r113, %r114;
$L__tmp21:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 1809 41 // standard.py:191:16
	shfl.sync.bfly.b32 	%r116, %r115, 8, 31, -1;
$L__tmp22:
	.loc	2 170 12, function_name $L__info_string0, inlined_at 2 191 16 // standard.py:170:12
	max.f32 	%r117, %r115, %r116;
$L__tmp23:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 1809 41 // standard.py:191:16
	shfl.sync.bfly.b32 	%r118, %r117, 4, 31, -1;
$L__tmp24:
	.loc	2 170 12, function_name $L__info_string0, inlined_at 2 191 16 // standard.py:170:12
	max.f32 	%r119, %r117, %r118;
$L__tmp25:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 1809 41 // standard.py:191:16
	shfl.sync.bfly.b32 	%r120, %r119, 2, 31, -1;
$L__tmp26:
	.loc	2 170 12, function_name $L__info_string0, inlined_at 2 191 16 // standard.py:170:12
	max.f32 	%r121, %r119, %r120;
$L__tmp27:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 1809 41 // standard.py:191:16
	shfl.sync.bfly.b32 	%r122, %r121, 1, 31, -1;
$L__tmp28:
	.loc	2 170 12, function_name $L__info_string0, inlined_at 2 191 16 // standard.py:170:12
	max.f32 	%r11, %r121, %r122;
$L__tmp29:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 1809 41 // standard.py:191:16
	// begin inline asm
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 1809 41 // standard.py:191:16
	st.shared::cta.b32 [ %r9 + 0 ], %r11;

	// end inline asm
	bar.sync 	0;
	ld.shared.b32 	%r123, [%r93];
	shfl.sync.bfly.b32 	%r124, %r123, 2, 31, -1;
$L__tmp30:
	.loc	2 170 12, function_name $L__info_string0, inlined_at 2 191 16 // standard.py:170:12
	max.f32 	%r125, %r123, %r124;
$L__tmp31:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 1809 41 // standard.py:191:16
	shfl.sync.bfly.b32 	%r126, %r125, 1, 31, -1;
$L__tmp32:
	.loc	2 170 12, function_name $L__info_string0, inlined_at 2 191 16 // standard.py:170:12
	max.f32 	%r127, %r125, %r126;
$L__tmp33:
	.loc	1 1809 17                       // ltxv_nodes.py:1809:17
	max.f32 	%r128, %r98, %r127;
	mov.b32 	%r129, 0f42FE0000;
	div.full.f32 	%r130, %r128, %r129;
	add.f32 	%r16, %r130, 0f33D6BF95;
	.loc	1 1810 19                       // ltxv_nodes.py:1810:19
	div.full.f32 	%r131, %r48, %r16;
	div.full.f32 	%r132, %r49, %r16;
	div.full.f32 	%r133, %r50, %r16;
	div.full.f32 	%r134, %r51, %r16;
	div.full.f32 	%r135, %r52, %r16;
	div.full.f32 	%r136, %r53, %r16;
	div.full.f32 	%r137, %r54, %r16;
	div.full.f32 	%r138, %r55, %r16;
	.loc	1 1811 19                       // ltxv_nodes.py:1811:19
	div.full.f32 	%r139, %r56, %r16;
	div.full.f32 	%r140, %r57, %r16;
	div.full.f32 	%r141, %r58, %r16;
	div.full.f32 	%r142, %r59, %r16;
	div.full.f32 	%r143, %r60, %r16;
	div.full.f32 	%r144, %r61, %r16;
	div.full.f32 	%r145, %r62, %r16;
	div.full.f32 	%r146, %r63, %r16;
	.loc	1 1812 35                       // ltxv_nodes.py:1812:35
	setp.ge.f32 	%p4, %r131, 0f00000000;
	setp.ge.f32 	%p5, %r132, 0f00000000;
	setp.ge.f32 	%p6, %r133, 0f00000000;
	setp.ge.f32 	%p7, %r134, 0f00000000;
	.loc	1 1812 26                       // ltxv_nodes.py:1812:26
	selp.f32 	%r147, 0f3F000000, 0fBF000000, %p7;
	selp.f32 	%r148, 0f3F000000, 0fBF000000, %p6;
	selp.f32 	%r149, 0f3F000000, 0fBF000000, %p5;
	selp.f32 	%r150, 0f3F000000, 0fBF000000, %p4;
	.loc	1 1812 9                        // ltxv_nodes.py:1812:9
	add.f32 	%r151, %r131, %r150;
	add.f32 	%r152, %r132, %r149;
	add.f32 	%r153, %r133, %r148;
	add.f32 	%r154, %r134, %r147;
	.loc	1 1814 19                       // ltxv_nodes.py:1814:19
	cvt.rzi.s16.f32 	%rs17, %r154;
	cvt.rzi.s16.f32 	%rs18, %r153;
	mov.b32 	%r155, {%rs18, %rs17};
	mov.b32 	{%rs19, %rs20}, %r155;
	cvt.u32.u16 	%r156, %rs20;
	cvt.u32.u16 	%r157, %rs19;
	prmt.b32 	%r158, %r157, %r156, 0x3340U;
	cvt.rzi.s16.f32 	%rs21, %r152;
	cvt.rzi.s16.f32 	%rs22, %r151;
	mov.b32 	%r159, {%rs22, %rs21};
	mov.b32 	{%rs23, %rs24}, %r159;
	cvt.u32.u16 	%r160, %rs24;
	cvt.u32.u16 	%r161, %rs23;
	prmt.b32 	%r162, %r161, %r160, 0x3340U;
	prmt.b32 	%r12, %r162, %r158, 0x5410U;
	.loc	1 1812 35                       // ltxv_nodes.py:1812:35
	setp.ge.f32 	%p8, %r135, 0f00000000;
	setp.ge.f32 	%p9, %r136, 0f00000000;
	setp.ge.f32 	%p10, %r137, 0f00000000;
	setp.ge.f32 	%p11, %r138, 0f00000000;
	.loc	1 1812 26                       // ltxv_nodes.py:1812:26
	selp.f32 	%r163, 0f3F000000, 0fBF000000, %p11;
	selp.f32 	%r164, 0f3F000000, 0fBF000000, %p10;
	selp.f32 	%r165, 0f3F000000, 0fBF000000, %p9;
	selp.f32 	%r166, 0f3F000000, 0fBF000000, %p8;
	.loc	1 1812 9                        // ltxv_nodes.py:1812:9
	add.f32 	%r167, %r135, %r166;
	add.f32 	%r168, %r136, %r165;
	add.f32 	%r169, %r137, %r164;
	add.f32 	%r170, %r138, %r163;
	.loc	1 1814 19                       // ltxv_nodes.py:1814:19
	cvt.rzi.s16.f32 	%rs25, %r170;
	cvt.rzi.s16.f32 	%rs26, %r169;
	mov.b32 	%r171, {%rs26, %rs25};
	mov.b32 	{%rs27, %rs28}, %r171;
	cvt.u32.u16 	%r172, %rs28;
	cvt.u32.u16 	%r173, %rs27;
	prmt.b32 	%r174, %r173, %r172, 0x3340U;
	cvt.rzi.s16.f32 	%rs29, %r168;
	cvt.rzi.s16.f32 	%rs30, %r167;
	mov.b32 	%r175, {%rs30, %rs29};
	mov.b32 	{%rs31, %rs32}, %r175;
	cvt.u32.u16 	%r176, %rs32;
	cvt.u32.u16 	%r177, %rs31;
	prmt.b32 	%r178, %r177, %r176, 0x3340U;
	prmt.b32 	%r13, %r178, %r174, 0x5410U;
	.loc	1 1813 35                       // ltxv_nodes.py:1813:35
	setp.ge.f32 	%p12, %r139, 0f00000000;
	setp.ge.f32 	%p13, %r140, 0f00000000;
	setp.ge.f32 	%p14, %r141, 0f00000000;
	setp.ge.f32 	%p15, %r142, 0f00000000;
	.loc	1 1813 26                       // ltxv_nodes.py:1813:26
	selp.f32 	%r179, 0f3F000000, 0fBF000000, %p15;
	selp.f32 	%r180, 0f3F000000, 0fBF000000, %p14;
	selp.f32 	%r181, 0f3F000000, 0fBF000000, %p13;
	selp.f32 	%r182, 0f3F000000, 0fBF000000, %p12;
	.loc	1 1813 9                        // ltxv_nodes.py:1813:9
	add.f32 	%r183, %r139, %r182;
	add.f32 	%r184, %r140, %r181;
	add.f32 	%r185, %r141, %r180;
	add.f32 	%r186, %r142, %r179;
	.loc	1 1815 19                       // ltxv_nodes.py:1815:19
	cvt.rzi.s16.f32 	%rs33, %r186;
	cvt.rzi.s16.f32 	%rs34, %r185;
	mov.b32 	%r187, {%rs34, %rs33};
	mov.b32 	{%rs35, %rs36}, %r187;
	cvt.u32.u16 	%r188, %rs36;
	cvt.u32.u16 	%r189, %rs35;
	prmt.b32 	%r190, %r189, %r188, 0x3340U;
	cvt.rzi.s16.f32 	%rs37, %r184;
	cvt.rzi.s16.f32 	%rs38, %r183;
	mov.b32 	%r191, {%rs38, %rs37};
	mov.b32 	{%rs39, %rs40}, %r191;
	cvt.u32.u16 	%r192, %rs40;
	cvt.u32.u16 	%r193, %rs39;
	prmt.b32 	%r194, %r193, %r192, 0x3340U;
	prmt.b32 	%r14, %r194, %r190, 0x5410U;
	.loc	1 1813 35                       // ltxv_nodes.py:1813:35
	setp.ge.f32 	%p16, %r143, 0f00000000;
	setp.ge.f32 	%p17, %r144, 0f00000000;
	setp.ge.f32 	%p18, %r145, 0f00000000;
	setp.ge.f32 	%p19, %r146, 0f00000000;
	.loc	1 1813 26                       // ltxv_nodes.py:1813:26
	selp.f32 	%r195, 0f3F000000, 0fBF000000, %p19;
	selp.f32 	%r196, 0f3F000000, 0fBF000000, %p18;
	selp.f32 	%r197, 0f3F000000, 0fBF000000, %p17;
	selp.f32 	%r198, 0f3F000000, 0fBF000000, %p16;
	.loc	1 1813 9                        // ltxv_nodes.py:1813:9
	add.f32 	%r199, %r143, %r198;
	add.f32 	%r200, %r144, %r197;
	add.f32 	%r201, %r145, %r196;
	add.f32 	%r202, %r146, %r195;
	.loc	1 1815 19                       // ltxv_nodes.py:1815:19
	cvt.rzi.s16.f32 	%rs41, %r202;
	cvt.rzi.s16.f32 	%rs42, %r201;
	mov.b32 	%r203, {%rs42, %rs41};
	mov.b32 	{%rs43, %rs44}, %r203;
	cvt.u32.u16 	%r204, %rs44;
	cvt.u32.u16 	%r205, %rs43;
	prmt.b32 	%r206, %r205, %r204, 0x3340U;
	cvt.rzi.s16.f32 	%rs45, %r200;
	cvt.rzi.s16.f32 	%rs46, %r199;
	mov.b32 	%r207, {%rs46, %rs45};
	mov.b32 	{%rs47, %rs48}, %r207;
	cvt.u32.u16 	%r208, %rs48;
	cvt.u32.u16 	%r209, %rs47;
	prmt.b32 	%r210, %r209, %r208, 0x3340U;
	prmt.b32 	%r15, %r210, %r206, 0x5410U;
	.loc	1 1816 9                        // ltxv_nodes.py:1816:9
	// begin inline asm
	@%p1 st.global.v2.b32 [ %rd3 + 0 ], { %r12, %r13 };

	// end inline asm
	.loc	1 1817 9                        // ltxv_nodes.py:1817:9
	// begin inline asm
	@%p2 st.global.v2.b32 [ %rd4 + 0 ], { %r14, %r15 };

	// end inline asm
	.loc	1 1818 9                        // ltxv_nodes.py:1818:9
	setp.eq.b32 	%p3, %r28, 0;
	// begin inline asm
	@%p3 st.global.b32 [ %rd5 + 0 ], { %r16 };

	// end inline asm
	.loc	1 1785 5                        // ltxv_nodes.py:1785:5
	ret;
$L__tmp34:
$L__func_end0:
                                        // -- End function
}
	.file	1 "D:\\Comfy\\App\\ComfyUI_windows_portable\\ComfyUI\\custom_nodes\\comfyui-kjnodes\\nodes\\ltxv_nodes.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 5                                   // DW_FORM_data2
.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 279                                // 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:0x110 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 108                                 // DW_AT_name
.b8 116
.b8 120
.b8 118
.b8 95
.b8 110
.b8 111
.b8 100
.b8 101
.b8 115
.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 67
.b8 111
.b8 109
.b8 102
.b8 121
.b8 85
.b8 73
.b8 92
.b8 99
.b8 117
.b8 115
.b8 116
.b8 111
.b8 109
.b8 95
.b8 110
.b8 111
.b8 100
.b8 101
.b8 115
.b8 92
.b8 99
.b8 111
.b8 109
.b8 102
.b8 121
.b8 117
.b8 105
.b8 45
.b8 107
.b8 106
.b8 110
.b8 111
.b8 100
.b8 101
.b8 115
.b8 92
.b8 110
.b8 111
.b8 100
.b8 101
.b8 115
.b8 0
.b8 2                                   // Abbrev [2] 0x78:0x28 DW_TAG_subprogram
.b8 95                                  // DW_AT_name
.b8 113
.b8 117
.b8 97
.b8 110
.b8 116
.b8 95
.b8 107
.b8 101
.b8 121
.b8 95
.b8 112
.b8 101
.b8 114
.b8 95
.b8 116
.b8 104
.b8 114
.b8 101
.b8 97
.b8 100
.b8 95
.b8 105
.b8 110
.b8 116
.b8 56
.b8 95
.b8 105
.b8 54
.b8 52
.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] 0xa0:0x7a DW_TAG_subprogram
.b64 $L__func_begin0                    // DW_AT_low_pc
.b64 $L__func_end0                      // DW_AT_high_pc
.b32 120                                // DW_AT_abstract_origin
.b8 4                                   // Abbrev [4] 0xb5:0x32 DW_TAG_inlined_subroutine
.b32 120                                // 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 17                                  // DW_AT_call_line
.b8 7
.b8 21                                  // DW_AT_call_column
.b8 5                                   // Abbrev [5] 0xce:0x18 DW_TAG_inlined_subroutine
.b32 120                                // 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 4                                   // Abbrev [4] 0xe7:0x32 DW_TAG_inlined_subroutine
.b32 120                                // DW_AT_abstract_origin
.b64 $L__tmp17                          // DW_AT_low_pc
.b64 $L__tmp33                          // DW_AT_high_pc
.b8 1                                   // DW_AT_call_file
.b8 17                                  // DW_AT_call_line
.b8 7
.b8 41                                  // DW_AT_call_column
.b8 5                                   // Abbrev [5] 0x100:0x18 DW_TAG_inlined_subroutine
.b32 120                                // DW_AT_abstract_origin
.b64 $L__tmp18                          // DW_AT_low_pc
.b64 $L__tmp33                          // 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 95                                  // string offset=0 ; _quant_key_per_thread_int8_i64_kernel
.b8 113
.b8 117
.b8 97
.b8 110
.b8 116
.b8 95
.b8 107
.b8 101
.b8 121
.b8 95
.b8 112
.b8 101
.b8 114
.b8 95
.b8 116
.b8 104
.b8 114
.b8 101
.b8 97
.b8 100
.b8 95
.b8 105
.b8 110
.b8 116
.b8 56
.b8 95
.b8 105
.b8 54
.b8 52
.b8 95
.b8 107
.b8 101
.b8 114
.b8 110
.b8 101
.b8 108
.b8 0
	}
	.section	.debug_macinfo	{	}