	.version 2.2
	.target sm_20
	// compiled with ../../../External/3rdParty/NVIDIA/CUDA/win/bin/../open64/lib//be.exe
	// nvopencc 3.2 built on 2010-11-04

	.visible .func (.param .s32 __cudaretf__Z15IntegerMultiplyii) _Z15IntegerMultiplyii (.param .s32 __cudaparmf1__Z15IntegerMultiplyii, .param .s32 __cudaparmf2__Z15IntegerMultiplyii)

	.visible .func (.param .s32 __cudaretf__Z17Standard2DKernelXv) _Z17Standard2DKernelXv ()

	.visible .func (.param .s32 __cudaretf__Z17Standard2DKernelYv) _Z17Standard2DKernelYv ()

	.visible .func (.param .align 16 .b8 __cudaretf__Z13Half4ToFloat47ushort4[16]) _Z13Half4ToFloat47ushort4 (.param .align 8 .b8 __cudaparmf1__Z13Half4ToFloat47ushort4[8])

	.visible .func (.param .align 8 .b8 __cudaretf__Z13Float4ToHalf46float4[8]) _Z13Float4ToHalf46float4 (.param .align 16 .b8 __cudaparmf1__Z13Float4ToHalf46float4[16])

	.visible .func (.param .u32 __cudaretf__Z4Mix3RjS_S_) _Z4Mix3RjS_S_ (.param .u64 __cudaparmf1__Z4Mix3RjS_S_, .param .u64 __cudaparmf2__Z4Mix3RjS_S_, .param .u64 __cudaparmf3__Z4Mix3RjS_S_)

	.visible .func (.param .s32 __cudaretf__Z4Randj) _Z4Randj (.param .u32 __cudaparmf1__Z4Randj)

	.visible .func (.param .s32 __cudaretf__Z6Rand2Djjj) _Z6Rand2Djjj (.param .u32 __cudaparmf1__Z6Rand2Djjj, .param .u32 __cudaparmf2__Z6Rand2Djjj, .param .u32 __cudaparmf3__Z6Rand2Djjj)

	.visible .func (.param .s32 __cudaretf__Z6Rand2Dj) _Z6Rand2Dj (.param .u32 __cudaparmf1__Z6Rand2Dj)

	.visible .func _Z7Write2DI7ushort4EvT_PS1_iii (.param .align 8 .b8 __cudaparmf1__Z7Write2DI7ushort4EvT_PS1_iii[8], .param .u64 __cudaparmf2__Z7Write2DI7ushort4EvT_PS1_iii, .param .s32 __cudaparmf3__Z7Write2DI7ushort4EvT_PS1_iii, .param .s32 __cudaparmf4__Z7Write2DI7ushort4EvT_PS1_iii, .param .s32 __cudaparmf5__Z7Write2DI7ushort4EvT_PS1_iii)

	.visible .func _Z7Write2DI6float4EvT_PS1_iii (.param .align 16 .b8 __cudaparmf1__Z7Write2DI6float4EvT_PS1_iii[16], .param .u64 __cudaparmf2__Z7Write2DI6float4EvT_PS1_iii, .param .s32 __cudaparmf3__Z7Write2DI6float4EvT_PS1_iii, .param .s32 __cudaparmf4__Z7Write2DI6float4EvT_PS1_iii, .param .s32 __cudaparmf5__Z7Write2DI6float4EvT_PS1_iii)

	.visible .func (.param .align 16 .b8 __cudaretf__Z18UnpremultiplyPixel8PixelRGB[16]) _Z18UnpremultiplyPixel8PixelRGB (.param .align 16 .b8 __cudaparmf1__Z18UnpremultiplyPixel8PixelRGB[16])

	.visible .func (.param .f32 __cudaretf__Z13ToLinearColorf) _Z13ToLinearColorf (.param .f32 __cudaparmf1__Z13ToLinearColorf)

	.visible .func (.param .f32 __cudaretf__Z15FromLinearColorf) _Z15FromLinearColorf (.param .f32 __cudaparmf1__Z15FromLinearColorf)

	.visible .func (.param .align 16 .b8 __cudaretf__Z25PremultiplyLinearizePixel8PixelRGB[16]) _Z25PremultiplyLinearizePixel8PixelRGB (.param .align 16 .b8 __cudaparmf1__Z25PremultiplyLinearizePixel8PixelRGB[16])

	.visible .func (.param .align 16 .b8 __cudaretf__Z29UnpremultiplyUnlinearizePixel8PixelRGB[16]) _Z29UnpremultiplyUnlinearizePixel8PixelRGB (.param .align 16 .b8 __cudaparmf1__Z29UnpremultiplyUnlinearizePixel8PixelRGB[16])

	.visible .func (.param .align 16 .b8 __cudaretf__Z20PremultiplyLinearize6float4[16]) _Z20PremultiplyLinearize6float4 (.param .align 16 .b8 __cudaparmf1__Z20PremultiplyLinearize6float4[16])

	.visible .func (.param .align 16 .b8 __cudaretf__Z24UnpremultiplyUnlinearize6float4[16]) _Z24UnpremultiplyUnlinearize6float4 (.param .align 16 .b8 __cudaparmf1__Z24UnpremultiplyUnlinearize6float4[16])

	.visible .func (.param .align 16 .b8 __cudaretf__Z17AverageComponentsI8PixelRGBET_RKS1_S3_[16]) _Z17AverageComponentsI8PixelRGBET_RKS1_S3_ (.param .u64 __cudaparmf1__Z17AverageComponentsI8PixelRGBET_RKS1_S3_, .param .u64 __cudaparmf2__Z17AverageComponentsI8PixelRGBET_RKS1_S3_)

	.visible .func (.param .align 16 .b8 __cudaretf__ZmlI8PixelRGBET_RKS1_f[16]) _ZmlI8PixelRGBET_RKS1_f (.param .u64 __cudaparmf1__ZmlI8PixelRGBET_RKS1_f, .param .f32 __cudaparmf2__ZmlI8PixelRGBET_RKS1_f)

	.visible .func (.param .align 16 .b8 __cudaretf__Z16MedianComponentsI8PixelRGBET_RKS1_S3_S3_[16]) _Z16MedianComponentsI8PixelRGBET_RKS1_S3_S3_ (.param .u64 __cudaparmf1__Z16MedianComponentsI8PixelRGBET_RKS1_S3_S3_, .param .u64 __cudaparmf2__Z16MedianComponentsI8PixelRGBET_RKS1_S3_S3_, .param .u64 __cudaparmf3__Z16MedianComponentsI8PixelRGBET_RKS1_S3_S3_)

	//-----------------------------------------------------------
	// Compiling C:/Users/dvaeng/AppData/Local/Temp/tmpxft_00003b38_00000000-11_Deinterlace.cpp3.i (C:/Users/dvaeng/AppData/Local/Temp/ccBI#.a15892)
	//-----------------------------------------------------------

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

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


	.visible .func (.param .s32 __cudaretf__Z15IntegerMultiplyii) _Z15IntegerMultiplyii (.param .s32 __cudaparmf1__Z15IntegerMultiplyii, .param .s32 __cudaparmf2__Z15IntegerMultiplyii)
	{
	.reg .u32 %r<7>;
	.loc	19	60	0
$LDWbegin__Z15IntegerMultiplyii:
	ld.param.u32 	%r1, [__cudaparmf1__Z15IntegerMultiplyii];
	mov.s32 	%r2, %r1;
	ld.param.u32 	%r3, [__cudaparmf2__Z15IntegerMultiplyii];
	mov.s32 	%r4, %r3;
	.loc	19	64	0
	mul.lo.s32 	%r5, %r2, %r4;
	st.param.s32 	[__cudaretf__Z15IntegerMultiplyii], %r5;
	ret;
$LDWend__Z15IntegerMultiplyii:
	} // _Z15IntegerMultiplyii

	.visible .func (.param .s32 __cudaretf__Z17Standard2DKernelXv) _Z17Standard2DKernelXv ()
	{
	.reg .u32 %r<7>;
	.loc	19	73	0
$LDWbegin__Z17Standard2DKernelXv:
	.loc	19	74	0
	mov.u32 	%r1, %tid.x;
	cvt.s32.u32 	%r2, %ctaid.x;
	cvt.s32.u32 	%r3, %ntid.x;
	mul.lo.s32 	%r4, %r2, %r3;
	add.u32 	%r5, %r1, %r4;
	st.param.s32 	[__cudaretf__Z17Standard2DKernelXv], %r5;
	ret;
$LDWend__Z17Standard2DKernelXv:
	} // _Z17Standard2DKernelXv

	.visible .func (.param .s32 __cudaretf__Z17Standard2DKernelYv) _Z17Standard2DKernelYv ()
	{
	.reg .u32 %r<7>;
	.loc	19	77	0
$LDWbegin__Z17Standard2DKernelYv:
	.loc	19	78	0
	mov.u32 	%r1, %tid.y;
	cvt.s32.u32 	%r2, %ctaid.y;
	cvt.s32.u32 	%r3, %ntid.y;
	mul.lo.s32 	%r4, %r2, %r3;
	add.u32 	%r5, %r1, %r4;
	st.param.s32 	[__cudaretf__Z17Standard2DKernelYv], %r5;
	ret;
$LDWend__Z17Standard2DKernelYv:
	} // _Z17Standard2DKernelYv

	.visible .func (.param .align 16 .b8 __cudaretf__Z13Half4ToFloat47ushort4[16]) _Z13Half4ToFloat47ushort4 (.param .align 8 .b8 __cudaparmf1__Z13Half4ToFloat47ushort4[8])
	{
	.reg .u32 %r<14>;
	.reg .f32 %f<9>;
	.loc	19	86	0
$LDWbegin__Z13Half4ToFloat47ushort4:
	ld.param.u16 	%r1, [__cudaparmf1__Z13Half4ToFloat47ushort4+0];
	mov.s32 	%r2, %r1;
	ld.param.u16 	%r3, [__cudaparmf1__Z13Half4ToFloat47ushort4+2];
	mov.s32 	%r4, %r3;
	ld.param.u16 	%r5, [__cudaparmf1__Z13Half4ToFloat47ushort4+4];
	mov.s32 	%r6, %r5;
	ld.param.u16 	%r7, [__cudaparmf1__Z13Half4ToFloat47ushort4+6];
	mov.s32 	%r8, %r7;
	.loc	19	87	0
	cvt.u16.u32 	%r9, %r4;
	{ .reg .b32 %b1;
	mov.b32		%b1, %r9;
	cvt.ftz.f32.f16	%f1, %b1; }
	cvt.u16.u32 	%r10, %r6;
	{ .reg .b32 %b1;
	mov.b32		%b1, %r10;
	cvt.ftz.f32.f16	%f2, %b1; }
	cvt.u16.u32 	%r11, %r8;
	{ .reg .b32 %b1;
	mov.b32		%b1, %r11;
	cvt.ftz.f32.f16	%f3, %b1; }
	cvt.u16.u32 	%r12, %r2;
	{ .reg .b32 %b1;
	mov.b32		%b1, %r12;
	cvt.ftz.f32.f16	%f4, %b1; }
	st.param.f32 	[__cudaretf__Z13Half4ToFloat47ushort4+0], %f4;
	mov.f32 	%f5, %f1;
	st.param.f32 	[__cudaretf__Z13Half4ToFloat47ushort4+4], %f5;
	mov.f32 	%f6, %f2;
	st.param.f32 	[__cudaretf__Z13Half4ToFloat47ushort4+8], %f6;
	mov.f32 	%f7, %f3;
	st.param.f32 	[__cudaretf__Z13Half4ToFloat47ushort4+12], %f7;
	ret;
$LDWend__Z13Half4ToFloat47ushort4:
	} // _Z13Half4ToFloat47ushort4

	.visible .func (.param .align 8 .b8 __cudaretf__Z13Float4ToHalf46float4[8]) _Z13Float4ToHalf46float4 (.param .align 16 .b8 __cudaparmf1__Z13Float4ToHalf46float4[16])
	{
	.reg .u32 %r<13>;
	.reg .f32 %f<10>;
	.loc	19	95	0
$LDWbegin__Z13Float4ToHalf46float4:
	ld.param.f32 	%f1, [__cudaparmf1__Z13Float4ToHalf46float4+0];
	mov.f32 	%f2, %f1;
	ld.param.f32 	%f3, [__cudaparmf1__Z13Float4ToHalf46float4+4];
	mov.f32 	%f4, %f3;
	ld.param.f32 	%f5, [__cudaparmf1__Z13Float4ToHalf46float4+8];
	mov.f32 	%f6, %f5;
	ld.param.f32 	%f7, [__cudaparmf1__Z13Float4ToHalf46float4+12];
	mov.f32 	%f8, %f7;
	.loc	19	96	0
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f4;
	mov.b32		%r1, %b1; }
	cvt.u16.u32 	%r2, %r1;
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f6;
	mov.b32		%r3, %b1; }
	cvt.u16.u32 	%r4, %r3;
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f8;
	mov.b32		%r5, %b1; }
	cvt.u16.u32 	%r6, %r5;
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f2;
	mov.b32		%r7, %b1; }
	cvt.u16.u32 	%r8, %r7;
	st.param.u16 	[__cudaretf__Z13Float4ToHalf46float4+0], %r8;
	mov.s32 	%r9, %r2;
	st.param.u16 	[__cudaretf__Z13Float4ToHalf46float4+2], %r9;
	mov.s32 	%r10, %r4;
	st.param.u16 	[__cudaretf__Z13Float4ToHalf46float4+4], %r10;
	mov.s32 	%r11, %r6;
	st.param.u16 	[__cudaretf__Z13Float4ToHalf46float4+6], %r11;
	ret;
