//
// Generated by NVIDIA NVVM Compiler
// Compiler built on Wed Jul 10 12:41:20 2013 (1373485280)
// Cuda compilation tools, release 5.5, V5.5.0
//

.version 3.2
.target sm_30
.address_size 64

	.file	1 "c:/caravan/releases/2014.09/shared/adobe/MediaCore/GPUFoundation/Src/ImageProcessing/Win/PixelFormatConvert_NV12Centered.cu", 1411781712, 5912
	.file	2 "c:\\caravan\\releases\\2014.09\\shared\\adobe\\MediaCore\\GPUFoundation\\Inc\\ImageProcessing/PixelFormatConvert_Common.h", 1411781711, 21667
	.file	3 "c:\\caravan\\releases\\2014.09\\shared\\adobe\\mediacore\\external\\3rdparty\\nvidia\\cuda\\win\\include\\device_functions.h", 1411781673, 191626
	.file	4 "c:\\caravan\\releases\\2014.09\\shared\\adobe\\MediaCore\\GPUFoundation\\API\\Inc\\GPUFoundation/KernelSupport/KernelCore.h", 1411781711, 7840
.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, 188, 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, 147, 24, 206, 61, 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 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};
.global .align 1 .b8 $str[11] = {95, 95, 67, 85, 68, 65, 95, 70, 84, 90, 0};

.visible .func  (.param .b32 func_retval0) _Z13MaxDepthValuej(
	.param .b32 _Z13MaxDepthValuej_param_0
)
{
	.reg .pred 	%p<4>;
	.reg .s32 	%r<3>;
	.reg .f32 	%f<6>;


	ld.param.u32 	%r2, [_Z13MaxDepthValuej_param_0];
	.loc 2 51 1
	and.b32  	%r1, %r2, 248;
	setp.ne.s32	%p1, %r1, 0;
	@%p1 bra 	BB0_2;

	mov.f32 	%f5, 0f437F0000;
	bra.uni 	BB0_5;

BB0_2:
	.loc 2 51 1
	setp.ne.s32	%p2, %r1, 8;
	@%p2 bra 	BB0_4;

	mov.f32 	%f5, 0f447FC000;
	bra.uni 	BB0_5;

BB0_4:
	.loc 2 51 1
	setp.eq.s32	%p3, %r1, 16;
	.loc 2 51 1
	selp.f32	%f5, 0f47000000, 0f3F800000, %p3;

BB0_5:
	st.param.f32	[func_retval0+0], %f5;
	.loc 2 51 1
	ret;
}

.visible .func  (.param .b64 func_retval0) _Z23ColorSpaceConvertMatrixjj(
	.param .b32 _Z23ColorSpaceConvertMatrixjj_param_0,
	.param .b32 _Z23ColorSpaceConvertMatrixjj_param_1
)
{
	.reg .pred 	%p<33>;
	.reg .s32 	%r<16>;
	.reg .s64 	%rd<32>;


	ld.param.u32 	%r9, [_Z23ColorSpaceConvertMatrixjj_param_0];
	ld.param.u32 	%r10, [_Z23ColorSpaceConvertMatrixjj_param_1];
	.loc 2 51 1
	and.b32  	%r11, %r9, 512;
	setp.eq.s32	%p1, %r11, 0;
	.loc 2 51 1
	and.b32  	%r1, %r10, 256;
	.loc 2 51 1
	@%p1 bra 	BB1_29;

	.loc 2 51 1
	and.b32  	%r12, %r9, 2048;
	setp.eq.s32	%p2, %r12, 0;
	.loc 2 51 1
	and.b32  	%r2, %r9, 4096;
	.loc 2 51 1
	and.b32  	%r3, %r10, 512;
	.loc 2 51 1
	@%p2 bra 	BB1_15;

	setp.eq.s32	%p3, %r2, 0;
	.loc 2 51 1
	@%p3 bra 	BB1_13;

	setp.eq.s32	%p4, %r3, 0;
	.loc 2 51 1
	and.b32  	%r4, %r9, 256;
	.loc 2 51 1
	@%p4 bra 	BB1_7;

	.loc 2 51 1
	setp.ne.s32	%p5, %r4, 0;
	mov.u64 	%rd31, 0;
	.loc 2 51 1
	@%p5 bra 	BB1_47;

	.loc 2 51 1
	and.b32  	%r13, %r10, 4096;
	setp.eq.s32	%p6, %r13, 0;
	.loc 2 51 1
	setp.ne.s32	%p7, %r1, 0;
	or.pred  	%p8, %p6, %p7;
	mov.u64 	%rd31, 0;
	.loc 2 51 1
	@%p8 bra 	BB1_47;

	.loc 2 51 1
	cvta.const.u64 	%rd31, k709YCbCr_To_601YCbCr;
	bra.uni 	BB1_47;

BB1_7:
	setp.eq.s32	%p9, %r4, 0;
	.loc 2 51 1
	@%p9 bra 	BB1_10;

	setp.eq.s32	%p10, %r1, 0;
	mov.u64 	%rd31, 0;
	.loc 2 51 1
	@%p10 bra 	BB1_47;

	.loc 2 51 1
	cvta.const.u64 	%rd31, k709YPbPr_To_RGB32f;
	bra.uni 	BB1_47;

BB1_10:
	setp.eq.s32	%p11, %r1, 0;
	.loc 2 51 1
	@%p11 bra 	BB1_12;

	.loc 2 51 1
	cvta.const.u64 	%rd31, k709YCbCr_To_RGB32f;
	bra.uni 	BB1_47;

BB1_12:
	.loc 2 51 1
	cvta.const.u64 	%rd31, k709YCbCr_To_RGB8u;
	bra.uni 	BB1_47;

BB1_13:
	.loc 2 51 1
	setp.ne.s32	%p12, %r3, 0;
	setp.eq.s32	%p13, %r1, 0;
	.loc 2 51 1
	or.pred  	%p14, %p13, %p12;
	mov.u64 	%rd31, 0;
	.loc 2 51 1
	@%p14 bra 	BB1_47;

	.loc 2 51 1
	cvta.const.u64 	%rd31, k709YCbCrFullRange_To_RGB32f;
	bra.uni 	BB1_47;

BB1_15:
	setp.eq.s32	%p15, %r2, 0;
	.loc 2 51 1
	@%p15 bra 	BB1_25;

	setp.eq.s32	%p16, %r3, 0;
	.loc 2 51 1
	and.b32  	%r5, %r9, 256;
	.loc 2 51 1
	@%p16 bra 	BB1_19;

	or.b32  	%r14, %r5, %r1;
	.loc 2 51 1
	setp.ne.s32	%p17, %r14, 0;
	mov.u64 	%rd31, 0;
	.loc 2 51 1
	@%p17 bra 	BB1_47;

	.loc 2 51 1
	cvta.const.u64 	%rd31, k601YCbCr_To_709YCbCr;
	bra.uni 	BB1_47;

BB1_19:
	setp.eq.s32	%p18, %r5, 0;
	.loc 2 51 1
	@%p18 bra 	BB1_22;

	setp.eq.s32	%p19, %r1, 0;
	mov.u64 	%rd31, 0;
	.loc 2 51 1
	@%p19 bra 	BB1_47;

	.loc 2 51 1
	cvta.const.u64 	%rd31, k601YPbPr_To_RGB32f;
	bra.uni 	BB1_47;

BB1_22:
	setp.eq.s32	%p20, %r1, 0;
	.loc 2 51 1
	@%p20 bra 	BB1_24;

	.loc 2 51 1
	cvta.const.u64 	%rd31, k601YCbCr_To_RGB32f;
	bra.uni 	BB1_47;

BB1_24:
	.loc 2 51 1
	cvta.const.u64 	%rd31, k601YCbCr_To_RGB8u;
	bra.uni 	BB1_47;

BB1_25:
	.loc 2 51 1
	setp.ne.s32	%p21, %r3, 0;
	mov.u64 	%rd31, 0;
	.loc 2 51 1
	@%p21 bra 	BB1_47;

	setp.eq.s32	%p22, %r1, 0;
	.loc 2 51 1
	@%p22 bra 	BB1_28;

	.loc 2 51 1
	cvta.const.u64 	%rd31, k601YCbCrFullRange_To_RGB32f;
	bra.uni 	BB1_47;

BB1_28:
	.loc 2 51 1
	cvta.const.u64 	%rd31, k601YCbCrFullRange_To_RGB8u;
	bra.uni 	BB1_47;

BB1_29:
	.loc 2 51 1
	and.b32  	%r15, %r9, 256;
	setp.eq.s32	%p23, %r15, 0;
	.loc 2 51 1
	and.b32  	%r6, %r10, 2048;
	.loc 2 51 1
	@%p23 bra 	BB1_40;

	setp.eq.s32	%p24, %r1, 0;
	.loc 2 51 1
	@%p24 bra 	BB1_34;

	setp.eq.s32	%p25, %r6, 0;
	.loc 2 51 1
	@%p25 bra 	BB1_33;

	.loc 2 51 1
	cvta.const.u64 	%rd31, kRGB32f_To_709YPbPr;
	bra.uni 	BB1_47;

BB1_33:
	.loc 2 51 1
	cvta.const.u64 	%rd31, kRGB32f_To_601YPbPr;
	bra.uni 	BB1_47;

BB1_34:
	setp.eq.s32	%p26, %r6, 0;
	.loc 2 51 1
	and.b32  	%r7, %r10, 4096;
	.loc 2 51 1
	@%p26 bra 	BB1_37;

	setp.eq.s32	%p27, %r7, 0;
	mov.u64 	%rd31, 0;
	.loc 2 51 1
	@%p27 bra 	BB1_47;

	.loc 2 51 1
	cvta.const.u64 	%rd31, kRGB32f_To_709YCbCr;
	bra.uni 	BB1_47;

BB1_37:
	setp.eq.s32	%p28, %r7, 0;
	.loc 2 51 1
	@%p28 bra 	BB1_39;

	.loc 2 51 1
	cvta.const.u64 	%rd31, kRGB32f_To_601YCbCr;
	bra.uni 	BB1_47;

BB1_39:
	.loc 2 51 1
	cvta.const.u64 	%rd31, kRGB32f_To_601YCbCrFullRange;
	bra.uni 	BB1_47;

BB1_40:
	.loc 2 51 1
	setp.ne.s32	%p29, %r1, 0;
	mov.u64 	%rd31, 0;
	.loc 2 51 1
	@%p29 bra 	BB1_47;

	setp.eq.s32	%p30, %r6, 0;
	.loc 2 51 1
	and.b32  	%r8, %r10, 4096;
	.loc 2 51 1
	@%p30 bra 	BB1_44;

	.loc 2 51 1
	setp.ne.s32	%p31, %r8, 0;
	mov.u64 	%rd31, 0;
	.loc 2 51 1
	@%p31 bra 	BB1_47;

	.loc 2 51 1
	cvta.const.u64 	%rd31, kRGB8u_To_709YCbCr;
	bra.uni 	BB1_47;

BB1_44:
	setp.eq.s32	%p32, %r8, 0;
	.loc 2 51 1
	@%p32 bra 	BB1_46;

	.loc 2 51 1
	cvta.const.u64 	%rd31, kRGB8u_To_601YCbCr;
	bra.uni 	BB1_47;

BB1_46:
	.loc 2 51 1
	cvta.const.u64 	%rd31, kRGB8u_To_601YCbCrFullRange;

BB1_47:
	st.param.b64	[func_retval0+0], %rd31;
	.loc 2 51 1
	ret;
}

