A   +       $      e垿Dec  3 202420:16:52HOST64sm_86//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-31833905
// Cuda compilation tools, release 11.8, V11.8.89
// Based on NVVM 7.0.1
//

.version 7.8
.target sm_50
.address_size 64

	// .globl	ShaderKernel_fxLegalizePAL
.const .align 4 .b8 kRGB32f_To_601YPbPr[36] = {135, 22, 153, 62, 162, 69, 22, 63, 213, 120, 233, 61, 33, 201, 44, 190, 111, 155, 169, 190, 0, 0, 0, 63, 0, 0, 0, 63, 70, 94, 214, 190, 232, 134, 166, 189};
.const .align 4 .b8 k601YPbPr_To_RGB32f[36] = {0, 0, 128, 63, 0, 0, 0, 0, 188, 116, 179, 63, 0, 0, 128, 63, 152, 50, 176, 190, 158, 209, 54, 191, 0, 0, 128, 63, 229, 208, 226, 63, 0, 0, 0, 0};
.const .align 4 .b8 kRGB32f_To_601YCbCr[36] = {70, 246, 130, 66, 145, 141, 0, 67, 94, 186, 199, 65, 33, 48, 23, 194, 240, 103, 148, 194, 0, 0, 224, 66, 0, 0, 224, 66, 111, 146, 187, 194, 70, 182, 145, 193};
.const .align 4 .b8 k601YCbCr_To_RGB32f[36] = {37, 160, 149, 59, 0, 0, 0, 0, 182, 23, 205, 59, 37, 160, 149, 59, 40, 15, 201, 186, 156, 239, 80, 187, 37, 160, 149, 59, 236, 155, 1, 60, 0, 0, 0, 0};
.const .align 4 .b8 kRGB8u_To_601YCbCr[36] = {219, 121, 131, 62, 152, 14, 1, 63, 18, 131, 200, 61, 174, 199, 23, 190, 238, 252, 148, 190, 197, 224, 224, 62, 197, 224, 224, 62, 217, 78, 188, 190, 174, 71, 146, 189};
.const .align 4 .b8 k601YCbCr_To_RGB8u[36] = {127, 10, 149, 63, 0, 0, 0, 0, 160, 74, 204, 63, 127, 10, 149, 63, 254, 148, 200, 190, 184, 30, 80, 191, 127, 10, 149, 63, 78, 26, 1, 64, 0, 0, 0, 0};
.const .align 4 .b8 kRGB8u_To_601YCbCrFullRange[36] = {135, 22, 153, 62, 162, 69, 22, 63, 213, 120, 233, 61, 166, 27, 44, 190, 39, 241, 168, 190, 250, 254, 254, 62, 250, 254, 254, 62, 43, 135, 213, 190, 59, 223, 165, 189};
.const .align 4 .b8 k601YCbCrFullRange_To_RGB8u[36] = {0, 0, 128, 63, 0, 0, 0, 0, 72, 193, 178, 63, 0, 0, 128, 63, 143, 130, 175, 190, 225, 26, 54, 191, 0, 0, 128, 63, 20, 238, 225, 63, 0, 0, 0, 0};
.const .align 4 .b8 kRGB32f_To_601YCbCrFullRange[36] = {113, 125, 152, 66, 92, 175, 21, 67, 92, 143, 232, 65, 158, 111, 43, 194, 49, 72, 168, 194, 0, 0, 254, 66, 0, 0, 254, 66, 170, 177, 212, 194, 88, 57, 165, 193};
.const .align 4 .b8 k601YCbCrFullRange_To_RGB32f[36] = {129, 128, 128, 59, 0, 0, 0, 0, 189, 116, 179, 59, 129, 128, 128, 59, 194, 50, 176, 186, 179, 209, 54, 187, 129, 128, 128, 59, 229, 208, 226, 59, 0, 0, 0, 0};
.const .align 4 .b8 kRGB32f_To_709YPbPr[36] = {208, 179, 89, 62, 89, 23, 55, 63, 152, 221, 147, 61, 186, 164, 234, 189, 210, 86, 197, 190, 0, 0, 0, 63, 0, 0, 0, 63, 190, 134, 232, 190, 16, 202, 59, 189};
.const .align 4 .b8 k709YPbPr_To_RGB32f[36] = {0, 0, 128, 63, 0, 0, 0, 0, 12, 147, 201, 63, 0, 0, 128, 63, 221, 209, 63, 190, 243, 173, 239, 190, 0, 0, 128, 63, 77, 132, 237, 63, 0, 0, 0, 0};
.const .align 4 .b8 kRGB32f_To_709YCbCr[36] = {106, 60, 58, 66, 6, 161, 28, 67, 244, 253, 124, 65, 223, 79, 205, 193, 8, 172, 172, 194, 0, 0, 224, 66, 0, 0, 224, 66, 195, 117, 203, 194, 236, 81, 36, 193};
.const .align 4 .b8 k709YCbCr_To_RGB32f[36] = {37, 160, 149, 59, 0, 0, 0, 0, 239, 94, 230, 59, 37, 160, 149, 59, 33, 57, 91, 186, 178, 245, 8, 187, 37, 160, 149, 59, 82, 185, 7, 60, 0, 0, 0, 0};
.const .align 4 .b8 k709YCbCrFullRange_To_RGB32f[36] = {131, 128, 128, 59, 0, 0, 0, 0, 28, 147, 201, 59, 131, 128, 128, 59, 61, 210, 63, 186, 248, 173, 239, 186, 131, 128, 128, 59, 82, 132, 237, 59, 0, 0, 0, 0};
.const .align 4 .b8 kRGB8u_To_709YCbCr[36] = {207, 247, 58, 62, 53, 62, 29, 63, 231, 251, 125, 61, 184, 30, 206, 189, 23, 89, 173, 190, 197, 224, 224, 62, 197, 224, 224, 62, 12, 66, 204, 190, 195, 245, 36, 189};
.const .align 4 .b8 k709YCbCr_To_RGB8u[36] = {127, 10, 149, 63, 0, 0, 0, 0, 147, 120, 229, 63, 127, 10, 149, 63, 53, 94, 90, 190, 205, 108, 8, 191, 127, 10, 149, 63, 154, 49, 7, 64, 0, 0, 0, 0};
.const .align 4 .b8 k709YCbCr_To_601YCbCr[36] = {0, 0, 128, 63, 23, 100, 203, 61, 1, 77, 68, 62, 0, 0, 0, 0, 18, 103, 125, 63, 10, 158, 226, 189, 0, 0, 0, 0, 61, 98, 148, 189, 249, 191, 123, 63};
.const .align 4 .b8 k601YCbCr_To_709YCbCr[36] = {0, 0, 128, 63, 122, 165, 236, 189, 179, 237, 84, 190, 0, 0, 0, 0, 204, 98, 130, 63, 216, 188, 234, 61, 0, 0, 0, 0, 74, 179, 153, 61, 234, 61, 131, 63};
.const .align 4 .b8 kZeroMatrix[36];
.const .align 4 .b8 kYCbCrOffset[12] = {0, 0, 128, 65, 0, 0, 0, 67, 0, 0, 0, 67};
.const .align 4 .b8 kYCbCrFullRangeOffset[12] = {0, 0, 0, 0, 0, 0, 0, 67, 0, 0, 0, 67};
.const .align 4 .f32 PQ_m1 = 0f3E232000;
.const .align 4 .f32 PQ_m1Inv = 0f40C8E06B;
.const .align 4 .f32 PQ_m2 = 0f429DB000;
.const .align 4 .f32 PQ_m2Inv = 0f3C4FCDAC;
.const .align 4 .f32 PQ_c1 = 0f3F560000;
.const .align 4 .f32 PQ_c2 = 0f4196D000;
.const .align 4 .f32 PQ_c3 = 0f41958000;
.const .align 4 .f32 Gamma1886 = 0f4019999A;
.const .align 4 .f32 PQ_Lpeak = 0f461C4000;
.const .align 4 .f32 PQ_a = 0f3F8CAC08;
.const .align 4 .f32 PQ_b = 0f426E1556;
.const .align 4 .f32 PQ_c = 0f39B033E5;
.const .align 4 .f32 PQ_s = 0f4385EB85;
.const .align 4 .f32 PQ_g = 0f3EE66666;
.const .align 4 .f32 scaleFD = 0f42C80000;
.const .align 4 .f32 HLG_a = 0f3E371FF0;
.const .align 4 .f32 HLG_b = 0f3E91C020;
.const .align 4 .f32 HLG_c = 0f3F0F564F;
.const .align 4 .f32 HLG_inva = 0f40B2F029;
.const .align 4 .f32 HLG_alpha = 0f41200000;
.const .align 4 .f32 HLG_invAlpha = 0f3DCCCCCD;
.const .align 4 .f32 HLG_gamma = 0f3F99999A;
.const .align 4 .f32 HLG_gammaM1 = 0f3E4CCCCD;
.const .align 4 .f32 HLG_gammaM1Dgamma = 0f3E2AAAAB;
.const .align 4 .f32 HLG_YR = 0f3E86809D;
.const .align 4 .f32 HLG_YG = 0f3F2D9168;
.const .align 4 .f32 HLG_YB = 0f3D72E48F;
.const .align 4 .f32 HLG_Lpeak = 0f447A0000;
.const .align 4 .f32 AppleLog_R0 = 0fBD670F18;
.const .align 4 .f32 AppleLog_Rt = 0f3C23D70A;
.const .align 4 .f32 AppleLog_c = 0f423D2601;
.const .align 4 .f32 AppleLog_beta = 0f3C1DF346;
.const .align 4 .f32 AppleLog_gamma = 0f3DAF1D23;
.const .align 4 .f32 AppleLog_delta = 0f3F3180A9;
.global .texref texture0_RECT;
// _ZZ35ShaderKernel_fxLegalizePAL_DelegateP6float4PKS_Pvi17DevicePixelFormatii5uint2E7p_local has been demoted