$LDWend__Z13Float4ToHalf46float4:
	} // _Z13Float4ToHalf46float4

	.visible .func (.param .u32 __cudaretf__Z4Mix3RjS_S_) _Z4Mix3RjS_S_ (.param .u64 __cudaparmf1__Z4Mix3RjS_S_, .param .u64 __cudaparmf2__Z4Mix3RjS_S_, .param .u64 __cudaparmf3__Z4Mix3RjS_S_)
	{
	.reg .u32 %r<75>;
	.reg .u64 %rd<8>;
	.loc	19	138	0
$LDWbegin__Z4Mix3RjS_S_:
	ld.param.u64 	%rd1, [__cudaparmf1__Z4Mix3RjS_S_];
	mov.s64 	%rd2, %rd1;
	ld.param.u64 	%rd3, [__cudaparmf2__Z4Mix3RjS_S_];
	mov.s64 	%rd4, %rd3;
	ld.param.u64 	%rd5, [__cudaparmf3__Z4Mix3RjS_S_];
	mov.s64 	%rd6, %rd5;
	.loc	19	139	0
	ld.u32 	%r1, [%rd2+0];
	ld.u32 	%r2, [%rd4+0];
	sub.u32 	%r3, %r1, %r2;
	st.u32 	[%rd2+0], %r3;
	ld.u32 	%r4, [%rd6+0];
	sub.u32 	%r5, %r3, %r4;
	st.u32 	[%rd2+0], %r5;
	ld.u32 	%r6, [%rd6+0];
	shr.u32 	%r7, %r6, 13;
	xor.b32 	%r8, %r5, %r7;
	st.u32 	[%rd2+0], %r8;
	.loc	19	140	0
	ld.u32 	%r9, [%rd4+0];
	ld.u32 	%r10, [%rd6+0];
	sub.u32 	%r11, %r9, %r10;
	st.u32 	[%rd4+0], %r11;
	ld.u32 	%r12, [%rd2+0];
	sub.u32 	%r13, %r11, %r12;
	st.u32 	[%rd4+0], %r13;
	ld.u32 	%r14, [%rd2+0];
	shl.b32 	%r15, %r14, 8;
	xor.b32 	%r16, %r13, %r15;
	st.u32 	[%rd4+0], %r16;
	.loc	19	141	0
	ld.u32 	%r17, [%rd6+0];
	ld.u32 	%r18, [%rd2+0];
	sub.u32 	%r19, %r17, %r18;
	st.u32 	[%rd6+0], %r19;
	ld.u32 	%r20, [%rd4+0];
	sub.u32 	%r21, %r19, %r20;
	st.u32 	[%rd6+0], %r21;
	ld.u32 	%r22, [%rd4+0];
	shr.u32 	%r23, %r22, 13;
	xor.b32 	%r24, %r21, %r23;
	st.u32 	[%rd6+0], %r24;
	.loc	19	142	0
	ld.u32 	%r25, [%rd2+0];
	ld.u32 	%r26, [%rd4+0];
	sub.u32 	%r27, %r25, %r26;
	st.u32 	[%rd2+0], %r27;
	ld.u32 	%r28, [%rd6+0];
	sub.u32 	%r29, %r27, %r28;
	st.u32 	[%rd2+0], %r29;
	ld.u32 	%r30, [%rd6+0];
	shr.u32 	%r31, %r30, 12;
	xor.b32 	%r32, %r29, %r31;
	st.u32 	[%rd2+0], %r32;
	.loc	19	143	0
	ld.u32 	%r33, [%rd4+0];
	ld.u32 	%r34, [%rd6+0];
	sub.u32 	%r35, %r33, %r34;
	st.u32 	[%rd4+0], %r35;
	ld.u32 	%r36, [%rd2+0];
	sub.u32 	%r37, %r35, %r36;
	st.u32 	[%rd4+0], %r37;
	ld.u32 	%r38, [%rd2+0];
	shl.b32 	%r39, %r38, 16;
	xor.b32 	%r40, %r37, %r39;
	st.u32 	[%rd4+0], %r40;
	.loc	19	144	0
	ld.u32 	%r41, [%rd6+0];
	ld.u32 	%r42, [%rd2+0];
	sub.u32 	%r43, %r41, %r42;
	st.u32 	[%rd6+0], %r43;
	ld.u32 	%r44, [%rd4+0];
	sub.u32 	%r45, %r43, %r44;
	st.u32 	[%rd6+0], %r45;
	ld.u32 	%r46, [%rd4+0];
	shr.u32 	%r47, %r46, 5;
	xor.b32 	%r48, %r45, %r47;
	st.u32 	[%rd6+0], %r48;
	.loc	19	145	0
	ld.u32 	%r49, [%rd2+0];
	ld.u32 	%r50, [%rd4+0];
	sub.u32 	%r51, %r49, %r50;
	st.u32 	[%rd2+0], %r51;
	ld.u32 	%r52, [%rd6+0];
	sub.u32 	%r53, %r51, %r52;
	st.u32 	[%rd2+0], %r53;
	ld.u32 	%r54, [%rd6+0];
	shr.u32 	%r55, %r54, 3;
	xor.b32 	%r56, %r53, %r55;
	st.u32 	[%rd2+0], %r56;
	.loc	19	146	0
	ld.u32 	%r57, [%rd4+0];
	ld.u32 	%r58, [%rd6+0];
	sub.u32 	%r59, %r57, %r58;
	st.u32 	[%rd4+0], %r59;
	ld.u32 	%r60, [%rd2+0];
	sub.u32 	%r61, %r59, %r60;
	st.u32 	[%rd4+0], %r61;
	ld.u32 	%r62, [%rd2+0];
	shl.b32 	%r63, %r62, 10;
	xor.b32 	%r64, %r61, %r63;
	st.u32 	[%rd4+0], %r64;
	.loc	19	147	0
	ld.u32 	%r65, [%rd6+0];
	ld.u32 	%r66, [%rd2+0];
	sub.u32 	%r67, %r65, %r66;
	st.u32 	[%rd6+0], %r67;
	ld.u32 	%r68, [%rd4+0];
	sub.u32 	%r69, %r67, %r68;
	st.u32 	[%rd6+0], %r69;
	ld.u32 	%r70, [%rd4+0];
	shr.u32 	%r71, %r70, 15;
	xor.b32 	%r72, %r69, %r71;
	st.u32 	[%rd6+0], %r72;
	.loc	19	148	0
	mov.s32 	%r73, %r72;
	st.param.u32 	[__cudaretf__Z4Mix3RjS_S_], %r73;
	ret;
$LDWend__Z4Mix3RjS_S_:
	} // _Z4Mix3RjS_S_

	.visible .func (.param .s32 __cudaretf__Z4Randj) _Z4Randj (.param .u32 __cudaparmf1__Z4Randj)
	{
	.reg .u32 %r<14>;
	.loc	19	152	0
$LDWbegin__Z4Randj:
	ld.param.u32 	%r1, [__cudaparmf1__Z4Randj];
	mov.s32 	%r2, %r1;
	.loc	19	163	0
	mul.lo.u32 	%r3, %r2, 1103515245;
	add.u32 	%r4, %r3, 12345;
	shr.u32 	%r5, %r4, 16;
	and.b32 	%r6, %r5, 255;
	shl.b32 	%r7, %r6, 7;
	mul.lo.u32 	%r8, %r2, -1029531031;
	sub.u32 	%r9, %r8, 740551042;
	shr.u32 	%r10, %r9, 16;
	and.b32 	%r11, %r10, 255;
	xor.b32 	%r12, %r7, %r11;
	st.param.s32 	[__cudaretf__Z4Randj], %r12;
	ret;
$LDWend__Z4Randj:
	} // _Z4Randj

	.visible .func (.param .s32 __cudaretf__Z6Rand2Djjj) _Z6Rand2Djjj (.param .u32 __cudaparmf1__Z6Rand2Djjj, .param .u32 __cudaparmf2__Z6Rand2Djjj, .param .u32 __cudaparmf3__Z6Rand2Djjj)
	{
	.reg .u32 %r<54>;
	.loc	19	169	0
$LDWbegin__Z6Rand2Djjj:
	ld.param.u32 	%r1, [__cudaparmf1__Z6Rand2Djjj];
	mov.s32 	%r2, %r1;
	ld.param.u32 	%r3, [__cudaparmf2__Z6Rand2Djjj];
	mov.s32 	%r4, %r3;
	ld.param.u32 	%r5, [__cudaparmf3__Z6Rand2Djjj];
	mov.s32 	%r6, %r5;
	.loc	19	139	0
	sub.u32 	%r7, %r2, %r4;
	sub.u32 	%r8, %r7, %r6;
	shr.u32 	%r9, %r6, 13;
	xor.b32 	%r10, %r8, %r9;
	.loc	19	140	0
	sub.u32 	%r11, %r4, %r6;
	sub.u32 	%r12, %r11, %r10;
	shl.b32 	%r13, %r10, 8;
	xor.b32 	%r14, %r12, %r13;
	.loc	19	141	0
	sub.u32 	%r15, %r6, %r10;
	sub.u32 	%r16, %r15, %r14;
	shr.u32 	%r17, %r14, 13;
	xor.b32 	%r18, %r16, %r17;
	.loc	19	142	0
	sub.u32 	%r19, %r10, %r14;
	sub.u32 	%r20, %r19, %r18;
	shr.u32 	%r21, %r18, 12;
	xor.b32 	%r22, %r20, %r21;
	.loc	19	143	0
	sub.u32 	%r23, %r14, %r18;
	sub.u32 	%r24, %r23, %r22;
	shl.b32 	%r25, %r22, 16;
	xor.b32 	%r26, %r24, %r25;
	.loc	19	144	0
	sub.u32 	%r27, %r18, %r22;
	sub.u32 	%r28, %r27, %r26;
	shr.u32 	%r29, %r26, 5;
	xor.b32 	%r30, %r28, %r29;
	.loc	19	145	0
	sub.u32 	%r31, %r22, %r26;
	sub.u32 	%r32, %r31, %r30;
	shr.u32 	%r33, %r30, 3;
	xor.b32 	%r34, %r32, %r33;
	.loc	19	146	0
	sub.u32 	%r35, %r26, %r30;
	sub.u32 	%r36, %r35, %r34;
	shl.b32 	%r37, %r34, 10;
	xor.b32 	%r38, %r36, %r37;
	.loc	19	147	0
	sub.u32 	%r39, %r30, %r34;
	sub.u32 	%r40, %r39, %r38;
	shr.u32 	%r41, %r38, 15;
	xor.b32 	%r42, %r40, %r41;
	.loc	19	170	0
	mul.lo.u32 	%r43, %r42, 1103515245;
	add.u32 	%r44, %r43, 12345;
	shr.u32 	%r45, %r44, 16;
	and.b32 	%r46, %r45, 255;
	shl.b32 	%r47, %r46, 7;
	mul.lo.u32 	%r48, %r42, -1029531031;
	sub.u32 	%r49, %r48, 740551042;
	shr.u32 	%r50, %r49, 16;
	and.b32 	%r51, %r50, 255;
	xor.b32 	%r52, %r47, %r51;
	st.param.s32 	[__cudaretf__Z6Rand2Djjj], %r52;
	ret;
$LDWend__Z6Rand2Djjj:
	} // _Z6Rand2Djjj

	.visible .func (.param .s32 __cudaretf__Z6Rand2Dj) _Z6Rand2Dj (.param .u32 __cudaparmf1__Z6Rand2Dj)
	{
	.reg .u32 %r<60>;
	.loc	19	175	0
$LDWbegin__Z6Rand2Dj:
	ld.param.u32 	%r1, [__cudaparmf1__Z6Rand2Dj];
	mov.s32 	%r2, %r1;
	.loc	19	143	0
	cvt.s32.u32 	%r3, %ctaid.y;
	cvt.s32.u32 	%r4, %ntid.y;
	mul.lo.s32 	%r5, %r3, %r4;
	cvt.s32.u32 	%r6, %ctaid.x;
	cvt.s32.u32 	%r7, %ntid.x;
	mul.lo.s32 	%r8, %r6, %r7;
	mov.u32 	%r9, %tid.y;
	add.u32 	%r10, %r5, %r9;
	mov.u32 	%r11, %tid.x;
	add.u32 	%r12, %r8, %r11;
	shr.u32 	%r13, %r10, 13;
	sub.u32 	%r14, %r2, %r12;
	sub.u32 	%r15, %r12, %r10;
	sub.u32 	%r16, %r14, %r10;
	xor.b32 	%r17, %r13, %r16;
	shl.b32 	%r18, %r17, 8;
	sub.u32 	%r19, %r15, %r17;
	sub.u32 	%r20, %r10, %r17;
	xor.b32 	%r21, %r18, %r19;
	shr.u32 	%r22, %r21, 13;
	sub.u32 	%r23, %r20, %r21;
	sub.u32 	%r24, %r17, %r21;
	xor.b32 	%r25, %r22, %r23;
	shr.u32 	%r26, %r25, 12;
	sub.u32 	%r27, %r24, %r25;
	xor.b32 	%r28, %r26, %r27;
	sub.u32 	%r29, %r21, %r25;
	sub.u32 	%r30, %r29, %r28;
	shl.b32 	%r31, %r28, 16;
	xor.b32 	%r32, %r30, %r31;
	.loc	19	144	0
	sub.u32 	%r33, %r25, %r28;
	sub.u32 	%r34, %r33, %r32;
	shr.u32 	%r35, %r32, 5;
	xor.b32 	%r36, %r34, %r35;
	.loc	19	145	0
	sub.u32 	%r37, %r28, %r32;
	sub.u32 	%r38, %r37, %r36;
	shr.u32 	%r39, %r36, 3;
	xor.b32 	%r40, %r38, %r39;
	.loc	19	146	0
	sub.u32 	%r41, %r32, %r36;
	sub.u32 	%r42, %r41, %r40;
	shl.b32 	%r43, %r40, 10;
	xor.b32 	%r44, %r42, %r43;
	.loc	19	147	0
	sub.u32 	%r45, %r36, %r40;
	sub.u32 	%r46, %r45, %r44;
	shr.u32 	%r47, %r44, 15;
	xor.b32 	%r48, %r46, %r47;
	.loc	19	176	0
	mul.lo.u32 	%r49, %r48, 1103515245;
	add.u32 	%r50, %r49, 12345;
	shr.u32 	%r51, %r50, 16;
	and.b32 	%r52, %r51, 255;
	shl.b32 	%r53, %r52, 7;
	mul.lo.u32 	%r54, %r48, -1029531031;
	sub.u32 	%r55, %r54, 740551042;
	shr.u32 	%r56, %r55, 16;
	and.b32 	%r57, %r56, 255;
	xor.b32 	%r58, %r53, %r57;
	st.param.s32 	[__cudaretf__Z6Rand2Dj], %r58;
	ret;
$LDWend__Z6Rand2Dj:
	} // _Z6Rand2Dj

	.visible .func _Z7Write2DI7ushort4EvT_PS1_iii (.param .align 8 .b8 __cudaparmf1__Z7Write2DI7ushort4EvT_PS1_iii[8], .param .u64 __cudaparmf2__Z7Write2DI7ushort4EvT_PS1_iii, .param .s32 __cudaparmf3__Z7Write2DI7ushort4EvT_PS1_iii, .param .s32 __cudaparmf4__Z7Write2DI7ushort4EvT_PS1_iii, .param .s32 __cudaparmf5__Z7Write2DI7ushort4EvT_PS1_iii)
	{
	.reg .u32 %r<18>;
	.reg .u64 %rd<7>;
	.loc	19	125	0
$LDWbegin__Z7Write2DI7ushort4EvT_PS1_iii:
	ld.param.u16 	%r1, [__cudaparmf1__Z7Write2DI7ushort4EvT_PS1_iii+0];
	mov.s32 	%r2, %r1;
	ld.param.u16 	%r3, [__cudaparmf1__Z7Write2DI7ushort4EvT_PS1_iii+2];
	mov.s32 	%r4, %r3;
	ld.param.u16 	%r5, [__cudaparmf1__Z7Write2DI7ushort4EvT_PS1_iii+4];
	mov.s32 	%r6, %r5;
	ld.param.u16 	%r7, [__cudaparmf1__Z7Write2DI7ushort4EvT_PS1_iii+6];
	mov.s32 	%r8, %r7;
	ld.param.u64 	%rd1, [__cudaparmf2__Z7Write2DI7ushort4EvT_PS1_iii];
	mov.s64 	%rd2, %rd1;
	ld.param.u32 	%r9, [__cudaparmf3__Z7Write2DI7ushort4EvT_PS1_iii];
	mov.s32 	%r10, %r9;
	ld.param.u32 	%r11, [__cudaparmf4__Z7Write2DI7ushort4EvT_PS1_iii];
	mov.s32 	%r12, %r11;
	ld.param.u32 	%r13, [__cudaparmf5__Z7Write2DI7ushort4EvT_PS1_iii];
	mov.s32 	%r14, %r13;
	.loc	19	126	0
	mul.lo.s32 	%r15, %r10, %r14;
	add.s32 	%r16, %r12, %r15;
	cvt.s64.s32 	%rd3, %r16;
	mul.wide.s32 	%rd4, %r16, 8;
	add.u64 	%rd5, %rd2, %rd4;
	st.v4.u16 	[%rd5+0], {%r2,%r4,%r6,%r8};
	.loc	19	127	0
	ret;
$LDWend__Z7Write2DI7ushort4EvT_PS1_iii:
	} // _Z7Write2DI7ushort4EvT_PS1_iii

	.visible .func _Z7Write2DI6float4EvT_PS1_iii (.param .align 16 .b8 __cudaparmf1__Z7Write2DI6float4EvT_PS1_iii[16], .param .u64 __cudaparmf2__Z7Write2DI6float4EvT_PS1_iii, .param .s32 __cudaparmf3__Z7Write2DI6float4EvT_PS1_iii, .param .s32 __cudaparmf4__Z7Write2DI6float4EvT_PS1_iii, .param .s32 __cudaparmf5__Z7Write2DI6float4EvT_PS1_iii)
	{
	.reg .u32 %r<10>;
	.reg .u64 %rd<7>;
	.reg .f32 %f<10>;
	.loc	19	125	0
$LDWbegin__Z7Write2DI6float4EvT_PS1_iii:
	ld.param.f32 	%f1, [__cudaparmf1__Z7Write2DI6float4EvT_PS1_iii+0];
	mov.f32 	%f2, %f1;
	ld.param.f32 	%f3, [__cudaparmf1__Z7Write2DI6float4EvT_PS1_iii+4];
	mov.f32 	%f4, %f3;
	ld.param.f32 	%f5, [__cudaparmf1__Z7Write2DI6float4EvT_PS1_iii+8];
	mov.f32 	%f6, %f5;
	ld.param.f32 	%f7, [__cudaparmf1__Z7Write2DI6float4EvT_PS1_iii+12];
	mov.f32 	%f8, %f7;
	ld.param.u64 	%rd1, [__cudaparmf2__Z7Write2DI6float4EvT_PS1_iii];
	mov.s64 	%rd2, %rd1;
	ld.param.u32 	%r1, [__cudaparmf3__Z7Write2DI6float4EvT_PS1_iii];
	mov.s32 	%r2, %r1;
	ld.param.u32 	%r3, [__cudaparmf4__Z7Write2DI6float4EvT_PS1_iii];
	mov.s32 	%r4, %r3;
	ld.param.u32 	%r5, [__cudaparmf5__Z7Write2DI6float4EvT_PS1_iii];
	mov.s32 	%r6, %r5;
	.loc	19	126	0
	mul.lo.s32 	%r7, %r2, %r6;
	add.s32 	%r8, %r4, %r7;
	cvt.s64.s32 	%rd3, %r8;
	mul.wide.s32 	%rd4, %r8, 16;
	add.u64 	%rd5, %rd2, %rd4;
	st.v4.f32 	[%rd5+0], {%f2,%f4,%f6,%f8};
	.loc	19	127	0
	ret;
$LDWend__Z7Write2DI6float4EvT_PS1_iii:
	} // _Z7Write2DI6float4EvT_PS1_iii

	.visible .func (.param .align 16 .b8 __cudaretf__Z18UnpremultiplyPixel8PixelRGB[16]) _Z18UnpremultiplyPixel8PixelRGB (.param .align 16 .b8 __cudaparmf1__Z18UnpremultiplyPixel8PixelRGB[16])
	{
	.reg .f32 %f<23>;
	.reg .pred %p<3>;
	.loc	4	206	0
$LDWbegin__Z18UnpremultiplyPixel8PixelRGB:
	ld.param.f32 	%f1, [__cudaparmf1__Z18UnpremultiplyPixel8PixelRGB+0];
	mov.f32 	%f2, %f1;
	ld.param.f32 	%f3, [__cudaparmf1__Z18UnpremultiplyPixel8PixelRGB+4];
	mov.f32 	%f4, %f3;
	ld.param.f32 	%f5, [__cudaparmf1__Z18UnpremultiplyPixel8PixelRGB+8];
	mov.f32 	%f6, %f5;
	ld.param.f32 	%f7, [__cudaparmf1__Z18UnpremultiplyPixel8PixelRGB+12];
	mov.f32 	%f8, %f7;
	.loc	4	208	0
	cvt.ftz.sat.f32.f32 	%f9, %f8;
	mov.f32 	%f10, %f9;
	mov.f32 	%f11, 0fb70637bd;    	// -8e-006
	add.ftz.f32 	%f12, %f9, %f11;
	mov.f32 	%f13, 0f00000000;    	// 0
	setp.le.ftz.f32 	%p1, %f12, %f13;
	@%p1 bra 	$Lt_11_1282;
	.loc	4	213	0
	rcp.approx.ftz.f32 	%f14, %f9;
	mul.ftz.f32 	%f15, %f14, %f6;
	.loc	4	214	0
	mul.ftz.f32 	%f16, %f14, %f4;
	.loc	4	215	0
	mul.ftz.f32 	%f17, %f14, %f2;
	bra.uni 	$Lt_11_1026;
$Lt_11_1282:
	.loc	4	219	0
	mov.f32 	%f15, 0f00000000;    	// 0
	mov.f32 	%f16, 0f00000000;    	// 0
	mov.f32 	%f17, 0f00000000;    	// 0
	mov.f32 	%f10, 0f00000000;    	// 0
$Lt_11_1026:
	.loc	4	224	0
	mov.f32 	%f18, %f17;
	st.param.f32 	[__cudaretf__Z18UnpremultiplyPixel8PixelRGB+0], %f18;
	mov.f32 	%f19, %f16;
	st.param.f32 	[__cudaretf__Z18UnpremultiplyPixel8PixelRGB+4], %f19;
	mov.f32 	%f20, %f15;
	st.param.f32 	[__cudaretf__Z18UnpremultiplyPixel8PixelRGB+8], %f20;
	mov.f32 	%f21, %f10;
	st.param.f32 	[__cudaretf__Z18UnpremultiplyPixel8PixelRGB+12], %f21;
	ret;
$LDWend__Z18UnpremultiplyPixel8PixelRGB:
	} // _Z18UnpremultiplyPixel8PixelRGB

	.visible .func (.param .f32 __cudaretf__Z13ToLinearColorf) _Z13ToLinearColorf (.param .f32 __cudaparmf1__Z13ToLinearColorf)
	{
	.reg .f32 %f<15>;
	.reg .pred %p<3>;
	.loc	4	231	0
$LDWbegin__Z13ToLinearColorf:
	ld.param.f32 	%f1, [__cudaparmf1__Z13ToLinearColorf];
	mov.f32 	%f2, %f1;
	mov.f32 	%f3, 0f00000000;     	// 0
	setp.lt.ftz.f32 	%p1, %f2, %f3;
	@!%p1 bra 	$Lt_12_1026;
	.loc	4	234	0
	neg.ftz.f32 	%f4, %f2;
	lg2.approx.ftz.f32 	%f5, %f4;
	mov.f32 	%f6, 0f400ccccd;     	// 2.2
	mul.ftz.f32 	%f7, %f5, %f6;
	ex2.approx.ftz.f32 	%f8, %f7;
	neg.ftz.f32 	%f9, %f8;
	bra.uni 	$LBB4__Z13ToLinearColorf;
$Lt_12_1026:
	.loc	4	236	0
	lg2.approx.ftz.f32 	%f10, %f2;
	mov.f32 	%f11, 0f400ccccd;    	// 2.2
	mul.ftz.f32 	%f12, %f10, %f11;
	ex2.approx.ftz.f32 	%f9, %f12;
$LBB4__Z13ToLinearColorf:
	mov.f32 	%f13, %f9;
	st.param.f32 	[__cudaretf__Z13ToLinearColorf], %f13;
	ret;
$LDWend__Z13ToLinearColorf:
	} // _Z13ToLinearColorf

	.visible .func (.param .f32 __cudaretf__Z15FromLinearColorf) _Z15FromLinearColorf (.param .f32 __cudaparmf1__Z15FromLinearColorf)
	{
	.reg .f32 %f<15>;
	.reg .pred %p<3>;
	.loc	4	239	0
$LDWbegin__Z15FromLinearColorf:
	ld.param.f32 	%f1, [__cudaparmf1__Z15FromLinearColorf];
	mov.f32 	%f2, %f1;
	mov.f32 	%f3, 0f00000000;     	// 0
	setp.lt.ftz.f32 	%p1, %f2, %f3;
	@!%p1 bra 	$Lt_13_1026;
	.loc	4	242	0
	neg.ftz.f32 	%f4, %f2;
	lg2.approx.ftz.f32 	%f5, %f4;
	mov.f32 	%f6, 0f3ee8ba2e;     	// 0.454545
	mul.ftz.f32 	%f7, %f5, %f6;
	ex2.approx.ftz.f32 	%f8, %f7;
	neg.ftz.f32 	%f9, %f8;
	bra.uni 	$LBB4__Z15FromLinearColorf;
$Lt_13_1026:
	.loc	4	244	0
	lg2.approx.ftz.f32 	%f10, %f2;
	mov.f32 	%f11, 0f3ee8ba2e;    	// 0.454545
	mul.ftz.f32 	%f12, %f10, %f11;
	ex2.approx.ftz.f32 	%f9, %f12;
$LBB4__Z15FromLinearColorf:
	mov.f32 	%f13, %f9;
	st.param.f32 	[__cudaretf__Z15FromLinearColorf], %f13;
	ret;
$LDWend__Z15FromLinearColorf:
	} // _Z15FromLinearColorf

	.visible .func (.param .align 16 .b8 __cudaretf__Z25PremultiplyLinearizePixel8PixelRGB[16]) _Z25PremultiplyLinearizePixel8PixelRGB (.param .align 16 .b8 __cudaparmf1__Z25PremultiplyLinearizePixel8PixelRGB[16])
	{
	.reg .f32 %f<47>;
	.reg .pred %p<5>;
	.loc	4	252	0
$LDWbegin__Z25PremultiplyLinearizePixel8PixelRGB:
	ld.param.f32 	%f1, [__cudaparmf1__Z25PremultiplyLinearizePixel8PixelRGB+0];
	mov.f32 	%f2, %f1;
	ld.param.f32 	%f3, [__cudaparmf1__Z25PremultiplyLinearizePixel8PixelRGB+4];
	mov.f32 	%f4, %f3;
	ld.param.f32 	%f5, [__cudaparmf1__Z25PremultiplyLinearizePixel8PixelRGB+8];
	mov.f32 	%f6, %f5;
	ld.param.f32 	%f7, [__cudaparmf1__Z25PremultiplyLinearizePixel8PixelRGB+12];
	mov.f32 	%f8, %f7;
	.loc	4	254	0
	cvt.ftz.sat.f32.f32 	%f9, %f8;
	.loc	4	255	0
	mov.f32 	%f10, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p1, %f2, %f10;
	@!%p1 bra 	$Lt_14_4098;
	.loc	4	234	0
	neg.ftz.f32 	%f11, %f2;
	lg2.approx.ftz.f32 	%f12, %f11;
	mov.f32 	%f13, 0f400ccccd;    	// 2.2
	mul.ftz.f32 	%f14, %f12, %f13;
	ex2.approx.ftz.f32 	%f15, %f14;
	neg.ftz.f32 	%f16, %f15;
	bra.uni 	$LDWendi___log2f_191_5;
$Lt_14_4098:
	.loc	4	236	0
	lg2.approx.ftz.f32 	%f17, %f2;
	mov.f32 	%f18, 0f400ccccd;    	// 2.2
	mul.ftz.f32 	%f19, %f17, %f18;
	ex2.approx.ftz.f32 	%f16, %f19;
$LDWendi___log2f_191_5:
	.loc	4	256	0
	mov.f32 	%f20, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p2, %f4, %f20;
	@!%p2 bra 	$Lt_14_4610;
	.loc	4	234	0
	neg.ftz.f32 	%f21, %f4;
	lg2.approx.ftz.f32 	%f22, %f21;
	mov.f32 	%f23, 0f400ccccd;    	// 2.2
	mul.ftz.f32 	%f24, %f22, %f23;
	ex2.approx.ftz.f32 	%f25, %f24;
	neg.ftz.f32 	%f26, %f25;
	bra.uni 	$LDWendi___log2f_191_3;
$Lt_14_4610:
	.loc	4	236	0
	lg2.approx.ftz.f32 	%f27, %f4;
	mov.f32 	%f28, 0f400ccccd;    	// 2.2
	mul.ftz.f32 	%f29, %f27, %f28;
	ex2.approx.ftz.f32 	%f26, %f29;
$LDWendi___log2f_191_3:
	.loc	4	257	0
	mov.f32 	%f30, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p3, %f6, %f30;
	@!%p3 bra 	$Lt_14_5122;
	.loc	4	234	0
	neg.ftz.f32 	%f31, %f6;
	lg2.approx.ftz.f32 	%f32, %f31;
	mov.f32 	%f33, 0f400ccccd;    	// 2.2
	mul.ftz.f32 	%f34, %f32, %f33;
	ex2.approx.ftz.f32 	%f35, %f34;
	neg.ftz.f32 	%f36, %f35;
	bra.uni 	$LDWendi___log2f_191_1;
$Lt_14_5122:
	.loc	4	236	0
	lg2.approx.ftz.f32 	%f37, %f6;
	mov.f32 	%f38, 0f400ccccd;    	// 2.2
	mul.ftz.f32 	%f39, %f37, %f38;
	ex2.approx.ftz.f32 	%f36, %f39;
$LDWendi___log2f_191_1:
	.loc	4	259	0
	mul.ftz.f32 	%f40, %f36, %f9;
	mul.ftz.f32 	%f41, %f26, %f9;
	mul.ftz.f32 	%f42, %f16, %f9;
	st.param.f32 	[__cudaretf__Z25PremultiplyLinearizePixel8PixelRGB+0], %f42;
	mov.f32 	%f43, %f41;
	st.param.f32 	[__cudaretf__Z25PremultiplyLinearizePixel8PixelRGB+4], %f43;
	mov.f32 	%f44, %f40;
	st.param.f32 	[__cudaretf__Z25PremultiplyLinearizePixel8PixelRGB+8], %f44;
	mov.f32 	%f45, %f9;
	st.param.f32 	[__cudaretf__Z25PremultiplyLinearizePixel8PixelRGB+12], %f45;
	ret;
$LDWend__Z25PremultiplyLinearizePixel8PixelRGB:
	} // _Z25PremultiplyLinearizePixel8PixelRGB

	.visible .func (.param .align 16 .b8 __cudaretf__Z29UnpremultiplyUnlinearizePixel8PixelRGB[16]) _Z29UnpremultiplyUnlinearizePixel8PixelRGB (.param .align 16 .b8 __cudaparmf1__Z29UnpremultiplyUnlinearizePixel8PixelRGB[16])
	{
	.reg .f32 %f<53>;
	.reg .pred %p<6>;
	.loc	4	263	0
$LDWbegin__Z29UnpremultiplyUnlinearizePixel8PixelRGB:
	ld.param.f32 	%f1, [__cudaparmf1__Z29UnpremultiplyUnlinearizePixel8PixelRGB+0];
	mov.f32 	%f2, %f1;
	ld.param.f32 	%f3, [__cudaparmf1__Z29UnpremultiplyUnlinearizePixel8PixelRGB+4];
	mov.f32 	%f4, %f3;
	ld.param.f32 	%f5, [__cudaparmf1__Z29UnpremultiplyUnlinearizePixel8PixelRGB+8];
	mov.f32 	%f6, %f5;
	ld.param.f32 	%f7, [__cudaparmf1__Z29UnpremultiplyUnlinearizePixel8PixelRGB+12];
	mov.f32 	%f8, %f7;
	.loc	4	208	0
	cvt.ftz.sat.f32.f32 	%f9, %f8;
	mov.f32 	%f10, %f9;
	mov.f32 	%f11, 0fb70637bd;    	// -8e-006
	add.ftz.f32 	%f12, %f9, %f11;
	mov.f32 	%f13, 0f00000000;    	// 0
	setp.le.ftz.f32 	%p1, %f12, %f13;
	@%p1 bra 	$Lt_15_5122;
	.loc	4	213	0
	rcp.approx.ftz.f32 	%f14, %f9;
	mul.ftz.f32 	%f15, %f14, %f6;
	.loc	4	214	0
	mul.ftz.f32 	%f16, %f14, %f4;
	.loc	4	215	0
	mul.ftz.f32 	%f17, %f14, %f2;
	bra.uni 	$Lt_15_4866;
$Lt_15_5122:
	.loc	4	219	0
	mov.f32 	%f15, 0f00000000;    	// 0
	mov.f32 	%f16, 0f00000000;    	// 0
	mov.f32 	%f17, 0f00000000;    	// 0
	mov.f32 	%f10, 0f00000000;    	// 0
$Lt_15_4866:
	.loc	4	266	0
	mov.f32 	%f18, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p2, %f17, %f18;
	@!%p2 bra 	$Lt_15_5378;
	.loc	4	242	0
	neg.ftz.f32 	%f19, %f17;
	lg2.approx.ftz.f32 	%f20, %f19;
	mov.f32 	%f21, 0f3ee8ba2e;    	// 0.454545
	mul.ftz.f32 	%f22, %f20, %f21;
	ex2.approx.ftz.f32 	%f23, %f22;
	neg.ftz.f32 	%f24, %f23;
	bra.uni 	$LDWendi___log2f_192_5;
$Lt_15_5378:
	.loc	4	244	0
	lg2.approx.ftz.f32 	%f25, %f17;
	mov.f32 	%f26, 0f3ee8ba2e;    	// 0.454545
	mul.ftz.f32 	%f27, %f25, %f26;
	ex2.approx.ftz.f32 	%f24, %f27;
$LDWendi___log2f_192_5:
	.loc	4	267	0
	mov.f32 	%f28, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p3, %f16, %f28;
	@!%p3 bra 	$Lt_15_5890;
	.loc	4	242	0
	neg.ftz.f32 	%f29, %f16;
	lg2.approx.ftz.f32 	%f30, %f29;
	mov.f32 	%f31, 0f3ee8ba2e;    	// 0.454545
	mul.ftz.f32 	%f32, %f30, %f31;
	ex2.approx.ftz.f32 	%f33, %f32;
	neg.ftz.f32 	%f34, %f33;
	bra.uni 	$LDWendi___log2f_192_3;
$Lt_15_5890:
	.loc	4	244	0
	lg2.approx.ftz.f32 	%f35, %f16;
	mov.f32 	%f36, 0f3ee8ba2e;    	// 0.454545
	mul.ftz.f32 	%f37, %f35, %f36;
	ex2.approx.ftz.f32 	%f34, %f37;
$LDWendi___log2f_192_3:
	.loc	4	268	0
	mov.f32 	%f38, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p4, %f15, %f38;
	@!%p4 bra 	$Lt_15_6402;
	.loc	4	242	0
	neg.ftz.f32 	%f39, %f15;
	lg2.approx.ftz.f32 	%f40, %f39;
	mov.f32 	%f41, 0f3ee8ba2e;    	// 0.454545
	mul.ftz.f32 	%f42, %f40, %f41;
	ex2.approx.ftz.f32 	%f43, %f42;
	neg.ftz.f32 	%f44, %f43;
	bra.uni 	$LDWendi___log2f_192_1;
$Lt_15_6402:
	.loc	4	244	0
	lg2.approx.ftz.f32 	%f45, %f15;
	mov.f32 	%f46, 0f3ee8ba2e;    	// 0.454545
	mul.ftz.f32 	%f47, %f45, %f46;
	ex2.approx.ftz.f32 	%f44, %f47;
$LDWendi___log2f_192_1:
	.loc	4	269	0
	mov.f32 	%f48, %f24;
	st.param.f32 	[__cudaretf__Z29UnpremultiplyUnlinearizePixel8PixelRGB+0], %f48;
	mov.f32 	%f49, %f34;
	st.param.f32 	[__cudaretf__Z29UnpremultiplyUnlinearizePixel8PixelRGB+4], %f49;
	mov.f32 	%f50, %f44;
	st.param.f32 	[__cudaretf__Z29UnpremultiplyUnlinearizePixel8PixelRGB+8], %f50;
	mov.f32 	%f51, %f10;
	st.param.f32 	[__cudaretf__Z29UnpremultiplyUnlinearizePixel8PixelRGB+12], %f51;
	ret;
$LDWend__Z29UnpremultiplyUnlinearizePixel8PixelRGB:
	} // _Z29UnpremultiplyUnlinearizePixel8PixelRGB

	.visible .func (.param .align 16 .b8 __cudaretf__Z20PremultiplyLinearize6float4[16]) _Z20PremultiplyLinearize6float4 (.param .align 16 .b8 __cudaparmf1__Z20PremultiplyLinearize6float4[16])
	{
	.reg .f32 %f<47>;
	.reg .pred %p<5>;
	.loc	4	277	0
$LDWbegin__Z20PremultiplyLinearize6float4:
	ld.param.f32 	%f1, [__cudaparmf1__Z20PremultiplyLinearize6float4+0];
	mov.f32 	%f2, %f1;
	ld.param.f32 	%f3, [__cudaparmf1__Z20PremultiplyLinearize6float4+4];
	mov.f32 	%f4, %f3;
	ld.param.f32 	%f5, [__cudaparmf1__Z20PremultiplyLinearize6float4+8];
	mov.f32 	%f6, %f5;
	ld.param.f32 	%f7, [__cudaparmf1__Z20PremultiplyLinearize6float4+12];
	mov.f32 	%f8, %f7;
	.loc	4	254	0
	cvt.ftz.sat.f32.f32 	%f9, %f8;
	.loc	4	255	0
	mov.f32 	%f10, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p1, %f2, %f10;
	@!%p1 bra 	$Lt_16_4098;
	.loc	4	234	0
	neg.ftz.f32 	%f11, %f2;
	lg2.approx.ftz.f32 	%f12, %f11;
	mov.f32 	%f13, 0f400ccccd;    	// 2.2
	mul.ftz.f32 	%f14, %f12, %f13;
	ex2.approx.ftz.f32 	%f15, %f14;
	neg.ftz.f32 	%f16, %f15;
	bra.uni 	$LDWendi___log2f_193_5;
$Lt_16_4098:
	.loc	4	236	0
	lg2.approx.ftz.f32 	%f17, %f2;
	mov.f32 	%f18, 0f400ccccd;    	// 2.2
	mul.ftz.f32 	%f19, %f17, %f18;
	ex2.approx.ftz.f32 	%f16, %f19;
$LDWendi___log2f_193_5:
	.loc	4	256	0
	mov.f32 	%f20, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p2, %f4, %f20;
	@!%p2 bra 	$Lt_16_4610;
	.loc	4	234	0
	neg.ftz.f32 	%f21, %f4;
	lg2.approx.ftz.f32 	%f22, %f21;
	mov.f32 	%f23, 0f400ccccd;    	// 2.2
	mul.ftz.f32 	%f24, %f22, %f23;
	ex2.approx.ftz.f32 	%f25, %f24;
	neg.ftz.f32 	%f26, %f25;
	bra.uni 	$LDWendi___log2f_193_3;
$Lt_16_4610:
	.loc	4	236	0
	lg2.approx.ftz.f32 	%f27, %f4;
	mov.f32 	%f28, 0f400ccccd;    	// 2.2
	mul.ftz.f32 	%f29, %f27, %f28;
	ex2.approx.ftz.f32 	%f26, %f29;
$LDWendi___log2f_193_3:
	.loc	4	257	0
	mov.f32 	%f30, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p3, %f6, %f30;
	@!%p3 bra 	$Lt_16_5122;
	.loc	4	234	0
	neg.ftz.f32 	%f31, %f6;
	lg2.approx.ftz.f32 	%f32, %f31;
	mov.f32 	%f33, 0f400ccccd;    	// 2.2
	mul.ftz.f32 	%f34, %f32, %f33;
	ex2.approx.ftz.f32 	%f35, %f34;
	neg.ftz.f32 	%f36, %f35;
	bra.uni 	$LDWendi___log2f_193_1;
$Lt_16_5122:
	.loc	4	236	0
	lg2.approx.ftz.f32 	%f37, %f6;
	mov.f32 	%f38, 0f400ccccd;    	// 2.2
	mul.ftz.f32 	%f39, %f37, %f38;
	ex2.approx.ftz.f32 	%f36, %f39;
$LDWendi___log2f_193_1:
	.loc	4	259	0
	mul.ftz.f32 	%f40, %f36, %f9;
	mul.ftz.f32 	%f41, %f26, %f9;
	.loc	4	278	0
	mul.ftz.f32 	%f42, %f16, %f9;
	st.param.f32 	[__cudaretf__Z20PremultiplyLinearize6float4+0], %f42;
	mov.f32 	%f43, %f41;
	st.param.f32 	[__cudaretf__Z20PremultiplyLinearize6float4+4], %f43;
	mov.f32 	%f44, %f40;
	st.param.f32 	[__cudaretf__Z20PremultiplyLinearize6float4+8], %f44;
	mov.f32 	%f45, %f9;
	st.param.f32 	[__cudaretf__Z20PremultiplyLinearize6float4+12], %f45;
	ret;
$LDWend__Z20PremultiplyLinearize6float4:
	} // _Z20PremultiplyLinearize6float4

	.visible .func (.param .align 16 .b8 __cudaretf__Z24UnpremultiplyUnlinearize6float4[16]) _Z24UnpremultiplyUnlinearize6float4 (.param .align 16 .b8 __cudaparmf1__Z24UnpremultiplyUnlinearize6float4[16])
	{
	.reg .f32 %f<53>;
	.reg .pred %p<6>;
	.loc	4	284	0
$LDWbegin__Z24UnpremultiplyUnlinearize6float4:
	ld.param.f32 	%f1, [__cudaparmf1__Z24UnpremultiplyUnlinearize6float4+0];
	mov.f32 	%f2, %f1;
	ld.param.f32 	%f3, [__cudaparmf1__Z24UnpremultiplyUnlinearize6float4+4];
	mov.f32 	%f4, %f3;
	ld.param.f32 	%f5, [__cudaparmf1__Z24UnpremultiplyUnlinearize6float4+8];
	mov.f32 	%f6, %f5;
	ld.param.f32 	%f7, [__cudaparmf1__Z24UnpremultiplyUnlinearize6float4+12];
	mov.f32 	%f8, %f7;
	.loc	4	208	0
	cvt.ftz.sat.f32.f32 	%f9, %f8;
	mov.f32 	%f10, %f9;
	mov.f32 	%f11, 0fb70637bd;    	// -8e-006
	add.ftz.f32 	%f12, %f9, %f11;
	mov.f32 	%f13, 0f00000000;    	// 0
	setp.le.ftz.f32 	%p1, %f12, %f13;
	@%p1 bra 	$Lt_17_5122;
	.loc	4	213	0
	rcp.approx.ftz.f32 	%f14, %f9;
	mul.ftz.f32 	%f15, %f14, %f6;
	.loc	4	214	0
	mul.ftz.f32 	%f16, %f14, %f4;
	.loc	4	215	0
	mul.ftz.f32 	%f17, %f14, %f2;
	bra.uni 	$Lt_17_4866;
$Lt_17_5122:
	.loc	4	219	0
	mov.f32 	%f15, 0f00000000;    	// 0
	mov.f32 	%f16, 0f00000000;    	// 0
	mov.f32 	%f17, 0f00000000;    	// 0
	mov.f32 	%f10, 0f00000000;    	// 0
$Lt_17_4866:
	.loc	4	266	0
	mov.f32 	%f18, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p2, %f17, %f18;
	@!%p2 bra 	$Lt_17_5378;
	.loc	4	242	0
	neg.ftz.f32 	%f19, %f17;
	lg2.approx.ftz.f32 	%f20, %f19;
	mov.f32 	%f21, 0f3ee8ba2e;    	// 0.454545
	mul.ftz.f32 	%f22, %f20, %f21;
	ex2.approx.ftz.f32 	%f23, %f22;
	neg.ftz.f32 	%f24, %f23;
	bra.uni 	$LDWendi___log2f_194_5;
$Lt_17_5378:
	.loc	4	244	0
	lg2.approx.ftz.f32 	%f25, %f17;
	mov.f32 	%f26, 0f3ee8ba2e;    	// 0.454545
	mul.ftz.f32 	%f27, %f25, %f26;
	ex2.approx.ftz.f32 	%f24, %f27;
$LDWendi___log2f_194_5:
	.loc	4	267	0
	mov.f32 	%f28, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p3, %f16, %f28;
	@!%p3 bra 	$Lt_17_5890;
	.loc	4	242	0
	neg.ftz.f32 	%f29, %f16;
	lg2.approx.ftz.f32 	%f30, %f29;
	mov.f32 	%f31, 0f3ee8ba2e;    	// 0.454545
	mul.ftz.f32 	%f32, %f30, %f31;
	ex2.approx.ftz.f32 	%f33, %f32;
	neg.ftz.f32 	%f34, %f33;
	bra.uni 	$LDWendi___log2f_194_3;
$Lt_17_5890:
	.loc	4	244	0
	lg2.approx.ftz.f32 	%f35, %f16;
	mov.f32 	%f36, 0f3ee8ba2e;    	// 0.454545
	mul.ftz.f32 	%f37, %f35, %f36;
	ex2.approx.ftz.f32 	%f34, %f37;
$LDWendi___log2f_194_3:
	.loc	4	268	0
	mov.f32 	%f38, 0f00000000;    	// 0
	setp.lt.ftz.f32 	%p4, %f15, %f38;
	@!%p4 bra 	$Lt_17_6402;
	.loc	4	242	0
	neg.ftz.f32 	%f39, %f15;
	lg2.approx.ftz.f32 	%f40, %f39;
	mov.f32 	%f41, 0f3ee8ba2e;    	// 0.454545
	mul.ftz.f32 	%f42, %f40, %f41;
	ex2.approx.ftz.f32 	%f43, %f42;
	neg.ftz.f32 	%f44, %f43;
	bra.uni 	$LDWendi___log2f_194_1;
$Lt_17_6402:
	.loc	4	244	0
	lg2.approx.ftz.f32 	%f45, %f15;
	mov.f32 	%f46, 0f3ee8ba2e;    	// 0.454545
	mul.ftz.f32 	%f47, %f45, %f46;
	ex2.approx.ftz.f32 	%f44, %f47;
$LDWendi___log2f_194_1:
	.loc	4	285	0
	mov.f32 	%f48, %f24;
	st.param.f32 	[__cudaretf__Z24UnpremultiplyUnlinearize6float4+0], %f48;
	mov.f32 	%f49, %f34;
	st.param.f32 	[__cudaretf__Z24UnpremultiplyUnlinearize6float4+4], %f49;
	mov.f32 	%f50, %f44;
	st.param.f32 	[__cudaretf__Z24UnpremultiplyUnlinearize6float4+8], %f50;
	mov.f32 	%f51, %f10;
	st.param.f32 	[__cudaretf__Z24UnpremultiplyUnlinearize6float4+12], %f51;
	ret;
$LDWend__Z24UnpremultiplyUnlinearize6float4:
	} // _Z24UnpremultiplyUnlinearize6float4

	.visible .func (.param .align 16 .b8 __cudaretf__Z17AverageComponentsI8PixelRGBET_RKS1_S3_[16]) _Z17AverageComponentsI8PixelRGBET_RKS1_S3_ (.param .u64 __cudaparmf1__Z17AverageComponentsI8PixelRGBET_RKS1_S3_, .param .u64 __cudaparmf2__Z17AverageComponentsI8PixelRGBET_RKS1_S3_)
	{
	.reg .u64 %rd<6>;
	.reg .f32 %f<22>;
	.loc	20	304	0
$LDWbegin__Z17AverageComponentsI8PixelRGBET_RKS1_S3_:
	ld.param.u64 	%rd1, [__cudaparmf1__Z17AverageComponentsI8PixelRGBET_RKS1_S3_];
	mov.s64 	%rd2, %rd1;
	ld.param.u64 	%rd3, [__cudaparmf2__Z17AverageComponentsI8PixelRGBET_RKS1_S3_];
	mov.s64 	%rd4, %rd3;
	ld.v4.f32 	{%f1,%f2,%f3,%f4}, [%rd2+0];
	ld.v4.f32 	{%f5,%f6,%f7,%f8}, [%rd4+0];
	.loc	20	305	0
	add.ftz.f32 	%f9, %f1, %f5;
	add.ftz.f32 	%f10, %f2, %f6;
	add.ftz.f32 	%f11, %f3, %f7;
	add.ftz.f32 	%f12, %f4, %f8;
	mov.f32 	%f13, 0f3f000000;    	// 0.5
	mul.ftz.f32 	%f14, %f9, %f13;
	st.param.f32 	[__cudaretf__Z17AverageComponentsI8PixelRGBET_RKS1_S3_+0], %f14;
	mov.f32 	%f15, 0f3f000000;    	// 0.5
	mul.ftz.f32 	%f16, %f10, %f15;
	st.param.f32 	[__cudaretf__Z17AverageComponentsI8PixelRGBET_RKS1_S3_+4], %f16;
	mov.f32 	%f17, 0f3f000000;    	// 0.5
	mul.ftz.f32 	%f18, %f11, %f17;
	st.param.f32 	[__cudaretf__Z17AverageComponentsI8PixelRGBET_RKS1_S3_+8], %f18;
	mov.f32 	%f19, 0f3f000000;    	// 0.5
	mul.ftz.f32 	%f20, %f12, %f19;
	st.param.f32 	[__cudaretf__Z17AverageComponentsI8PixelRGBET_RKS1_S3_+12], %f20;
	ret;
$LDWend__Z17AverageComponentsI8PixelRGBET_RKS1_S3_:
	} // _Z17AverageComponentsI8PixelRGBET_RKS1_S3_

	.visible .func (.param .align 16 .b8 __cudaretf__Z16MedianComponentsI8PixelRGBET_RKS1_S3_S3_[16]) _Z16MedianComponentsI8PixelRGBET_RKS1_S3_S3_ (.param .u64 __cudaparmf1__Z16MedianComponentsI8PixelRGBET_RKS1_S3_S3_, .param .u64 __cudaparmf2__Z16MedianComponentsI8PixelRGBET_RKS1_S3_S3_, .param .u64 __cudaparmf3__Z16MedianComponentsI8PixelRGBET_RKS1_S3_S3_)
	{
	.reg .u64 %rd<8>;
	.reg .f32 %f<30>;
	.reg .pred %p<18>;
	.loc	20	313	0
$LDWbegin__Z16MedianComponentsI8PixelRGBET_RKS1_S3_S3_:
	ld.param.u64 	%rd1, [__cudaparmf1__Z16MedianComponentsI8PixelRGBET_RKS1_S3_S3_];
	mov.s64 	%rd2, %rd1;
	ld.param.u64 	%rd3, [__cudaparmf2__Z16MedianComponentsI8PixelRGBET_RKS1_S3_S3_];
	mov.s64 	%rd4, %rd3;
	ld.param.u64 	%rd5, [__cudaparmf3__Z16MedianComponentsI8PixelRGBET_RKS1_S3_S3_];
	mov.s64 	%rd6, %rd5;
	.loc	20	315	0
	ld.f32 	%f1, [%rd2+0];
	ld.f32 	%f2, [%rd4+0];
	ld.f32 	%f3, [%rd6+0];
	setp.gt.ftz.f32 	%p1, %f1, %f2;
	@!%p1 bra 	$Lt_19_14850;
	setp.gt.ftz.f32 	%p2, %f2, %f3;
	@!%p2 bra 	$Lt_19_15362;
	mov.f32 	%f4, %f2;
	bra.uni 	$Lt_19_15106;
$Lt_19_15362:
	setp.gt.ftz.f32 	%p3, %f1, %f3;
	selp.f32 	%f4, %f3, %f1, %p3;
$Lt_19_15106:
	mov.f32 	%f5, %f4;
	bra.uni 	$Lt_19_14594;
$Lt_19_14850:
	setp.lt.ftz.f32 	%p4, %f2, %f3;
	@!%p4 bra 	$Lt_19_15874;
	mov.f32 	%f6, %f2;
	bra.uni 	$Lt_19_15618;
$Lt_19_15874:
	max.ftz.f32 	%f6, %f1, %f3;
$Lt_19_15618:
	mov.f32 	%f5, %f6;
$Lt_19_14594:
	.loc	20	316	0
	ld.f32 	%f7, [%rd2+4];
	ld.f32 	%f8, [%rd4+4];
	ld.f32 	%f9, [%rd6+4];
	setp.gt.ftz.f32 	%p5, %f7, %f8;
	@!%p5 bra 	$Lt_19_16386;
	setp.gt.ftz.f32 	%p6, %f8, %f9;
	@!%p6 bra 	$Lt_19_16898;
	mov.f32 	%f10, %f8;
	bra.uni 	$Lt_19_16642;
$Lt_19_16898:
	setp.gt.ftz.f32 	%p7, %f7, %f9;
	selp.f32 	%f10, %f9, %f7, %p7;
$Lt_19_16642:
	mov.f32 	%f11, %f10;
	bra.uni 	$Lt_19_16130;
$Lt_19_16386:
	setp.lt.ftz.f32 	%p8, %f8, %f9;
	@!%p8 bra 	$Lt_19_17410;
	mov.f32 	%f12, %f8;
	bra.uni 	$Lt_19_17154;
$Lt_19_17410:
	max.ftz.f32 	%f12, %f7, %f9;
$Lt_19_17154:
	mov.f32 	%f11, %f12;
$Lt_19_16130:
	.loc	20	317	0
	ld.f32 	%f13, [%rd2+8];
	ld.f32 	%f14, [%rd4+8];
	ld.f32 	%f15, [%rd6+8];
	setp.gt.ftz.f32 	%p9, %f13, %f14;
	@!%p9 bra 	$Lt_19_17922;
	setp.gt.ftz.f32 	%p10, %f14, %f15;
	@!%p10 bra 	$Lt_19_18434;
	mov.f32 	%f16, %f14;
	bra.uni 	$Lt_19_18178;
$Lt_19_18434:
	setp.gt.ftz.f32 	%p11, %f13, %f15;
	selp.f32 	%f16, %f15, %f13, %p11;
$Lt_19_18178:
	mov.f32 	%f17, %f16;
	bra.uni 	$Lt_19_17666;
$Lt_19_17922:
	setp.lt.ftz.f32 	%p12, %f14, %f15;
	@!%p12 bra 	$Lt_19_18946;
	mov.f32 	%f18, %f14;
	bra.uni 	$Lt_19_18690;
$Lt_19_18946:
	max.ftz.f32 	%f18, %f13, %f15;
$Lt_19_18690:
	mov.f32 	%f17, %f18;
$Lt_19_17666:
	.loc	20	318	0
	ld.f32 	%f19, [%rd2+12];
	ld.f32 	%f20, [%rd4+12];
	ld.f32 	%f21, [%rd6+12];
	setp.gt.ftz.f32 	%p13, %f19, %f20;
	@!%p13 bra 	$Lt_19_19458;
	setp.gt.ftz.f32 	%p14, %f20, %f21;
	@!%p14 bra 	$Lt_19_19970;
	mov.f32 	%f22, %f20;
	bra.uni 	$Lt_19_19714;
$Lt_19_19970:
	setp.gt.ftz.f32 	%p15, %f19, %f21;
	selp.f32 	%f22, %f21, %f19, %p15;
$Lt_19_19714:
	mov.f32 	%f23, %f22;
	bra.uni 	$Lt_19_19202;
$Lt_19_19458:
	setp.lt.ftz.f32 	%p16, %f20, %f21;
	@!%p16 bra 	$Lt_19_20482;
	mov.f32 	%f24, %f20;
	bra.uni 	$Lt_19_20226;
$Lt_19_20482:
	max.ftz.f32 	%f24, %f19, %f21;
$Lt_19_20226:
	mov.f32 	%f23, %f24;
$Lt_19_19202:
	.loc	20	319	0
	mov.f32 	%f25, %f5;
	st.param.f32 	[__cudaretf__Z16MedianComponentsI8PixelRGBET_RKS1_S3_S3_+0], %f25;
	mov.f32 	%f26, %f11;
	st.param.f32 	[__cudaretf__Z16MedianComponentsI8PixelRGBET_RKS1_S3_S3_+4], %f26;
	mov.f32 	%f27, %f17;
	st.param.f32 	[__cudaretf__Z16MedianComponentsI8PixelRGBET_RKS1_S3_S3_+8], %f27;
	mov.f32 	%f28, %f23;
	st.param.f32 	[__cudaretf__Z16MedianComponentsI8PixelRGBET_RKS1_S3_S3_+12], %f28;
	ret;
$LDWend__Z16MedianComponentsI8PixelRGBET_RKS1_S3_S3_:
	} // _Z16MedianComponentsI8PixelRGBET_RKS1_S3_S3_

	.visible .func (.param .align 16 .b8 __cudaretf__ZmlI8PixelRGBET_RKS1_f[16]) _ZmlI8PixelRGBET_RKS1_f (.param .u64 __cudaparmf1__ZmlI8PixelRGBET_RKS1_f, .param .f32 __cudaparmf2__ZmlI8PixelRGBET_RKS1_f)
	{
	.reg .u64 %rd<4>;
	.reg .f32 %f<12>;
	.loc	20	150	0
$LDWbegin__ZmlI8PixelRGBET_RKS1_f:
	ld.param.u64 	%rd1, [__cudaparmf1__ZmlI8PixelRGBET_RKS1_f];
	mov.s64 	%rd2, %rd1;
	ld.param.f32 	%f1, [__cudaparmf2__ZmlI8PixelRGBET_RKS1_f];
	mov.f32 	%f2, %f1;
	ld.v4.f32 	{%f3,%f4,%f5,%f6}, [%rd2+0];
	.loc	20	153	0
	mul.ftz.f32 	%f7, %f3, %f2;
	st.param.f32 	[__cudaretf__ZmlI8PixelRGBET_RKS1_f+0], %f7;
	mul.ftz.f32 	%f8, %f4, %f2;
	st.param.f32 	[__cudaretf__ZmlI8PixelRGBET_RKS1_f+4], %f8;
	mul.ftz.f32 	%f9, %f5, %f2;
	st.param.f32 	[__cudaretf__ZmlI8PixelRGBET_RKS1_f+8], %f9;
	mul.ftz.f32 	%f10, %f6, %f2;
	st.param.f32 	[__cudaretf__ZmlI8PixelRGBET_RKS1_f+12], %f10;
	ret;
$LDWend__ZmlI8PixelRGBET_RKS1_f:
	} // _ZmlI8PixelRGBET_RKS1_f
	.global .texref sTexture;

	.entry DeinterlaceKernel (
		.param .u64 __cudaparm_DeinterlaceKernel_inImage,
		.param .s32 __cudaparm_DeinterlaceKernel_inPitch,
		.param .u32 __cudaparm_DeinterlaceKernel_inDeviceFormat,
		.param .s32 __cudaparm_DeinterlaceKernel_inWidth,
		.param .s32 __cudaparm_DeinterlaceKernel_inHeight,
		.param .u32 __cudaparm_DeinterlaceKernel_inWhichField)
	{
	.reg .u32 %r<83>;
	.reg .u64 %rd<26>;
	.reg .f32 %f<284>;
	.reg .pred %p<51>;
	.loc	21	54	0
$LDWbegin_DeinterlaceKernel:
	.loc	21	57	0
	cvt.s32.u32 	%r1, %ctaid.x;
	cvt.s32.u32 	%r2, %ntid.x;
	mul.lo.s32 	%r3, %r1, %r2;
	mov.u32 	%r4, %tid.x;
	add.u32 	%r5, %r3, %r4;
	ld.param.s32 	%r6, [__cudaparm_DeinterlaceKernel_inWidth];
	setp.le.s32 	%p1, %r6, %r5;
	@%p1 bra 	$Lt_21_60674;
	.loc	21	68	0
	ld.param.s32 	%r7, [__cudaparm_DeinterlaceKernel_inWhichField];
	mov.s32 	%r8, 0;
	setp.eq.s32 	%p2, %r7, %r8;
	cvt.s32.u32 	%r9, %ctaid.y;
	cvt.s32.u32 	%r10, %ntid.y;
	mul.lo.s32 	%r11, %r9, %r10;
	mov.u32 	%r12, %tid.y;
	add.u32 	%r13, %r11, %r12;
	shl.b32 	%r14, %r13, 1;
	add.s32 	%r15, %r14, 1;
	selp.s32 	%r16, %r15, %r14, %p2;
	mov.s32 	%r17, 0;
	setp.gt.s32 	%p3, %r16, %r17;
	add.s32 	%r18, %r16, 1;
	ld.param.s32 	%r19, [__cudaparm_DeinterlaceKernel_inHeight];
	setp.lt.s32 	%p4, %r18, %r19;
	@!%p3 bra 	$Lt_21_61698;
	@!%p4 bra 	$Lt_21_61698;
	mov.u32 	%r20, 0;
	setp.le.s32 	%p5, %r5, %r20;
	@%p5 bra 	$Lt_21_61698;
	add.s32 	%r21, %r5, 1;
	setp.le.s32 	%p6, %r6, %r21;
	@%p6 bra 	$Lt_21_61698;
	cvt.rn.f32.s32 	%f1, %r5;
	mov.f32 	%f2, 0fbf000000;     	// -0.5
	add.ftz.f32 	%f3, %f1, %f2;
	cvt.rn.f32.s32 	%f4, %r16;
	mov.f32 	%f5, 0fbf000000;     	// -0.5
	add.ftz.f32 	%f6, %f4, %f5;
	mov.f32 	%f7, %f3;
	mov.f32 	%f8, %f6;
	mov.f32 	%f9, 0f00000000;     	// 0
	mov.f32 	%f10, %f9;
	mov.f32 	%f11, 0f00000000;    	// 0
	mov.f32 	%f12, %f11;
	tex.2d.v4.f32.f32 {%f13,%f14,%f15,%f16},[sTexture,{%f7,%f8,%f10,%f12}];
	.loc	21	70	0
	mov.f32 	%f17, %f13;
	mov.f32 	%f18, %f14;
	mov.f32 	%f19, %f15;
	mov.f32 	%f20, %f16;
	mov.f32 	%f21, 0f3f000000;    	// 0.5
	add.ftz.f32 	%f22, %f1, %f21;
	mov.f32 	%f23, %f22;
	mov.f32 	%f24, %f6;
	mov.f32 	%f25, 0f00000000;    	// 0
	mov.f32 	%f26, %f25;
	mov.f32 	%f27, 0f00000000;    	// 0
	mov.f32 	%f28, %f27;
	tex.2d.v4.f32.f32 {%f29,%f30,%f31,%f32},[sTexture,{%f23,%f24,%f26,%f28}];
	.loc	21	71	0
	mov.f32 	%f33, %f29;
	mov.f32 	%f34, %f30;
	mov.f32 	%f35, %f31;
	mov.f32 	%f36, %f32;
	mov.f32 	%f37, 0f3fc00000;    	// 1.5
	add.ftz.f32 	%f38, %f1, %f37;
	mov.f32 	%f39, %f38;
	mov.f32 	%f40, %f6;
	mov.f32 	%f41, 0f00000000;    	// 0
	mov.f32 	%f42, %f41;
	mov.f32 	%f43, 0f00000000;    	// 0
	mov.f32 	%f44, %f43;
	tex.2d.v4.f32.f32 {%f45,%f46,%f47,%f48},[sTexture,{%f39,%f40,%f42,%f44}];
	.loc	21	72	0
	mov.f32 	%f49, %f45;
	mov.f32 	%f50, %f46;
	mov.f32 	%f51, %f47;
	mov.f32 	%f52, %f48;
	mov.f32 	%f53, 0f3fc00000;    	// 1.5
	add.ftz.f32 	%f54, %f4, %f53;
	mov.f32 	%f55, %f3;
	mov.f32 	%f56, %f54;
	mov.f32 	%f57, 0f00000000;    	// 0
	mov.f32 	%f58, %f57;
	mov.f32 	%f59, 0f00000000;    	// 0
	mov.f32 	%f60, %f59;
	tex.2d.v4.f32.f32 {%f61,%f62,%f63,%f64},[sTexture,{%f55,%f56,%f58,%f60}];
	.loc	21	74	0
	mov.f32 	%f65, %f61;
	mov.f32 	%f66, %f62;
	mov.f32 	%f67, %f63;
	mov.f32 	%f68, %f64;
	mov.f32 	%f69, %f22;
	mov.f32 	%f70, %f54;
	mov.f32 	%f71, 0f00000000;    	// 0
	mov.f32 	%f72, %f71;
	mov.f32 	%f73, 0f00000000;    	// 0
	mov.f32 	%f74, %f73;
	tex.2d.v4.f32.f32 {%f75,%f76,%f77,%f78},[sTexture,{%f69,%f70,%f72,%f74}];
	.loc	21	75	0
	mov.f32 	%f79, %f75;
	mov.f32 	%f80, %f76;
	mov.f32 	%f81, %f77;
	mov.f32 	%f82, %f78;
	mov.f32 	%f83, %f38;
	mov.f32 	%f84, %f54;
	mov.f32 	%f85, 0f00000000;    	// 0
	mov.f32 	%f86, %f85;
	mov.f32 	%f87, 0f00000000;    	// 0
	mov.f32 	%f88, %f87;
	tex.2d.v4.f32.f32 {%f89,%f90,%f91,%f92},[sTexture,{%f83,%f84,%f86,%f88}];
	.loc	21	76	0
	mov.f32 	%f93, %f89;
	mov.f32 	%f94, %f90;
	mov.f32 	%f95, %f91;
	mov.f32 	%f96, %f92;
	mul.ftz.f32 	%f97, %f19, %f20;
	mul.ftz.f32 	%f98, %f95, %f96;
	mul.ftz.f32 	%f99, %f18, %f20;
	mul.ftz.f32 	%f100, %f94, %f96;
	mul.ftz.f32 	%f101, %f17, %f20;
	mul.ftz.f32 	%f102, %f93, %f96;
	mul.ftz.f32 	%f103, %f35, %f36;
	mul.ftz.f32 	%f104, %f81, %f82;
	mul.ftz.f32 	%f105, %f34, %f36;
	mul.ftz.f32 	%f106, %f80, %f82;
	mul.ftz.f32 	%f107, %f33, %f36;
	mul.ftz.f32 	%f108, %f79, %f82;
	mul.ftz.f32 	%f109, %f51, %f52;
	mul.ftz.f32 	%f110, %f67, %f68;
	mul.ftz.f32 	%f111, %f50, %f52;
	mul.ftz.f32 	%f112, %f66, %f68;
	mul.ftz.f32 	%f113, %f49, %f52;
	mul.ftz.f32 	%f114, %f65, %f68;
	sub.ftz.f32 	%f115, %f97, %f98;
	sub.ftz.f32 	%f116, %f99, %f100;
	sub.ftz.f32 	%f117, %f101, %f102;
	sub.ftz.f32 	%f118, %f103, %f104;
	sub.ftz.f32 	%f119, %f105, %f106;
	sub.ftz.f32 	%f120, %f107, %f108;
	sub.ftz.f32 	%f121, %f109, %f110;
	sub.ftz.f32 	%f122, %f111, %f112;
	sub.ftz.f32 	%f123, %f113, %f114;
	mul.ftz.f32 	%f124, %f115, %f115;
	mul.ftz.f32 	%f125, %f118, %f118;
	mul.ftz.f32 	%f126, %f121, %f121;
	fma.rn.ftz.f32 	%f127, %f116, %f116, %f124;
	fma.rn.ftz.f32 	%f128, %f119, %f119, %f125;
	fma.rn.ftz.f32 	%f129, %f122, %f122, %f126;
	fma.rn.ftz.f32 	%f130, %f117, %f117, %f127;
	fma.rn.ftz.f32 	%f131, %f120, %f120, %f128;
	fma.rn.ftz.f32 	%f132, %f123, %f123, %f129;
	set.lt.ftz.u32.f32 	%r22, %f130, %f131;
	neg.s32 	%r23, %r22;
	set.lt.ftz.u32.f32 	%r24, %f130, %f132;
	neg.s32 	%r25, %r24;
	and.b32 	%r26, %r23, %r25;
	mov.u32 	%r27, 0;
	setp.eq.s32 	%p7, %r26, %r27;
	@%p7 bra 	$Lt_21_43522;
	.loc	21	87	0
	mov.f32 	%f133, %f101;
	mov.f32 	%f134, %f99;
	mov.f32 	%f135, %f97;
	mov.f32 	%f136, %f20;
	.loc	21	88	0
	mov.f32 	%f137, %f102;
	mov.f32 	%f138, %f100;
	mov.f32 	%f139, %f98;
	mov.f32 	%f140, %f96;
	bra.uni 	$Lt_21_43778;
$Lt_21_43522:
	setp.lt.ftz.f32 	%p8, %f131, %f132;
	@!%p8 bra 	$Lt_21_44034;
	.loc	21	92	0
	mov.f32 	%f133, %f107;
	mov.f32 	%f134, %f105;
	mov.f32 	%f135, %f103;
	mov.f32 	%f136, %f36;
	.loc	21	93	0
	mov.f32 	%f137, %f108;
	mov.f32 	%f138, %f106;
	mov.f32 	%f139, %f104;
	mov.f32 	%f140, %f82;
	bra.uni 	$Lt_21_43778;
$Lt_21_44034:
	.loc	21	97	0
	mov.f32 	%f133, %f113;
	mov.f32 	%f134, %f111;
	mov.f32 	%f135, %f109;
	mov.f32 	%f136, %f52;
	.loc	21	98	0
	mov.f32 	%f137, %f114;
	mov.f32 	%f138, %f112;
	mov.f32 	%f139, %f110;
	mov.f32 	%f140, %f68;
$Lt_21_43778:
$Lt_21_43266:
	.loc	21	103	0
	add.ftz.f32 	%f141, %f133, %f137;
	mov.f32 	%f142, 0f3f000000;   	// 0.5
	mul.ftz.f32 	%f143, %f141, %f142;
	setp.gt.ftz.f32 	%p9, %f107, %f108;
	@!%p9 bra 	$Lt_21_44546;
	setp.gt.ftz.f32 	%p10, %f108, %f143;
	@!%p10 bra 	$Lt_21_45058;
	.loc	20	315	0
	mov.f32 	%f144, %f108;
	bra.uni 	$Lt_21_44802;
$Lt_21_45058:
	setp.gt.ftz.f32 	%p11, %f107, %f143;
	selp.f32 	%f144, %f143, %f107, %p11;
$Lt_21_44802:
	mov.f32 	%f145, %f144;
	bra.uni 	$Lt_21_44290;
$Lt_21_44546:
	setp.lt.ftz.f32 	%p12, %f108, %f143;
	@!%p12 bra 	$Lt_21_45570;
	mov.f32 	%f146, %f108;
	bra.uni 	$Lt_21_45314;
$Lt_21_45570:
	max.ftz.f32 	%f146, %f107, %f143;
$Lt_21_45314:
	mov.f32 	%f145, %f146;
$Lt_21_44290:
	add.ftz.f32 	%f147, %f134, %f138;
	mov.f32 	%f148, 0f3f000000;   	// 0.5
	mul.ftz.f32 	%f149, %f147, %f148;
	setp.gt.ftz.f32 	%p13, %f105, %f106;
	@!%p13 bra 	$Lt_21_46082;
	setp.gt.ftz.f32 	%p14, %f106, %f149;
	@!%p14 bra 	$Lt_21_46594;
	.loc	20	316	0
	mov.f32 	%f150, %f106;
	bra.uni 	$Lt_21_46338;
$Lt_21_46594:
	setp.gt.ftz.f32 	%p15, %f105, %f149;
	selp.f32 	%f150, %f149, %f105, %p15;
$Lt_21_46338:
	mov.f32 	%f151, %f150;
	bra.uni 	$Lt_21_45826;
$Lt_21_46082:
	setp.lt.ftz.f32 	%p16, %f106, %f149;
	@!%p16 bra 	$Lt_21_47106;
	mov.f32 	%f152, %f106;
	bra.uni 	$Lt_21_46850;
$Lt_21_47106:
	max.ftz.f32 	%f152, %f105, %f149;
$Lt_21_46850:
	mov.f32 	%f151, %f152;
$Lt_21_45826:
	add.ftz.f32 	%f153, %f135, %f139;
	mov.f32 	%f154, 0f3f000000;   	// 0.5
	mul.ftz.f32 	%f155, %f153, %f154;
	setp.gt.ftz.f32 	%p17, %f103, %f104;
	@!%p17 bra 	$Lt_21_47618;
	setp.gt.ftz.f32 	%p18, %f104, %f155;
	@!%p18 bra 	$Lt_21_48130;
	.loc	20	317	0
	mov.f32 	%f156, %f104;
	bra.uni 	$Lt_21_47874;
$Lt_21_48130:
	setp.gt.ftz.f32 	%p19, %f103, %f155;
	selp.f32 	%f156, %f155, %f103, %p19;
$Lt_21_47874:
	mov.f32 	%f157, %f156;
	bra.uni 	$Lt_21_47362;
$Lt_21_47618:
	setp.lt.ftz.f32 	%p20, %f104, %f155;
	@!%p20 bra 	$Lt_21_48642;
	mov.f32 	%f158, %f104;
	bra.uni 	$Lt_21_48386;
$Lt_21_48642:
	max.ftz.f32 	%f158, %f103, %f155;
$Lt_21_48386:
	mov.f32 	%f157, %f158;
$Lt_21_47362:
	add.ftz.f32 	%f159, %f136, %f140;
	mov.f32 	%f160, 0f3f000000;   	// 0.5
	mul.ftz.f32 	%f161, %f159, %f160;
	setp.gt.ftz.f32 	%p21, %f36, %f82;
	@!%p21 bra 	$Lt_21_49154;
	setp.lt.ftz.f32 	%p22, %f161, %f82;
	@!%p22 bra 	$Lt_21_49666;
	.loc	20	318	0
	mov.f32 	%f162, %f82;
	bra.uni 	$Lt_21_49410;
$Lt_21_49666:
	setp.lt.ftz.f32 	%p23, %f161, %f36;
	selp.f32 	%f162, %f161, %f36, %p23;
$Lt_21_49410:
	mov.f32 	%f163, %f162;
	bra.uni 	$Lt_21_48898;
$Lt_21_49154:
	setp.gt.ftz.f32 	%p24, %f161, %f82;
	@!%p24 bra 	$Lt_21_50178;
	mov.f32 	%f164, %f82;
	bra.uni 	$Lt_21_49922;
$Lt_21_50178:
	max.ftz.f32 	%f164, %f161, %f36;
$Lt_21_49922:
	mov.f32 	%f163, %f164;
$Lt_21_48898:
	.loc	4	208	0
	cvt.ftz.sat.f32.f32 	%f165, %f163;
	mov.f32 	%f166, %f165;
	mov.f32 	%f167, 0fb70637bd;   	// -8e-006
	add.ftz.f32 	%f168, %f165, %f167;
	mov.f32 	%f169, 0f00000000;   	// 0
	setp.le.ftz.f32 	%p25, %f168, %f169;
	@%p25 bra 	$Lt_21_50690;
	.loc	4	213	0
	rcp.approx.ftz.f32 	%f170, %f165;
	mul.ftz.f32 	%f171, %f157, %f170;
	.loc	4	214	0
	mul.ftz.f32 	%f172, %f151, %f170;
	.loc	4	215	0
	mul.ftz.f32 	%f173, %f145, %f170;
	bra.uni 	$Lt_21_50434;
$Lt_21_50690:
	.loc	4	219	0
	mov.f32 	%f171, 0f00000000;   	// 0
	mov.f32 	%f172, 0f00000000;   	// 0
	mov.f32 	%f173, 0f00000000;   	// 0
	mov.f32 	%f166, 0f00000000;   	// 0
$Lt_21_50434:
	.loc	21	104	0
	ld.param.s32 	%r28, [__cudaparm_DeinterlaceKernel_inPitch];
	mul.lo.s32 	%r29, %r28, %r16;
	add.s32 	%r30, %r5, %r29;
	cvt.s64.s32 	%rd1, %r30;
	ld.param.u64 	%rd2, [__cudaparm_DeinterlaceKernel_inImage];
	ld.param.s32 	%r31, [__cudaparm_DeinterlaceKernel_inDeviceFormat];
	mov.u32 	%r32, 0;
	setp.ne.s32 	%p26, %r31, %r32;
	@%p26 bra 	$Lt_21_51202;
	.loc	19	126	0
	mul.lo.u64 	%rd3, %rd1, 8;
	add.u64 	%rd4, %rd2, %rd3;
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f173;
	mov.b32		%r33, %b1; }
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f172;
	mov.b32		%r34, %b1; }
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f171;
	mov.b32		%r35, %b1; }
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f166;
	mov.b32		%r36, %b1; }
	st.global.v4.u16 	[%rd4+0], {%r33,%r34,%r35,%r36};
	.loc	21	104	0
	bra.uni 	$Lt_21_60674;