.visible .func  (.param .b32 func_retval0) _Z5clampIfET_S0_S0_S0_(
	.param .b32 _Z5clampIfET_S0_S0_S0__param_0,
	.param .b32 _Z5clampIfET_S0_S0_S0__param_1,
	.param .b32 _Z5clampIfET_S0_S0_S0__param_2
)
{
	.reg .f32 	%f<6>;


	ld.param.f32 	%f1, [_Z5clampIfET_S0_S0_S0__param_0];
	ld.param.f32 	%f2, [_Z5clampIfET_S0_S0_S0__param_1];
	ld.param.f32 	%f3, [_Z5clampIfET_S0_S0_S0__param_2];
	.loc 3 2770 10
	max.ftz.f32 	%f4, %f1, %f2;
	.loc 3 2765 10
	min.ftz.f32 	%f5, %f4, %f3;
	st.param.f32	[func_retval0+0], %f5;
	.loc 4 146 39
	ret;
}

.visible .func  (.param .align 16 .b8 func_retval0[16]) _Z23UnpremultiplyComponents6float4j(
	.param .align 16 .b8 _Z23UnpremultiplyComponents6float4j_param_0[16],
	.param .b32 _Z23UnpremultiplyComponents6float4j_param_1
)
{
	.reg .pred 	%p<15>;
	.reg .s32 	%r<8>;
	.reg .f32 	%f<77>;


	ld.param.f32 	%f70, [_Z23UnpremultiplyComponents6float4j_param_0+12];
	ld.param.f32 	%f69, [_Z23UnpremultiplyComponents6float4j_param_0+8];
	ld.param.f32 	%f68, [_Z23UnpremultiplyComponents6float4j_param_0+4];
	ld.param.f32 	%f72, [_Z23UnpremultiplyComponents6float4j_param_0];
	ld.param.u32 	%r5, [_Z23UnpremultiplyComponents6float4j_param_1];
	.loc 2 51 1
	and.b32  	%r1, %r5, 768;
	setp.ne.s32	%p1, %r1, 512;
	@%p1 bra 	BB3_7;

	.loc 2 51 1
	and.b32  	%r2, %r5, 248;
	setp.ne.s32	%p2, %r2, 0;
	@%p2 bra 	BB3_3;

	mov.f32 	%f67, 0f437F0000;
	bra.uni 	BB3_6;

BB3_3:
	.loc 2 51 1
	setp.ne.s32	%p3, %r2, 8;
	@%p3 bra 	BB3_5;

	mov.f32 	%f67, 0f447FC000;
	bra.uni 	BB3_6;

BB3_5:
	.loc 2 51 1
	setp.eq.s32	%p4, %r2, 16;
	.loc 2 51 1
	selp.f32	%f67, 0f47000000, 0f3F800000, %p4;

BB3_6:
	.loc 2 51 119
	and.b32  	%r6, %r5, 4096;
	setp.eq.s32	%p5, %r6, 0;
	ld.const.f32 	%f32, [kYCbCrOffset];
	ld.const.f32 	%f33, [kYCbCrFullRangeOffset];
	selp.f32	%f34, %f33, %f32, %p5;
	mov.f32 	%f35, 0f437F0000;
	.loc 3 3606 10
	div.approx.ftz.f32 	%f36, %f67, %f35;
	.loc 2 51 119
	mul.ftz.f32 	%f37, %f34, %f36;
	sub.ftz.f32 	%f68, %f68, %f37;
	ld.const.f32 	%f38, [kYCbCrOffset+4];
	ld.const.f32 	%f39, [kYCbCrFullRangeOffset+4];
	selp.f32	%f40, %f39, %f38, %p5;
	mul.ftz.f32 	%f41, %f40, %f36;
	sub.ftz.f32 	%f69, %f69, %f41;
	ld.const.f32 	%f42, [kYCbCrOffset+8];
	ld.const.f32 	%f43, [kYCbCrFullRangeOffset+8];
	selp.f32	%f44, %f43, %f42, %p5;
	mul.ftz.f32 	%f45, %f44, %f36;
	sub.ftz.f32 	%f70, %f70, %f45;

BB3_7:
	.loc 2 51 1
	add.ftz.f32 	%f46, %f72, 0fB70637BD;
	setp.gtu.ftz.f32	%p6, %f46, 0f00000000;
	@%p6 bra 	BB3_9;

	mov.f32 	%f76, 0f00000000;
	mov.f32 	%f75, %f76;
	mov.f32 	%f74, %f76;
	mov.f32 	%f72, %f76;
	bra.uni 	BB3_15;

BB3_9:
	.loc 2 51 1
	and.b32  	%r3, %r5, 248;
	setp.ne.s32	%p7, %r3, 0;
	@%p7 bra 	BB3_11;

	mov.f32 	%f71, 0f437F0000;
	bra.uni 	BB3_14;

BB3_11:
	.loc 2 51 1
	setp.ne.s32	%p8, %r3, 8;
	@%p8 bra 	BB3_13;

	mov.f32 	%f71, 0f447FC000;
	bra.uni 	BB3_14;

BB3_13:
	.loc 2 51 1
	setp.eq.s32	%p9, %r3, 16;
	.loc 2 51 1
	selp.f32	%f71, 0f47000000, 0f3F800000, %p9;

BB3_14:
	.loc 3 3606 10
	div.approx.ftz.f32 	%f53, %f71, %f72;
	.loc 2 51 1
	mul.ftz.f32 	%f74, %f68, %f53;
	mul.ftz.f32 	%f75, %f69, %f53;
	mul.ftz.f32 	%f76, %f70, %f53;

BB3_15:
	.loc 2 51 1
	@%p1 bra 	BB3_22;

	.loc 2 51 1
	and.b32  	%r4, %r5, 248;
	setp.eq.s32	%p11, %r4, 8;
	@%p11 bra 	BB3_19;

	setp.ne.s32	%p12, %r4, 0;
	@%p12 bra 	BB3_20;

	mov.f32 	%f73, 0f437F0000;
	bra.uni 	BB3_21;

BB3_19:
	mov.f32 	%f73, 0f447FC000;
	bra.uni 	BB3_21;

BB3_20:
	.loc 2 51 1
	setp.eq.s32	%p13, %r4, 16;
	.loc 2 51 1
	selp.f32	%f73, 0f47000000, 0f3F800000, %p13;

BB3_21:
	.loc 2 51 120
	and.b32  	%r7, %r5, 4096;
	setp.eq.s32	%p14, %r7, 0;
	ld.const.f32 	%f56, [kYCbCrOffset];
	ld.const.f32 	%f57, [kYCbCrFullRangeOffset];
	selp.f32	%f58, %f57, %f56, %p14;
	mov.f32 	%f59, 0f437F0000;
	.loc 3 3606 10
	div.approx.ftz.f32 	%f60, %f73, %f59;
	.loc 2 51 120
	fma.rn.ftz.f32 	%f74, %f58, %f60, %f74;
	ld.const.f32 	%f61, [kYCbCrOffset+4];
	ld.const.f32 	%f62, [kYCbCrFullRangeOffset+4];
	selp.f32	%f63, %f62, %f61, %p14;
	fma.rn.ftz.f32 	%f75, %f63, %f60, %f75;
	ld.const.f32 	%f64, [kYCbCrOffset+8];
	ld.const.f32 	%f65, [kYCbCrFullRangeOffset+8];
	selp.f32	%f66, %f65, %f64, %p14;
	fma.rn.ftz.f32 	%f76, %f66, %f60, %f76;

BB3_22:
	st.param.f32	[func_retval0+0], %f72;
	st.param.f32	[func_retval0+4], %f74;
	st.param.f32	[func_retval0+8], %f75;
	st.param.f32	[func_retval0+12], %f76;
	.loc 2 51 1
	ret;
}