.visible .entry ShaderKernel_fxLegalizePAL(
	.param .u64 ShaderKernel_fxLegalizePAL_param_0,
	.param .u64 ShaderKernel_fxLegalizePAL_param_1,
	.param .u64 ShaderKernel_fxLegalizePAL_param_2,
	.param .u32 ShaderKernel_fxLegalizePAL_param_3,
	.param .u32 ShaderKernel_fxLegalizePAL_param_4,
	.param .u32 ShaderKernel_fxLegalizePAL_param_5,
	.param .u32 ShaderKernel_fxLegalizePAL_param_6
)
{
	.reg .pred 	%p<10>;
	.reg .b16 	%rs<5>;
	.reg .f32 	%f<73>;
	.reg .b32 	%r<23>;
	.reg .b64 	%rd<11>;
	// demoted variable
	.shared .align 16 .b8 _ZZ35ShaderKernel_fxLegalizePAL_DelegateP6float4PKS_Pvi17DevicePixelFormatii5uint2E7p_local[16];

	ld.param.u64 	%rd4, [ShaderKernel_fxLegalizePAL_param_0];
	ld.param.u64 	%rd3, [ShaderKernel_fxLegalizePAL_param_1];
	ld.param.u32 	%r4, [ShaderKernel_fxLegalizePAL_param_3];
	ld.param.u32 	%r5, [ShaderKernel_fxLegalizePAL_param_4];
	ld.param.u32 	%r6, [ShaderKernel_fxLegalizePAL_param_5];
	ld.param.u32 	%r7, [ShaderKernel_fxLegalizePAL_param_6];
	cvta.to.global.u64 	%rd1, %rd4;
	mov.u32 	%r1, %ntid.x;
	mov.u32 	%r8, %ctaid.x;
	mov.u32 	%r9, %tid.x;
	mad.lo.s32 	%r2, %r8, %r1, %r9;
	mov.u32 	%r10, %ntid.y;
	mov.u32 	%r11, %ctaid.y;
	mov.u32 	%r12, %tid.y;
	mad.lo.s32 	%r3, %r11, %r10, %r12;
	or.b32  	%r13, %r12, %r9;
	setp.ne.s32 	%p1, %r13, 0;
	@%p1 bra 	$L__BB0_3;

	setp.eq.s32 	%p2, %r1, 0;
	@%p2 bra 	$L__BB0_9;

	cvta.to.global.u64 	%rd5, %rd3;
	ld.global.v4.u32 	{%r14, %r15, %r16, %r17}, [%rd5];
	st.shared.v4.u32 	[_ZZ35ShaderKernel_fxLegalizePAL_DelegateP6float4PKS_Pvi17DevicePixelFormatii5uint2E7p_local], {%r14, %r15, %r16, %r17};

$L__BB0_3:
	bar.sync 	0;
	setp.ge.s32 	%p3, %r3, %r7;
	setp.ge.s32 	%p4, %r2, %r6;
	or.pred  	%p5, %p4, %p3;
	@%p5 bra 	$L__BB0_7;

	cvt.rn.f32.u32 	%f5, %r2;
	add.ftz.f32 	%f6, %f5, 0f3F000000;
	cvt.rn.f32.u32 	%f7, %r3;
	add.ftz.f32 	%f8, %f7, 0f3F000000;
	tex.2d.v4.f32.f32 	{%f9, %f10, %f11, %f1}, [texture0_RECT, {%f6, %f8}];
	cvt.ftz.sat.f32.f32 	%f12, %f11;
	cvt.ftz.sat.f32.f32 	%f13, %f10;
	cvt.ftz.sat.f32.f32 	%f14, %f9;
	mul.ftz.f32 	%f15, %f13, 0f3F1645A2;
	fma.rn.ftz.f32 	%f16, %f12, 0f3E991687, %f15;
	fma.rn.ftz.f32 	%f17, %f14, 0f3DE978D5, %f16;
	mul.ftz.f32 	%f18, %f12, 0fBE16872B;
	fma.rn.ftz.f32 	%f19, %f13, 0fBE93F7CF, %f18;
	fma.rn.ftz.f32 	%f20, %f14, 0f3EDF3B64, %f19;
	mul.ftz.f32 	%f21, %f13, 0fBF03DD98;
	fma.rn.ftz.f32 	%f22, %f12, 0f3F1D70A4, %f21;
	fma.rn.ftz.f32 	%f23, %f14, 0fBDCCCCCD, %f22;
	mul.ftz.f32 	%f24, %f23, %f23;
	fma.rn.ftz.f32 	%f25, %f20, %f20, %f24;
	add.ftz.f32 	%f26, %f25, 0f00000000;
	rsqrt.approx.ftz.f32 	%f27, %f26;
	mul.ftz.f32 	%f28, %f27, %f26;
	add.ftz.f32 	%f29, %f17, %f28;
	sub.ftz.f32 	%f30, %f17, %f28;
	fma.rn.ftz.f32 	%f31, %f28, 0f3FA9CA57, 0f00000000;
	mov.f32 	%f32, 0f3F11F682;
	fma.rn.ftz.f32 	%f33, %f29, 0f3F11F682, 0f00000000;
	mul.ftz.f32 	%f34, %f30, 0f3F11F682;
	sub.ftz.f32 	%f35, %f32, %f34;
	ld.shared.v4.f32 	{%f36, %f37, %f38, %f39}, [_ZZ35ShaderKernel_fxLegalizePAL_DelegateP6float4PKS_Pvi17DevicePixelFormatii5uint2E7p_local];
	mul.ftz.f32 	%f44, %f37, %f33;
	fma.rn.ftz.f32 	%f45, %f36, %f31, %f44;
	fma.rn.ftz.f32 	%f46, %f38, %f35, %f45;
	add.ftz.f32 	%f47, %f39, %f46;
	mul.ftz.f32 	%f48, %f28, %f47;
	sub.ftz.f32 	%f49, %f28, %f48;
	add.ftz.f32 	%f50, %f17, %f49;
	sub.ftz.f32 	%f51, %f17, %f49;
	max.ftz.f32 	%f52, %f51, 0fBE5B22D1;
	setp.gt.ftz.f32 	%p6, %f50, 0f3F8A1CAC;
	selp.f32 	%f53, 0f3F8A1CAC, %f50, %p6;
	sub.ftz.f32 	%f54, %f53, %f17;
	sub.ftz.f32 	%f55, %f17, %f52;
	setp.gt.ftz.f32 	%p7, %f55, %f54;
	selp.f32 	%f56, %f54, %f55, %p7;
	setp.gt.ftz.f32 	%p8, %f56, 0f3F8A1CAC;
	selp.f32 	%f57, 0f3F8A1CAC, %f56, %p8;
	mov.f32 	%f58, 0f3F800000;
	div.rn.ftz.f32 	%f59, %f58, %f28;
	mul.ftz.f32 	%f60, %f59, %f57;
	mul.ftz.f32 	%f61, %f20, %f60;
	mul.ftz.f32 	%f62, %f23, %f60;
	fma.rn.ftz.f32 	%f63, %f61, 0f38257EA2, %f17;
	fma.rn.ftz.f32 	%f64, %f62, 0f3F91E5E2, %f63;
	fma.rn.ftz.f32 	%f65, %f61, 0fBECA0A58, %f17;
	fma.rn.ftz.f32 	%f66, %f62, 0fBF149BAB, %f65;
	fma.rn.ftz.f32 	%f67, %f61, 0f40020C48, %f17;
	fma.rn.ftz.f32 	%f68, %f62, 0f39FC6131, %f67;
	max.ftz.f32 	%f2, %f64, 0f00000000;
	max.ftz.f32 	%f3, %f66, 0f00000000;
	max.ftz.f32 	%f4, %f68, 0f00000000;
	mad.lo.s32 	%r22, %r3, %r4, %r2;
	cvt.s64.s32 	%rd2, %r22;
	setp.eq.s32 	%p9, %r5, 0;
	@%p9 bra 	$L__BB0_6;

	shl.b64 	%rd7, %rd2, 4;
	add.s64 	%rd8, %rd1, %rd7;
	st.global.v4.f32 	[%rd8], {%f4, %f3, %f2, %f1};
	bra.uni 	$L__BB0_7;

$L__BB0_9:
	bra.uni 	$L__BB0_9;

$L__BB0_6:
	// begin inline asm
	{  cvt.rn.f16.f32 %rs1, %f4;}

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs2, %f3;}

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs3, %f2;}

	// end inline asm
	// begin inline asm
	{  cvt.rn.f16.f32 %rs4, %f1;}

	// end inline asm
	shl.b64 	%rd9, %rd2, 3;
	add.s64 	%rd10, %rd1, %rd9;
	st.global.u16 	[%rd10], %rs1;
	st.global.u16 	[%rd10+2], %rs2;
	st.global.u16 	[%rd10+4], %rs3;
	st.global.u16 	[%rd10+6], %rs4;