$Lt_21_51202:
	.loc	19	126	0
	mul.lo.u64 	%rd5, %rd1, 16;
	add.u64 	%rd6, %rd2, %rd5;
	st.global.v4.f32 	[%rd6+0], {%f173,%f172,%f171,%f166};
	bra.uni 	$Lt_21_60674;
$Lt_21_61698:
$L_21_41730:
	.loc	21	112	0
	selp.s32 	%r37, 1, 0, %p3;
	selp.s32 	%r38, 1, 0, %p4;
	and.b32 	%r39, %r37, %r38;
	mov.u32 	%r40, 0;
	setp.eq.s32 	%p27, %r39, %r40;
	@%p27 bra 	$Lt_21_51714;
	cvt.rn.f32.s32 	%f174, %r5;
	mov.f32 	%f175, 0f3f000000;   	// 0.5
	add.ftz.f32 	%f176, %f174, %f175;
	cvt.rn.f32.s32 	%f177, %r16;
	mov.f32 	%f178, %f176;
	mov.f32 	%f179, 0fbf000000;   	// -0.5
	add.ftz.f32 	%f180, %f177, %f179;
	mov.f32 	%f181, %f180;
	mov.f32 	%f182, 0f00000000;   	// 0
	mov.f32 	%f183, %f182;
	mov.f32 	%f184, 0f00000000;   	// 0
	mov.f32 	%f185, %f184;
	tex.2d.v4.f32.f32 {%f186,%f187,%f188,%f189},[sTexture,{%f178,%f181,%f183,%f185}];
	.loc	21	115	0
	mov.f32 	%f190, %f186;
	mov.f32 	%f191, %f187;
	mov.f32 	%f192, %f188;
	mov.f32 	%f193, %f189;
	mov.f32 	%f194, %f176;
	mov.f32 	%f195, 0f3fc00000;   	// 1.5
	add.ftz.f32 	%f196, %f177, %f195;
	mov.f32 	%f197, %f196;
	mov.f32 	%f198, 0f00000000;   	// 0
	mov.f32 	%f199, %f198;
	mov.f32 	%f200, 0f00000000;   	// 0
	mov.f32 	%f201, %f200;
	tex.2d.v4.f32.f32 {%f202,%f203,%f204,%f205},[sTexture,{%f194,%f197,%f199,%f201}];
	.loc	21	116	0
	mov.f32 	%f206, %f202;
	mov.f32 	%f207, %f203;
	mov.f32 	%f208, %f204;
	mov.f32 	%f209, %f205;
	.loc	21	119	0
	mul.ftz.f32 	%f210, %f190, %f193;
	mul.ftz.f32 	%f211, %f206, %f209;
	add.ftz.f32 	%f212, %f210, %f211;
	mov.f32 	%f213, 0f3f000000;   	// 0.5
	mul.ftz.f32 	%f214, %f212, %f213;
	setp.gt.ftz.f32 	%p28, %f210, %f211;
	@!%p28 bra 	$Lt_21_52226;
	setp.gt.ftz.f32 	%p29, %f211, %f214;
	@!%p29 bra 	$Lt_21_52738;
	.loc	20	315	0
	mov.f32 	%f215, %f211;
	bra.uni 	$Lt_21_52482;