.visible .func _Z7DoPixelbiiPvi17DevicePixelFormatPKh6float2(
	.param .b32 _Z7DoPixelbiiPvi17DevicePixelFormatPKh6float2_param_0,
	.param .b32 _Z7DoPixelbiiPvi17DevicePixelFormatPKh6float2_param_1,
	.param .b32 _Z7DoPixelbiiPvi17DevicePixelFormatPKh6float2_param_2,
	.param .b64 _Z7DoPixelbiiPvi17DevicePixelFormatPKh6float2_param_3,
	.param .b32 _Z7DoPixelbiiPvi17DevicePixelFormatPKh6float2_param_4,
	.param .b32 _Z7DoPixelbiiPvi17DevicePixelFormatPKh6float2_param_5,
	.param .b64 _Z7DoPixelbiiPvi17DevicePixelFormatPKh6float2_param_6,
	.param .align 8 .b8 _Z7DoPixelbiiPvi17DevicePixelFormatPKh6float2_param_7[8]
)
{
	.reg .pred 	%p<3>;
	.reg .s16 	%rs<8>;
	.reg .s32 	%r<20>;
	.reg .f32 	%f<38>;
	.reg .s64 	%rd<9>;


	ld.param.u32 	%r2, [_Z7DoPixelbiiPvi17DevicePixelFormatPKh6float2_param_1];
	ld.param.u32 	%r3, [_Z7DoPixelbiiPvi17DevicePixelFormatPKh6float2_param_2];
	ld.param.u64 	%rd1, [_Z7DoPixelbiiPvi17DevicePixelFormatPKh6float2_param_3];
	ld.param.u32 	%r4, [_Z7DoPixelbiiPvi17DevicePixelFormatPKh6float2_param_4];
	ld.param.u32 	%r5, [_Z7DoPixelbiiPvi17DevicePixelFormatPKh6float2_param_5];
	ld.param.u64 	%rd2, [_Z7DoPixelbiiPvi17DevicePixelFormatPKh6float2_param_6];
	ld.param.f32 	%f2, [_Z7DoPixelbiiPvi17DevicePixelFormatPKh6float2_param_7+4];
	ld.param.f32 	%f1, [_Z7DoPixelbiiPvi17DevicePixelFormatPKh6float2_param_7];
	ld.param.s8 	%rs1, [_Z7DoPixelbiiPvi17DevicePixelFormatPKh6float2_param_0];
	and.b16  	%rs2, %rs1, 255;
	.loc 1 83 1
	setp.eq.s16	%p1, %rs2, 0;
	@%p1 bra 	BB4_4;

	cvt.s64.s32	%rd3, %r3;
	add.s64 	%rd4, %rd2, %rd3;
	.loc 1 85 1
	ld.u8 	%rs3, [%rd4];
	cvt.rn.f32.u16	%f10, %rs3;
	ld.const.f32 	%f11, [kYCbCrOffset];
	mov.f32 	%f12, 0f437F0000;
	.loc 3 3606 10
	div.approx.ftz.f32 	%f13, %f12, %f12;
	mul.ftz.f32 	%f14, %f11, %f13;
	sub.ftz.f32 	%f15, %f10, %f14;
	ld.const.f32 	%f16, [kYCbCrOffset+4];
	mul.ftz.f32 	%f17, %f16, %f13;
	sub.ftz.f32 	%f18, %f1, %f17;
	ld.const.f32 	%f19, [kYCbCrOffset+8];
	mul.ftz.f32 	%f20, %f19, %f13;
	sub.ftz.f32 	%f21, %f2, %f20;
	ld.const.f32 	%f22, [k601YCbCr_To_RGB32f];
	ld.const.f32 	%f23, [k601YCbCr_To_RGB32f+4];
	mul.ftz.f32 	%f24, %f18, %f23;
	fma.rn.ftz.f32 	%f25, %f15, %f22, %f24;
	ld.const.f32 	%f26, [k601YCbCr_To_RGB32f+8];
	fma.rn.ftz.f32 	%f3, %f21, %f26, %f25;
	ld.const.f32 	%f27, [k601YCbCr_To_RGB32f+12];
	ld.const.f32 	%f28, [k601YCbCr_To_RGB32f+16];
	mul.ftz.f32 	%f29, %f18, %f28;
	fma.rn.ftz.f32 	%f30, %f15, %f27, %f29;
	ld.const.f32 	%f31, [k601YCbCr_To_RGB32f+20];
	fma.rn.ftz.f32 	%f4, %f21, %f31, %f30;
	ld.const.f32 	%f32, [k601YCbCr_To_RGB32f+24];
	ld.const.f32 	%f33, [k601YCbCr_To_RGB32f+28];
	mul.ftz.f32 	%f34, %f18, %f33;
	fma.rn.ftz.f32 	%f35, %f15, %f32, %f34;
	ld.const.f32 	%f36, [k601YCbCr_To_RGB32f+32];
	fma.rn.ftz.f32 	%f5, %f21, %f36, %f35;
	mov.f32 	%f9, 0f3F800000;
	.loc 1 87 1
	mov.u32 	%r6, %ntid.x;
	mov.u32 	%r7, %ctaid.x;
	mov.u32 	%r8, %tid.x;
	mad.lo.s32 	%r9, %r6, %r7, %r8;
	shl.b32 	%r10, %r9, 1;
	mov.u32 	%r11, %ntid.y;
	mov.u32 	%r12, %ctaid.y;
	mov.u32 	%r13, %tid.y;
	mad.lo.s32 	%r14, %r11, %r12, %r13;
	shl.b32 	%r15, %r14, 1;
	add.s32 	%r16, %r2, %r15;
	add.s32 	%r17, %r16, -1;
	add.s32 	%r18, %r3, %r10;
	add.s32 	%r19, %r18, -1;
	.loc 1 87 1
	mad.lo.s32 	%r1, %r17, %r4, %r19;
	.loc 1 87 1
	setp.eq.s32	%p2, %r5, 0;
	@%p2 bra 	BB4_3;

	mul.wide.s32 	%rd5, %r1, 16;
	add.s64 	%rd6, %rd1, %rd5;
	.loc 1 87 1
	st.v4.f32 	[%rd6], {%f5, %f4, %f3, %f9};
	bra.uni 	BB4_4;

BB4_3:
	mul.wide.s32 	%rd7, %r1, 8;
	add.s64 	%rd8, %rd1, %rd7;
	.loc 3 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f5;
	mov.b16 	%rs4, %temp;
}
	.loc 3 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f4;
	mov.b16 	%rs5, %temp;
}
	.loc 3 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f3;
	mov.b16 	%rs6, %temp;
}
	.loc 3 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f9;
	mov.b16 	%rs7, %temp;
}
	.loc 1 87 250
	st.v4.u16 	[%rd8], {%rs4, %rs5, %rs6, %rs7};

