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

.version 8.8
.target sm_89
.address_size 64

	// .globl	_rms_norm_dual_scale_shift_kernel // -- Begin function _rms_norm_dual_scale_shift_kernel
.extern .shared .align 16 .b8 global_smem[];
.global .align 1 .b8 _$_str[11] = {95, 95, 67, 85, 68, 65, 95, 70, 84, 90};
                                        // @_rms_norm_dual_scale_shift_kernel
.visible .entry _rms_norm_dual_scale_shift_kernel(
	.param .u64 .ptr .global .align 1 _rms_norm_dual_scale_shift_kernel_param_0,
	.param .u64 .ptr .global .align 1 _rms_norm_dual_scale_shift_kernel_param_1,
	.param .u64 .ptr .global .align 1 _rms_norm_dual_scale_shift_kernel_param_2,
	.param .u64 .ptr .global .align 1 _rms_norm_dual_scale_shift_kernel_param_3,
	.param .u64 .ptr .global .align 1 _rms_norm_dual_scale_shift_kernel_param_4,
	.param .u64 .ptr .global .align 1 _rms_norm_dual_scale_shift_kernel_param_5,
	.param .u64 .ptr .global .align 1 _rms_norm_dual_scale_shift_kernel_param_6,
	.param .u32 _rms_norm_dual_scale_shift_kernel_param_7,
	.param .u32 _rms_norm_dual_scale_shift_kernel_param_8,
	.param .u64 .ptr .global .align 1 _rms_norm_dual_scale_shift_kernel_param_9,
	.param .u64 .ptr .global .align 1 _rms_norm_dual_scale_shift_kernel_param_10
)
.reqntid 512
{
	.reg .pred 	%p<2>;
	.reg .b16 	%rs<41>;
	.reg .b32 	%r<157>;
	.reg .b64 	%rd<16>;
	.loc	1 1273 0                        // ltxv_nodes.py:1273:0
$L__func_begin0:
	.loc	1 1273 0                        // ltxv_nodes.py:1273:0

// %bb.0:                               // %__nv_rsqrtf.exit
	ld.param.b64 	%rd8, [_rms_norm_dual_scale_shift_kernel_param_0];
	ld.param.b64 	%rd9, [_rms_norm_dual_scale_shift_kernel_param_1];
$L__tmp0:
	.loc	1 1283 15                       // ltxv_nodes.py:1283:15
	mov.u32 	%r32, %ctaid.x;
	ld.param.b64 	%rd10, [_rms_norm_dual_scale_shift_kernel_param_2];
	.loc	1 1284 16                       // ltxv_nodes.py:1284:16
	mov.u32 	%r33, %tid.x;
	shl.b32 	%r34, %r33, 3;
	ld.param.b64 	%rd11, [_rms_norm_dual_scale_shift_kernel_param_3];
	and.b32 	%r35, %r34, 4088;
	ld.param.b64 	%rd12, [_rms_norm_dual_scale_shift_kernel_param_4];
	ld.param.b32 	%r36, [_rms_norm_dual_scale_shift_kernel_param_7];
	.loc	1 1285 16                       // ltxv_nodes.py:1285:16
	setp.lt.s32 	%p1, %r35, %r36;
	ld.param.b64 	%rd13, [_rms_norm_dual_scale_shift_kernel_param_5];
	ld.param.b64 	%rd14, [_rms_norm_dual_scale_shift_kernel_param_6];
	.loc	1 1286 15                       // ltxv_nodes.py:1286:15
	mad.lo.s32 	%r37, %r36, %r32, %r35;
	.loc	1 1288 21                       // ltxv_nodes.py:1288:21
	mul.wide.s32 	%rd15, %r37, 2;
	add.s64 	%rd1, %rd8, %rd15;
	mov.b32 	%r5, 0;
	.loc	1 1288 13                       // ltxv_nodes.py:1288:13
	// begin inline asm
	.loc	1 1288 13                       // ltxv_nodes.py:1288:13
	mov.u32 %r1, %r5;
	.loc	1 1289 13                       // ltxv_nodes.py:1289:13
	mov.u32 %r2, %r5;
	.loc	1 1290 13                       // ltxv_nodes.py:1290:13
	mov.u32 %r3, %r5;
	.loc	1 1291 13                       // ltxv_nodes.py:1291:13
	mov.u32 %r4, %r5;
	@%p1 ld.global.v4.b32 { %r1, %r2, %r3, %r4 }, [ %rd1 + 0 ];

	// end inline asm
	mov.b32 	{%rs1, %rs2}, %r1;
	cvt.f32.bf16 	%r38, %rs1;
	cvt.f32.bf16 	%r39, %rs2;
	mov.b32 	{%rs3, %rs4}, %r2;
	cvt.f32.bf16 	%r40, %rs3;
	cvt.f32.bf16 	%r41, %rs4;
	mov.b32 	{%rs5, %rs6}, %r3;
	cvt.f32.bf16 	%r42, %rs5;
	cvt.f32.bf16 	%r43, %rs6;
	mov.b32 	{%rs7, %rs8}, %r4;
	cvt.f32.bf16 	%r44, %rs7;
	cvt.f32.bf16 	%r45, %rs8;
	.loc	1 1289 26                       // ltxv_nodes.py:1289:26
	mul.f32 	%r46, %r41, %r41;
	mul.f32 	%r47, %r43, %r43;
	mul.f32 	%r48, %r45, %r45;
$L__tmp1:
	.loc	1 1289 19                       // ltxv_nodes.py:1289:19
	.loc	2 293 12, function_name $L__info_string0, inlined_at 1 1289 19 // standard.py:293:12
	.loc	2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
	fma.rn.f32 	%r49, %r42, %r42, %r47;
	fma.rn.f32 	%r50, %r44, %r44, %r48;
$L__tmp2:
	.loc	1 1289 26                       // ltxv_nodes.py:1289:26
	mul.f32 	%r51, %r39, %r39;
	add.f32 	%r52, %r49, %r50;
	fma.rn.f32 	%r53, %r40, %r40, %r46;
$L__tmp3:
	.loc	2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
	fma.rn.f32 	%r54, %r38, %r38, %r51;
	add.f32 	%r55, %r54, %r53;
	add.f32 	%r56, %r55, %r52;
$L__tmp4:
	.loc	2 293 12, function_name $L__info_string0, inlined_at 1 1289 19 // standard.py:293:12
	shfl.sync.bfly.b32 	%r57, %r56, 16, 31, -1;
$L__tmp5:
	.loc	2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
	add.f32 	%r58, %r56, %r57;
$L__tmp6:
	.loc	2 293 12, function_name $L__info_string0, inlined_at 1 1289 19 // standard.py:293:12
	shfl.sync.bfly.b32 	%r59, %r58, 8, 31, -1;
$L__tmp7:
	.loc	2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
	add.f32 	%r60, %r58, %r59;
$L__tmp8:
	.loc	2 293 12, function_name $L__info_string0, inlined_at 1 1289 19 // standard.py:293:12
	shfl.sync.bfly.b32 	%r61, %r60, 4, 31, -1;
$L__tmp9:
	.loc	2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
	add.f32 	%r62, %r60, %r61;
$L__tmp10:
	.loc	2 293 12, function_name $L__info_string0, inlined_at 1 1289 19 // standard.py:293:12
	shfl.sync.bfly.b32 	%r63, %r62, 2, 31, -1;
$L__tmp11:
	.loc	2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
	add.f32 	%r64, %r62, %r63;
$L__tmp12:
	.loc	2 293 12, function_name $L__info_string0, inlined_at 1 1289 19 // standard.py:293:12
	shfl.sync.bfly.b32 	%r65, %r64, 1, 31, -1;
$L__tmp13:
	.loc	2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
	add.f32 	%r7, %r64, %r65;
$L__tmp14:
	.loc	2 293 12, function_name $L__info_string0, inlined_at 1 1289 19 // standard.py:293:12
	shr.u32 	%r66, %r33, 3;
	and.b32 	%r67, %r66, 60;
	mov.b32 	%r68, global_smem;
	add.s32 	%r6, %r68, %r67;
	// begin inline asm
	.loc	2 293 12, function_name $L__info_string0, inlined_at 1 1289 19 // standard.py:293:12
	st.shared::cta.b32 [ %r6 + 0 ], %r7;

	// end inline asm
	bar.sync 	0;
	shl.b32 	%r69, %r33, 2;
	and.b32 	%r70, %r69, 60;
	add.s32 	%r71, %r68, %r70;
	ld.shared.b32 	%r72, [%r71];
	shfl.sync.bfly.b32 	%r73, %r72, 8, 31, -1;
$L__tmp15:
	.loc	2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
	add.f32 	%r74, %r72, %r73;
$L__tmp16:
	.loc	2 293 12, function_name $L__info_string0, inlined_at 1 1289 19 // standard.py:293:12
	shfl.sync.bfly.b32 	%r75, %r74, 4, 31, -1;
$L__tmp17:
	.loc	2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
	add.f32 	%r76, %r74, %r75;
$L__tmp18:
	.loc	2 293 12, function_name $L__info_string0, inlined_at 1 1289 19 // standard.py:293:12
	shfl.sync.bfly.b32 	%r77, %r76, 2, 31, -1;
$L__tmp19:
	.loc	2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
	add.f32 	%r78, %r76, %r77;
$L__tmp20:
	.loc	2 293 12, function_name $L__info_string0, inlined_at 1 1289 19 // standard.py:293:12
	shfl.sync.bfly.b32 	%r79, %r78, 1, 31, -1;
$L__tmp21:
	.loc	2 263 12, function_name $L__info_string0, inlined_at 2 293 12 // standard.py:263:12
	add.f32 	%r80, %r78, %r79;
$L__tmp22:
	.loc	1 1289 19                       // ltxv_nodes.py:1289:19
	cvt.rn.f32.s32 	%r81, %r36;
	div.full.f32 	%r82, %r80, %r81;
	.loc	1 1290 36                       // ltxv_nodes.py:1290:36
	add.f32 	%r83, %r82, 0f358637BD;
	.loc	1 1290 22                       // ltxv_nodes.py:1290:22
	rsqrt.approx.ftz.f32 	%r84, %r83;
	.loc	1 1290 18                       // ltxv_nodes.py:1290:18
	mul.f32 	%r85, %r84, %r39;
	mul.f32 	%r86, %r84, %r38;
	mul.f32 	%r87, %r84, %r41;
	mul.f32 	%r88, %r84, %r40;
	mul.f32 	%r89, %r84, %r43;
	mul.f32 	%r90, %r84, %r42;
	mul.f32 	%r91, %r84, %r45;
	mul.f32 	%r92, %r84, %r44;
	.loc	1 1298 26                       // ltxv_nodes.py:1298:26
	add.s64 	%rd2, %rd11, %rd15;
	.loc	1 1298 18                       // ltxv_nodes.py:1298:18
	// begin inline asm
	.loc	1 1298 18                       // ltxv_nodes.py:1298:18
	mov.u32 %r8, %r5;
	.loc	1 1299 18                       // ltxv_nodes.py:1299:18
	mov.u32 %r9, %r5;
	.loc	1 1300 18                       // ltxv_nodes.py:1300:18
	mov.u32 %r10, %r5;
	.loc	1 1301 18                       // ltxv_nodes.py:1301:18
	mov.u32 %r11, %r5;
	@%p1 ld.global.v4.b32 { %r8, %r9, %r10, %r11 }, [ %rd2 + 0 ];

	// end inline asm
	.loc	1 1299 26                       // ltxv_nodes.py:1299:26
	add.s64 	%rd3, %rd12, %rd15;
	.loc	1 1299 18                       // ltxv_nodes.py:1299:18
	// begin inline asm
	.loc	1 1299 18                       // ltxv_nodes.py:1299:18
	mov.u32 %r12, %r5;
	.loc	1 1300 18                       // ltxv_nodes.py:1300:18
	mov.u32 %r13, %r5;
	.loc	1 1301 18                       // ltxv_nodes.py:1301:18
	mov.u32 %r14, %r5;
	.loc	1 1302 18                       // ltxv_nodes.py:1302:18
	mov.u32 %r15, %r5;
	@%p1 ld.global.v4.b32 { %r12, %r13, %r14, %r15 }, [ %rd3 + 0 ];

	// end inline asm
	.loc	1 1300 26                       // ltxv_nodes.py:1300:26
	add.s64 	%rd4, %rd13, %rd15;
	.loc	1 1300 18                       // ltxv_nodes.py:1300:18
	// begin inline asm
	.loc	1 1300 18                       // ltxv_nodes.py:1300:18
	mov.u32 %r16, %r5;
	.loc	1 1301 18                       // ltxv_nodes.py:1301:18
	mov.u32 %r17, %r5;
	.loc	1 1302 18                       // ltxv_nodes.py:1302:18
	mov.u32 %r18, %r5;
	.loc	1 1303 18                       // ltxv_nodes.py:1303:18
	mov.u32 %r19, %r5;
	@%p1 ld.global.v4.b32 { %r16, %r17, %r18, %r19 }, [ %rd4 + 0 ];

	// end inline asm
	.loc	1 1301 26                       // ltxv_nodes.py:1301:26
	add.s64 	%rd5, %rd14, %rd15;
	.loc	1 1301 18                       // ltxv_nodes.py:1301:18
	// begin inline asm
	.loc	1 1301 18                       // ltxv_nodes.py:1301:18
	mov.u32 %r20, %r5;
	.loc	1 1302 18                       // ltxv_nodes.py:1302:18
	mov.u32 %r21, %r5;
	.loc	1 1303 18                       // ltxv_nodes.py:1303:18
	mov.u32 %r22, %r5;
	.loc	1 1304 18                       // ltxv_nodes.py:1304:18
	mov.u32 %r23, %r5;
	@%p1 ld.global.v4.b32 { %r20, %r21, %r22, %r23 }, [ %rd5 + 0 ];

	// end inline asm
	.loc	1 1307 22                       // ltxv_nodes.py:1307:22
	add.s64 	%rd6, %rd9, %rd15;
	.loc	1 1298 18                       // ltxv_nodes.py:1298:18
	mov.b32 	{%rs9, %rs10}, %r8;
	cvt.f32.bf16 	%r93, %rs9;
	cvt.f32.bf16 	%r94, %rs10;
	.loc	1 1299 18                       // ltxv_nodes.py:1299:18
	mov.b32 	{%rs11, %rs12}, %r12;
	cvt.f32.bf16 	%r95, %rs12;
	cvt.f32.bf16 	%r96, %rs11;
	.loc	1 1303 26                       // ltxv_nodes.py:1303:26
	add.f32 	%r97, %r94, 0f3F800000;
	add.f32 	%r98, %r93, 0f3F800000;
	.loc	1 1303 16                       // ltxv_nodes.py:1303:16
	fma.rn.f32 	%r99, %r86, %r98, %r96;
	fma.rn.f32 	%r100, %r85, %r97, %r95;
	.loc	1 1307 38                       // ltxv_nodes.py:1307:38
	cvt.rn.bf16x2.f32 	%r24, %r100, %r99;
	.loc	1 1298 18                       // ltxv_nodes.py:1298:18
	mov.b32 	{%rs13, %rs14}, %r9;
	cvt.f32.bf16 	%r101, %rs13;
	cvt.f32.bf16 	%r102, %rs14;
	.loc	1 1299 18                       // ltxv_nodes.py:1299:18
	mov.b32 	{%rs15, %rs16}, %r13;
	cvt.f32.bf16 	%r103, %rs16;
	cvt.f32.bf16 	%r104, %rs15;
	.loc	1 1303 26                       // ltxv_nodes.py:1303:26
	add.f32 	%r105, %r102, 0f3F800000;
	add.f32 	%r106, %r101, 0f3F800000;
	.loc	1 1303 16                       // ltxv_nodes.py:1303:16
	fma.rn.f32 	%r107, %r88, %r106, %r104;
	fma.rn.f32 	%r108, %r87, %r105, %r103;
	.loc	1 1307 38                       // ltxv_nodes.py:1307:38
	cvt.rn.bf16x2.f32 	%r25, %r108, %r107;
	.loc	1 1298 18                       // ltxv_nodes.py:1298:18
	mov.b32 	{%rs17, %rs18}, %r10;
	cvt.f32.bf16 	%r109, %rs17;
	cvt.f32.bf16 	%r110, %rs18;
	.loc	1 1299 18                       // ltxv_nodes.py:1299:18
	mov.b32 	{%rs19, %rs20}, %r14;
	cvt.f32.bf16 	%r111, %rs20;
	cvt.f32.bf16 	%r112, %rs19;
	.loc	1 1303 26                       // ltxv_nodes.py:1303:26
	add.f32 	%r113, %r110, 0f3F800000;
	add.f32 	%r114, %r109, 0f3F800000;
	.loc	1 1303 16                       // ltxv_nodes.py:1303:16
	fma.rn.f32 	%r115, %r90, %r114, %r112;
	fma.rn.f32 	%r116, %r89, %r113, %r111;
	.loc	1 1307 38                       // ltxv_nodes.py:1307:38
	cvt.rn.bf16x2.f32 	%r26, %r116, %r115;
	.loc	1 1298 18                       // ltxv_nodes.py:1298:18
	mov.b32 	{%rs21, %rs22}, %r11;
	cvt.f32.bf16 	%r117, %rs21;
	cvt.f32.bf16 	%r118, %rs22;
	.loc	1 1299 18                       // ltxv_nodes.py:1299:18
	mov.b32 	{%rs23, %rs24}, %r15;
	cvt.f32.bf16 	%r119, %rs24;
	cvt.f32.bf16 	%r120, %rs23;
	.loc	1 1303 26                       // ltxv_nodes.py:1303:26
	add.f32 	%r121, %r118, 0f3F800000;
	add.f32 	%r122, %r117, 0f3F800000;
	.loc	1 1303 16                       // ltxv_nodes.py:1303:16
	fma.rn.f32 	%r123, %r92, %r122, %r120;
	fma.rn.f32 	%r124, %r91, %r121, %r119;
	.loc	1 1307 38                       // ltxv_nodes.py:1307:38
	cvt.rn.bf16x2.f32 	%r27, %r124, %r123;
	.loc	1 1307 13                       // ltxv_nodes.py:1307:13
	// begin inline asm
	@%p1 st.global.v4.b32 [ %rd6 + 0 ], { %r24, %r25, %r26, %r27 };

	// end inline asm
	.loc	1 1308 22                       // ltxv_nodes.py:1308:22
	add.s64 	%rd7, %rd10, %rd15;
	.loc	1 1300 18                       // ltxv_nodes.py:1300:18
	mov.b32 	{%rs25, %rs26}, %r16;
	cvt.f32.bf16 	%r125, %rs25;
	cvt.f32.bf16 	%r126, %rs26;
	.loc	1 1301 18                       // ltxv_nodes.py:1301:18
	mov.b32 	{%rs27, %rs28}, %r20;
	cvt.f32.bf16 	%r127, %rs28;
	cvt.f32.bf16 	%r128, %rs27;
	.loc	1 1304 26                       // ltxv_nodes.py:1304:26
	add.f32 	%r129, %r126, 0f3F800000;
	add.f32 	%r130, %r125, 0f3F800000;
	.loc	1 1304 16                       // ltxv_nodes.py:1304:16
	fma.rn.f32 	%r131, %r86, %r130, %r128;
	fma.rn.f32 	%r132, %r85, %r129, %r127;
	.loc	1 1308 38                       // ltxv_nodes.py:1308:38
	cvt.rn.bf16x2.f32 	%r28, %r132, %r131;
	.loc	1 1300 18                       // ltxv_nodes.py:1300:18
	mov.b32 	{%rs29, %rs30}, %r17;
	cvt.f32.bf16 	%r133, %rs29;
	cvt.f32.bf16 	%r134, %rs30;
	.loc	1 1301 18                       // ltxv_nodes.py:1301:18
	mov.b32 	{%rs31, %rs32}, %r21;
	cvt.f32.bf16 	%r135, %rs32;
	cvt.f32.bf16 	%r136, %rs31;
	.loc	1 1304 26                       // ltxv_nodes.py:1304:26
	add.f32 	%r137, %r134, 0f3F800000;
	add.f32 	%r138, %r133, 0f3F800000;
	.loc	1 1304 16                       // ltxv_nodes.py:1304:16
	fma.rn.f32 	%r139, %r88, %r138, %r136;
	fma.rn.f32 	%r140, %r87, %r137, %r135;
	.loc	1 1308 38                       // ltxv_nodes.py:1308:38
	cvt.rn.bf16x2.f32 	%r29, %r140, %r139;
	.loc	1 1300 18                       // ltxv_nodes.py:1300:18
	mov.b32 	{%rs33, %rs34}, %r18;
	cvt.f32.bf16 	%r141, %rs33;
	cvt.f32.bf16 	%r142, %rs34;
	.loc	1 1301 18                       // ltxv_nodes.py:1301:18
	mov.b32 	{%rs35, %rs36}, %r22;
	cvt.f32.bf16 	%r143, %rs36;
	cvt.f32.bf16 	%r144, %rs35;
	.loc	1 1304 26                       // ltxv_nodes.py:1304:26
	add.f32 	%r145, %r142, 0f3F800000;
	add.f32 	%r146, %r141, 0f3F800000;
	.loc	1 1304 16                       // ltxv_nodes.py:1304:16
	fma.rn.f32 	%r147, %r90, %r146, %r144;
	fma.rn.f32 	%r148, %r89, %r145, %r143;
	.loc	1 1308 38                       // ltxv_nodes.py:1308:38
	cvt.rn.bf16x2.f32 	%r30, %r148, %r147;
	.loc	1 1300 18                       // ltxv_nodes.py:1300:18
	mov.b32 	{%rs37, %rs38}, %r19;
	cvt.f32.bf16 	%r149, %rs37;
	cvt.f32.bf16 	%r150, %rs38;
	.loc	1 1301 18                       // ltxv_nodes.py:1301:18
	mov.b32 	{%rs39, %rs40}, %r23;
	cvt.f32.bf16 	%r151, %rs40;
	cvt.f32.bf16 	%r152, %rs39;
	.loc	1 1304 26                       // ltxv_nodes.py:1304:26
	add.f32 	%r153, %r150, 0f3F800000;
	add.f32 	%r154, %r149, 0f3F800000;
	.loc	1 1304 16                       // ltxv_nodes.py:1304:16
	fma.rn.f32 	%r155, %r92, %r154, %r152;
	fma.rn.f32 	%r156, %r91, %r153, %r151;
	.loc	1 1308 38                       // ltxv_nodes.py:1308:38
	cvt.rn.bf16x2.f32 	%r31, %r156, %r155;
	.loc	1 1308 13                       // ltxv_nodes.py:1308:13
	// begin inline asm
	@%p1 st.global.v4.b32 [ %rd7 + 0 ], { %r28, %r29, %r30, %r31 };

	// end inline asm
	.loc	1 1273 5                        // ltxv_nodes.py:1273:5
	ret;
$L__tmp23:
$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 5                                   // DW_FORM_data2
.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 226                                // 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:0xdb 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:0x24 DW_TAG_subprogram
.b8 95                                  // DW_AT_name
.b8 114
.b8 109
.b8 115
.b8 95
.b8 110
.b8 111
.b8 114
.b8 109
.b8 95
.b8 100
.b8 117
.b8 97
.b8 108
.b8 95
.b8 115
.b8 99
.b8 97
.b8 108
.b8 101
.b8 95
.b8 115
.b8 104
.b8 105
.b8 102
.b8 116
.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] 0x9c:0x49 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] 0xb1:0x33 DW_TAG_inlined_subroutine
.b32 120                                // DW_AT_abstract_origin
.b64 $L__tmp1                           // DW_AT_low_pc
.b64 $L__tmp22                          // DW_AT_high_pc
.b8 1                                   // DW_AT_call_file
.b8 9                                   // DW_AT_call_line
.b8 5
.b8 19                                  // DW_AT_call_column
.b8 5                                   // Abbrev [5] 0xca:0x19 DW_TAG_inlined_subroutine
.b32 120                                // DW_AT_abstract_origin
.b64 $L__tmp1                           // DW_AT_low_pc
.b64 $L__tmp22                          // DW_AT_high_pc
.b8 2                                   // DW_AT_call_file
.b8 37                                  // DW_AT_call_line
.b8 1
.b8 12                                  // 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 ; _rms_norm_dual_scale_shift_kernel
.b8 114
.b8 109
.b8 115
.b8 95
.b8 110
.b8 111
.b8 114
.b8 109
.b8 95
.b8 100
.b8 117
.b8 97
.b8 108
.b8 95
.b8 115
.b8 99
.b8 97
.b8 108
.b8 101
.b8 95
.b8 115
.b8 104
.b8 105
.b8 102
.b8 116
.b8 95
.b8 107
.b8 101
.b8 114
.b8 110
.b8 101
.b8 108
.b8 0
	}
	.section	.debug_macinfo	{	}