$Lt_21_52738:
	setp.gt.ftz.f32 	%p30, %f210, %f214;
	selp.f32 	%f215, %f214, %f210, %p30;
$Lt_21_52482:
	mov.f32 	%f216, %f215;
	bra.uni 	$Lt_21_51970;
$Lt_21_52226:
	setp.lt.ftz.f32 	%p31, %f211, %f214;
	@!%p31 bra 	$Lt_21_53250;
	mov.f32 	%f217, %f211;
	bra.uni 	$Lt_21_52994;
$Lt_21_53250:
	max.ftz.f32 	%f217, %f210, %f214;
$Lt_21_52994:
	mov.f32 	%f216, %f217;
$Lt_21_51970:
	mul.ftz.f32 	%f218, %f191, %f193;
	mul.ftz.f32 	%f219, %f207, %f209;
	add.ftz.f32 	%f220, %f218, %f219;
	mov.f32 	%f221, 0f3f000000;   	// 0.5
	mul.ftz.f32 	%f222, %f220, %f221;
	setp.gt.ftz.f32 	%p32, %f218, %f219;
	@!%p32 bra 	$Lt_21_53762;
	setp.gt.ftz.f32 	%p33, %f219, %f222;
	@!%p33 bra 	$Lt_21_54274;
	.loc	20	316	0
	mov.f32 	%f223, %f219;
	bra.uni 	$Lt_21_54018;