BB4_4:
	.loc 1 89 2
	ret;
}

.visible .entry PixelFormatConvert_NV12Centered_To_BGRA_4444_Kernel(
	.param .u64 PixelFormatConvert_NV12Centered_To_BGRA_4444_Kernel_param_0,
	.param .u32 PixelFormatConvert_NV12Centered_To_BGRA_4444_Kernel_param_1,
	.param .u64 PixelFormatConvert_NV12Centered_To_BGRA_4444_Kernel_param_2,
	.param .u32 PixelFormatConvert_NV12Centered_To_BGRA_4444_Kernel_param_3,
	.param .u32 PixelFormatConvert_NV12Centered_To_BGRA_4444_Kernel_param_4,
	.param .u32 PixelFormatConvert_NV12Centered_To_BGRA_4444_Kernel_param_5,
	.param .u32 PixelFormatConvert_NV12Centered_To_BGRA_4444_Kernel_param_6
)
{
	.reg .pred 	%p<22>;
	.reg .s16 	%rs<61>;
	.reg .s32 	%r<103>;
	.reg .f32 	%f<189>;
	.reg .s64 	%rd<67>;


	ld.param.u64 	%rd7, [PixelFormatConvert_NV12Centered_To_BGRA_4444_Kernel_param_0];
	ld.param.u32 	%r7, [PixelFormatConvert_NV12Centered_To_BGRA_4444_Kernel_param_1];
	ld.param.u64 	%rd8, [PixelFormatConvert_NV12Centered_To_BGRA_4444_Kernel_param_2];
	ld.param.u32 	%r8, [PixelFormatConvert_NV12Centered_To_BGRA_4444_Kernel_param_3];
	ld.param.u32 	%r9, [PixelFormatConvert_NV12Centered_To_BGRA_4444_Kernel_param_4];
	ld.param.u32 	%r10, [PixelFormatConvert_NV12Centered_To_BGRA_4444_Kernel_param_5];
	ld.param.u32 	%r11, [PixelFormatConvert_NV12Centered_To_BGRA_4444_Kernel_param_6];
	cvta.to.global.u64 	%rd1, %rd7;
	.loc 1 100 1
	mov.u32 	%r12, %ntid.x;
	mov.u32 	%r13, %ctaid.x;
	mul.lo.s32 	%r1, %r12, %r13;
	mov.u32 	%r2, %tid.x;
	add.s32 	%r14, %r1, %r2;
	shl.b32 	%r3, %r14, 1;
	setp.gt.u32	%p1, %r3, %r10;
	mov.u32 	%r15, %ntid.y;
	mov.u32 	%r16, %ctaid.y;
	mul.lo.s32 	%r4, %r15, %r16;
	mov.u32 	%r5, %tid.y;
	add.s32 	%r17, %r4, %r5;
	shl.b32 	%r6, %r17, 1;
	setp.gt.u32	%p2, %r6, %r11;
	or.pred  	%p3, %p1, %p2;
	@%p3 bra 	BB5_17;

	.loc 1 105 1
	neg.s32 	%r18, %r2;
	setp.ne.s32	%p4, %r1, %r18;
	.loc 1 106 1
	neg.s32 	%r19, %r5;
	setp.ne.s32	%p5, %r4, %r19;
	.loc 1 111 1
	add.s32 	%r20, %r5, %r11;
	add.s32 	%r21, %r20, %r4;
	mad.lo.s32 	%r22, %r21, %r7, %r3;
	.loc 1 112 1
	selp.b64	%rd9, -2, 0, %p4;
	.loc 1 107 1
	setp.lt.u32	%p6, %r3, %r10;
	.loc 1 112 1
	selp.b64	%rd10, 0, -2, %p6;
	.loc 1 115 1
	neg.s32 	%r23, %r7;
	cvt.s64.s32	%rd11, %r23;
	selp.b64	%rd12, %rd11, 0, %p5;
	.loc 1 108 1
	setp.lt.u32	%p7, %r6, %r11;
	.loc 1 115 1
	selp.b64	%rd13, 0, %rd11, %p7;
	.loc 1 119 1
	cvt.s64.s32	%rd14, %r22;
	add.s64 	%rd15, %rd12, %rd9;
	add.s64 	%rd16, %rd15, %rd14;
	add.s64 	%rd17, %rd1, %rd16;
	ld.global.v2.u8 	{%rs1, %rs2}, [%rd17];
	and.b16  	%rs4, %rs1, 255;
	cvt.rn.f32.u16	%f37, %rs4;
	and.b16  	%rs6, %rs2, 255;
	cvt.rn.f32.u16	%f38, %rs6;
	add.s64 	%rd18, %rd10, %rd12;
	add.s64 	%rd19, %rd18, %rd14;
	add.s64 	%rd20, %rd1, %rd19;
	ld.global.v2.u8 	{%rs7, %rs8}, [%rd20];
	and.b16  	%rs10, %rs7, 255;
	cvt.rn.f32.u16	%f39, %rs10;
	and.b16  	%rs12, %rs8, 255;
	cvt.rn.f32.u16	%f40, %rs12;
	add.s64 	%rd21, %rd13, %rd9;
	add.s64 	%rd22, %rd21, %rd14;
	add.s64 	%rd23, %rd1, %rd22;
	ld.global.v2.u8 	{%rs13, %rs14}, [%rd23];
	and.b16  	%rs16, %rs13, 255;
	cvt.rn.f32.u16	%f41, %rs16;
	and.b16  	%rs18, %rs14, 255;
	cvt.rn.f32.u16	%f42, %rs18;
	add.s64 	%rd24, %rd13, %rd10;
	add.s64 	%rd25, %rd24, %rd14;
	add.s64 	%rd26, %rd1, %rd25;
	ld.global.v2.u8 	{%rs19, %rs20}, [%rd26];
	and.b16  	%rs22, %rs19, 255;
	cvt.rn.f32.u16	%f43, %rs22;
	and.b16  	%rs24, %rs20, 255;
	cvt.rn.f32.u16	%f44, %rs24;
	.loc 3 3272 10
	cvt.rzi.ftz.u32.f32	%r24, %f37;
	cvt.u16.u32	%rs25, %r24;
	.loc 3 3272 10
	cvt.rzi.ftz.u32.f32	%r25, %f39;
	cvt.u16.u32	%rs26, %r25;
	.loc 3 3272 10
	cvt.rzi.ftz.u32.f32	%r26, %f41;
	cvt.u16.u32	%rs27, %r26;
	.loc 3 3272 10
	cvt.rzi.ftz.u32.f32	%r27, %f43;
	cvt.u16.u32	%rs28, %r27;
	and.b16  	%rs29, %rs25, 255;
	cvt.rn.f32.u16	%f45, %rs29;
	and.b16  	%rs30, %rs26, 255;
	cvt.rn.f32.u16	%f46, %rs30;
	sub.ftz.f32 	%f47, %f46, %f45;
	fma.rn.ftz.f32 	%f48, %f47, 0f3E158106, %f45;
	and.b16  	%rs31, %rs27, 255;
	cvt.rn.f32.u16	%f49, %rs31;
	and.b16  	%rs32, %rs28, 255;
	cvt.rn.f32.u16	%f50, %rs32;
	sub.ftz.f32 	%f51, %f50, %f49;
	fma.rn.ftz.f32 	%f52, %f51, 0f3E158106, %f49;
	sub.ftz.f32 	%f53, %f52, %f48;
	fma.rn.ftz.f32 	%f1, %f53, 0f3E158106, %f48;
	.loc 3 3272 10
	cvt.rzi.ftz.u32.f32	%r28, %f38;
	cvt.u16.u32	%rs33, %r28;
	.loc 3 3272 10
	cvt.rzi.ftz.u32.f32	%r29, %f40;
	cvt.u16.u32	%rs34, %r29;
	.loc 3 3272 10
	cvt.rzi.ftz.u32.f32	%r30, %f42;
	cvt.u16.u32	%rs35, %r30;
	.loc 3 3272 10
	cvt.rzi.ftz.u32.f32	%r31, %f44;
	cvt.u16.u32	%rs36, %r31;
	and.b16  	%rs37, %rs33, 255;
	cvt.rn.f32.u16	%f54, %rs37;
	and.b16  	%rs38, %rs34, 255;
	cvt.rn.f32.u16	%f55, %rs38;
	sub.ftz.f32 	%f56, %f55, %f54;
	fma.rn.ftz.f32 	%f57, %f56, 0f3E158106, %f54;
	and.b16  	%rs39, %rs35, 255;
	cvt.rn.f32.u16	%f58, %rs39;
	and.b16  	%rs40, %rs36, 255;
	cvt.rn.f32.u16	%f59, %rs40;
	sub.ftz.f32 	%f60, %f59, %f58;
	fma.rn.ftz.f32 	%f61, %f60, 0f3E158106, %f58;
	sub.ftz.f32 	%f62, %f61, %f57;
	fma.rn.ftz.f32 	%f2, %f62, 0f3E158106, %f57;
	sub.ftz.f32 	%f63, %f45, %f46;
	fma.rn.ftz.f32 	%f64, %f63, 0f3E158106, %f46;
	sub.ftz.f32 	%f65, %f49, %f50;
	fma.rn.ftz.f32 	%f66, %f65, 0f3E158106, %f50;
	sub.ftz.f32 	%f67, %f66, %f64;
	fma.rn.ftz.f32 	%f3, %f67, 0f3E158106, %f64;
	sub.ftz.f32 	%f68, %f54, %f55;
	fma.rn.ftz.f32 	%f69, %f68, 0f3E158106, %f55;
	sub.ftz.f32 	%f70, %f58, %f59;
	fma.rn.ftz.f32 	%f71, %f70, 0f3E158106, %f59;
	sub.ftz.f32 	%f72, %f71, %f69;
	fma.rn.ftz.f32 	%f4, %f72, 0f3E158106, %f69;
	sub.ftz.f32 	%f73, %f48, %f52;
	fma.rn.ftz.f32 	%f5, %f73, 0f3E158106, %f52;
	sub.ftz.f32 	%f74, %f57, %f61;
	fma.rn.ftz.f32 	%f6, %f74, 0f3E158106, %f61;
	sub.ftz.f32 	%f75, %f64, %f66;
	fma.rn.ftz.f32 	%f7, %f75, 0f3E158106, %f66;
	sub.ftz.f32 	%f76, %f69, %f71;
	fma.rn.ftz.f32 	%f8, %f76, 0f3E158106, %f71;
	.loc 1 133 1
	and.pred  	%p8, %p5, %p4;
	.loc 1 83 1
	@!%p8 bra 	BB5_5;
	bra.uni 	BB5_2;

BB5_2:
	.loc 1 131 1
	add.s32 	%r32, %r6, -1;
	mul.lo.s32 	%r33, %r32, %r7;
	cvt.s64.s32	%rd28, %r33;
	cvt.s64.s32	%rd29, %r3;
	add.s64 	%rd30, %rd29, %rd28;
	.loc 1 85 1
	add.s64 	%rd31, %rd1, %rd30;
	.loc 1 85 1
	ld.global.u8 	%rs41, [%rd31+-1];
	cvt.rn.f32.u16	%f77, %rs41;
	ld.const.f32 	%f78, [kYCbCrOffset];
	mov.f32 	%f79, 0f437F0000;
	.loc 3 3606 10
	div.approx.ftz.f32 	%f80, %f79, %f79;
	mul.ftz.f32 	%f81, %f78, %f80;
	sub.ftz.f32 	%f82, %f77, %f81;
	ld.const.f32 	%f83, [kYCbCrOffset+4];
	mul.ftz.f32 	%f84, %f83, %f80;
	sub.ftz.f32 	%f85, %f1, %f84;
	ld.const.f32 	%f86, [kYCbCrOffset+8];
	mul.ftz.f32 	%f87, %f86, %f80;
	sub.ftz.f32 	%f88, %f2, %f87;
	ld.const.f32 	%f89, [k601YCbCr_To_RGB32f];
	ld.const.f32 	%f90, [k601YCbCr_To_RGB32f+4];
	mul.ftz.f32 	%f91, %f85, %f90;
	fma.rn.ftz.f32 	%f92, %f82, %f89, %f91;
	ld.const.f32 	%f93, [k601YCbCr_To_RGB32f+8];
	fma.rn.ftz.f32 	%f9, %f88, %f93, %f92;
	ld.const.f32 	%f94, [k601YCbCr_To_RGB32f+12];
	ld.const.f32 	%f95, [k601YCbCr_To_RGB32f+16];
	mul.ftz.f32 	%f96, %f85, %f95;
	fma.rn.ftz.f32 	%f97, %f82, %f94, %f96;
	ld.const.f32 	%f98, [k601YCbCr_To_RGB32f+20];
	fma.rn.ftz.f32 	%f10, %f88, %f98, %f97;
	ld.const.f32 	%f99, [k601YCbCr_To_RGB32f+24];
	ld.const.f32 	%f100, [k601YCbCr_To_RGB32f+28];
	mul.ftz.f32 	%f101, %f85, %f100;
	fma.rn.ftz.f32 	%f102, %f82, %f99, %f101;
	ld.const.f32 	%f103, [k601YCbCr_To_RGB32f+32];
	fma.rn.ftz.f32 	%f11, %f88, %f103, %f102;
	mov.f32 	%f15, 0f3F800000;
	.loc 1 87 1
	mad.lo.s32 	%r34, %r32, %r8, %r3;
	.loc 1 87 1
	add.s32 	%r35, %r34, -1;
	.loc 1 87 1
	cvt.s64.s32	%rd2, %r35;
	.loc 1 87 1
	setp.eq.s32	%p9, %r9, 0;
	@%p9 bra 	BB5_4;

	cvta.to.global.u64 	%rd32, %rd8;
	.loc 1 87 1
	shl.b64 	%rd33, %rd2, 4;
	add.s64 	%rd34, %rd32, %rd33;
	st.global.v4.f32 	[%rd34], {%f11, %f10, %f9, %f15};
	bra.uni 	BB5_5;

BB5_4:
	cvta.to.global.u64 	%rd35, %rd8;
	.loc 1 100 1
	mad.lo.s32 	%r39, %r12, %r13, %r2;
	shl.b32 	%r40, %r39, 1;
	mad.lo.s32 	%r44, %r15, %r16, %r5;
	shl.b32 	%r45, %r44, 1;
	.loc 1 131 1
	add.s32 	%r46, %r45, -1;
	.loc 1 87 1
	mad.lo.s32 	%r47, %r46, %r8, %r40;
	.loc 1 87 1
	add.s32 	%r48, %r47, -1;
	.loc 1 87 1
	mul.wide.s32 	%rd36, %r48, 8;
	add.s64 	%rd37, %rd35, %rd36;
	.loc 3 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f11;
	mov.b16 	%rs42, %temp;
}
	.loc 3 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f10;
	mov.b16 	%rs43, %temp;
}
	.loc 3 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f9;
	mov.b16 	%rs44, %temp;
}
	.loc 3 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f15;
	mov.b16 	%rs45, %temp;
}
	.loc 1 87 250
	st.global.v4.u16 	[%rd37], {%rs42, %rs43, %rs44, %rs45};

