	.version 1.4
	.target sm_11, map_f64_to_f32
	// compiled with ../../../External/3rdParty/NVIDIA/CUDA/win/64/bin64/../open64/lib//be.exe
	// nvopencc 2.3 built on 2009-07-14

	//-----------------------------------------------------------
	// Compiling C:\Users\dvaeng\AppData\Local\Temp/tmpxft_0001da9c_00000000-9_Deinterlace.cpp3.i (C:/Users/dvaeng/AppData/Local/Temp/ccBI#.a21356)
	//-----------------------------------------------------------

	//-----------------------------------------------------------
	// Options:
	//-----------------------------------------------------------
	//  Target:ptx, ISA:sm_11, Endian:little, Pointer Size:64
	//  -O3	(Optimization level)
	//  -g0	(Debug level)
	//  -m2	(Report advisories)
	//-----------------------------------------------------------

	.file	1	"C:\Users\dvaeng\AppData\Local\Temp/tmpxft_0001da9c_00000000-8_Deinterlace.cudafe2.gpu"
	.file	2	"c:\scully64\shared\adobe\MediaCore\ImageRenderer\API\Inc\ImageRenderer.h"
	.file	3	"c:\scully64\shared\adobe\MediaCore\GPUFoundation\API\Inc\GPUFoundation/KernelSupport/PixelRGB.h"
	.file	4	"C:\Program Files (x86)\Microsoft Visual Studio 9.0\VC\include\crtdefs.h"
	.file	5	"c:\scully64\shared\adobe\MediaCore\External\3rdParty\NVIDIA\CUDA\win\64\include\crt/device_runtime.h"
	.file	6	"c:\scully64\shared\adobe\MediaCore\External\3rdParty\NVIDIA\CUDA\win\64\include\host_defines.h"
	.file	7	"c:\scully64\shared\adobe\MediaCore\External\3rdParty\NVIDIA\CUDA\win\64\include\builtin_types.h"
	.file	8	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\device_types.h"
	.file	9	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\driver_types.h"
	.file	10	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\texture_types.h"
	.file	11	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\vector_types.h"
	.file	12	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\host_defines.h"
	.file	13	"c:\scully64\shared\adobe\MediaCore\External\3rdParty\NVIDIA\CUDA\win\64\include\device_launch_parameters.h"
	.file	14	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\crt\storage_class.h"
	.file	15	"C:\Program Files (x86)\Microsoft Visual Studio 9.0\VC\include\time.h"
	.file	16	"c:\scully64\shared\adobe\MediaCore\GPUFoundation\API\Inc\GPUFoundation/KernelSupport/Utils.h"
	.file	17	"c:/scully64/shared/adobe/MediaCore/GPUFoundation/Src/ImageProcessing/Deinterlace.cu"
	.file	18	"c:\scully64\shared\adobe\MediaCore\GPUFoundation\API\Inc\GPUFoundation/KernelSupport/VectorUtils.h"
	.file	19	"c:\scully64\shared\adobe\MediaCore\External\3rdParty\NVIDIA\CUDA\win\64\include\common_functions.h"
	.file	20	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\crt/func_macro.h"
	.file	21	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\math_functions.h"
	.file	22	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\device_functions.h"
	.file	23	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\math_constants.h"
	.file	24	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\sm_11_atomic_functions.h"
	.file	25	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\sm_12_atomic_functions.h"
	.file	26	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\sm_13_double_functions.h"
	.file	27	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\common_types.h"
	.file	28	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\texture_fetch_functions.h"
	.file	29	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\math_functions_dbl_ptx1.h"


	.entry DeinterlaceKernel (
		.param .u64 __cudaparm_DeinterlaceKernel_inImage,
		.param .s32 __cudaparm_DeinterlaceKernel_inPitch,
		.param .s32 __cudaparm_DeinterlaceKernel_inWidth,
		.param .s32 __cudaparm_DeinterlaceKernel_inHeight,
		.param .u32 __cudaparm_DeinterlaceKernel_inWhichField)
	{
	.reg .u32 %r<129>;
	.reg .u64 %rd<94>;
	.reg .f32 %f<237>;
	.reg .pred %p<98>;
	.shared .align 4 .b8 __cuda_sGPlane24[3072];
	.shared .align 4 .b8 __cuda_sAPlane3096[3072];
	.shared .align 4 .b8 __cuda_sBPlane6168[3072];
	.shared .align 4 .b8 __cuda_sRPlane9240[3072];
	.loc	17	166	0
$LBB1_DeinterlaceKernel:
	.loc	17	176	0
	cvt.s32.u16 	%r1, %ctaid.x;
	cvt.s32.u16 	%r2, %ntid.x;
	mul24.lo.s32 	%r3, %r1, %r2;
	shl.b32 	%r4, %r1, 1;
	ld.param.s32 	%r5, [__cudaparm_DeinterlaceKernel_inWhichField];
	mov.s32 	%r6, 0;
	setp.eq.s32 	%p1, %r5, %r6;
	cvt.s32.u16 	%r7, %ctaid.y;
	cvt.s32.u16 	%r8, %ntid.y;
	mul24.lo.s32 	%r9, %r7, %r8;
	cvt.u32.u16 	%r10, %tid.x;
	add.u32 	%r11, %r3, %r10;
	cvt.u32.u16 	%r12, %tid.y;
	add.u32 	%r13, %r9, %r12;
	sub.u32 	%r14, %r11, %r4;
	mul.lo.s32 	%r15, %r13, 4;
	ld.param.s32 	%r16, [__cudaparm_DeinterlaceKernel_inWidth];
	setp.gt.s32 	%p2, %r16, %r14;
	add.s32 	%r17, %r15, 1;
	selp.s32 	%r18, 1, 0, %p2;
	selp.s32 	%r19, %r17, %r15, %p1;
	mov.s32 	%r20, 0;
	setp.gt.s32 	%p3, %r19, %r20;
	selp.s32 	%r21, 1, 0, %p3;
	and.b32 	%r22, %r18, %r21;
	mov.u32 	%r23, 0;
	setp.eq.s32 	%p4, %r22, %r23;
	@%p4 bra 	$Lt_0_84226;
	.loc	17	207	0
	mov.u64 	%rd1, __cuda_sGPlane24;
	mov.u64 	%rd2, __cuda_sAPlane3096;
	mov.u64 	%rd3, __cuda_sBPlane6168;
	mov.u64 	%rd4, __cuda_sRPlane9240;
	sub.s32 	%r24, %r19, 1;
	ld.param.s32 	%r25, [__cudaparm_DeinterlaceKernel_inPitch];
	mul24.lo.s32 	%r26, %r24, %r25;
	ld.param.u64 	%rd5, [__cudaparm_DeinterlaceKernel_inImage];
	add.s32 	%r27, %r26, %r14;
	cvt.s64.s32 	%rd6, %r27;
	mul.lo.u64 	%rd7, %rd6, 16;
	add.u64 	%rd8, %rd5, %rd7;
	ld.global.v4.f32 	{%f1,%f2,%f3,%f4}, [%rd8+0];
	.loc	17	208	0
	cvt.u64.u32 	%rd9, %r10;
	mul.lo.u64 	%rd10, %rd9, 4;
	mul.f32 	%f5, %f1, %f4;
	add.u64 	%rd11, %rd10, %rd3;
	st.shared.f32 	[%rd11+0], %f5;
	.loc	17	209	0
	mul.f32 	%f6, %f2, %f4;
	add.u64 	%rd12, %rd10, %rd1;
	st.shared.f32 	[%rd12+0], %f6;
	.loc	17	210	0
	mul.f32 	%f7, %f3, %f4;
	add.u64 	%rd13, %rd10, %rd4;
	st.shared.f32 	[%rd13+0], %f7;
	.loc	17	211	0
	add.u64 	%rd14, %rd10, %rd2;
	st.shared.f32 	[%rd14+0], %f4;
$Lt_0_84226:
	mov.u64 	%rd1, __cuda_sGPlane24;
	mov.u64 	%rd2, __cuda_sAPlane3096;
	mov.u64 	%rd3, __cuda_sBPlane6168;
	mov.u64 	%rd4, __cuda_sRPlane9240;
	add.s32 	%r28, %r19, 1;
	ld.param.s32 	%r29, [__cudaparm_DeinterlaceKernel_inHeight];
	setp.lt.s32 	%p5, %r28, %r29;
	selp.s32 	%r30, 1, 0, %p5;
	and.b32 	%r31, %r18, %r30;
	mov.u32 	%r32, 0;
	setp.eq.s32 	%p6, %r31, %r32;
	@%p6 bra 	$Lt_0_84738;
	.loc	16	89	0
	ld.param.s32 	%r33, [__cudaparm_DeinterlaceKernel_inPitch];
	mul24.lo.s32 	%r34, %r28, %r33;
	ld.param.u64 	%rd15, [__cudaparm_DeinterlaceKernel_inImage];
	add.s32 	%r35, %r34, %r14;
	cvt.s64.s32 	%rd16, %r35;
	mul.lo.u64 	%rd17, %rd16, 16;
	add.u64 	%rd18, %rd15, %rd17;
	ld.global.v4.f32 	{%f8,%f9,%f10,%f11}, [%rd18+0];
	.loc	17	216	0
	cvt.u64.u32 	%rd19, %r10;
	mul.lo.u64 	%rd10, %rd19, 4;
	mul.f32 	%f12, %f8, %f11;
	add.u64 	%rd20, %rd10, %rd3;
	st.shared.f32 	[%rd20+1024], %f12;
	.loc	17	217	0
	mul.f32 	%f13, %f9, %f11;
	add.u64 	%rd21, %rd10, %rd1;
	st.shared.f32 	[%rd21+1024], %f13;
	.loc	17	218	0
	mul.f32 	%f14, %f10, %f11;
	add.u64 	%rd22, %rd10, %rd4;
	st.shared.f32 	[%rd22+1024], %f14;
	.loc	17	219	0
	add.u64 	%rd23, %rd10, %rd2;
	st.shared.f32 	[%rd23+1024], %f11;
$Lt_0_84738:
	add.s32 	%r36, %r19, 3;
	setp.lt.s32 	%p7, %r36, %r29;
	selp.s32 	%r37, 1, 0, %p7;
	and.b32 	%r38, %r18, %r37;
	mov.u32 	%r39, 0;
	setp.eq.s32 	%p8, %r38, %r39;
	@%p8 bra 	$Lt_0_85250;
	.loc	16	89	0
	ld.param.s32 	%r40, [__cudaparm_DeinterlaceKernel_inPitch];
	mul24.lo.s32 	%r41, %r36, %r40;
	ld.param.u64 	%rd24, [__cudaparm_DeinterlaceKernel_inImage];
	add.s32 	%r42, %r41, %r14;
	cvt.s64.s32 	%rd25, %r42;
	mul.lo.u64 	%rd26, %rd25, 16;
	add.u64 	%rd27, %rd24, %rd26;
	ld.global.v4.f32 	{%f15,%f16,%f17,%f18}, [%rd27+0];
	.loc	17	224	0
	cvt.u64.u32 	%rd28, %r10;
	mul.lo.u64 	%rd10, %rd28, 4;
	mul.f32 	%f19, %f15, %f18;
	add.u64 	%rd29, %rd10, %rd3;
	st.shared.f32 	[%rd29+2048], %f19;
	.loc	17	225	0
	mul.f32 	%f20, %f16, %f18;
	add.u64 	%rd30, %rd10, %rd1;
	st.shared.f32 	[%rd30+2048], %f20;
	.loc	17	226	0
	mul.f32 	%f21, %f17, %f18;
	add.u64 	%rd31, %rd10, %rd4;
	st.shared.f32 	[%rd31+2048], %f21;
	.loc	17	227	0
	add.u64 	%rd32, %rd10, %rd2;
	st.shared.f32 	[%rd32+2048], %f18;
$Lt_0_85250:
	.loc	17	231	0
	bar.sync 	0;
	mov.s32 	%r43, 0;
	set.eq.u32.s32 	%r44, %r14, %r43;
	neg.s32 	%r45, %r44;
	sub.u32 	%r46, %r10, 1;
	mov.u32 	%r47, 253;
	set.le.u32.u32 	%r48, %r46, %r47;
	neg.s32 	%r49, %r48;
	or.b32 	%r50, %r45, %r49;
	and.b32 	%r51, %r18, %r50;
	mov.u32 	%r52, 0;
	setp.eq.s32 	%p9, %r51, %r52;
	@%p9 bra 	$Lt_0_119042;
	.loc	17	237	0
	@!%p3 bra 	$Lt_0_120578;
	@!%p5 bra 	$Lt_0_120578;
	mov.u32 	%r53, 0;
	setp.le.s32 	%p10, %r14, %r53;
	@%p10 bra 	$Lt_0_120578;
	add.s32 	%r54, %r14, 1;
	setp.le.s32 	%p11, %r16, %r54;
	@%p11 bra 	$Lt_0_120578;
	.loc	17	78	0
	cvt.u64.u32 	%rd33, %r10;
	mul.lo.u64 	%rd10, %rd33, 4;
	add.u64 	%rd34, %rd10, %rd3;
	ld.shared.f32 	%f22, [%rd34+-4];
	add.u64 	%rd35, %rd10, %rd1;
	ld.shared.f32 	%f23, [%rd35+-4];
	add.u64 	%rd36, %rd10, %rd4;
	ld.shared.f32 	%f24, [%rd36+-4];
	add.u64 	%rd37, %rd10, %rd2;
	ld.shared.f32 	%f25, [%rd37+-4];
	.loc	17	79	0
	ld.shared.f32 	%f26, [%rd34+0];
	ld.shared.f32 	%f27, [%rd35+0];
	ld.shared.f32 	%f28, [%rd36+0];
	ld.shared.f32 	%f29, [%rd37+0];
	.loc	17	80	0
	ld.shared.f32 	%f30, [%rd34+4];
	ld.shared.f32 	%f31, [%rd35+4];
	ld.shared.f32 	%f32, [%rd36+4];
	ld.shared.f32 	%f33, [%rd37+4];
	.loc	17	82	0
	ld.shared.f32 	%f34, [%rd34+1020];
	ld.shared.f32 	%f35, [%rd35+1020];
	ld.shared.f32 	%f36, [%rd36+1020];
	ld.shared.f32 	%f37, [%rd37+1020];
	.loc	17	83	0
	ld.shared.f32 	%f38, [%rd34+1024];
	ld.shared.f32 	%f39, [%rd35+1024];
	ld.shared.f32 	%f40, [%rd36+1024];
	ld.shared.f32 	%f41, [%rd37+1024];
	.loc	17	84	0
	ld.shared.f32 	%f42, [%rd34+1028];
	ld.shared.f32 	%f43, [%rd35+1028];
	ld.shared.f32 	%f44, [%rd36+1028];
	ld.shared.f32 	%f45, [%rd37+1028];
	sub.f32 	%f46, %f24, %f44;
	sub.f32 	%f47, %f23, %f43;
	sub.f32 	%f48, %f22, %f42;
	sub.f32 	%f49, %f28, %f40;
	sub.f32 	%f50, %f27, %f39;
	sub.f32 	%f51, %f26, %f38;
	sub.f32 	%f52, %f32, %f36;
	sub.f32 	%f53, %f31, %f35;
	sub.f32 	%f54, %f30, %f34;
	mul.f32 	%f55, %f46, %f46;
	mul.f32 	%f56, %f49, %f49;
	mul.f32 	%f57, %f52, %f52;
	mad.f32 	%f58, %f47, %f47, %f55;
	mad.f32 	%f59, %f50, %f50, %f56;
	mad.f32 	%f60, %f53, %f53, %f57;
	mad.f32 	%f61, %f48, %f48, %f58;
	mad.f32 	%f62, %f51, %f51, %f59;
	mad.f32 	%f63, %f54, %f54, %f60;
	set.lt.u32.f32 	%r55, %f61, %f62;
	neg.s32 	%r56, %r55;
	set.lt.u32.f32 	%r57, %f61, %f63;
	neg.s32 	%r58, %r57;
	and.b32 	%r59, %r56, %r58;
	mov.u32 	%r60, 0;
	setp.eq.s32 	%p12, %r59, %r60;
	@%p12 bra 	$Lt_0_86530;
	.loc	17	95	0
	mov.f32 	%f64, %f22;
	mov.f32 	%f65, %f23;
	mov.f32 	%f66, %f24;
	mov.f32 	%f67, %f25;
	.loc	17	96	0
	mov.f32 	%f68, %f42;
	mov.f32 	%f69, %f43;
	mov.f32 	%f70, %f44;
	mov.f32 	%f71, %f45;
	bra.uni 	$Lt_0_86786;
$Lt_0_86530:
	setp.lt.f32 	%p13, %f62, %f63;
	@!%p13 bra 	$Lt_0_87042;
	.loc	17	100	0
	mov.f32 	%f64, %f26;
	mov.f32 	%f65, %f27;
	mov.f32 	%f66, %f28;
	mov.f32 	%f67, %f29;
	.loc	17	101	0
	mov.f32 	%f68, %f38;
	mov.f32 	%f69, %f39;
	mov.f32 	%f70, %f40;
	mov.f32 	%f71, %f41;
	bra.uni 	$Lt_0_86786;
$Lt_0_87042:
	.loc	17	105	0
	mov.f32 	%f64, %f30;
	mov.f32 	%f65, %f31;
	mov.f32 	%f66, %f32;
	mov.f32 	%f67, %f33;
	.loc	17	106	0
	mov.f32 	%f68, %f34;
	mov.f32 	%f69, %f35;
	mov.f32 	%f70, %f36;
	mov.f32 	%f71, %f37;
$Lt_0_86786:
$Lt_0_86274:
	.loc	17	111	0
	add.f32 	%f72, %f64, %f68;
	mov.f32 	%f73, 0f3f000000;    	// 0.5
	mul.f32 	%f74, %f72, %f73;
	setp.lt.f32 	%p14, %f38, %f26;
	@!%p14 bra 	$Lt_0_87554;
	setp.lt.f32 	%p15, %f74, %f38;
	@!%p15 bra 	$Lt_0_88066;
	.loc	18	315	0
	mov.f32 	%f75, %f38;
	bra.uni 	$Lt_0_87810;
$Lt_0_88066:
	setp.lt.f32 	%p16, %f74, %f26;
	selp.f32 	%f75, %f74, %f26, %p16;
$Lt_0_87810:
	mov.f32 	%f76, %f75;
	bra.uni 	$Lt_0_87298;
$Lt_0_87554:
	setp.gt.f32 	%p17, %f74, %f38;
	@!%p17 bra 	$Lt_0_88578;
	mov.f32 	%f77, %f38;
	bra.uni 	$Lt_0_88322;
$Lt_0_88578:
	max.f32 	%f77, %f74, %f26;
$Lt_0_88322:
	mov.f32 	%f76, %f77;
$Lt_0_87298:
	add.f32 	%f78, %f65, %f69;
	mov.f32 	%f79, 0f3f000000;    	// 0.5
	mul.f32 	%f80, %f78, %f79;
	setp.lt.f32 	%p18, %f39, %f27;
	@!%p18 bra 	$Lt_0_89090;
	setp.lt.f32 	%p19, %f80, %f39;
	@!%p19 bra 	$Lt_0_89602;
	.loc	18	316	0
	mov.f32 	%f81, %f39;
	bra.uni 	$Lt_0_89346;
$Lt_0_89602:
	setp.lt.f32 	%p20, %f80, %f27;
	selp.f32 	%f81, %f80, %f27, %p20;
$Lt_0_89346:
	mov.f32 	%f82, %f81;
	bra.uni 	$Lt_0_88834;
$Lt_0_89090:
	setp.gt.f32 	%p21, %f80, %f39;
	@!%p21 bra 	$Lt_0_90114;
	mov.f32 	%f83, %f39;
	bra.uni 	$Lt_0_89858;
$Lt_0_90114:
	max.f32 	%f83, %f80, %f27;
$Lt_0_89858:
	mov.f32 	%f82, %f83;
$Lt_0_88834:
	add.f32 	%f84, %f66, %f70;
	mov.f32 	%f85, 0f3f000000;    	// 0.5
	mul.f32 	%f86, %f84, %f85;
	setp.lt.f32 	%p22, %f40, %f28;
	@!%p22 bra 	$Lt_0_90626;
	setp.lt.f32 	%p23, %f86, %f40;
	@!%p23 bra 	$Lt_0_91138;
	.loc	18	317	0
	mov.f32 	%f87, %f40;
	bra.uni 	$Lt_0_90882;
$Lt_0_91138:
	setp.lt.f32 	%p24, %f86, %f28;
	selp.f32 	%f87, %f86, %f28, %p24;
$Lt_0_90882:
	mov.f32 	%f88, %f87;
	bra.uni 	$Lt_0_90370;
$Lt_0_90626:
	setp.gt.f32 	%p25, %f86, %f40;
	@!%p25 bra 	$Lt_0_91650;
	mov.f32 	%f89, %f40;
	bra.uni 	$Lt_0_91394;
$Lt_0_91650:
	max.f32 	%f89, %f86, %f28;
$Lt_0_91394:
	mov.f32 	%f88, %f89;
$Lt_0_90370:
	add.f32 	%f90, %f67, %f71;
	mov.f32 	%f91, 0f3f000000;    	// 0.5
	mul.f32 	%f92, %f90, %f91;
	setp.lt.f32 	%p26, %f41, %f29;
	@!%p26 bra 	$Lt_0_92162;
	setp.lt.f32 	%p27, %f92, %f41;
	@!%p27 bra 	$Lt_0_92674;
	.loc	18	318	0
	mov.f32 	%f93, %f41;
	bra.uni 	$Lt_0_92418;
$Lt_0_92674:
	setp.lt.f32 	%p28, %f92, %f29;
	selp.f32 	%f93, %f92, %f29, %p28;
$Lt_0_92418:
	mov.f32 	%f94, %f93;
	bra.uni 	$Lt_0_91906;
$Lt_0_92162:
	setp.gt.f32 	%p29, %f92, %f41;
	@!%p29 bra 	$Lt_0_93186;
	mov.f32 	%f95, %f41;
	bra.uni 	$Lt_0_92930;
$Lt_0_93186:
	max.f32 	%f95, %f92, %f29;
$Lt_0_92930:
	mov.f32 	%f94, %f95;
$Lt_0_91906:
	.loc	17	112	0
	mov.f32 	%f96, %f76;
	mov.f32 	%f97, %f82;
	mov.f32 	%f98, %f88;
	mov.f32 	%f99, 0f00000000;    	// 0
	setp.neu.f32 	%p30, %f94, %f99;
	@!%p30 bra 	$Lt_0_93442;
	div.full.f32 	%f98, %f98, %f94;
	div.full.f32 	%f97, %f97, %f94;
	div.full.f32 	%f96, %f96, %f94;
$Lt_0_93442:
	.loc	16	100	0
	ld.param.s32 	%r61, [__cudaparm_DeinterlaceKernel_inPitch];
	mul24.lo.s32 	%r62, %r19, %r61;
	ld.param.u64 	%rd38, [__cudaparm_DeinterlaceKernel_inImage];
	add.s32 	%r63, %r62, %r14;
	cvt.s64.s32 	%rd39, %r63;
	mul.lo.u64 	%rd40, %rd39, 16;
	add.u64 	%rd41, %rd38, %rd40;
	st.global.v4.f32 	[%rd41+0], {%f96,%f97,%f98,%f94};
	.loc	17	112	0
	bra.uni 	$Lt_0_102146;
$Lt_0_120578:
$L_0_81922:
	.loc	17	119	0
	and.b32 	%r64, %r21, %r30;
	mov.u32 	%r65, 0;
	setp.eq.s32 	%p31, %r64, %r65;
	@%p31 bra 	$Lt_0_94210;
	.loc	17	122	0
	cvt.u64.u32 	%rd42, %r10;
	mul.lo.u64 	%rd10, %rd42, 4;
	add.u64 	%rd34, %rd10, %rd3;
	ld.shared.f32 	%f100, [%rd34+0];
	add.u64 	%rd35, %rd10, %rd1;
	ld.shared.f32 	%f101, [%rd35+0];
	add.u64 	%rd36, %rd10, %rd4;
	ld.shared.f32 	%f102, [%rd36+0];
	add.u64 	%rd37, %rd10, %rd2;
	ld.shared.f32 	%f103, [%rd37+0];
	.loc	17	123	0
	ld.shared.f32 	%f104, [%rd34+1024];
	ld.shared.f32 	%f105, [%rd35+1024];
	ld.shared.f32 	%f106, [%rd36+1024];
	ld.shared.f32 	%f107, [%rd37+1024];
	.loc	17	126	0
	add.f32 	%f108, %f104, %f100;
	mov.f32 	%f109, 0f3f000000;   	// 0.5
	mul.f32 	%f110, %f108, %f109;
	setp.lt.f32 	%p32, %f104, %f100;
	@!%p32 bra 	$Lt_0_94722;
	setp.lt.f32 	%p33, %f110, %f104;
	@!%p33 bra 	$Lt_0_95234;
	.loc	18	315	0
	mov.f32 	%f111, %f104;
	bra.uni 	$Lt_0_94978;
$Lt_0_95234:
	setp.lt.f32 	%p34, %f110, %f100;
	selp.f32 	%f111, %f110, %f100, %p34;
$Lt_0_94978:
	mov.f32 	%f112, %f111;
	bra.uni 	$Lt_0_94466;
$Lt_0_94722:
	setp.gt.f32 	%p35, %f110, %f104;
	@!%p35 bra 	$Lt_0_95746;
	mov.f32 	%f113, %f104;
	bra.uni 	$Lt_0_95490;
$Lt_0_95746:
	max.f32 	%f113, %f110, %f100;
$Lt_0_95490:
	mov.f32 	%f112, %f113;
$Lt_0_94466:
	add.f32 	%f114, %f105, %f101;
	mov.f32 	%f115, 0f3f000000;   	// 0.5
	mul.f32 	%f116, %f114, %f115;
	setp.lt.f32 	%p36, %f105, %f101;
	@!%p36 bra 	$Lt_0_96258;
	setp.lt.f32 	%p37, %f116, %f105;
	@!%p37 bra 	$Lt_0_96770;
	.loc	18	316	0
	mov.f32 	%f117, %f105;
	bra.uni 	$Lt_0_96514;
$Lt_0_96770:
	setp.lt.f32 	%p38, %f116, %f101;
	selp.f32 	%f117, %f116, %f101, %p38;
$Lt_0_96514:
	mov.f32 	%f118, %f117;
	bra.uni 	$Lt_0_96002;
$Lt_0_96258:
	setp.gt.f32 	%p39, %f116, %f105;
	@!%p39 bra 	$Lt_0_97282;
	mov.f32 	%f119, %f105;
	bra.uni 	$Lt_0_97026;
$Lt_0_97282:
	max.f32 	%f119, %f116, %f101;
$Lt_0_97026:
	mov.f32 	%f118, %f119;
$Lt_0_96002:
	add.f32 	%f120, %f106, %f102;
	mov.f32 	%f121, 0f3f000000;   	// 0.5
	mul.f32 	%f122, %f120, %f121;
	setp.lt.f32 	%p40, %f106, %f102;
	@!%p40 bra 	$Lt_0_97794;
	setp.lt.f32 	%p41, %f122, %f106;
	@!%p41 bra 	$Lt_0_98306;
	.loc	18	317	0
	mov.f32 	%f123, %f106;
	bra.uni 	$Lt_0_98050;
$Lt_0_98306:
	setp.lt.f32 	%p42, %f122, %f102;
	selp.f32 	%f123, %f122, %f102, %p42;
$Lt_0_98050:
	mov.f32 	%f124, %f123;
	bra.uni 	$Lt_0_97538;
$Lt_0_97794:
	setp.gt.f32 	%p43, %f122, %f106;
	@!%p43 bra 	$Lt_0_98818;
	mov.f32 	%f125, %f106;
	bra.uni 	$Lt_0_98562;
$Lt_0_98818:
	max.f32 	%f125, %f122, %f102;
$Lt_0_98562:
	mov.f32 	%f124, %f125;
$Lt_0_97538:
	add.f32 	%f126, %f107, %f103;
	mov.f32 	%f127, 0f3f000000;   	// 0.5
	mul.f32 	%f128, %f126, %f127;
	setp.lt.f32 	%p44, %f107, %f103;
	@!%p44 bra 	$Lt_0_99330;
	setp.lt.f32 	%p45, %f128, %f107;
	@!%p45 bra 	$Lt_0_99842;
	.loc	18	318	0
	mov.f32 	%f129, %f107;
	bra.uni 	$Lt_0_99586;
$Lt_0_99842:
	setp.lt.f32 	%p46, %f128, %f103;
	selp.f32 	%f129, %f128, %f103, %p46;
$Lt_0_99586:
	mov.f32 	%f130, %f129;
	bra.uni 	$Lt_0_99074;
$Lt_0_99330:
	setp.gt.f32 	%p47, %f128, %f107;
	@!%p47 bra 	$Lt_0_100354;
	mov.f32 	%f131, %f107;
	bra.uni 	$Lt_0_100098;
$Lt_0_100354:
	max.f32 	%f131, %f128, %f103;
$Lt_0_100098:
	mov.f32 	%f130, %f131;
$Lt_0_99074:
	.loc	17	127	0
	mov.f32 	%f132, %f112;
	mov.f32 	%f133, %f118;
	mov.f32 	%f134, %f124;
	mov.f32 	%f135, 0f00000000;   	// 0
	setp.neu.f32 	%p48, %f130, %f135;
	@!%p48 bra 	$Lt_0_100610;
	div.full.f32 	%f134, %f134, %f130;
	div.full.f32 	%f133, %f133, %f130;
	div.full.f32 	%f132, %f132, %f130;
$Lt_0_100610:
	.loc	16	100	0
	ld.param.s32 	%r66, [__cudaparm_DeinterlaceKernel_inPitch];
	mul24.lo.s32 	%r67, %r19, %r66;
	ld.param.u64 	%rd43, [__cudaparm_DeinterlaceKernel_inImage];
	add.s32 	%r68, %r67, %r14;
	cvt.s64.s32 	%rd44, %r68;
	mul.lo.u64 	%rd45, %rd44, 16;
	add.u64 	%rd46, %rd43, %rd45;
	st.global.v4.f32 	[%rd46+0], {%f132,%f133,%f134,%f130};
	.loc	17	127	0
	bra.uni 	$Lt_0_102146;
$Lt_0_94210:
	mov.s32 	%r69, 0;
	set.eq.u32.s32 	%r70, %r19, %r69;
	neg.s32 	%r71, %r70;
	mov.s32 	%r72, 1;
	set.eq.u32.s32 	%r73, %r5, %r72;
	neg.s32 	%r74, %r73;
	and.b32 	%r75, %r71, %r74;
	mov.u32 	%r76, 0;
	setp.eq.s32 	%p49, %r75, %r76;
	@%p49 bra 	$Lt_0_101378;
	.loc	17	137	0
	cvt.u64.u32 	%rd47, %r10;
	mul.lo.u64 	%rd10, %rd47, 4;
	add.u64 	%rd48, %rd10, %rd3;
	ld.shared.f32 	%f136, [%rd48+1024];
	add.u64 	%rd49, %rd10, %rd1;
	ld.shared.f32 	%f137, [%rd49+1024];
	add.u64 	%rd50, %rd10, %rd4;
	ld.shared.f32 	%f138, [%rd50+1024];
	add.u64 	%rd51, %rd10, %rd2;
	ld.shared.f32 	%f139, [%rd51+1024];
	.loc	17	138	0
	mov.f32 	%f140, %f136;
	mov.f32 	%f141, %f137;
	mov.f32 	%f142, %f138;
	mov.f32 	%f143, 0f00000000;   	// 0
	setp.neu.f32 	%p50, %f139, %f143;
	@!%p50 bra 	$Lt_0_101634;
	.loc	17	137	0
	ld.shared.f32 	%f138, [%rd50+1024];
	ld.shared.f32 	%f139, [%rd51+1024];
	.loc	17	138	0
	div.full.f32 	%f142, %f138, %f139;
	.loc	17	137	0
	ld.shared.f32 	%f137, [%rd49+1024];
	.loc	17	138	0
	div.full.f32 	%f141, %f137, %f139;
	.loc	17	137	0
	ld.shared.f32 	%f136, [%rd48+1024];
	.loc	17	138	0
	div.full.f32 	%f140, %f136, %f139;
$Lt_0_101634:
	.loc	16	100	0
	ld.param.s32 	%r77, [__cudaparm_DeinterlaceKernel_inPitch];
	mul24.lo.s32 	%r78, %r19, %r77;
	ld.param.u64 	%rd52, [__cudaparm_DeinterlaceKernel_inImage];
	add.s32 	%r79, %r78, %r14;
	cvt.s64.s32 	%rd53, %r79;
	mul.lo.u64 	%rd54, %rd53, 16;
	add.u64 	%rd55, %rd52, %rd54;
	.loc	17	137	0
	ld.shared.f32 	%f139, [%rd51+1024];
	st.global.v4.f32 	[%rd55+0], {%f140,%f141,%f142,%f139};
	.loc	17	138	0
	bra.uni 	$Lt_0_102146;
$Lt_0_101378:
	selp.s32 	%r80, 1, 0, %p1;
	sub.s32 	%r81, %r29, 1;
	set.eq.u32.s32 	%r82, %r19, %r81;
	neg.s32 	%r83, %r82;
	and.b32 	%r84, %r80, %r83;
	mov.u32 	%r85, 0;
	setp.eq.s32 	%p51, %r84, %r85;
	@%p51 bra 	$Lt_0_102146;
	.loc	17	148	0
	cvt.u64.u32 	%rd56, %r10;
	mul.lo.u64 	%rd10, %rd56, 4;
	add.u64 	%rd57, %rd10, %rd3;
	ld.shared.f32 	%f144, [%rd57+0];
	add.u64 	%rd58, %rd10, %rd1;
	ld.shared.f32 	%f145, [%rd58+0];
	add.u64 	%rd59, %rd10, %rd4;
	ld.shared.f32 	%f146, [%rd59+0];
	add.u64 	%rd60, %rd10, %rd2;
	ld.shared.f32 	%f147, [%rd60+0];
	.loc	17	149	0
	mov.f32 	%f148, %f144;
	mov.f32 	%f149, %f145;
	mov.f32 	%f150, %f146;
	mov.f32 	%f151, 0f00000000;   	// 0
	setp.neu.f32 	%p52, %f147, %f151;
	@!%p52 bra 	$Lt_0_102658;
	.loc	17	148	0
	ld.shared.f32 	%f146, [%rd59+0];
	ld.shared.f32 	%f147, [%rd60+0];
	.loc	17	149	0
	div.full.f32 	%f150, %f146, %f147;
	.loc	17	148	0
	ld.shared.f32 	%f145, [%rd58+0];
	.loc	17	149	0
	div.full.f32 	%f149, %f145, %f147;
	.loc	17	148	0
	ld.shared.f32 	%f144, [%rd57+0];
	.loc	17	149	0
	div.full.f32 	%f148, %f144, %f147;
$Lt_0_102658:
	.loc	16	100	0
	ld.param.s32 	%r86, [__cudaparm_DeinterlaceKernel_inPitch];
	mul24.lo.s32 	%r87, %r19, %r86;
	ld.param.u64 	%rd61, [__cudaparm_DeinterlaceKernel_inImage];
	add.s32 	%r88, %r87, %r14;
	cvt.s64.s32 	%rd62, %r88;
	mul.lo.u64 	%rd63, %rd62, 16;
	add.u64 	%rd64, %rd61, %rd63;
	.loc	17	148	0
	ld.shared.f32 	%f147, [%rd60+0];
	st.global.v4.f32 	[%rd64+0], {%f148,%f149,%f150,%f147};
$Lt_0_102146:
$Lt_0_101122:
$Lt_0_93954:
$L_0_81666:
	.loc	17	237	0
	add.s32 	%r89, %r19, 2;
	mov.s32 	%r90, 0;
	setp.gt.s32 	%p53, %r89, %r90;
	@!%p53 bra 	$Lt_0_121602;
	@!%p7 bra 	$Lt_0_121602;
	mov.u32 	%r91, 0;
	setp.le.s32 	%p54, %r14, %r91;
	@%p54 bra 	$Lt_0_121602;
	add.s32 	%r92, %r14, 1;
	setp.le.s32 	%p55, %r16, %r92;
	@%p55 bra 	$Lt_0_121602;
	.loc	17	78	0
	cvt.u64.u32 	%rd65, %r10;
	mul.lo.u64 	%rd10, %rd65, 4;
	add.u64 	%rd34, %rd10, %rd3;
	ld.shared.f32 	%f22, [%rd34+1020];
	add.u64 	%rd35, %rd10, %rd1;
	ld.shared.f32 	%f23, [%rd35+1020];
	add.u64 	%rd36, %rd10, %rd4;
	ld.shared.f32 	%f24, [%rd36+1020];
	add.u64 	%rd37, %rd10, %rd2;
	ld.shared.f32 	%f25, [%rd37+1020];
	.loc	17	79	0
	ld.shared.f32 	%f152, [%rd34+1024];
	ld.shared.f32 	%f153, [%rd35+1024];
	ld.shared.f32 	%f154, [%rd36+1024];
	ld.shared.f32 	%f155, [%rd37+1024];
	.loc	17	80	0
	ld.shared.f32 	%f30, [%rd34+1028];
	ld.shared.f32 	%f31, [%rd35+1028];
	ld.shared.f32 	%f32, [%rd36+1028];
	ld.shared.f32 	%f33, [%rd37+1028];
	.loc	17	82	0
	ld.shared.f32 	%f34, [%rd34+2044];
	ld.shared.f32 	%f35, [%rd35+2044];
	ld.shared.f32 	%f36, [%rd36+2044];
	ld.shared.f32 	%f37, [%rd37+2044];
	.loc	17	83	0
	ld.shared.f32 	%f156, [%rd34+2048];
	ld.shared.f32 	%f157, [%rd35+2048];
	ld.shared.f32 	%f158, [%rd36+2048];
	ld.shared.f32 	%f159, [%rd37+2048];
	.loc	17	84	0
	ld.shared.f32 	%f42, [%rd34+2052];
	ld.shared.f32 	%f43, [%rd35+2052];
	ld.shared.f32 	%f44, [%rd36+2052];
	ld.shared.f32 	%f45, [%rd37+2052];
	sub.f32 	%f160, %f24, %f44;
	sub.f32 	%f161, %f23, %f43;
	sub.f32 	%f162, %f22, %f42;
	sub.f32 	%f163, %f154, %f158;
	sub.f32 	%f164, %f153, %f157;
	sub.f32 	%f165, %f152, %f156;
	sub.f32 	%f166, %f32, %f36;
	sub.f32 	%f167, %f31, %f35;
	sub.f32 	%f168, %f30, %f34;
	mul.f32 	%f169, %f160, %f160;
	mul.f32 	%f170, %f163, %f163;
	mul.f32 	%f171, %f166, %f166;
	mad.f32 	%f172, %f161, %f161, %f169;
	mad.f32 	%f173, %f164, %f164, %f170;
	mad.f32 	%f174, %f167, %f167, %f171;
	mad.f32 	%f175, %f162, %f162, %f172;
	mad.f32 	%f62, %f165, %f165, %f173;
	mad.f32 	%f63, %f168, %f168, %f174;
	set.lt.u32.f32 	%r93, %f175, %f62;
	neg.s32 	%r94, %r93;
	set.lt.u32.f32 	%r95, %f175, %f63;
	neg.s32 	%r96, %r95;
	and.b32 	%r97, %r94, %r96;
	mov.u32 	%r98, 0;
	setp.eq.s32 	%p56, %r97, %r98;
	@%p56 bra 	$Lt_0_103426;
	.loc	17	95	0
	mov.f32 	%f64, %f22;
	mov.f32 	%f65, %f23;
	mov.f32 	%f66, %f24;
	mov.f32 	%f67, %f25;
	.loc	17	96	0
	mov.f32 	%f68, %f42;
	mov.f32 	%f69, %f43;
	mov.f32 	%f70, %f44;
	mov.f32 	%f71, %f45;
	bra.uni 	$Lt_0_103682;
$Lt_0_103426:
	setp.lt.f32 	%p57, %f62, %f63;
	@!%p57 bra 	$Lt_0_103938;
	.loc	17	100	0
	mov.f32 	%f64, %f152;
	mov.f32 	%f65, %f153;
	mov.f32 	%f66, %f154;
	mov.f32 	%f67, %f155;
	.loc	17	101	0
	mov.f32 	%f68, %f156;
	mov.f32 	%f69, %f157;
	mov.f32 	%f70, %f158;
	mov.f32 	%f71, %f159;
	bra.uni 	$Lt_0_103682;
$Lt_0_103938:
	.loc	17	105	0
	mov.f32 	%f64, %f30;
	mov.f32 	%f65, %f31;
	mov.f32 	%f66, %f32;
	mov.f32 	%f67, %f33;
	.loc	17	106	0
	mov.f32 	%f68, %f34;
	mov.f32 	%f69, %f35;
	mov.f32 	%f70, %f36;
	mov.f32 	%f71, %f37;
$Lt_0_103682:
$Lt_0_103170:
	.loc	17	111	0
	add.f32 	%f176, %f64, %f68;
	mov.f32 	%f177, 0f3f000000;   	// 0.5
	mul.f32 	%f74, %f176, %f177;
	setp.lt.f32 	%p58, %f156, %f152;
	@!%p58 bra 	$Lt_0_104450;
	setp.lt.f32 	%p59, %f74, %f156;
	@!%p59 bra 	$Lt_0_104962;
	.loc	18	315	0
	mov.f32 	%f178, %f156;
	bra.uni 	$Lt_0_104706;
$Lt_0_104962:
	setp.lt.f32 	%p60, %f74, %f152;
	selp.f32 	%f178, %f74, %f152, %p60;
$Lt_0_104706:
	mov.f32 	%f179, %f178;
	bra.uni 	$Lt_0_104194;
$Lt_0_104450:
	setp.gt.f32 	%p61, %f74, %f156;
	@!%p61 bra 	$Lt_0_105474;
	mov.f32 	%f180, %f156;
	bra.uni 	$Lt_0_105218;
$Lt_0_105474:
	max.f32 	%f180, %f74, %f152;
$Lt_0_105218:
	mov.f32 	%f179, %f180;
$Lt_0_104194:
	add.f32 	%f181, %f65, %f69;
	mov.f32 	%f182, 0f3f000000;   	// 0.5
	mul.f32 	%f80, %f181, %f182;
	setp.lt.f32 	%p62, %f157, %f153;
	@!%p62 bra 	$Lt_0_105986;
	setp.lt.f32 	%p63, %f80, %f157;
	@!%p63 bra 	$Lt_0_106498;
	.loc	18	316	0
	mov.f32 	%f183, %f157;
	bra.uni 	$Lt_0_106242;
$Lt_0_106498:
	setp.lt.f32 	%p64, %f80, %f153;
	selp.f32 	%f183, %f80, %f153, %p64;
$Lt_0_106242:
	mov.f32 	%f184, %f183;
	bra.uni 	$Lt_0_105730;
$Lt_0_105986:
	setp.gt.f32 	%p65, %f80, %f157;
	@!%p65 bra 	$Lt_0_107010;
	mov.f32 	%f185, %f157;
	bra.uni 	$Lt_0_106754;
$Lt_0_107010:
	max.f32 	%f185, %f80, %f153;
$Lt_0_106754:
	mov.f32 	%f184, %f185;
$Lt_0_105730:
	add.f32 	%f186, %f66, %f70;
	mov.f32 	%f187, 0f3f000000;   	// 0.5
	mul.f32 	%f86, %f186, %f187;
	setp.lt.f32 	%p66, %f158, %f154;
	@!%p66 bra 	$Lt_0_107522;
	setp.lt.f32 	%p67, %f86, %f158;
	@!%p67 bra 	$Lt_0_108034;
	.loc	18	317	0
	mov.f32 	%f188, %f158;
	bra.uni 	$Lt_0_107778;
$Lt_0_108034:
	setp.lt.f32 	%p68, %f86, %f154;
	selp.f32 	%f188, %f86, %f154, %p68;
$Lt_0_107778:
	mov.f32 	%f189, %f188;
	bra.uni 	$Lt_0_107266;
$Lt_0_107522:
	setp.gt.f32 	%p69, %f86, %f158;
	@!%p69 bra 	$Lt_0_108546;
	mov.f32 	%f190, %f158;
	bra.uni 	$Lt_0_108290;
$Lt_0_108546:
	max.f32 	%f190, %f86, %f154;
$Lt_0_108290:
	mov.f32 	%f189, %f190;
$Lt_0_107266:
	add.f32 	%f191, %f67, %f71;
	mov.f32 	%f192, 0f3f000000;   	// 0.5
	mul.f32 	%f92, %f191, %f192;
	setp.lt.f32 	%p70, %f159, %f155;
	@!%p70 bra 	$Lt_0_109058;
	setp.lt.f32 	%p71, %f92, %f159;
	@!%p71 bra 	$Lt_0_109570;
	.loc	18	318	0
	mov.f32 	%f193, %f159;
	bra.uni 	$Lt_0_109314;
$Lt_0_109570:
	setp.lt.f32 	%p72, %f92, %f155;
	selp.f32 	%f193, %f92, %f155, %p72;
$Lt_0_109314:
	mov.f32 	%f194, %f193;
	bra.uni 	$Lt_0_108802;
$Lt_0_109058:
	setp.gt.f32 	%p73, %f92, %f159;
	@!%p73 bra 	$Lt_0_110082;
	mov.f32 	%f195, %f159;
	bra.uni 	$Lt_0_109826;
$Lt_0_110082:
	max.f32 	%f195, %f92, %f155;
$Lt_0_109826:
	mov.f32 	%f194, %f195;
$Lt_0_108802:
	.loc	17	112	0
	mov.f32 	%f96, %f179;
	mov.f32 	%f97, %f184;
	mov.f32 	%f98, %f189;
	mov.f32 	%f196, 0f00000000;   	// 0
	setp.neu.f32 	%p74, %f194, %f196;
	@!%p74 bra 	$Lt_0_110338;
	div.full.f32 	%f98, %f98, %f194;
	div.full.f32 	%f97, %f97, %f194;
	div.full.f32 	%f96, %f96, %f194;
$Lt_0_110338:
	.loc	16	100	0
	ld.param.s32 	%r99, [__cudaparm_DeinterlaceKernel_inPitch];
	mul24.lo.s32 	%r100, %r89, %r99;
	ld.param.u64 	%rd66, [__cudaparm_DeinterlaceKernel_inImage];
	add.s32 	%r101, %r100, %r14;
	cvt.s64.s32 	%rd67, %r101;
	mul.lo.u64 	%rd68, %rd67, 16;
	add.u64 	%rd69, %rd66, %rd68;
	st.global.v4.f32 	[%rd69+0], {%f96,%f97,%f98,%f194};
	.loc	17	112	0
	bra.uni 	$Lt_0_119042;
$Lt_0_121602:
$L_0_83202:
	.loc	17	119	0
	selp.s32 	%r102, 1, 0, %p53;
	and.b32 	%r103, %r102, %r37;
	mov.u32 	%r104, 0;
	setp.eq.s32 	%p75, %r103, %r104;
	@%p75 bra 	$Lt_0_111106;
	.loc	17	122	0
	cvt.u64.u32 	%rd70, %r10;
	mul.lo.u64 	%rd10, %rd70, 4;
	add.u64 	%rd34, %rd10, %rd3;
	ld.shared.f32 	%f197, [%rd34+1024];
	add.u64 	%rd35, %rd10, %rd1;
	ld.shared.f32 	%f198, [%rd35+1024];
	add.u64 	%rd36, %rd10, %rd4;
	ld.shared.f32 	%f199, [%rd36+1024];
	add.u64 	%rd37, %rd10, %rd2;
	ld.shared.f32 	%f200, [%rd37+1024];
	.loc	17	123	0
	ld.shared.f32 	%f201, [%rd34+2048];
	ld.shared.f32 	%f202, [%rd35+2048];
	ld.shared.f32 	%f203, [%rd36+2048];
	ld.shared.f32 	%f204, [%rd37+2048];
	.loc	17	126	0
	add.f32 	%f205, %f201, %f197;
	mov.f32 	%f206, 0f3f000000;   	// 0.5
	mul.f32 	%f110, %f205, %f206;
	setp.lt.f32 	%p76, %f201, %f197;
	@!%p76 bra 	$Lt_0_111618;
	setp.lt.f32 	%p77, %f110, %f201;
	@!%p77 bra 	$Lt_0_112130;
	.loc	18	315	0
	mov.f32 	%f207, %f201;
	bra.uni 	$Lt_0_111874;
$Lt_0_112130:
	setp.lt.f32 	%p78, %f110, %f197;
	selp.f32 	%f207, %f110, %f197, %p78;
$Lt_0_111874:
	mov.f32 	%f208, %f207;
	bra.uni 	$Lt_0_111362;
$Lt_0_111618:
	setp.gt.f32 	%p79, %f110, %f201;
	@!%p79 bra 	$Lt_0_112642;
	mov.f32 	%f209, %f201;
	bra.uni 	$Lt_0_112386;
$Lt_0_112642:
	max.f32 	%f209, %f110, %f197;
$Lt_0_112386:
	mov.f32 	%f208, %f209;
$Lt_0_111362:
	add.f32 	%f210, %f202, %f198;
	mov.f32 	%f211, 0f3f000000;   	// 0.5
	mul.f32 	%f116, %f210, %f211;
	setp.lt.f32 	%p80, %f202, %f198;
	@!%p80 bra 	$Lt_0_113154;
	setp.lt.f32 	%p81, %f116, %f202;
	@!%p81 bra 	$Lt_0_113666;
	.loc	18	316	0
	mov.f32 	%f212, %f202;
	bra.uni 	$Lt_0_113410;
$Lt_0_113666:
	setp.lt.f32 	%p82, %f116, %f198;
	selp.f32 	%f212, %f116, %f198, %p82;
$Lt_0_113410:
	mov.f32 	%f213, %f212;
	bra.uni 	$Lt_0_112898;
$Lt_0_113154:
	setp.gt.f32 	%p83, %f116, %f202;
	@!%p83 bra 	$Lt_0_114178;
	mov.f32 	%f214, %f202;
	bra.uni 	$Lt_0_113922;
$Lt_0_114178:
	max.f32 	%f214, %f116, %f198;
$Lt_0_113922:
	mov.f32 	%f213, %f214;
$Lt_0_112898:
	add.f32 	%f215, %f203, %f199;
	mov.f32 	%f216, 0f3f000000;   	// 0.5
	mul.f32 	%f122, %f215, %f216;
	setp.lt.f32 	%p84, %f203, %f199;
	@!%p84 bra 	$Lt_0_114690;
	setp.lt.f32 	%p85, %f122, %f203;
	@!%p85 bra 	$Lt_0_115202;
	.loc	18	317	0
	mov.f32 	%f217, %f203;
	bra.uni 	$Lt_0_114946;
$Lt_0_115202:
	setp.lt.f32 	%p86, %f122, %f199;
	selp.f32 	%f217, %f122, %f199, %p86;
$Lt_0_114946:
	mov.f32 	%f218, %f217;
	bra.uni 	$Lt_0_114434;
$Lt_0_114690:
	setp.gt.f32 	%p87, %f122, %f203;
	@!%p87 bra 	$Lt_0_115714;
	mov.f32 	%f219, %f203;
	bra.uni 	$Lt_0_115458;
$Lt_0_115714:
	max.f32 	%f219, %f122, %f199;
$Lt_0_115458:
	mov.f32 	%f218, %f219;
$Lt_0_114434:
	add.f32 	%f220, %f204, %f200;
	mov.f32 	%f221, 0f3f000000;   	// 0.5
	mul.f32 	%f128, %f220, %f221;
	setp.lt.f32 	%p88, %f204, %f200;
	@!%p88 bra 	$Lt_0_116226;
	setp.lt.f32 	%p89, %f128, %f204;
	@!%p89 bra 	$Lt_0_116738;
	.loc	18	318	0
	mov.f32 	%f222, %f204;
	bra.uni 	$Lt_0_116482;
$Lt_0_116738:
	setp.lt.f32 	%p90, %f128, %f200;
	selp.f32 	%f222, %f128, %f200, %p90;
$Lt_0_116482:
	mov.f32 	%f223, %f222;
	bra.uni 	$Lt_0_115970;
$Lt_0_116226:
	setp.gt.f32 	%p91, %f128, %f204;
	@!%p91 bra 	$Lt_0_117250;
	mov.f32 	%f224, %f204;
	bra.uni 	$Lt_0_116994;
$Lt_0_117250:
	max.f32 	%f224, %f128, %f200;
$Lt_0_116994:
	mov.f32 	%f223, %f224;
$Lt_0_115970:
	.loc	17	127	0
	mov.f32 	%f132, %f208;
	mov.f32 	%f133, %f213;
	mov.f32 	%f134, %f218;
	mov.f32 	%f225, 0f00000000;   	// 0
	setp.neu.f32 	%p92, %f223, %f225;
	@!%p92 bra 	$Lt_0_117506;
	div.full.f32 	%f134, %f134, %f223;
	div.full.f32 	%f133, %f133, %f223;
	div.full.f32 	%f132, %f132, %f223;
$Lt_0_117506:
	.loc	16	100	0
	ld.param.s32 	%r105, [__cudaparm_DeinterlaceKernel_inPitch];
	mul24.lo.s32 	%r106, %r89, %r105;
	ld.param.u64 	%rd71, [__cudaparm_DeinterlaceKernel_inImage];
	add.s32 	%r107, %r106, %r14;
	cvt.s64.s32 	%rd72, %r107;
	mul.lo.u64 	%rd73, %rd72, 16;
	add.u64 	%rd74, %rd71, %rd73;
	st.global.v4.f32 	[%rd74+0], {%f132,%f133,%f134,%f223};
	.loc	17	127	0
	bra.uni 	$Lt_0_119042;
$Lt_0_111106:
	mov.s32 	%r108, -2;
	set.eq.u32.s32 	%r109, %r19, %r108;
	neg.s32 	%r110, %r109;
	mov.s32 	%r111, 1;
	set.eq.u32.s32 	%r112, %r5, %r111;
	neg.s32 	%r113, %r112;
	and.b32 	%r114, %r110, %r113;
	mov.u32 	%r115, 0;
	setp.eq.s32 	%p93, %r114, %r115;
	@%p93 bra 	$Lt_0_118274;
	.loc	17	137	0
	cvt.u64.u32 	%rd75, %r10;
	mul.lo.u64 	%rd10, %rd75, 4;
	add.u64 	%rd76, %rd10, %rd3;
	ld.shared.f32 	%f226, [%rd76+2048];
	add.u64 	%rd77, %rd10, %rd1;
	ld.shared.f32 	%f227, [%rd77+2048];
	add.u64 	%rd78, %rd10, %rd4;
	ld.shared.f32 	%f228, [%rd78+2048];
	add.u64 	%rd79, %rd10, %rd2;
	ld.shared.f32 	%f229, [%rd79+2048];
	.loc	17	138	0
	mov.f32 	%f140, %f226;
	mov.f32 	%f141, %f227;
	mov.f32 	%f142, %f228;
	mov.f32 	%f230, 0f00000000;   	// 0
	setp.neu.f32 	%p94, %f229, %f230;
	@!%p94 bra 	$Lt_0_118530;
	.loc	17	137	0
	ld.shared.f32 	%f228, [%rd78+2048];
	ld.shared.f32 	%f229, [%rd79+2048];
	.loc	17	138	0
	div.full.f32 	%f142, %f228, %f229;
	.loc	17	137	0
	ld.shared.f32 	%f227, [%rd77+2048];
	.loc	17	138	0
	div.full.f32 	%f141, %f227, %f229;
	.loc	17	137	0
	ld.shared.f32 	%f226, [%rd76+2048];
	.loc	17	138	0
	div.full.f32 	%f140, %f226, %f229;
$Lt_0_118530:
	.loc	16	100	0
	ld.param.s32 	%r116, [__cudaparm_DeinterlaceKernel_inPitch];
	mul24.lo.s32 	%r117, %r89, %r116;
	ld.param.u64 	%rd80, [__cudaparm_DeinterlaceKernel_inImage];
	add.s32 	%r118, %r117, %r14;
	cvt.s64.s32 	%rd81, %r118;
	mul.lo.u64 	%rd82, %rd81, 16;
	add.u64 	%rd83, %rd80, %rd82;
	.loc	17	137	0
	ld.shared.f32 	%f229, [%rd79+2048];
	st.global.v4.f32 	[%rd83+0], {%f140,%f141,%f142,%f229};
	.loc	17	138	0
	bra.uni 	$Lt_0_119042;
$Lt_0_118274:
	selp.s32 	%r119, 1, 0, %p1;
	sub.s32 	%r120, %r29, 1;
	set.eq.u32.s32 	%r121, %r89, %r120;
	neg.s32 	%r122, %r121;
	and.b32 	%r123, %r119, %r122;
	mov.u32 	%r124, 0;
	setp.eq.s32 	%p95, %r123, %r124;
	@%p95 bra 	$Lt_0_119042;
	.loc	17	148	0
	cvt.u64.u32 	%rd84, %r10;
	mul.lo.u64 	%rd10, %rd84, 4;
	add.u64 	%rd85, %rd10, %rd3;
	ld.shared.f32 	%f231, [%rd85+1024];
	add.u64 	%rd86, %rd10, %rd1;
	ld.shared.f32 	%f232, [%rd86+1024];
	add.u64 	%rd87, %rd10, %rd4;
	ld.shared.f32 	%f233, [%rd87+1024];
	add.u64 	%rd88, %rd10, %rd2;
	ld.shared.f32 	%f234, [%rd88+1024];
	.loc	17	149	0
	mov.f32 	%f148, %f231;
	mov.f32 	%f149, %f232;
	mov.f32 	%f150, %f233;
	mov.f32 	%f235, 0f00000000;   	// 0
	setp.neu.f32 	%p96, %f234, %f235;
	@!%p96 bra 	$Lt_0_119554;
	.loc	17	148	0
	ld.shared.f32 	%f233, [%rd87+1024];
	ld.shared.f32 	%f234, [%rd88+1024];
	.loc	17	149	0
	div.full.f32 	%f150, %f233, %f234;
	.loc	17	148	0
	ld.shared.f32 	%f232, [%rd86+1024];
	.loc	17	149	0
	div.full.f32 	%f149, %f232, %f234;
	.loc	17	148	0
	ld.shared.f32 	%f231, [%rd85+1024];
	.loc	17	149	0
	div.full.f32 	%f148, %f231, %f234;
$Lt_0_119554:
	.loc	16	100	0
	ld.param.s32 	%r125, [__cudaparm_DeinterlaceKernel_inPitch];
	mul24.lo.s32 	%r126, %r89, %r125;
	ld.param.u64 	%rd89, [__cudaparm_DeinterlaceKernel_inImage];
	add.s32 	%r127, %r126, %r14;
	cvt.s64.s32 	%rd90, %r127;
	mul.lo.u64 	%rd91, %rd90, 16;
	add.u64 	%rd92, %rd89, %rd91;
	.loc	17	148	0
	ld.shared.f32 	%f234, [%rd88+1024];
	st.global.v4.f32 	[%rd92+0], {%f148,%f149,%f150,%f234};
$Lt_0_119042:
$Lt_0_118018:
$Lt_0_110850:
$L_0_82946:
$Lt_0_85762:
	.loc	17	271	0
	exit;
$LDWend_DeinterlaceKernel:
	} // DeinterlaceKernel