$Lt_21_54274:
	setp.gt.ftz.f32 	%p34, %f218, %f222;
	selp.f32 	%f223, %f222, %f218, %p34;
$Lt_21_54018:
	mov.f32 	%f224, %f223;
	bra.uni 	$Lt_21_53506;
$Lt_21_53762:
	setp.lt.ftz.f32 	%p35, %f219, %f222;
	@!%p35 bra 	$Lt_21_54786;
	mov.f32 	%f225, %f219;
	bra.uni 	$Lt_21_54530;
$Lt_21_54786:
	max.ftz.f32 	%f225, %f218, %f222;
$Lt_21_54530:
	mov.f32 	%f224, %f225;
$Lt_21_53506:
	mul.ftz.f32 	%f226, %f192, %f193;
	mul.ftz.f32 	%f227, %f208, %f209;
	add.ftz.f32 	%f228, %f226, %f227;
	mov.f32 	%f229, 0f3f000000;   	// 0.5
	mul.ftz.f32 	%f230, %f228, %f229;
	setp.gt.ftz.f32 	%p36, %f226, %f227;
	@!%p36 bra 	$Lt_21_55298;
	setp.gt.ftz.f32 	%p37, %f227, %f230;
	@!%p37 bra 	$Lt_21_55810;
	.loc	20	317	0
	mov.f32 	%f231, %f227;
	bra.uni 	$Lt_21_55554;