BB5_5:
	.loc 1 134 1
	and.pred  	%p12, %p5, %p6;
	.loc 1 83 1
	@!%p12 bra 	BB5_9;
	bra.uni 	BB5_6;

BB5_6:
	.loc 1 100 1
	mad.lo.s32 	%r57, %r12, %r13, %r2;
	shl.b32 	%r58, %r57, 1;
	.loc 1 131 1
	cvt.s64.s32	%rd39, %r58;
	.loc 1 100 1
	mad.lo.s32 	%r62, %r15, %r16, %r5;
	shl.b32 	%r63, %r62, 1;
	.loc 1 131 1
	add.s32 	%r64, %r63, -1;
	mul.lo.s32 	%r65, %r64, %r7;
	cvt.s64.s32	%rd40, %r65;
	add.s64 	%rd41, %rd39, %rd40;
	.loc 1 85 1
	add.s64 	%rd42, %rd1, %rd41;
	ld.global.u8 	%rs46, [%rd42];
	cvt.rn.f32.u16	%f105, %rs46;
	ld.const.f32 	%f106, [kYCbCrOffset];
	mov.f32 	%f107, 0f437F0000;
	.loc 3 3606 10
	div.approx.ftz.f32 	%f108, %f107, %f107;
	mul.ftz.f32 	%f109, %f106, %f108;
	sub.ftz.f32 	%f110, %f105, %f109;
	ld.const.f32 	%f111, [kYCbCrOffset+4];
	mul.ftz.f32 	%f112, %f111, %f108;
	sub.ftz.f32 	%f113, %f3, %f112;
	ld.const.f32 	%f114, [kYCbCrOffset+8];
	mul.ftz.f32 	%f115, %f114, %f108;
	sub.ftz.f32 	%f116, %f4, %f115;
	ld.const.f32 	%f117, [k601YCbCr_To_RGB32f];
	ld.const.f32 	%f118, [k601YCbCr_To_RGB32f+4];
	mul.ftz.f32 	%f119, %f113, %f118;
	fma.rn.ftz.f32 	%f120, %f110, %f117, %f119;
	ld.const.f32 	%f121, [k601YCbCr_To_RGB32f+8];
	fma.rn.ftz.f32 	%f16, %f116, %f121, %f120;
	ld.const.f32 	%f122, [k601YCbCr_To_RGB32f+12];
	ld.const.f32 	%f123, [k601YCbCr_To_RGB32f+16];
	mul.ftz.f32 	%f124, %f113, %f123;
	fma.rn.ftz.f32 	%f125, %f110, %f122, %f124;
	ld.const.f32 	%f126, [k601YCbCr_To_RGB32f+20];
	fma.rn.ftz.f32 	%f17, %f116, %f126, %f125;
	ld.const.f32 	%f127, [k601YCbCr_To_RGB32f+24];
	ld.const.f32 	%f128, [k601YCbCr_To_RGB32f+28];
	mul.ftz.f32 	%f129, %f113, %f128;
	fma.rn.ftz.f32 	%f130, %f110, %f127, %f129;
	ld.const.f32 	%f131, [k601YCbCr_To_RGB32f+32];
	fma.rn.ftz.f32 	%f18, %f116, %f131, %f130;
	mov.f32 	%f22, 0f3F800000;
	.loc 1 87 1
	mad.lo.s32 	%r66, %r64, %r8, %r58;
	.loc 1 87 1
	cvt.s64.s32	%rd3, %r66;
	.loc 1 87 1
	setp.eq.s32	%p13, %r9, 0;
	@%p13 bra 	BB5_8;

	cvta.to.global.u64 	%rd43, %rd8;
	.loc 1 87 1
	shl.b64 	%rd44, %rd3, 4;
	add.s64 	%rd45, %rd43, %rd44;
	st.global.v4.f32 	[%rd45], {%f18, %f17, %f16, %f22};
	bra.uni 	BB5_9;