$L__BB0_7:
	ret;

}

  ELF3                    @#            V2 @ 8  @    .shstrtab .strtab .symtab .symtab_shndx .nv.info .text.ShaderKernel_fxLegalizePAL .nv.info.ShaderKernel_fxLegalizePAL .nv.shared.ShaderKernel_fxLegalizePAL .nv.constant3 .nv.constant0.ShaderKernel_fxLegalizePAL .rel.nv.constant0.ShaderKernel_fxLegalizePAL .debug_frame .rel.debug_frame .rela.debug_frame .nv.callgraph .nv.prototype .nv.rel.action  .shstrtab .strtab .symtab .symtab_shndx .nv.info .text.ShaderKernel_fxLegalizePAL .nv.info.ShaderKernel_fxLegalizePAL .nv.shared.ShaderKernel_fxLegalizePAL .nv.constant3 kRGB32f_To_601YPbPr k601YPbPr_To_RGB32f kRGB32f_To_601YCbCr k601YCbCr_To_RGB32f kRGB8u_To_601YCbCr k601YCbCr_To_RGB8u kRGB8u_To_601YCbCrFullRange k601YCbCrFullRange_To_RGB8u kRGB32f_To_601YCbCrFullRange k601YCbCrFullRange_To_RGB32f kRGB32f_To_709YPbPr k709YPbPr_To_RGB32f kRGB32f_To_709YCbCr k709YCbCr_To_RGB32f k709YCbCrFullRange_To_RGB32f kRGB8u_To_709YCbCr k709YCbCr_To_RGB8u k709YCbCr_To_601YCbCr k601YCbCr_To_709YCbCr kZeroMatrix kYCbCrOffset kYCbCrFullRangeOffset PQ_m1 PQ_m1Inv PQ_m2 PQ_m2Inv PQ_c1 PQ_c2 PQ_c3 Gamma1886 PQ_Lpeak PQ_a PQ_b PQ_c PQ_s PQ_g scaleFD HLG_a HLG_b HLG_c HLG_inva HLG_alpha HLG_invAlpha HLG_gamma HLG_gammaM1 HLG_gammaM1Dgamma HLG_YR HLG_YG HLG_YB HLG_Lpeak AppleLog_R0 AppleLog_Rt AppleLog_c AppleLog_beta AppleLog_gamma AppleLog_delta $__internal_0_$__cuda_sm3x_div_rn_ftz_f32_slowpath .rel.nv.constant0.ShaderKernel_fxLegalizePAL .nv.constant0.ShaderKernel_fxLegalizePAL .debug_frame .rel.debug_frame .rela.debug_frame .nv.callgraph .nv.prototype .nv.rel.action ShaderKernel_fxLegalizePAL texture0_RECT                           2                     w                                                       $            $       $            H       $            l       $                   $                  $       !           $       =           $       Y           $       v    D      $           h      $                 $                 $                 $                 $                  $           @      $       &    d      $       <          $       R          $       ^                 k                                                                                                                                                                                                                                                                                      $                 (                 ,                 0                  4             
    8                 <             !    @             -    D             ?    H             F    L             M    P             T    T             ^    X             j    \             v    `                 d                 h                 l               "                                   6                    f                                                                              $        |( ((   4                                 $   (          <        |(( ((((        ( " ,       p                       	((      / A       <        A        A       7    5  
 =   `( (       $                                       !         !           !     @  0                             s          % 6         B             A   D          A             A   >E?x=!,o   ?   ?F^־膦  ?    t?  ?26  ??    FB C^A!0g  B  BoF%;    ;%;(ɺP%;<    y>?=>>NG
?    J?
?ȾP
?N@    >E?x=,'>>+վ;ߥ  ?    H?  ?6  ??    q}B\C\Ao+1H  B  BX9;    t;;26;;    гY>Y7?ݓ=Vž   ?   ?;  ?    ?  ??  ?M?    j<:BC|AO  B  BuQ$%;    ^;%;!9[%;R<    ;    ;;=?ﺃ;R;    :>5>?}=νY>>B̾$
?    x?
?5^Zl
?1@      ?d=MD>    g}?
    =b{?  ?z콳T    b?ؼ=    J==?                                      A   C   C       C   C  #>k@ BO<  V? ЖA A@ @F?VnB39Cff>  B7> >OV?)@   A=?L>*>>h-?r=  zDg
#<&=BF<#=1?                                                                                                                                                                                                                                                                                                                                                                                                                                                z  
       y      !   " z  F   
    y       "   ( y      %   h y	      &    r      $z     / $z		     O G	  `       z     pR  G  0      z  Z       z  [       y      s        A Hy      {           z 	 a  pb  z  `  pfp   M	          Er          Er 	        x         !t    ?     !t   ?     `{b 0 B x   ?      Ey 0      !r       !r         !r         x  A    x +  A    xE?  A   #x
 p?    #x    #x  >    y        " #x
̽
   #xd;>   #xx=     x ?      r

     A   #r       !r        s        d  r      A  / !r         #x?    !r         r     A   #x Wʩ?    #t?   #r       s         $ #r       |        !r        s         d #r        !r       !r       #t   ?    	x?    	x"[    #r        !r      !r       #x  ?     r     @  #t   ?    r         #r        	x?    G  0      / x        Dy         r         Ay          r     A   z  _  pR   r      A   r
     A   $z	 ^    #x ~%8   #x X
ʾ   #x H@    x
      #x?   #x    #x1a9    
 X     	r       r         
	 Y  
$   	r        	r       	      M	          z X    >r        >r        z Y  
   x 2v        x2v       y      y      y     y     My          Gy    x       Ey  0      Ey        x     x  x    p@  G         x       B	         G	        x       B         G  0      r        B         G        Ay         x  Ey       $x     s         " !r      #t   ?     #r           #x   ?     #t  ?     #r         #t  ?     #r        x       x       r       x   x    p`  G  `       x     pB  x          pb            Gy         $x    Ay         Gy         x    ?<  x       x     Gy  @       x    ?<  x      Gy         !t   ?     Ay          r         x          Py   Gy    y            y            y            y            y            y            y            y            y            y            y            y            y            y            y                                                                                                  @       \                                                                                                            X      H         A                                                                          )      p                      0                             S      p@                                                 1    p                l                                   M    p                                                      	   @                                                  	   @                     0                                                      p                                   B               P                                  2                                    A                 w      C                                                        @#                                                                                                                                                  @#                                           