$Lt_21_55810:
	setp.gt.ftz.f32 	%p38, %f226, %f230;
	selp.f32 	%f231, %f230, %f226, %p38;
$Lt_21_55554:
	mov.f32 	%f232, %f231;
	bra.uni 	$Lt_21_55042;
$Lt_21_55298:
	setp.lt.ftz.f32 	%p39, %f227, %f230;
	@!%p39 bra 	$Lt_21_56322;
	mov.f32 	%f233, %f227;
	bra.uni 	$Lt_21_56066;
$Lt_21_56322:
	max.ftz.f32 	%f233, %f226, %f230;
$Lt_21_56066:
	mov.f32 	%f232, %f233;
$Lt_21_55042:
	add.ftz.f32 	%f234, %f193, %f209;
	mov.f32 	%f235, 0f3f000000;   	// 0.5
	mul.ftz.f32 	%f236, %f234, %f235;
	setp.gt.ftz.f32 	%p40, %f193, %f209;
	@!%p40 bra 	$Lt_21_56834;
	setp.lt.ftz.f32 	%p41, %f236, %f209;
	@!%p41 bra 	$Lt_21_57346;
	.loc	20	318	0
	mov.f32 	%f237, %f209;
	bra.uni 	$Lt_21_57090;
$Lt_21_57346:
	setp.lt.ftz.f32 	%p42, %f236, %f193;
	selp.f32 	%f237, %f236, %f193, %p42;