BB5_8:
	cvta.to.global.u64 	%rd46, %rd8;
	.loc 1 87 1
	shl.b64 	%rd47, %rd3, 3;
	add.s64 	%rd48, %rd46, %rd47;
	.loc 3 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f18;
	mov.b16 	%rs47, %temp;
}
	.loc 3 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f17;
	mov.b16 	%rs48, %temp;
}
	.loc 3 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f16;
	mov.b16 	%rs49, %temp;
}
	.loc 3 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f22;
	mov.b16 	%rs50, %temp;
}
	.loc 1 87 250
	st.global.v4.u16 	[%rd48], {%rs47, %rs48, %rs49, %rs50};

BB5_9:
	.loc 1 100 1
	mad.lo.s32 	%r70, %r15, %r16, %r5;
	shl.b32 	%r71, %r70, 1;
	setp.lt.u32	%p14, %r71, %r11;
	.loc 1 135 1
	and.pred  	%p16, %p14, %p4;
	.loc 1 131 1
	cvt.s64.s32	%rd49, %r3;
	add.s32 	%r79, %r71, -1;
	mul.lo.s32 	%r80, %r79, %r7;
	cvt.s64.s32	%rd50, %r80;
	add.s64 	%rd51, %rd49, %rd50;
	.loc 1 132 1
	cvt.u64.u32	%rd52, %r7;
	.loc 1 85 1
	add.s64 	%rd53, %rd51, %rd52;
	add.s64 	%rd4, %rd1, %rd53;
	.loc 1 83 1
	@!%p16 bra 	BB5_13;
	bra.uni 	BB5_10;

