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

.version 8.8
.target sm_89
.address_size 64

	// .globl	_quant_query_per_thread_int8_i64_kernel // -- Begin function _quant_query_per_thread_int8_i64_kernel
.extern .shared .align 16 .b8 global_smem[];
                                        // @_quant_query_per_thread_int8_i64_kernel
.visible .entry _quant_query_per_thread_int8_i64_kernel(
	.param .u64 .ptr .global .align 1 _quant_query_per_thread_int8_i64_kernel_param_0,
	.param .u64 .ptr .global .align 1 _quant_query_per_thread_int8_i64_kernel_param_1,
	.param .u64 .ptr .global .align 1 _quant_query_per_thread_int8_i64_kernel_param_2,
	.param .u32 _quant_query_per_thread_int8_i64_kernel_param_3,
	.param .u32 _quant_query_per_thread_int8_i64_kernel_param_4,
	.param .u32 _quant_query_per_thread_int8_i64_kernel_param_5,
	.param .u32 _quant_query_per_thread_int8_i64_kernel_param_6,
	.param .u32 _quant_query_per_thread_int8_i64_kernel_param_7,
	.param .u32 _quant_query_per_thread_int8_i64_kernel_param_8,
	.param .u32 _quant_query_per_thread_int8_i64_kernel_param_9,
	.param .u32 _quant_query_per_thread_int8_i64_kernel_param_10,
	.param .u32 _quant_query_per_thread_int8_i64_kernel_param_11,
	.param .u64 .ptr .global .align 1 _quant_query_per_thread_int8_i64_kernel_param_12,
	.param .u64 .ptr .global .align 1 _quant_query_per_thread_int8_i64_kernel_param_13
)
.reqntid 128
{
	.reg .pred 	%p<7>;
	.reg .b16 	%rs<13>;
	.reg .b32 	%r<87>;
	.reg .b64 	%rd<21>;
	.loc	1 1758 0                        // ltxv_nodes.py:1758:0
$L__func_begin0:
	.loc	1 1758 0                        // ltxv_nodes.py:1758:0

// %bb.0:
	ld.param.b64 	%rd4, [_quant_query_per_thread_int8_i64_kernel_param_0];
	ld.param.b64 	%rd5, [_quant_query_per_thread_int8_i64_kernel_param_1];
$L__tmp0:
	.loc	1 1763 19                       // ltxv_nodes.py:1763:19
	mov.u32 	%r7, %ctaid.x;
	.loc	1 1764 19                       // ltxv_nodes.py:1764:19
	and.b32 	%r8, %r7, 7;
	ld.param.b64 	%rd6, [_quant_query_per_thread_int8_i64_kernel_param_2];
	ld.param.b32 	%r9, [_quant_query_per_thread_int8_i64_kernel_param_3];
	.loc	1 1765 17                       // ltxv_nodes.py:1765:17
	mov.u32 	%r10, %ctaid.y;
	.loc	1 1766 17                       // ltxv_nodes.py:1766:17
	mov.u32 	%r11, %ctaid.z;
	ld.param.b32 	%r12, [_quant_query_per_thread_int8_i64_kernel_param_4];
	.loc	1 1768 18                       // ltxv_nodes.py:1768:18
	shl.b32 	%r13, %r7, 2;
	ld.param.b32 	%r14, [_quant_query_per_thread_int8_i64_kernel_param_5];
	and.b32 	%r15, %r13, -32;
	ld.param.b32 	%r16, [_quant_query_per_thread_int8_i64_kernel_param_6];
	ld.param.b32 	%r17, [_quant_query_per_thread_int8_i64_kernel_param_7];
	.loc	1 1768 34                       // ltxv_nodes.py:1768:34
	mov.u32 	%r18, %tid.x;
	and.b32 	%r19, %r18, 96;
	ld.param.b32 	%r20, [_quant_query_per_thread_int8_i64_kernel_param_8];
	shr.u32 	%r21, %r19, 2;
	ld.param.b32 	%r22, [_quant_query_per_thread_int8_i64_kernel_param_9];
	.loc	1 1768 18                       // ltxv_nodes.py:1768:18
	or.b32 	%r23, %r21, %r15;
	ld.param.b32 	%r24, [_quant_query_per_thread_int8_i64_kernel_param_10];
	or.b32 	%r25, %r23, %r8;
	ld.param.b32 	%r26, [_quant_query_per_thread_int8_i64_kernel_param_11];
	.loc	1 1771 30                       // ltxv_nodes.py:1771:30
	mul.lo.s32 	%r27, %r12, %r11;
	.loc	1 1771 22                       // ltxv_nodes.py:1771:22
	mad.wide.s32 	%rd7, %r27, 2, %rd4;
	.loc	1 1771 50                       // ltxv_nodes.py:1771:50
	mul.lo.s32 	%r28, %r14, %r10;
	.loc	1 1771 22                       // ltxv_nodes.py:1771:22
	mad.wide.s32 	%rd8, %r28, 2, %rd7;
	.loc	1 1771 70                       // ltxv_nodes.py:1771:70
	mul.wide.s32 	%rd9, %r16, %r25;
	.loc	1 1771 22                       // ltxv_nodes.py:1771:22
	shl.b64 	%rd10, %rd9, 1;
	add.s64 	%rd11, %rd8, %rd10;
	.loc	1 1771 113                      // ltxv_nodes.py:1771:113
	shl.b32 	%r29, %r18, 2;
	and.b32 	%r30, %r29, 124;
	.loc	1 1771 22                       // ltxv_nodes.py:1771:22
	cvt.u64.u32 	%rd12, %r30;
	mad.wide.u32 	%rd1, %r30, 2, %rd11;
	.loc	1 1772 32                       // ltxv_nodes.py:1772:32
	mul.lo.s32 	%r31, %r17, %r11;
	.loc	1 1772 23                       // ltxv_nodes.py:1772:23
	cvt.s64.s32 	%rd13, %r31;
	add.s64 	%rd14, %rd5, %rd13;
	.loc	1 1772 52                       // ltxv_nodes.py:1772:52
	mul.lo.s32 	%r32, %r20, %r10;
	.loc	1 1772 23                       // ltxv_nodes.py:1772:23
	cvt.s64.s32 	%rd15, %r32;
	add.s64 	%rd16, %rd14, %rd15;
	mad.wide.s32 	%rd17, %r22, %r25, %rd16;
	add.s64 	%rd2, %rd17, %rd12;
	.loc	1 1773 30                       // ltxv_nodes.py:1773:30
	mul.lo.s32 	%r33, %r24, %r11;
	.loc	1 1773 22                       // ltxv_nodes.py:1773:22
	mad.wide.s32 	%rd18, %r33, 4, %rd6;
	.loc	1 1773 50                       // ltxv_nodes.py:1773:50
	mul.lo.s32 	%r34, %r26, %r10;
	.loc	1 1773 22                       // ltxv_nodes.py:1773:22
	mad.wide.s32 	%rd19, %r34, 4, %rd18;
	.loc	1 1773 70                       // ltxv_nodes.py:1773:70
	and.b32 	%r35, %r7, 2147483640;
	.loc	1 1773 22                       // ltxv_nodes.py:1773:22
	mad.wide.u32 	%rd20, %r35, 4, %rd19;
	mad.wide.u32 	%rd3, %r8, 4, %rd20;
	.loc	1 1775 38                       // ltxv_nodes.py:1775:38
	setp.lt.s32 	%p1, %r25, %r9;
	.loc	1 1775 13                       // ltxv_nodes.py:1775:13
	// begin inline asm
	.loc	1 1775 13                       // ltxv_nodes.py:1775:13
	mov.u32 %r1, 0x0;
	.loc	1 1776 13                       // ltxv_nodes.py:1776:13
	mov.u32 %r2, 0x0;
	@%p1 ld.global.v2.b32 { %r1, %r2 }, [ %rd1 + 0 ];

	// end inline asm
	mov.b32 	{%rs1, %rs2}, %r1;
	mov.b32 	{%rs3, %rs4}, %r2;
	.loc	1 1776 13                       // ltxv_nodes.py:1776:13
	cvt.f32.bf16 	%r36, %rs1;
	cvt.f32.bf16 	%r37, %rs2;
	cvt.f32.bf16 	%r38, %rs3;
	cvt.f32.bf16 	%r39, %rs4;
	.loc	1 1777 24                       // ltxv_nodes.py:1777:24
	abs.f32 	%r40, %r36;
	abs.f32 	%r41, %r37;
	abs.f32 	%r42, %r38;
	abs.f32 	%r43, %r39;
$L__tmp1:
	.loc	1 1777 17                       // ltxv_nodes.py:1777:17
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 1777 17 // standard.py:191:16
	.loc	2 170 12, function_name $L__info_string0, inlined_at 2 191 16 // standard.py:170:12
	max.f32 	%r44, %r40, %r41;
	max.f32 	%r45, %r42, %r43;
	max.f32 	%r46, %r44, %r45;
$L__tmp2:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 1777 17 // standard.py:191:16
	shfl.sync.bfly.b32 	%r47, %r46, 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 	%r48, %r46, %r47;
$L__tmp4:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 1777 17 // standard.py:191:16
	shfl.sync.bfly.b32 	%r49, %r48, 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 	%r50, %r48, %r49;
$L__tmp6:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 1777 17 // standard.py:191:16
	shfl.sync.bfly.b32 	%r51, %r50, 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 	%r52, %r50, %r51;
$L__tmp8:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 1777 17 // standard.py:191:16
	shfl.sync.bfly.b32 	%r53, %r52, 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 	%r54, %r52, %r53;
$L__tmp10:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 1777 17 // standard.py:191:16
	shfl.sync.bfly.b32 	%r55, %r54, 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 	%r4, %r54, %r55;
$L__tmp12:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 1777 17 // standard.py:191:16
	shr.u32 	%r56, %r19, 3;
	mov.b32 	%r57, global_smem;
	add.s32 	%r3, %r57, %r56;
	// begin inline asm
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 1777 17 // standard.py:191:16
	st.shared::cta.b32 [ %r3 + 0 ], %r4;

	// end inline asm
	bar.sync 	0;
	and.b32 	%r58, %r29, 12;
	add.s32 	%r59, %r57, %r58;
	ld.shared.b32 	%r60, [%r59];
	shfl.sync.bfly.b32 	%r61, %r60, 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 	%r62, %r60, %r61;
$L__tmp14:
	.loc	2 191 16, function_name $L__info_string0, inlined_at 1 1777 17 // standard.py:191:16
	shfl.sync.bfly.b32 	%r63, %r62, 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 	%r64, %r62, %r63;
	mov.b32 	%r65, 0f42FE0000;
$L__tmp16:
	.loc	1 1777 17                       // ltxv_nodes.py:1777:17
	div.full.f32 	%r66, %r64, %r65;
	add.f32 	%r6, %r66, 0f33D6BF95;
	.loc	1 1778 18                       // ltxv_nodes.py:1778:18
	div.full.f32 	%r67, %r36, %r6;
	div.full.f32 	%r68, %r37, %r6;
	div.full.f32 	%r69, %r38, %r6;
	div.full.f32 	%r70, %r39, %r6;
	.loc	1 1779 34                       // ltxv_nodes.py:1779:34
	setp.ge.f32 	%p3, %r67, 0f00000000;
	setp.ge.f32 	%p4, %r68, 0f00000000;
	setp.ge.f32 	%p5, %r69, 0f00000000;
	setp.ge.f32 	%p6, %r70, 0f00000000;
	.loc	1 1779 25                       // ltxv_nodes.py:1779:25
	selp.f32 	%r71, 0f3F000000, 0fBF000000, %p6;
	selp.f32 	%r72, 0f3F000000, 0fBF000000, %p5;
	selp.f32 	%r73, 0f3F000000, 0fBF000000, %p4;
	selp.f32 	%r74, 0f3F000000, 0fBF000000, %p3;
	.loc	1 1779 9                        // ltxv_nodes.py:1779:9
	add.f32 	%r75, %r67, %r74;
	add.f32 	%r76, %r68, %r73;
	add.f32 	%r77, %r69, %r72;
	add.f32 	%r78, %r70, %r71;
	.loc	1 1780 18                       // ltxv_nodes.py:1780:18
	cvt.rzi.s16.f32 	%rs5, %r78;
	cvt.rzi.s16.f32 	%rs6, %r77;
	mov.b32 	%r79, {%rs6, %rs5};
	mov.b32 	{%rs7, %rs8}, %r79;
	cvt.u32.u16 	%r80, %rs8;
	cvt.u32.u16 	%r81, %rs7;
	prmt.b32 	%r82, %r81, %r80, 0x3340U;
	cvt.rzi.s16.f32 	%rs9, %r76;
	cvt.rzi.s16.f32 	%rs10, %r75;
	mov.b32 	%r83, {%rs10, %rs9};
	mov.b32 	{%rs11, %rs12}, %r83;
	cvt.u32.u16 	%r84, %rs12;
	cvt.u32.u16 	%r85, %rs11;
	prmt.b32 	%r86, %r85, %r84, 0x3340U;
	prmt.b32 	%r5, %r86, %r82, 0x5410U;
	.loc	1 1781 9                        // ltxv_nodes.py:1781:9
	// begin inline asm
	@%p1 st.global.b32 [ %rd2 + 0 ], { %r5 };

	// end inline asm
	.loc	1 1782 9                        // ltxv_nodes.py:1782:9
	setp.eq.b32 	%p2, %r18, 0;
	// begin inline asm
	@%p2 st.global.b32 [ %rd3 + 0 ], { %r6 };

	// end inline asm
	.loc	1 1758 5                        // ltxv_nodes.py:1758:5
	ret;
$L__tmp17:
$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 231                                // 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:0xe0 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:0x2a DW_TAG_subprogram
.b8 95                                  // DW_AT_name
.b8 113
.b8 117
.b8 97
.b8 110
.b8 116
.b8 95
.b8 113
.b8 117
.b8 101
.b8 114
.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] 0xa2:0x48 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] 0xb7: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 241                                 // DW_AT_call_line
.b8 6
.b8 17                                  // DW_AT_call_column
.b8 5                                   // Abbrev [5] 0xd0: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 0                                   // End Of Children Mark
.b8 0                                   // End Of Children Mark
	}
	.section	.debug_str
	{
$L__info_string0:
.b8 95                                  // string offset=0 ; _quant_query_per_thread_int8_i64_kernel
.b8 113
.b8 117
.b8 97
.b8 110
.b8 116
.b8 95
.b8 113
.b8 117
.b8 101
.b8 114
.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	{	}