$Lt_21_57090:
	mov.f32 	%f238, %f237;
	bra.uni 	$Lt_21_56578;
$Lt_21_56834:
	setp.gt.ftz.f32 	%p43, %f236, %f209;
	@!%p43 bra 	$Lt_21_57858;
	mov.f32 	%f239, %f209;
	bra.uni 	$Lt_21_57602;
$Lt_21_57858:
	max.ftz.f32 	%f239, %f236, %f193;
$Lt_21_57602:
	mov.f32 	%f238, %f239;
$Lt_21_56578:
	.loc	4	208	0
	cvt.ftz.sat.f32.f32 	%f240, %f238;
	mov.f32 	%f166, %f240;
	mov.f32 	%f241, 0fb70637bd;   	// -8e-006
	add.ftz.f32 	%f242, %f240, %f241;
	mov.f32 	%f243, 0f00000000;   	// 0
	setp.le.ftz.f32 	%p44, %f242, %f243;
	@%p44 bra 	$Lt_21_58370;
	.loc	4	213	0
	rcp.approx.ftz.f32 	%f244, %f240;
	mul.ftz.f32 	%f171, %f232, %f244;
	.loc	4	214	0
	mul.ftz.f32 	%f172, %f224, %f244;
	.loc	4	215	0
	mul.ftz.f32 	%f173, %f216, %f244;
	bra.uni 	$Lt_21_58114;