BB5_10:
	.loc 1 85 1
	ld.global.u8 	%rs51, [%rd4+-1];
	cvt.rn.f32.u16	%f133, %rs51;
	ld.const.f32 	%f134, [kYCbCrOffset];
	mov.f32 	%f135, 0f437F0000;
	.loc 3 3606 10
	div.approx.ftz.f32 	%f136, %f135, %f135;
	mul.ftz.f32 	%f137, %f134, %f136;
	sub.ftz.f32 	%f138, %f133, %f137;
	ld.const.f32 	%f139, [kYCbCrOffset+4];
	mul.ftz.f32 	%f140, %f139, %f136;
	sub.ftz.f32 	%f141, %f5, %f140;
	ld.const.f32 	%f142, [kYCbCrOffset+8];
	mul.ftz.f32 	%f143, %f142, %f136;
	sub.ftz.f32 	%f144, %f6, %f143;
	ld.const.f32 	%f145, [k601YCbCr_To_RGB32f];
	ld.const.f32 	%f146, [k601YCbCr_To_RGB32f+4];
	mul.ftz.f32 	%f147, %f141, %f146;
	fma.rn.ftz.f32 	%f148, %f138, %f145, %f147;
	ld.const.f32 	%f149, [k601YCbCr_To_RGB32f+8];
	fma.rn.ftz.f32 	%f23, %f144, %f149, %f148;
	ld.const.f32 	%f150, [k601YCbCr_To_RGB32f+12];
	ld.const.f32 	%f151, [k601YCbCr_To_RGB32f+16];
	mul.ftz.f32 	%f152, %f141, %f151;
	fma.rn.ftz.f32 	%f153, %f138, %f150, %f152;
	ld.const.f32 	%f154, [k601YCbCr_To_RGB32f+20];
	fma.rn.ftz.f32 	%f24, %f144, %f154, %f153;
	ld.const.f32 	%f155, [k601YCbCr_To_RGB32f+24];
	ld.const.f32 	%f156, [k601YCbCr_To_RGB32f+28];
	mul.ftz.f32 	%f157, %f141, %f156;
	fma.rn.ftz.f32 	%f158, %f138, %f155, %f157;
	ld.const.f32 	%f159, [k601YCbCr_To_RGB32f+32];
	fma.rn.ftz.f32 	%f25, %f144, %f159, %f158;
	mov.f32 	%f29, 0f3F800000;
	.loc 1 100 1
	mad.lo.s32 	%r88, %r12, %r13, %r2;
	mad.lo.s32 	%r89, %r70, %r8, %r88;
	shl.b32 	%r90, %r89, 1;
	.loc 1 87 1
	add.s32 	%r91, %r90, -1;
	.loc 1 87 1
	cvt.s64.s32	%rd5, %r91;
	.loc 1 87 1
	setp.eq.s32	%p17, %r9, 0;
	@%p17 bra 	BB5_12;

	cvta.to.global.u64 	%rd55, %rd8;
	.loc 1 87 1
	shl.b64 	%rd56, %rd5, 4;
	add.s64 	%rd57, %rd55, %rd56;
	st.global.v4.f32 	[%rd57], {%f25, %f24, %f23, %f29};
	bra.uni 	BB5_13;

BB5_12:
	cvta.to.global.u64 	%rd58, %rd8;
	.loc 1 87 1
	shl.b64 	%rd59, %rd5, 3;
	add.s64 	%rd60, %rd58, %rd59;
	.loc 3 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f25;
	mov.b16 	%rs52, %temp;
}
	.loc 3 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f24;
	mov.b16 	%rs53, %temp;
}
	.loc 3 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f23;
	mov.b16 	%rs54, %temp;
}
	.loc 3 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f29;
	mov.b16 	%rs55, %temp;
}
	.loc 1 87 250
	st.global.v4.u16 	[%rd60], {%rs52, %rs53, %rs54, %rs55};

BB5_13:
	.loc 1 136 1
	and.pred  	%p20, %p14, %p6;
	.loc 1 83 1
	@!%p20 bra 	BB5_17;
	bra.uni 	BB5_14;

BB5_14:
	.loc 1 85 1
	ld.global.u8 	%rs56, [%rd4];
	cvt.rn.f32.u16	%f161, %rs56;
	ld.const.f32 	%f162, [kYCbCrOffset];
	mov.f32 	%f163, 0f437F0000;
	.loc 3 3606 10
	div.approx.ftz.f32 	%f164, %f163, %f163;
	mul.ftz.f32 	%f165, %f162, %f164;
	sub.ftz.f32 	%f166, %f161, %f165;
	ld.const.f32 	%f167, [kYCbCrOffset+4];
	mul.ftz.f32 	%f168, %f167, %f164;
	sub.ftz.f32 	%f169, %f7, %f168;
	ld.const.f32 	%f170, [kYCbCrOffset+8];
	mul.ftz.f32 	%f171, %f170, %f164;
	sub.ftz.f32 	%f172, %f8, %f171;
	ld.const.f32 	%f173, [k601YCbCr_To_RGB32f];
	ld.const.f32 	%f174, [k601YCbCr_To_RGB32f+4];
	mul.ftz.f32 	%f175, %f169, %f174;
	fma.rn.ftz.f32 	%f176, %f166, %f173, %f175;
	ld.const.f32 	%f177, [k601YCbCr_To_RGB32f+8];
	fma.rn.ftz.f32 	%f30, %f172, %f177, %f176;
	ld.const.f32 	%f178, [k601YCbCr_To_RGB32f+12];
	ld.const.f32 	%f179, [k601YCbCr_To_RGB32f+16];
	mul.ftz.f32 	%f180, %f169, %f179;
	fma.rn.ftz.f32 	%f181, %f166, %f178, %f180;
	ld.const.f32 	%f182, [k601YCbCr_To_RGB32f+20];
	fma.rn.ftz.f32 	%f31, %f172, %f182, %f181;
	ld.const.f32 	%f183, [k601YCbCr_To_RGB32f+24];
	ld.const.f32 	%f184, [k601YCbCr_To_RGB32f+28];
	mul.ftz.f32 	%f185, %f169, %f184;
	fma.rn.ftz.f32 	%f186, %f166, %f183, %f185;
	ld.const.f32 	%f187, [k601YCbCr_To_RGB32f+32];
	fma.rn.ftz.f32 	%f32, %f172, %f187, %f186;
	mov.f32 	%f36, 0f3F800000;
	.loc 1 87 1
	shl.b32 	%r101, %r8, 1;
	.loc 1 87 1
	mad.lo.s32 	%r102, %r101, %r70, %r3;
	.loc 1 87 1
	cvt.s64.s32	%rd6, %r102;
	.loc 1 87 1
	setp.eq.s32	%p21, %r9, 0;
	@%p21 bra 	BB5_16;

	cvta.to.global.u64 	%rd61, %rd8;
	.loc 1 87 1
	shl.b64 	%rd62, %rd6, 4;
	add.s64 	%rd63, %rd61, %rd62;
	st.global.v4.f32 	[%rd63], {%f32, %f31, %f30, %f36};
	bra.uni 	BB5_17;

BB5_16:
	cvta.to.global.u64 	%rd64, %rd8;
	.loc 1 87 1
	shl.b64 	%rd65, %rd6, 3;
	add.s64 	%rd66, %rd64, %rd65;
	.loc 3 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f32;
	mov.b16 	%rs57, %temp;
}
	.loc 3 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f31;
	mov.b16 	%rs58, %temp;
}
	.loc 3 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f30;
	mov.b16 	%rs59, %temp;
}
	.loc 3 3513 10
	{
	.reg .b16 %temp;
	cvt.rn.ftz.f16.f32 	%temp, %f36;
	mov.b16 	%rs60, %temp;
}
	.loc 1 87 250
	st.global.v4.u16 	[%rd66], {%rs57, %rs58, %rs59, %rs60};

BB5_17:
	.loc 1 137 2
	ret;
}

.visible .entry CopyNV12_Kernel(
	.param .u64 CopyNV12_Kernel_param_0,
	.param .u32 CopyNV12_Kernel_param_1,
	.param .u64 CopyNV12_Kernel_param_2,
	.param .u32 CopyNV12_Kernel_param_3,
	.param .u32 CopyNV12_Kernel_param_4
)
{
	.reg .pred 	%p<4>;
	.reg .s32 	%r<18>;
	.reg .s64 	%rd<9>;


	ld.param.u64 	%rd3, [CopyNV12_Kernel_param_0];
	ld.param.u32 	%r4, [CopyNV12_Kernel_param_1];
	ld.param.u64 	%rd4, [CopyNV12_Kernel_param_2];
	ld.param.u32 	%r5, [CopyNV12_Kernel_param_3];
	ld.param.u32 	%r6, [CopyNV12_Kernel_param_4];
	cvta.to.global.u64 	%rd1, %rd4;
	cvta.to.global.u64 	%rd2, %rd3;
	.loc 1 146 1
	mov.u32 	%r7, %ntid.x;
	mov.u32 	%r8, %ctaid.x;
	mul.lo.s32 	%r1, %r7, %r8;
	mov.u32 	%r2, %tid.x;
	add.s32 	%r9, %r1, %r2;
	setp.lt.s32	%p1, %r9, %r5;
	mov.u32 	%r10, %ntid.y;
	mov.u32 	%r11, %ctaid.y;
	mov.u32 	%r12, %tid.y;
	mad.lo.s32 	%r3, %r10, %r11, %r12;
	setp.lt.s32	%p2, %r3, %r6;
	and.pred  	%p3, %p1, %p2;
	@!%p3 bra 	BB6_2;
	bra.uni 	BB6_1;

BB6_1:
	.loc 1 148 1
	mad.lo.s32 	%r13, %r3, %r4, %r2;
	add.s32 	%r14, %r13, %r1;
	mul.wide.s32 	%rd5, %r14, 4;
	add.s64 	%rd6, %rd2, %rd5;
	.loc 1 148 1
	mad.lo.s32 	%r15, %r3, %r5, %r2;
	add.s32 	%r16, %r15, %r1;
	mul.wide.s32 	%rd7, %r16, 4;
	add.s64 	%rd8, %rd1, %rd7;
	.loc 1 148 1
	ld.global.u32 	%r17, [%rd6];
	st.global.u32 	[%rd8], %r17;

BB6_2:
	.loc 1 151 2
	ret;
}