$Lt_21_58370:
	.loc	4	219	0
	mov.f32 	%f171, 0f00000000;   	// 0
	mov.f32 	%f172, 0f00000000;   	// 0
	mov.f32 	%f173, 0f00000000;   	// 0
	mov.f32 	%f166, 0f00000000;   	// 0
$Lt_21_58114:
	.loc	21	120	0
	ld.param.s32 	%r41, [__cudaparm_DeinterlaceKernel_inPitch];
	mul.lo.s32 	%r42, %r41, %r16;
	add.s32 	%r43, %r5, %r42;
	cvt.s64.s32 	%rd7, %r43;
	ld.param.u64 	%rd8, [__cudaparm_DeinterlaceKernel_inImage];
	ld.param.s32 	%r44, [__cudaparm_DeinterlaceKernel_inDeviceFormat];
	mov.u32 	%r45, 0;
	setp.ne.s32 	%p45, %r44, %r45;
	@%p45 bra 	$Lt_21_58882;
	.loc	19	126	0
	mul.lo.u64 	%rd9, %rd7, 8;
	add.u64 	%rd10, %rd8, %rd9;
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f173;
	mov.b32		%r46, %b1; }
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f172;
	mov.b32		%r47, %b1; }
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f171;
	mov.b32		%r48, %b1; }
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f166;
	mov.b32		%r49, %b1; }
	st.global.v4.u16 	[%rd10+0], {%r46,%r47,%r48,%r49};
	.loc	21	120	0
	bra.uni 	$Lt_21_60674;
$Lt_21_58882:
	.loc	19	126	0
	mul.lo.u64 	%rd11, %rd7, 16;
	add.u64 	%rd12, %rd8, %rd11;
	st.global.v4.f32 	[%rd12+0], {%f173,%f172,%f171,%f166};
	bra.uni 	$Lt_21_60674;
$Lt_21_51714:
	.loc	21	120	0
	mov.s32 	%r50, 0;
	set.eq.u32.s32 	%r51, %r16, %r50;
	neg.s32 	%r52, %r51;
	mov.s32 	%r53, 1;
	set.eq.u32.s32 	%r54, %r7, %r53;
	neg.s32 	%r55, %r54;
	and.b32 	%r56, %r52, %r55;
	mov.u32 	%r57, 0;
	setp.eq.s32 	%p46, %r56, %r57;
	@%p46 bra 	$Lt_21_59394;
	cvt.rn.f32.s32 	%f245, %r5;
	mov.f32 	%f246, 0f3f000000;   	// 0.5
	add.ftz.f32 	%f247, %f245, %f246;
	mov.f32 	%f248, %f247;
	mov.f32 	%f249, 0f3fc00000;   	// 1.5
	mov.f32 	%f250, %f249;
	mov.f32 	%f251, 0f00000000;   	// 0
	mov.f32 	%f252, %f251;
	mov.f32 	%f253, 0f00000000;   	// 0
	mov.f32 	%f254, %f253;
	tex.2d.v4.f32.f32 {%f255,%f256,%f257,%f258},[sTexture,{%f248,%f250,%f252,%f254}];
	.loc	21	131	0
	mov.f32 	%f259, %f255;
	mov.f32 	%f260, %f256;
	mov.f32 	%f261, %f257;
	mov.f32 	%f262, %f258;
	ld.param.s32 	%r58, [__cudaparm_DeinterlaceKernel_inPitch];
	mul.lo.s32 	%r59, %r58, %r16;
	add.s32 	%r60, %r5, %r59;
	cvt.s64.s32 	%rd13, %r60;
	ld.param.u64 	%rd14, [__cudaparm_DeinterlaceKernel_inImage];
	ld.param.s32 	%r61, [__cudaparm_DeinterlaceKernel_inDeviceFormat];
	mov.u32 	%r62, 0;
	setp.ne.s32 	%p47, %r61, %r62;
	@%p47 bra 	$Lt_21_59906;
	.loc	19	126	0
	mul.lo.u64 	%rd15, %rd13, 8;
	add.u64 	%rd16, %rd14, %rd15;
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f259;
	mov.b32		%r63, %b1; }
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f260;
	mov.b32		%r64, %b1; }
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f261;
	mov.b32		%r65, %b1; }
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f262;
	mov.b32		%r66, %b1; }
	st.global.v4.u16 	[%rd16+0], {%r63,%r64,%r65,%r66};
	.loc	21	131	0
	bra.uni 	$Lt_21_60674;
$Lt_21_59906:
	.loc	19	126	0
	mul.lo.u64 	%rd17, %rd13, 16;
	add.u64 	%rd18, %rd14, %rd17;
	st.global.v4.f32 	[%rd18+0], {%f259,%f260,%f261,%f262};
	bra.uni 	$Lt_21_60674;
$Lt_21_59394:
	.loc	21	131	0
	selp.s32 	%r67, 1, 0, %p2;
	sub.s32 	%r68, %r19, 1;
	set.eq.u32.s32 	%r69, %r16, %r68;
	neg.s32 	%r70, %r69;
	and.b32 	%r71, %r67, %r70;
	mov.u32 	%r72, 0;
	setp.eq.s32 	%p48, %r71, %r72;
	@%p48 bra 	$Lt_21_60674;
	cvt.rn.f32.s32 	%f263, %r5;
	mov.f32 	%f264, 0f3f000000;   	// 0.5
	add.ftz.f32 	%f265, %f263, %f264;
	mov.f32 	%f266, %f265;
	cvt.rn.f32.s32 	%f267, %r16;
	mov.f32 	%f268, 0fbf000000;   	// -0.5
	add.ftz.f32 	%f269, %f267, %f268;
	mov.f32 	%f270, %f269;
	mov.f32 	%f271, 0f00000000;   	// 0
	mov.f32 	%f272, %f271;
	mov.f32 	%f273, 0f00000000;   	// 0
	mov.f32 	%f274, %f273;
	tex.2d.v4.f32.f32 {%f275,%f276,%f277,%f278},[sTexture,{%f266,%f270,%f272,%f274}];
	.loc	21	142	0
	mov.f32 	%f279, %f275;
	mov.f32 	%f280, %f276;
	mov.f32 	%f281, %f277;
	mov.f32 	%f282, %f278;
	ld.param.s32 	%r73, [__cudaparm_DeinterlaceKernel_inPitch];
	mul.lo.s32 	%r74, %r73, %r16;
	add.s32 	%r75, %r5, %r74;
	cvt.s64.s32 	%rd19, %r75;
	ld.param.u64 	%rd20, [__cudaparm_DeinterlaceKernel_inImage];
	ld.param.s32 	%r76, [__cudaparm_DeinterlaceKernel_inDeviceFormat];
	mov.u32 	%r77, 0;
	setp.ne.s32 	%p49, %r76, %r77;
	@%p49 bra 	$Lt_21_60930;
	.loc	19	126	0
	mul.lo.u64 	%rd21, %rd19, 8;
	add.u64 	%rd22, %rd20, %rd21;
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f279;
	mov.b32		%r78, %b1; }
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f280;
	mov.b32		%r79, %b1; }
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f281;
	mov.b32		%r80, %b1; }
	{ .reg .b32 %b1;
	cvt.rn.ftz.f16.f32	%b1, %f282;
	mov.b32		%r81, %b1; }
	st.global.v4.u16 	[%rd22+0], {%r78,%r79,%r80,%r81};
	.loc	21	142	0
	bra.uni 	$Lt_21_60674;
$Lt_21_60930:
	.loc	19	126	0
	mul.lo.u64 	%rd23, %rd19, 16;
	add.u64 	%rd24, %rd20, %rd23;
	st.global.v4.f32 	[%rd24+0], {%f279,%f280,%f281,%f282};
$Lt_21_60674:
$Lt_21_60162:
$Lt_21_59138:
$Lt_21_51458:
$L_21_41474:
$Lt_21_42754:
	.loc	21	151	0
	exit;
$LDWend_DeinterlaceKernel:
	} // DeinterlaceKernel

