	.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_0001d5a4_00000000-9_keyer_Kernels.cpp3.i (C:/Users/dvaeng/AppData/Local/Temp/ccBI#.a21076)
	//-----------------------------------------------------------

	//-----------------------------------------------------------
	// 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_0001d5a4_00000000-8_keyer_Kernels.cudafe2.gpu"
	.file	2	"c:\scully64\shared\adobe\MediaCore\Display\Inc\CUDA/Effects/Keyer_Kernels.h"
	.file	3	"C:\Program Files (x86)\Microsoft Visual Studio 9.0\VC\include\crtdefs.h"
	.file	4	"c:\scully64\shared\adobe\MediaCore\External\3rdParty\NVIDIA\CUDA\win\64\include\crt/device_runtime.h"
	.file	5	"c:\scully64\shared\adobe\MediaCore\External\3rdParty\NVIDIA\CUDA\win\64\include\host_defines.h"
	.file	6	"c:\scully64\shared\adobe\MediaCore\External\3rdParty\NVIDIA\CUDA\win\64\include\builtin_types.h"
	.file	7	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\device_types.h"
	.file	8	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\driver_types.h"
	.file	9	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\texture_types.h"
	.file	10	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\vector_types.h"
	.file	11	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\host_defines.h"
	.file	12	"c:\scully64\shared\adobe\MediaCore\External\3rdParty\NVIDIA\CUDA\win\64\include\device_launch_parameters.h"
	.file	13	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\crt\storage_class.h"
	.file	14	"C:\Program Files (x86)\Microsoft Visual Studio 9.0\VC\include\time.h"
	.file	15	"c:\scully64\shared\adobe\MediaCore\GPUFoundation\API\Inc\GPUFoundation/KernelSupport/Utils.h"
	.file	16	"c:/scully64/shared/adobe/MediaCore/Display/Src/CUDA/Effects/keyer_Kernels.cu"
	.file	17	"c:\scully64\shared\adobe\MediaCore\External\3rdParty\NVIDIA\CUDA\win\64\include\common_functions.h"
	.file	18	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\crt/func_macro.h"
	.file	19	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\math_functions.h"
	.file	20	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\device_functions.h"
	.file	21	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\math_constants.h"
	.file	22	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\sm_11_atomic_functions.h"
	.file	23	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\sm_12_atomic_functions.h"
	.file	24	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\sm_13_double_functions.h"
	.file	25	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\common_types.h"
	.file	26	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\texture_fetch_functions.h"
	.file	27	"c:\scully64\shared\adobe\mediacore\external\3rdparty\nvidia\cuda\win\64\include\math_functions_dbl_ptx1.h"

	.const .align 4 .b8 p[164];

	.entry cuda_kernel_chokefilter_x (
		.param .u64 __cudaparm_cuda_kernel_chokefilter_x_id,
		.param .u64 __cudaparm_cuda_kernel_chokefilter_x_od,
		.param .s32 __cudaparm_cuda_kernel_chokefilter_x_idPitch,
		.param .s32 __cudaparm_cuda_kernel_chokefilter_x_odPitch,
		.param .s32 __cudaparm_cuda_kernel_chokefilter_x_w,
		.param .s32 __cudaparm_cuda_kernel_chokefilter_x_h,
		.param .s32 __cudaparm_cuda_kernel_chokefilter_x_r)
	{
	.reg .u32 %r<38>;
	.reg .u64 %rd<37>;
	.reg .f32 %f<54>;
	.reg .pred %p<13>;
	.shared .align 4 .b8 __cuda_data204[1152];
	.loc	16	163	0
$LBB1_cuda_kernel_chokefilter_x:
	.loc	16	166	0
	cvt.s32.u16 	%r1, %ntid.x;
	cvt.s32.u16 	%r2, %ctaid.x;
	mul24.lo.s32 	%r3, %r2, %r1;
	cvt.u32.u16 	%r4, %tid.x;
	add.u32 	%r5, %r3, %r4;
	ld.param.s32 	%r6, [__cudaparm_cuda_kernel_chokefilter_x_w];
	setp.gt.s32 	%p1, %r6, %r5;
	@%p1 bra 	$Lt_0_12290;
	bra.uni 	$LBB26_cuda_kernel_chokefilter_x;
$Lt_0_12290:
	.loc	16	167	0
	mov.u64 	%rd1, __cuda_data204;
	.loc	16	176	0
	cvt.u32.u16 	%r7, %ctaid.y;
	cvt.u64.u32 	%rd2, %r4;
	mul.lo.u64 	%rd3, %rd2, 4;
	cvt.u64.s32 	%rd4, %r5;
	add.u64 	%rd5, %rd3, %rd1;
	ld.param.u64 	%rd6, [__cudaparm_cuda_kernel_chokefilter_x_id];
	ld.param.u32 	%r8, [__cudaparm_cuda_kernel_chokefilter_x_idPitch];
	mul.lo.u32 	%r9, %r8, %r7;
	cvt.u64.u32 	%rd7, %r9;
	mul.lo.u64 	%rd8, %rd4, 16;
	add.u64 	%rd9, %rd7, %rd8;
	add.u64 	%rd10, %rd6, %rd9;
	ld.global.f32 	%f1, [%rd10+12];
	st.shared.f32 	[%rd5+64], %f1;
	ld.param.s32 	%r10, [__cudaparm_cuda_kernel_chokefilter_x_r];
	add.s32 	%r11, %r10, 1;
	setp.le.u32 	%p2, %r11, %r4;
	@%p2 bra 	$Lt_0_13058;
	neg.s32 	%r12, %r10;
	cvt.s64.s32 	%rd11, %r12;
	mul.lo.u64 	%rd12, %rd11, 4;
	add.u64 	%rd13, %rd5, %rd12;
	setp.le.s32 	%p3, %r11, %r5;
	@%p3 bra 	$Lt_0_13570;
	.loc	16	180	0
	st.shared.f32 	[%rd13+60], %f1;
	bra.uni 	$Lt_0_14338;
$Lt_0_13570:
	.loc	16	181	0
	mul.lo.s32 	%r13, %r10, -4;
	cvt.s64.s32 	%rd14, %r13;
	mul.lo.u64 	%rd15, %rd14, 4;
	add.u64 	%rd16, %rd10, %rd15;
	ld.global.f32 	%f2, [%rd16+-4];
	st.shared.f32 	[%rd13+60], %f2;
	bra.uni 	$Lt_0_14338;
$Lt_0_13058:
	sub.u32 	%r14, %r1, %r10;
	sub.u32 	%r15, %r14, 1;
	setp.lt.u32 	%p4, %r4, %r15;
	@%p4 bra 	$Lt_0_14338;
	cvt.u64.s32 	%rd17, %r10;
	mul.lo.u64 	%rd18, %rd17, 4;
	add.u64 	%rd19, %rd5, %rd18;
	sub.s32 	%r16, %r6, %r10;
	sub.s32 	%r17, %r16, 1;
	setp.lt.s32 	%p5, %r5, %r17;
	@%p5 bra 	$Lt_0_14594;
	.loc	16	183	0
	st.shared.f32 	[%rd19+68], %f1;
	bra.uni 	$Lt_0_14338;
$Lt_0_14594:
	.loc	16	184	0
	mul.lo.s32 	%r18, %r10, 4;
	cvt.s64.s32 	%rd20, %r18;
	mul.lo.u64 	%rd21, %rd20, 4;
	add.u64 	%rd22, %rd10, %rd21;
	ld.global.f32 	%f3, [%rd22+28];
	st.shared.f32 	[%rd19+68], %f3;
$Lt_0_14338:
$Lt_0_13826:
$Lt_0_12802:
	.loc	16	187	0
	bar.sync 	0;
	ld.param.u64 	%rd23, [__cudaparm_cuda_kernel_chokefilter_x_od];
	ld.param.u32 	%r19, [__cudaparm_cuda_kernel_chokefilter_x_odPitch];
	mul.lo.u32 	%r20, %r19, %r7;
	cvt.u64.u32 	%rd24, %r20;
	mul.lo.u64 	%rd25, %rd4, 8;
	add.u64 	%rd26, %rd24, %rd25;
	add.u64 	%rd27, %rd23, %rd26;
	mov.u32 	%r21, 3;
	setp.ne.s32 	%p6, %r10, %r21;
	@%p6 bra 	$Lt_0_15106;
	.loc	16	151	0
	ld.shared.f32 	%f4, [%rd5+48];
	.loc	16	152	0
	ld.shared.f32 	%f5, [%rd5+52];
	.loc	16	154	0
	ld.shared.f32 	%f6, [%rd5+56];
	.loc	16	155	0
	ld.shared.f32 	%f7, [%rd5+60];
	min.f32 	%f8, %f6, %f5;
	min.f32 	%f9, %f7, %f8;
	.loc	16	154	0
	ld.shared.f32 	%f10, [%rd5+64];
	.loc	16	155	0
	ld.shared.f32 	%f11, [%rd5+68];
	min.f32 	%f12, %f10, %f9;
	min.f32 	%f9, %f11, %f12;
	.loc	16	154	0
	ld.shared.f32 	%f13, [%rd5+72];
	.loc	16	155	0
	ld.shared.f32 	%f14, [%rd5+76];
	min.f32 	%f15, %f13, %f9;
	min.f32 	%f9, %f14, %f15;
	.loc	16	157	0
	ld.shared.f32 	%f16, [%rd5+80];
	.loc	16	159	0
	min.f32 	%f17, %f4, %f16;
	min.f32 	%f18, %f9, %f17;
	st.global.v2.f32 	[%rd27+0], {%f18,%f9};
	.loc	16	190	0
	bra.uni 	$LBB26_cuda_kernel_chokefilter_x;
$Lt_0_15106:
	mov.u32 	%r22, 2;
	setp.ne.s32 	%p7, %r10, %r22;
	@%p7 bra 	$Lt_0_15618;
	.loc	16	151	0
	ld.shared.f32 	%f19, [%rd5+52];
	.loc	16	152	0
	ld.shared.f32 	%f20, [%rd5+56];
	.loc	16	154	0
	ld.shared.f32 	%f21, [%rd5+60];
	.loc	16	155	0
	ld.shared.f32 	%f22, [%rd5+64];
	min.f32 	%f23, %f21, %f20;
	min.f32 	%f24, %f22, %f23;
	.loc	16	154	0
	ld.shared.f32 	%f25, [%rd5+68];
	.loc	16	155	0
	ld.shared.f32 	%f26, [%rd5+72];
	min.f32 	%f27, %f25, %f24;
	min.f32 	%f24, %f26, %f27;
	.loc	16	157	0
	ld.shared.f32 	%f28, [%rd5+76];
	.loc	16	159	0
	min.f32 	%f29, %f28, %f19;
	min.f32 	%f30, %f24, %f29;
	st.global.v2.f32 	[%rd27+0], {%f30,%f24};
	.loc	16	191	0
	bra.uni 	$LBB26_cuda_kernel_chokefilter_x;
$Lt_0_15618:
	mov.u32 	%r23, 1;
	setp.ne.s32 	%p8, %r10, %r23;
	@%p8 bra 	$Lt_0_16130;
	.loc	16	151	0
	ld.shared.f32 	%f31, [%rd5+56];
	.loc	16	152	0
	ld.shared.f32 	%f32, [%rd5+60];
	.loc	16	154	0
	ld.shared.f32 	%f33, [%rd5+64];
	.loc	16	155	0
	ld.shared.f32 	%f34, [%rd5+68];
	min.f32 	%f35, %f33, %f32;
	min.f32 	%f36, %f34, %f35;
	.loc	16	157	0
	ld.shared.f32 	%f37, [%rd5+72];
	.loc	16	159	0
	min.f32 	%f38, %f37, %f31;
	min.f32 	%f39, %f36, %f38;
	st.global.v2.f32 	[%rd27+0], {%f39,%f36};
	.loc	16	192	0
	bra.uni 	$LBB26_cuda_kernel_chokefilter_x;
$Lt_0_16130:
	mov.u32 	%r24, 0;
	setp.ne.s32 	%p9, %r10, %r24;
	@%p9 bra 	$Lt_0_16642;
	.loc	16	151	0
	ld.shared.f32 	%f40, [%rd5+60];
	.loc	16	152	0
	ld.shared.f32 	%f41, [%rd5+64];
	.loc	16	157	0
	ld.shared.f32 	%f42, [%rd5+68];
	.loc	16	159	0
	min.f32 	%f43, %f42, %f40;
	min.f32 	%f44, %f41, %f43;
	st.global.v2.f32 	[%rd27+0], {%f44,%f41};
	.loc	16	193	0
	bra.uni 	$LBB26_cuda_kernel_chokefilter_x;
$Lt_0_16642:
	.loc	16	195	0
	neg.s32 	%r25, %r10;
	cvt.s64.s32 	%rd28, %r25;
	mul.lo.u64 	%rd29, %rd28, 4;
	add.u64 	%rd30, %rd5, %rd29;
	ld.shared.f32 	%f45, [%rd30+60];
	.loc	16	196	0
	ld.shared.f32 	%f46, [%rd30+64];
	mul.lo.s32 	%r26, %r10, 2;
	mov.u32 	%r27, 0;
	setp.le.s32 	%p10, %r26, %r27;
	@%p10 bra 	$Lt_0_17922;
	add.s32 	%r28, %r26, 1;
	shr.s32 	%r29, %r28, 31;
	mov.s32 	%r30, 1;
	and.b32 	%r31, %r29, %r30;
	add.s32 	%r32, %r31, %r28;
	shr.s32 	%r33, %r32, 1;
	mov.s32 	%r34, %r25;
	sub.s32 	%r35, %r26, %r10;
	add.s64 	%rd31, %rd3, %rd29;
	add.u64 	%rd32, %rd1, %rd31;
	mov.s32 	%r36, %r33;
$Lt_0_18434:
 //<loop> Loop body line 196, nesting depth: 1, estimated iterations: unknown
	.loc	16	198	0
	ld.shared.f32 	%f47, [%rd32+68];
	.loc	16	199	0
	ld.shared.f32 	%f48, [%rd32+72];
	min.f32 	%f49, %f46, %f47;
	min.f32 	%f46, %f48, %f49;
	add.s32 	%r34, %r34, 2;
	add.u64 	%rd32, %rd32, 8;
	setp.lt.s32 	%p11, %r34, %r35;
	@%p11 bra 	$Lt_0_18434;
$Lt_0_17922:
	.loc	16	201	0
	cvt.s64.s32 	%rd33, %r10;
	mul.lo.u64 	%rd34, %rd33, 4;
	add.u64 	%rd35, %rd5, %rd34;
	ld.shared.f32 	%f50, [%rd35+68];
	.loc	16	203	0
	min.f32 	%f51, %f45, %f50;
	min.f32 	%f52, %f46, %f51;
	st.global.v2.f32 	[%rd27+0], {%f52,%f46};
$LBB26_cuda_kernel_chokefilter_x:
	.loc	16	205	0
	exit;
$LDWend_cuda_kernel_chokefilter_x:
	} // cuda_kernel_chokefilter_x

	.entry cuda_kernel_chokefilter_y (
		.param .u64 __cudaparm_cuda_kernel_chokefilter_y_id,
		.param .u64 __cudaparm_cuda_kernel_chokefilter_y_od,
		.param .s32 __cudaparm_cuda_kernel_chokefilter_y_idPitch,
		.param .s32 __cudaparm_cuda_kernel_chokefilter_y_odPitch,
		.param .s32 __cudaparm_cuda_kernel_chokefilter_y_w,
		.param .s32 __cudaparm_cuda_kernel_chokefilter_y_h,
		.param .s32 __cudaparm_cuda_kernel_chokefilter_y_r,
		.param .f32 __cudaparm_cuda_kernel_chokefilter_y_f,
		.param .s8 __cudaparm_cuda_kernel_chokefilter_y_remap)
	{
	.reg .u32 %r<68>;
	.reg .u64 %rd<32>;
	.reg .f32 %f<90>;
	.reg .pred %p<14>;
	.shared .align 8 .b8 __cuda_data1408[3200];
	.loc	16	224	0
$LBB1_cuda_kernel_chokefilter_y:
	.loc	16	228	0
	cvt.s32.u16 	%r1, %ctaid.x;
	cvt.s32.u16 	%r2, %ntid.x;
	mul24.lo.s32 	%r3, %r1, %r2;
	cvt.s32.u16 	%r4, %ntid.y;
	cvt.s32.u16 	%r5, %ctaid.y;
	mul24.lo.s32 	%r6, %r5, %r4;
	cvt.u32.u16 	%r7, %tid.x;
	add.u32 	%r8, %r3, %r7;
	cvt.u32.u16 	%r9, %tid.y;
	add.u32 	%r10, %r6, %r9;
	ld.param.s32 	%r11, [__cudaparm_cuda_kernel_chokefilter_y_h];
	ld.param.s32 	%r12, [__cudaparm_cuda_kernel_chokefilter_y_w];
	set.le.u32.s32 	%r13, %r12, %r8;
	neg.s32 	%r14, %r13;
	set.le.u32.s32 	%r15, %r11, %r10;
	neg.s32 	%r16, %r15;
	or.b32 	%r17, %r14, %r16;
	mov.u32 	%r18, 0;
	setp.eq.s32 	%p1, %r17, %r18;
	@%p1 bra 	$Lt_1_13314;
	bra.uni 	$LBB27_cuda_kernel_chokefilter_y;
$Lt_1_13314:
	.loc	16	229	0
	mov.u64 	%rd1, __cuda_data1408;
	.loc	16	237	0
	mul24.lo.u32 	%r19, %r7, 25;
	ld.param.s32 	%r20, [__cudaparm_cuda_kernel_chokefilter_y_idPitch];
	add.u32 	%r21, %r19, %r9;
	cvt.u64.u32 	%rd2, %r21;
	mul.lo.u64 	%rd3, %rd2, 8;
	add.u64 	%rd4, %rd3, %rd1;
	mul.lo.s32 	%r22, %r20, %r10;
	mul.lo.s32 	%r23, %r8, 8;
	add.s32 	%r24, %r22, %r23;
	cvt.u64.s32 	%rd5, %r24;
	ld.param.u64 	%rd6, [__cudaparm_cuda_kernel_chokefilter_y_id];
	add.u64 	%rd7, %rd5, %rd6;
	ld.global.v2.f32 	{%f1,%f2}, [%rd7+0];
	st.shared.f32 	[%rd4+64], %f1;
	st.shared.f32 	[%rd4+68], %f2;
	ld.param.s32 	%r25, [__cudaparm_cuda_kernel_chokefilter_y_r];
	add.s32 	%r26, %r25, 1;
	setp.le.u32 	%p2, %r26, %r9;
	@%p2 bra 	$Lt_1_14338;
	neg.s32 	%r27, %r25;
	cvt.s64.s32 	%rd8, %r27;
	mul.lo.u64 	%rd9, %rd8, 8;
	add.u64 	%rd10, %rd4, %rd9;
	setp.le.s32 	%p3, %r26, %r10;
	@%p3 bra 	$Lt_1_14594;
	.loc	16	239	0
	st.shared.f32 	[%rd10+56], %f1;
	ld.shared.f32 	%f3, [%rd4+68];
	st.shared.f32 	[%rd10+60], %f3;
	bra.uni 	$Lt_1_14338;
$Lt_1_14594:
	.loc	16	240	0
	mul.lo.s32 	%r28, %r26, %r20;
	neg.s32 	%r29, %r28;
	shr.s32 	%r30, %r29, 31;
	mov.s32 	%r31, 7;
	and.b32 	%r32, %r30, %r31;
	add.s32 	%r33, %r32, %r29;
	shr.s32 	%r34, %r33, 3;
	cvt.s64.s32 	%rd11, %r34;
	mul.lo.u64 	%rd12, %rd11, 8;
	add.u64 	%rd13, %rd7, %rd12;
	ld.global.v2.f32 	{%f4,%f5}, [%rd13+0];
	st.shared.f32 	[%rd10+56], %f4;
	st.shared.f32 	[%rd10+60], %f5;
$Lt_1_14338:
$Lt_1_13826:
	sub.u32 	%r35, %r4, %r25;
	sub.u32 	%r36, %r35, 1;
	setp.lt.u32 	%p4, %r9, %r36;
	@%p4 bra 	$Lt_1_15362;
	cvt.u64.s32 	%rd14, %r25;
	mul.lo.u64 	%rd15, %rd14, 8;
	add.u64 	%rd16, %rd4, %rd15;
	sub.s32 	%r37, %r11, %r25;
	sub.s32 	%r38, %r37, 1;
	setp.lt.s32 	%p5, %r10, %r38;
	@%p5 bra 	$Lt_1_15618;
	.loc	16	243	0
	ld.shared.f32 	%f6, [%rd4+64];
	st.shared.f32 	[%rd16+72], %f6;
	ld.shared.f32 	%f7, [%rd4+68];
	st.shared.f32 	[%rd16+76], %f7;
	bra.uni 	$Lt_1_15362;
$Lt_1_15618:
	.loc	16	244	0
	mul.lo.s32 	%r39, %r26, %r20;
	shr.s32 	%r40, %r39, 31;
	mov.s32 	%r41, 7;
	and.b32 	%r42, %r40, %r41;
	add.s32 	%r43, %r42, %r39;
	shr.s32 	%r44, %r43, 3;
	cvt.s64.s32 	%rd17, %r44;
	mul.lo.u64 	%rd18, %rd17, 8;
	add.u64 	%rd19, %rd7, %rd18;
	ld.global.v2.f32 	{%f8,%f9}, [%rd19+0];
	st.shared.f32 	[%rd16+72], %f8;
	st.shared.f32 	[%rd16+76], %f9;
$Lt_1_15362:
$Lt_1_14850:
	.loc	16	247	0
	bar.sync 	0;
	mov.u32 	%r45, 3;
	setp.ne.s32 	%p6, %r25, %r45;
	@%p6 bra 	$Lt_1_16130;
	.loc	16	212	0
	ld.shared.f32 	%f10, [%rd4+32];
	ld.shared.f32 	%f11, [%rd4+40];
	.loc	16	213	0
	ld.shared.f32 	%f12, [%rd4+44];
	.loc	16	215	0
	ld.shared.f32 	%f13, [%rd4+48];
	.loc	16	216	0
	ld.shared.f32 	%f14, [%rd4+52];
	.loc	16	217	0
	ld.shared.f32 	%f15, [%rd4+56];
	min.f32 	%f16, %f11, %f10;
	min.f32 	%f17, %f13, %f16;
	min.f32 	%f18, %f15, %f17;
	.loc	16	218	0
	ld.shared.f32 	%f19, [%rd4+60];
	min.f32 	%f20, %f14, %f12;
	min.f32 	%f12, %f19, %f20;
	.loc	16	215	0
	ld.shared.f32 	%f21, [%rd4+64];
	.loc	16	216	0
	ld.shared.f32 	%f22, [%rd4+68];
	.loc	16	217	0
	ld.shared.f32 	%f23, [%rd4+72];
	min.f32 	%f24, %f21, %f18;
	min.f32 	%f18, %f23, %f24;
	.loc	16	218	0
	ld.shared.f32 	%f25, [%rd4+76];
	min.f32 	%f26, %f22, %f12;
	min.f32 	%f12, %f25, %f26;
	.loc	16	215	0
	ld.shared.f32 	%f27, [%rd4+80];
	.loc	16	216	0
	ld.shared.f32 	%f28, [%rd4+84];
	.loc	16	217	0
	ld.shared.f32 	%f29, [%rd4+88];
	min.f32 	%f30, %f27, %f18;
	min.f32 	%f18, %f29, %f30;
	.loc	16	218	0
	ld.shared.f32 	%f31, [%rd4+92];
	min.f32 	%f32, %f28, %f12;
	min.f32 	%f12, %f31, %f32;
	.loc	16	220	0
	ld.shared.f32 	%f33, [%rd4+96];
	.loc	16	250	0
	min.f32 	%f34, %f18, %f33;
	mov.f32 	%f35, %f12;
	bra.uni 	$Lt_1_17410;
$Lt_1_16130:
	mov.u32 	%r46, 2;
	setp.ne.s32 	%p7, %r25, %r46;
	@%p7 bra 	$Lt_1_16642;
	.loc	16	212	0
	ld.shared.f32 	%f36, [%rd4+40];
	ld.shared.f32 	%f37, [%rd4+48];
	.loc	16	213	0
	ld.shared.f32 	%f38, [%rd4+52];
	.loc	16	215	0
	ld.shared.f32 	%f39, [%rd4+56];
	.loc	16	216	0
	ld.shared.f32 	%f40, [%rd4+60];
	.loc	16	217	0
	ld.shared.f32 	%f41, [%rd4+64];
	min.f32 	%f42, %f37, %f36;
	min.f32 	%f43, %f39, %f42;
	min.f32 	%f44, %f41, %f43;
	.loc	16	218	0
	ld.shared.f32 	%f45, [%rd4+68];
	min.f32 	%f46, %f40, %f38;
	min.f32 	%f47, %f45, %f46;
	.loc	16	215	0
	ld.shared.f32 	%f48, [%rd4+72];
	.loc	16	216	0
	ld.shared.f32 	%f49, [%rd4+76];
	.loc	16	217	0
	ld.shared.f32 	%f50, [%rd4+80];
	min.f32 	%f51, %f48, %f44;
	min.f32 	%f44, %f50, %f51;
	.loc	16	218	0
	ld.shared.f32 	%f52, [%rd4+84];
	min.f32 	%f53, %f49, %f47;
	min.f32 	%f47, %f52, %f53;
	.loc	16	220	0
	ld.shared.f32 	%f54, [%rd4+88];
	.loc	16	251	0
	min.f32 	%f34, %f54, %f44;
	mov.f32 	%f35, %f47;
	bra.uni 	$Lt_1_17410;
$Lt_1_16642:
	mov.u32 	%r47, 1;
	setp.ne.s32 	%p8, %r25, %r47;
	@%p8 bra 	$Lt_1_17154;
	.loc	16	212	0
	ld.shared.f32 	%f55, [%rd4+48];
	ld.shared.f32 	%f56, [%rd4+56];
	.loc	16	213	0
	ld.shared.f32 	%f57, [%rd4+60];
	.loc	16	215	0
	ld.shared.f32 	%f58, [%rd4+64];
	.loc	16	216	0
	ld.shared.f32 	%f59, [%rd4+68];
	.loc	16	217	0
	ld.shared.f32 	%f60, [%rd4+72];
	.loc	16	218	0
	ld.shared.f32 	%f61, [%rd4+76];
	min.f32 	%f62, %f59, %f57;
	min.f32 	%f63, %f61, %f62;
	.loc	16	220	0
	ld.shared.f32 	%f64, [%rd4+80];
	.loc	16	252	0
	min.f32 	%f65, %f56, %f55;
	min.f32 	%f66, %f58, %f65;
	min.f32 	%f67, %f60, %f66;
	min.f32 	%f34, %f64, %f67;
	mov.f32 	%f35, %f63;
	bra.uni 	$Lt_1_17410;
$Lt_1_17154:
	mov.u32 	%r48, 0;
	setp.ne.s32 	%p9, %r25, %r48;
	@%p9 bra 	$Lt_1_17666;
	.loc	16	212	0
	ld.shared.f32 	%f68, [%rd4+56];
	ld.shared.f32 	%f69, [%rd4+64];
	.loc	16	213	0
	ld.shared.f32 	%f70, [%rd4+68];
	.loc	16	220	0
	ld.shared.f32 	%f71, [%rd4+72];
	.loc	16	253	0
	min.f32 	%f72, %f69, %f68;
	min.f32 	%f34, %f71, %f72;
	mov.f32 	%f35, %f70;
	bra.uni 	$Lt_1_17410;
$Lt_1_17666:
	.loc	16	255	0
	neg.s32 	%r49, %r25;
	cvt.s64.s32 	%rd20, %r49;
	mul.lo.u64 	%rd21, %rd20, 8;
	add.u64 	%rd22, %rd4, %rd21;
	ld.shared.f32 	%f73, [%rd22+56];
	ld.shared.f32 	%f74, [%rd22+64];
	min.f32 	%f34, %f73, %f74;
	.loc	16	256	0
	ld.shared.f32 	%f35, [%rd22+68];
	mul.lo.s32 	%r50, %r25, 2;
	mov.u32 	%r51, 0;
	setp.le.s32 	%p10, %r50, %r51;
	@%p10 bra 	$Lt_1_18946;
	add.s32 	%r52, %r50, 1;
	shr.s32 	%r53, %r52, 31;
	mov.s32 	%r54, 1;
	and.b32 	%r55, %r53, %r54;
	add.s32 	%r56, %r55, %r52;
	shr.s32 	%r57, %r56, 1;
	mov.s32 	%r58, %r49;
	sub.s32 	%r59, %r50, %r25;
	add.s64 	%rd23, %rd3, %rd21;
	add.u64 	%rd24, %rd1, %rd23;
	mov.s32 	%r60, %r57;
$Lt_1_19458:
 //<loop> Loop body line 256, nesting depth: 1, estimated iterations: unknown
	.loc	16	258	0
	ld.shared.f32 	%f75, [%rd24+72];
	.loc	16	259	0
	ld.shared.f32 	%f76, [%rd24+76];
	.loc	16	260	0
	ld.shared.f32 	%f77, [%rd24+80];
	min.f32 	%f78, %f34, %f75;
	min.f32 	%f34, %f77, %f78;
	.loc	16	261	0
	ld.shared.f32 	%f79, [%rd24+84];
	min.f32 	%f80, %f76, %f35;
	min.f32 	%f35, %f79, %f80;
	add.s32 	%r58, %r58, 2;
	add.u64 	%rd24, %rd24, 16;
	setp.lt.s32 	%p11, %r58, %r59;
	@%p11 bra 	$Lt_1_19458;
$Lt_1_18946:
	.loc	16	263	0
	cvt.s64.s32 	%rd25, %r25;
	mul.lo.u64 	%rd26, %rd25, 8;
	add.u64 	%rd27, %rd4, %rd26;
	ld.shared.f32 	%f81, [%rd27+72];
	min.f32 	%f34, %f34, %f81;
$Lt_1_17410:
$Lt_1_16898:
$Lt_1_16386:
$Lt_1_15874:
	.loc	16	265	0
	ld.param.f32 	%f82, [__cudaparm_cuda_kernel_chokefilter_y_f];
	mov.f32 	%f83, 0f3f800000;    	// 1
	sub.f32 	%f84, %f83, %f82;
	mul.f32 	%f85, %f35, %f84;
	mad.f32 	%f34, %f82, %f34, %f85;
	ld.param.s8 	%r61, [__cudaparm_cuda_kernel_chokefilter_y_remap];
	mov.u32 	%r62, 0;
	setp.eq.s32 	%p12, %r61, %r62;
	@%p12 bra 	$Lt_1_19970;
	.loc	16	266	0
	ld.const.f32 	%f86, [p+144];
	ld.const.f32 	%f87, [p+140];
	mad.f32 	%f88, %f34, %f87, %f86;
	cvt.sat.f32.f32 	%f34, %f88;
$Lt_1_19970:
	.loc	16	267	0
	ld.param.s32 	%r63, [__cudaparm_cuda_kernel_chokefilter_y_odPitch];
	mul.lo.s32 	%r64, %r63, %r10;
	mul.lo.s32 	%r65, %r8, 16;
	add.s32 	%r66, %r64, %r65;
	cvt.u64.s32 	%rd28, %r66;
	ld.param.u64 	%rd29, [__cudaparm_cuda_kernel_chokefilter_y_od];
	add.u64 	%rd30, %rd28, %rd29;
	st.global.f32 	[%rd30+12], %f34;
$LBB27_cuda_kernel_chokefilter_y:
	.loc	16	268	0
	exit;
$LDWend_cuda_kernel_chokefilter_y:
	} // cuda_kernel_chokefilter_y

	.entry cuda_kernel_boxfilter_x (
		.param .u64 __cudaparm_cuda_kernel_boxfilter_x_id,
		.param .u64 __cudaparm_cuda_kernel_boxfilter_x_od,
		.param .s32 __cudaparm_cuda_kernel_boxfilter_x_idPitch,
		.param .s32 __cudaparm_cuda_kernel_boxfilter_x_odPitch,
		.param .s32 __cudaparm_cuda_kernel_boxfilter_x_w,
		.param .s32 __cudaparm_cuda_kernel_boxfilter_x_h,
		.param .s32 __cudaparm_cuda_kernel_boxfilter_x_r)
	{
	.reg .u32 %r<29>;
	.reg .u64 %rd<24>;
	.reg .f32 %f<30>;
	.reg .pred %p<10>;
	.shared .align 4 .b8 __cuda_data4644[1152];
	.loc	16	285	0
$LBB1_cuda_kernel_boxfilter_x:
	.loc	16	288	0
	cvt.s32.u16 	%r1, %ntid.x;
	cvt.s32.u16 	%r2, %ctaid.x;
	mul24.lo.s32 	%r3, %r2, %r1;
	cvt.u32.u16 	%r4, %tid.x;
	add.u32 	%r5, %r3, %r4;
	ld.param.s32 	%r6, [__cudaparm_cuda_kernel_boxfilter_x_w];
	setp.gt.s32 	%p1, %r6, %r5;
	@%p1 bra 	$Lt_2_10498;
	bra.uni 	$LBB20_cuda_kernel_boxfilter_x;
$Lt_2_10498:
	.loc	16	297	0
	cvt.u32.u16 	%r7, %ctaid.y;
	ld.param.u32 	%r8, [__cudaparm_cuda_kernel_boxfilter_x_idPitch];
	mul.lo.u32 	%r9, %r8, %r7;
	mul.lo.s32 	%r10, %r5, 16;
	mov.u64 	%rd1, __cuda_data4644;
	cvt.u64.u32 	%rd2, %r4;
	mul.lo.u64 	%rd3, %rd2, 4;
	add.u64 	%rd4, %rd1, %rd3;
	add.u32 	%r11, %r10, %r9;
	cvt.u64.u32 	%rd5, %r11;
	ld.param.u64 	%rd6, [__cudaparm_cuda_kernel_boxfilter_x_id];
	add.u64 	%rd7, %rd5, %rd6;
	ld.global.f32 	%f1, [%rd7+12];
	st.shared.f32 	[%rd4+64], %f1;
	ld.param.s32 	%r12, [__cudaparm_cuda_kernel_boxfilter_x_r];
	add.s32 	%r13, %r12, 1;
	setp.le.u32 	%p2, %r13, %r4;
	@%p2 bra 	$Lt_2_11266;
	neg.s32 	%r14, %r12;
	cvt.s64.s32 	%rd8, %r14;
	mul.lo.u64 	%rd9, %rd8, 4;
	add.u64 	%rd10, %rd4, %rd9;
	setp.le.s32 	%p3, %r13, %r5;
	@%p3 bra 	$Lt_2_11778;
	.loc	16	299	0
	st.shared.f32 	[%rd10+60], %f1;
	bra.uni 	$Lt_2_12546;
$Lt_2_11778:
	.loc	16	300	0
	mul.lo.s32 	%r15, %r12, -4;
	cvt.s64.s32 	%rd11, %r15;
	mul.lo.u64 	%rd12, %rd11, 4;
	add.u64 	%rd13, %rd7, %rd12;
	ld.global.f32 	%f2, [%rd13+-4];
	st.shared.f32 	[%rd10+60], %f2;
	bra.uni 	$Lt_2_12546;
$Lt_2_11266:
	sub.u32 	%r16, %r1, %r12;
	sub.u32 	%r17, %r16, 1;
	setp.lt.u32 	%p4, %r4, %r17;
	@%p4 bra 	$Lt_2_12546;
	cvt.u64.s32 	%rd14, %r12;
	mul.lo.u64 	%rd15, %rd14, 4;
	add.u64 	%rd16, %rd4, %rd15;
	sub.s32 	%r18, %r6, %r12;
	sub.s32 	%r19, %r18, 1;
	setp.lt.s32 	%p5, %r5, %r19;
	@%p5 bra 	$Lt_2_12802;
	.loc	16	302	0
	st.shared.f32 	[%rd16+68], %f1;
	bra.uni 	$Lt_2_12546;
$Lt_2_12802:
	.loc	16	303	0
	mul.lo.s32 	%r20, %r12, 4;
	cvt.s64.s32 	%rd17, %r20;
	mul.lo.u64 	%rd18, %rd17, 4;
	add.u64 	%rd19, %rd7, %rd18;
	ld.global.f32 	%f3, [%rd19+28];
	st.shared.f32 	[%rd16+68], %f3;
$Lt_2_12546:
$Lt_2_12034:
$Lt_2_11010:
	.loc	16	306	0
	bar.sync 	0;
	ld.param.u32 	%r21, [__cudaparm_cuda_kernel_boxfilter_x_odPitch];
	mul.lo.u32 	%r22, %r21, %r7;
	mul.lo.s32 	%r23, %r5, 8;
	ld.shared.f32 	%f4, [%rd4+64];
	ld.shared.f32 	%f5, [%rd4+60];
	ld.shared.f32 	%f6, [%rd4+68];
	add.u32 	%r24, %r23, %r22;
	cvt.u64.u32 	%rd20, %r24;
	ld.param.u64 	%rd21, [__cudaparm_cuda_kernel_boxfilter_x_od];
	add.u64 	%rd22, %rd20, %rd21;
	mov.u32 	%r25, 3;
	setp.ne.s32 	%p6, %r12, %r25;
	@%p6 bra 	$Lt_2_13314;
	.loc	16	274	0
	ld.shared.f32 	%f7, [%rd4+48];
	.loc	16	275	0
	ld.shared.f32 	%f8, [%rd4+52];
	.loc	16	277	0
	ld.shared.f32 	%f9, [%rd4+56];
	add.f32 	%f10, %f9, %f8;
	.loc	16	306	0
	ld.shared.f32 	%f5, [%rd4+60];
	.loc	16	278	0
	add.f32 	%f10, %f5, %f10;
	.loc	16	306	0
	ld.shared.f32 	%f4, [%rd4+64];
	.loc	16	277	0
	add.f32 	%f10, %f4, %f10;
	.loc	16	306	0
	ld.shared.f32 	%f6, [%rd4+68];
	.loc	16	278	0
	add.f32 	%f10, %f6, %f10;
	.loc	16	277	0
	ld.shared.f32 	%f11, [%rd4+72];
	add.f32 	%f10, %f11, %f10;
	.loc	16	278	0
	ld.shared.f32 	%f12, [%rd4+76];
	add.f32 	%f10, %f12, %f10;
	.loc	16	280	0
	ld.shared.f32 	%f13, [%rd4+80];
	add.f32 	%f7, %f13, %f7;
	.loc	16	282	0
	add.f32 	%f14, %f7, %f10;
	st.global.v2.f32 	[%rd22+0], {%f14,%f10};
	.loc	16	308	0
	bra.uni 	$LBB20_cuda_kernel_boxfilter_x;
$Lt_2_13314:
	mov.u32 	%r26, 2;
	setp.ne.s32 	%p7, %r12, %r26;
	@%p7 bra 	$Lt_2_13826;
	.loc	16	274	0
	ld.shared.f32 	%f15, [%rd4+52];
	.loc	16	275	0
	ld.shared.f32 	%f16, [%rd4+56];
	.loc	16	306	0
	ld.shared.f32 	%f5, [%rd4+60];
	.loc	16	277	0
	add.f32 	%f17, %f16, %f5;
	.loc	16	306	0
	ld.shared.f32 	%f4, [%rd4+64];
	.loc	16	278	0
	add.f32 	%f17, %f4, %f17;
	.loc	16	306	0
	ld.shared.f32 	%f6, [%rd4+68];
	.loc	16	277	0
	add.f32 	%f17, %f6, %f17;
	.loc	16	278	0
	ld.shared.f32 	%f18, [%rd4+72];
	add.f32 	%f17, %f18, %f17;
	.loc	16	280	0
	ld.shared.f32 	%f19, [%rd4+76];
	add.f32 	%f20, %f19, %f15;
	.loc	16	282	0
	add.f32 	%f21, %f20, %f17;
	st.global.v2.f32 	[%rd22+0], {%f21,%f17};
	.loc	16	309	0
	bra.uni 	$LBB20_cuda_kernel_boxfilter_x;
$Lt_2_13826:
	mov.u32 	%r27, 1;
	setp.ne.s32 	%p8, %r12, %r27;
	@%p8 bra 	$Lt_2_14338;
	.loc	16	274	0
	ld.shared.f32 	%f22, [%rd4+56];
	.loc	16	306	0
	ld.shared.f32 	%f5, [%rd4+60];
	ld.shared.f32 	%f4, [%rd4+64];
	.loc	16	277	0
	add.f32 	%f23, %f5, %f4;
	.loc	16	306	0
	ld.shared.f32 	%f6, [%rd4+68];
	.loc	16	278	0
	add.f32 	%f23, %f6, %f23;
	.loc	16	280	0
	ld.shared.f32 	%f24, [%rd4+72];
	add.f32 	%f25, %f24, %f22;
	.loc	16	282	0
	add.f32 	%f26, %f25, %f23;
	st.global.v2.f32 	[%rd22+0], {%f26,%f23};
	.loc	16	310	0
	bra.uni 	$LBB20_cuda_kernel_boxfilter_x;
$Lt_2_14338:
	.loc	16	306	0
	ld.shared.f32 	%f6, [%rd4+68];
	ld.shared.f32 	%f5, [%rd4+60];
	.loc	16	280	0
	add.f32 	%f27, %f6, %f5;
	.loc	16	306	0
	ld.shared.f32 	%f4, [%rd4+64];
	.loc	16	282	0
	add.f32 	%f28, %f4, %f27;
	st.global.v2.f32 	[%rd22+0], {%f28,%f4};
$LBB20_cuda_kernel_boxfilter_x:
	.loc	16	313	0
	exit;
$LDWend_cuda_kernel_boxfilter_x:
	} // cuda_kernel_boxfilter_x

	.entry cuda_kernel_boxfilter_y (
		.param .u64 __cudaparm_cuda_kernel_boxfilter_y_id,
		.param .u64 __cudaparm_cuda_kernel_boxfilter_y_od,
		.param .s32 __cudaparm_cuda_kernel_boxfilter_y_idPitch,
		.param .s32 __cudaparm_cuda_kernel_boxfilter_y_odPitch,
		.param .s32 __cudaparm_cuda_kernel_boxfilter_y_w,
		.param .s32 __cudaparm_cuda_kernel_boxfilter_y_h,
		.param .s32 __cudaparm_cuda_kernel_boxfilter_y_r,
		.param .f32 __cudaparm_cuda_kernel_boxfilter_y_f,
		.param .s8 __cudaparm_cuda_kernel_boxfilter_y_remap)
	{
	.reg .u32 %r<55>;
	.reg .u64 %rd<24>;
	.reg .f32 %f<69>;
	.reg .pred %p<11>;
	.shared .align 8 .b8 __cuda_data5848[3200];
	.loc	16	336	0
$LBB1_cuda_kernel_boxfilter_y:
	.loc	16	340	0
	cvt.s32.u16 	%r1, %ctaid.x;
	cvt.s32.u16 	%r2, %ntid.x;
	mul24.lo.s32 	%r3, %r1, %r2;
	cvt.s32.u16 	%r4, %ntid.y;
	cvt.s32.u16 	%r5, %ctaid.y;
	mul24.lo.s32 	%r6, %r5, %r4;
	cvt.u32.u16 	%r7, %tid.x;
	add.u32 	%r8, %r3, %r7;
	cvt.u32.u16 	%r9, %tid.y;
	add.u32 	%r10, %r6, %r9;
	ld.param.s32 	%r11, [__cudaparm_cuda_kernel_boxfilter_y_h];
	ld.param.s32 	%r12, [__cudaparm_cuda_kernel_boxfilter_y_w];
	set.le.u32.s32 	%r13, %r12, %r8;
	neg.s32 	%r14, %r13;
	set.le.u32.s32 	%r15, %r11, %r10;
	neg.s32 	%r16, %r15;
	or.b32 	%r17, %r14, %r16;
	mov.u32 	%r18, 0;
	setp.eq.s32 	%p1, %r17, %r18;
	@%p1 bra 	$Lt_3_11522;
	bra.uni 	$LBB22_cuda_kernel_boxfilter_y;
$Lt_3_11522:
	.loc	16	349	0
	mul24.lo.u32 	%r19, %r7, 25;
	ld.param.s32 	%r20, [__cudaparm_cuda_kernel_boxfilter_y_idPitch];
	mov.u64 	%rd1, __cuda_data5848;
	add.u32 	%r21, %r19, %r9;
	cvt.u64.u32 	%rd2, %r21;
	mul.lo.u64 	%rd3, %rd2, 8;
	add.u64 	%rd4, %rd1, %rd3;
	mul.lo.s32 	%r22, %r20, %r10;
	mul.lo.s32 	%r23, %r8, 8;
	add.s32 	%r24, %r22, %r23;
	cvt.u64.s32 	%rd5, %r24;
	ld.param.u64 	%rd6, [__cudaparm_cuda_kernel_boxfilter_y_id];
	add.u64 	%rd7, %rd5, %rd6;
	ld.global.v2.f32 	{%f1,%f2}, [%rd7+0];
	st.shared.f32 	[%rd4+64], %f1;
	st.shared.f32 	[%rd4+68], %f2;
	ld.param.s32 	%r25, [__cudaparm_cuda_kernel_boxfilter_y_r];
	add.s32 	%r26, %r25, 1;
	setp.le.u32 	%p2, %r26, %r9;
	@%p2 bra 	$Lt_3_12290;
	neg.s32 	%r27, %r25;
	cvt.s64.s32 	%rd8, %r27;
	mul.lo.u64 	%rd9, %rd8, 8;
	add.u64 	%rd10, %rd4, %rd9;
	setp.le.s32 	%p3, %r26, %r10;
	@%p3 bra 	$Lt_3_12802;
	.loc	16	351	0
	st.shared.f32 	[%rd10+56], %f1;
	ld.shared.f32 	%f3, [%rd4+68];
	st.shared.f32 	[%rd10+60], %f3;
	bra.uni 	$Lt_3_13570;
$Lt_3_12802:
	.loc	16	352	0
	mul.lo.s32 	%r28, %r26, %r20;
	neg.s32 	%r29, %r28;
	shr.s32 	%r30, %r29, 31;
	mov.s32 	%r31, 7;
	and.b32 	%r32, %r30, %r31;
	add.s32 	%r33, %r32, %r29;
	shr.s32 	%r34, %r33, 3;
	cvt.s64.s32 	%rd11, %r34;
	mul.lo.u64 	%rd12, %rd11, 8;
	add.u64 	%rd13, %rd7, %rd12;
	ld.global.v2.f32 	{%f4,%f5}, [%rd13+0];
	st.shared.f32 	[%rd10+56], %f4;
	st.shared.f32 	[%rd10+60], %f5;
	bra.uni 	$Lt_3_13570;
$Lt_3_12290:
	sub.u32 	%r35, %r4, %r25;
	sub.u32 	%r36, %r35, 1;
	setp.lt.u32 	%p4, %r9, %r36;
	@%p4 bra 	$Lt_3_13570;
	cvt.u64.s32 	%rd14, %r25;
	mul.lo.u64 	%rd15, %rd14, 8;
	add.u64 	%rd16, %rd4, %rd15;
	sub.s32 	%r37, %r11, %r25;
	sub.s32 	%r38, %r37, 1;
	setp.lt.s32 	%p5, %r10, %r38;
	@%p5 bra 	$Lt_3_13826;
	.loc	16	354	0
	st.shared.f32 	[%rd16+72], %f1;
	ld.shared.f32 	%f6, [%rd4+68];
	st.shared.f32 	[%rd16+76], %f6;
	bra.uni 	$Lt_3_13570;
$Lt_3_13826:
	.loc	16	355	0
	mul.lo.s32 	%r39, %r26, %r20;
	shr.s32 	%r40, %r39, 31;
	mov.s32 	%r41, 7;
	and.b32 	%r42, %r40, %r41;
	add.s32 	%r43, %r42, %r39;
	shr.s32 	%r44, %r43, 3;
	cvt.s64.s32 	%rd17, %r44;
	mul.lo.u64 	%rd18, %rd17, 8;
	add.u64 	%rd19, %rd7, %rd18;
	ld.global.v2.f32 	{%f7,%f8}, [%rd19+0];
	st.shared.f32 	[%rd16+72], %f7;
	st.shared.f32 	[%rd16+76], %f8;
$Lt_3_13570:
$Lt_3_13058:
$Lt_3_12034:
	.loc	16	358	0
	bar.sync 	0;
	ld.param.f32 	%f9, [__cudaparm_cuda_kernel_boxfilter_y_f];
	mov.f32 	%f10, 0f3f800000;    	// 1
	sub.f32 	%f11, %f10, %f9;
	ld.shared.f32 	%f12, [%rd4+64];
	ld.shared.f32 	%f13, [%rd4+68];
	ld.shared.f32 	%f14, [%rd4+56];
	ld.shared.f32 	%f15, [%rd4+72];
	mov.u32 	%r45, 3;
	setp.ne.s32 	%p6, %r25, %r45;
	@%p6 bra 	$Lt_3_14338;
	.loc	16	322	0
	ld.shared.f32 	%f16, [%rd4+32];
	.loc	16	323	0
	ld.shared.f32 	%f17, [%rd4+40];
	add.f32 	%f16, %f17, %f16;
	.loc	16	324	0
	ld.shared.f32 	%f18, [%rd4+44];
	.loc	16	326	0
	ld.shared.f32 	%f19, [%rd4+48];
	add.f32 	%f16, %f19, %f16;
	.loc	16	327	0
	ld.shared.f32 	%f20, [%rd4+52];
	add.f32 	%f18, %f20, %f18;
	.loc	16	358	0
	ld.shared.f32 	%f14, [%rd4+56];
	.loc	16	328	0
	add.f32 	%f16, %f14, %f16;
	.loc	16	329	0
	ld.shared.f32 	%f21, [%rd4+60];
	add.f32 	%f18, %f21, %f18;
	.loc	16	358	0
	ld.shared.f32 	%f12, [%rd4+64];
	.loc	16	326	0
	add.f32 	%f16, %f12, %f16;
	.loc	16	358	0
	ld.shared.f32 	%f13, [%rd4+68];
	.loc	16	327	0
	add.f32 	%f18, %f13, %f18;
	.loc	16	358	0
	ld.shared.f32 	%f15, [%rd4+72];
	.loc	16	328	0
	add.f32 	%f16, %f15, %f16;
	.loc	16	329	0
	ld.shared.f32 	%f22, [%rd4+76];
	add.f32 	%f18, %f22, %f18;
	.loc	16	326	0
	ld.shared.f32 	%f23, [%rd4+80];
	add.f32 	%f16, %f23, %f16;
	.loc	16	327	0
	ld.shared.f32 	%f24, [%rd4+84];
	add.f32 	%f18, %f24, %f18;
	.loc	16	328	0
	ld.shared.f32 	%f25, [%rd4+88];
	add.f32 	%f16, %f25, %f16;
	.loc	16	329	0
	ld.shared.f32 	%f26, [%rd4+92];
	add.f32 	%f18, %f26, %f18;
	.loc	16	331	0
	ld.shared.f32 	%f27, [%rd4+96];
	add.f32 	%f16, %f27, %f16;
	.loc	16	332	0
	mov.f32 	%f28, 0f42a20000;    	// 81
	div.full.f32 	%f29, %f9, %f28;
	mul.f32 	%f30, %f16, %f29;
	mov.f32 	%f31, 0f42440000;    	// 49
	div.full.f32 	%f32, %f11, %f31;
	mad.f32 	%f18, %f18, %f32, %f30;
	.loc	16	361	0
	mov.f32 	%f33, %f18;
	bra.uni 	$Lt_3_15106;
$Lt_3_14338:
	mov.u32 	%r46, 2;
	setp.ne.s32 	%p7, %r25, %r46;
	@%p7 bra 	$Lt_3_14850;
	.loc	16	322	0
	ld.shared.f32 	%f34, [%rd4+40];
	.loc	16	323	0
	ld.shared.f32 	%f35, [%rd4+48];
	add.f32 	%f36, %f35, %f34;
	.loc	16	324	0
	ld.shared.f32 	%f37, [%rd4+52];
	.loc	16	358	0
	ld.shared.f32 	%f14, [%rd4+56];
	.loc	16	326	0
	add.f32 	%f36, %f14, %f36;
	.loc	16	327	0
	ld.shared.f32 	%f38, [%rd4+60];
	add.f32 	%f39, %f38, %f37;
	.loc	16	358	0
	ld.shared.f32 	%f12, [%rd4+64];
	.loc	16	328	0
	add.f32 	%f36, %f12, %f36;
	.loc	16	358	0
	ld.shared.f32 	%f13, [%rd4+68];
	.loc	16	329	0
	add.f32 	%f39, %f13, %f39;
	.loc	16	358	0
	ld.shared.f32 	%f15, [%rd4+72];
	.loc	16	326	0
	add.f32 	%f36, %f15, %f36;
	.loc	16	327	0
	ld.shared.f32 	%f40, [%rd4+76];
	add.f32 	%f39, %f40, %f39;
	.loc	16	328	0
	ld.shared.f32 	%f41, [%rd4+80];
	add.f32 	%f36, %f41, %f36;
	.loc	16	329	0
	ld.shared.f32 	%f42, [%rd4+84];
	add.f32 	%f39, %f42, %f39;
	.loc	16	331	0
	ld.shared.f32 	%f43, [%rd4+88];
	add.f32 	%f36, %f43, %f36;
	.loc	16	332	0
	mov.f32 	%f44, 0f42440000;    	// 49
	div.full.f32 	%f45, %f9, %f44;
	mul.f32 	%f46, %f36, %f45;
	mov.f32 	%f47, 0f41c80000;    	// 25
	div.full.f32 	%f48, %f11, %f47;
	mad.f32 	%f39, %f39, %f48, %f46;
	.loc	16	362	0
	mov.f32 	%f33, %f39;
	bra.uni 	$Lt_3_15106;
$Lt_3_14850:
	mov.u32 	%r47, 1;
	setp.ne.s32 	%p8, %r25, %r47;
	@%p8 bra 	$Lt_3_15362;
	.loc	16	322	0
	ld.shared.f32 	%f49, [%rd4+48];
	.loc	16	358	0
	ld.shared.f32 	%f14, [%rd4+56];
	.loc	16	323	0
	add.f32 	%f50, %f49, %f14;
	.loc	16	324	0
	ld.shared.f32 	%f51, [%rd4+60];
	.loc	16	358	0
	ld.shared.f32 	%f12, [%rd4+64];
	.loc	16	326	0
	add.f32 	%f50, %f12, %f50;
	.loc	16	358	0
	ld.shared.f32 	%f13, [%rd4+68];
	.loc	16	327	0
	add.f32 	%f52, %f51, %f13;
	.loc	16	358	0
	ld.shared.f32 	%f15, [%rd4+72];
	.loc	16	328	0
	add.f32 	%f50, %f15, %f50;
	.loc	16	329	0
	ld.shared.f32 	%f53, [%rd4+76];
	add.f32 	%f52, %f53, %f52;
	.loc	16	331	0
	ld.shared.f32 	%f54, [%rd4+80];
	add.f32 	%f50, %f54, %f50;
	.loc	16	332	0
	mov.f32 	%f55, 0f41c80000;    	// 25
	div.full.f32 	%f56, %f9, %f55;
	mul.f32 	%f57, %f50, %f56;
	mov.f32 	%f58, 0f41100000;    	// 9
	div.full.f32 	%f59, %f11, %f58;
	mad.f32 	%f52, %f52, %f59, %f57;
	.loc	16	363	0
	mov.f32 	%f33, %f52;
	bra.uni 	$Lt_3_15106;
$Lt_3_15362:
	.loc	16	358	0
	ld.shared.f32 	%f12, [%rd4+64];
	ld.shared.f32 	%f14, [%rd4+56];
	.loc	16	323	0
	add.f32 	%f60, %f12, %f14;
	.loc	16	358	0
	ld.shared.f32 	%f15, [%rd4+72];
	.loc	16	331	0
	add.f32 	%f60, %f15, %f60;
	.loc	16	332	0
	mov.f32 	%f61, 0f41100000;    	// 9
	div.full.f32 	%f62, %f9, %f61;
	mul.f32 	%f63, %f60, %f62;
	.loc	16	358	0
	ld.shared.f32 	%f13, [%rd4+68];
	.loc	16	332	0
	mad.f32 	%f64, %f13, %f11, %f63;
	.loc	16	364	0
	mov.f32 	%f33, %f64;
$Lt_3_15106:
$Lt_3_14594:
$Lt_3_14082:
	ld.param.s8 	%r48, [__cudaparm_cuda_kernel_boxfilter_y_remap];
	mov.u32 	%r49, 0;
	setp.eq.s32 	%p9, %r48, %r49;
	@%p9 bra 	$Lt_3_16642;
	.loc	16	366	0
	ld.const.f32 	%f65, [p+144];
	ld.const.f32 	%f66, [p+140];
	mad.f32 	%f67, %f33, %f66, %f65;
	cvt.sat.f32.f32 	%f33, %f67;
$Lt_3_16642:
	.loc	16	367	0
	ld.param.s32 	%r50, [__cudaparm_cuda_kernel_boxfilter_y_odPitch];
	mul.lo.s32 	%r51, %r50, %r10;
	mul.lo.s32 	%r52, %r8, 16;
	add.s32 	%r53, %r51, %r52;
	cvt.u64.s32 	%rd20, %r53;
	ld.param.u64 	%rd21, [__cudaparm_cuda_kernel_boxfilter_y_od];
	add.u64 	%rd22, %rd20, %rd21;
	st.global.f32 	[%rd22+12], %f33;
$LBB22_cuda_kernel_boxfilter_y:
	.loc	16	368	0
	exit;
$LDWend_cuda_kernel_boxfilter_y:
	} // cuda_kernel_boxfilter_y

	.entry cuda_kernel_composite (
		.param .u64 __cudaparm_cuda_kernel_composite_srcFrame,
		.param .u64 __cudaparm_cuda_kernel_composite_dstFrame,
		.param .s32 __cudaparm_cuda_kernel_composite_width,
		.param .s32 __cudaparm_cuda_kernel_composite_height,
		.param .s32 __cudaparm_cuda_kernel_composite_srcPitch,
		.param .s32 __cudaparm_cuda_kernel_composite_dstPitch,
		.param .s8 __cudaparm_cuda_kernel_composite_remap)
	{
	.reg .u32 %r<28>;
	.reg .u64 %rd<11>;
	.reg .f32 %f<187>;
	.reg .f64 %fd<5>;
	.reg .pred %p<12>;
	.loc	16	371	0
$LBB1_cuda_kernel_composite:
	.loc	16	374	0
	cvt.s32.u16 	%r1, %ctaid.x;
	cvt.s32.u16 	%r2, %ntid.x;
	mul24.lo.s32 	%r3, %r1, %r2;
	cvt.s32.u16 	%r4, %ctaid.y;
	cvt.s32.u16 	%r5, %ntid.y;
	mul24.lo.s32 	%r6, %r4, %r5;
	cvt.u32.u16 	%r7, %tid.x;
	add.u32 	%r8, %r3, %r7;
	cvt.u32.u16 	%r9, %tid.y;
	add.u32 	%r10, %r6, %r9;
	ld.param.s32 	%r11, [__cudaparm_cuda_kernel_composite_width];
	set.le.u32.s32 	%r12, %r11, %r8;
	neg.s32 	%r13, %r12;
	ld.param.s32 	%r14, [__cudaparm_cuda_kernel_composite_height];
	set.le.u32.s32 	%r15, %r14, %r10;
	neg.s32 	%r16, %r15;
	or.b32 	%r17, %r13, %r16;
	mov.u32 	%r18, 0;
	setp.eq.s32 	%p1, %r17, %r18;
	@%p1 bra 	$Lt_4_8194;
	bra.uni 	$LBB27_cuda_kernel_composite;
$Lt_4_8194:
	.loc	16	78	0
	mov.s32 	%r19, 16;
	mul24.lo.s32 	%r20, %r8, %r19;
	ld.param.s32 	%r21, [__cudaparm_cuda_kernel_composite_srcPitch];
	mul24.lo.s32 	%r22, %r10, %r21;
	cvt.u64.s32 	%rd1, %r20;
	cvt.u64.s32 	%rd2, %r22;
	ld.param.u64 	%rd3, [__cudaparm_cuda_kernel_composite_srcFrame];
	add.u64 	%rd4, %rd2, %rd3;
	add.u64 	%rd5, %rd1, %rd4;
	ld.global.v4.f32 	{%f1,%f2,%f3,%f4}, [%rd5+0];
	.loc	16	79	0
	cvt.sat.f32.f32 	%f5, %f1;
	mov.f32 	%f6, 0f3d25aee6;     	// 0.04045
	setp.gt.f32 	%p2, %f5, %f6;
	@!%p2 bra 	$Lt_4_8962;
	.loc	16	35	0
	mov.f32 	%f7, 0f3d6147ae;     	// 0.055
	add.f32 	%f8, %f5, %f7;
	mov.f32 	%f9, 0f3f870a3d;     	// 1.055
	div.full.f32 	%f10, %f8, %f9;
	lg2.approx.f32 	%f11, %f10;
	mov.f32 	%f12, 0f4019999a;    	// 2.4
	mul.f32 	%f13, %f11, %f12;
	ex2.approx.f32 	%f14, %f13;
	bra.uni 	$Lt_4_8706;
$Lt_4_8962:
	.loc	16	37	0
	mov.f32 	%f15, 0f414eb852;    	// 12.92
	div.full.f32 	%f14, %f5, %f15;
$Lt_4_8706:
	.loc	16	80	0
	cvt.sat.f32.f32 	%f16, %f2;
	mov.f32 	%f17, 0f3d25aee6;    	// 0.04045
	setp.gt.f32 	%p3, %f16, %f17;
	@!%p3 bra 	$Lt_4_9474;
	.loc	16	35	0
	mov.f32 	%f18, 0f3d6147ae;    	// 0.055
	add.f32 	%f19, %f16, %f18;
	mov.f32 	%f20, 0f3f870a3d;    	// 1.055
	div.full.f32 	%f21, %f19, %f20;
	lg2.approx.f32 	%f22, %f21;
	mov.f32 	%f23, 0f4019999a;    	// 2.4
	mul.f32 	%f24, %f22, %f23;
	ex2.approx.f32 	%f25, %f24;
	bra.uni 	$Lt_4_9218;
$Lt_4_9474:
	.loc	16	37	0
	mov.f32 	%f26, 0f414eb852;    	// 12.92
	div.full.f32 	%f25, %f16, %f26;
$Lt_4_9218:
	.loc	16	81	0
	cvt.sat.f32.f32 	%f27, %f3;
	mov.f32 	%f28, 0f3d25aee6;    	// 0.04045
	setp.gt.f32 	%p4, %f27, %f28;
	@!%p4 bra 	$Lt_4_9986;
	.loc	16	35	0
	mov.f32 	%f29, 0f3d6147ae;    	// 0.055
	add.f32 	%f30, %f27, %f29;
	mov.f32 	%f31, 0f3f870a3d;    	// 1.055
	div.full.f32 	%f32, %f30, %f31;
	lg2.approx.f32 	%f33, %f32;
	mov.f32 	%f34, 0f4019999a;    	// 2.4
	mul.f32 	%f35, %f33, %f34;
	ex2.approx.f32 	%f36, %f35;
	bra.uni 	$Lt_4_9730;
$Lt_4_9986:
	.loc	16	37	0
	mov.f32 	%f37, 0f414eb852;    	// 12.92
	div.full.f32 	%f36, %f27, %f37;
$Lt_4_9730:
	.loc	16	95	0
	ld.const.f32 	%f38, [p+56];
	mul.f32 	%f39, %f38, %f25;
	ld.const.f32 	%f40, [p+68];
	mul.f32 	%f41, %f40, %f25;
	ld.const.f32 	%f42, [p+44];
	mul.f32 	%f43, %f42, %f25;
	ld.const.f32 	%f44, [p+52];
	mad.f32 	%f45, %f36, %f44, %f39;
	ld.const.f32 	%f46, [p+64];
	mad.f32 	%f47, %f36, %f46, %f41;
	ld.const.f32 	%f48, [p+40];
	mad.f32 	%f49, %f36, %f48, %f43;
	ld.const.f32 	%f50, [p+60];
	mad.f32 	%f51, %f14, %f50, %f45;
	ld.const.f32 	%f52, [p+72];
	mad.f32 	%f53, %f14, %f52, %f47;
	ld.const.f32 	%f54, [p+48];
	mad.f32 	%f55, %f14, %f54, %f49;
	mov.f32 	%f56, 0f02081cea;    	// 1e-037
	max.f32 	%f57, %f53, %f56;
	div.approx.f32 	%f58, %f51, %f57;
	div.approx.f32 	%f59, %f55, %f57;
	ld.const.f32 	%f60, [p+152];
	sub.f32 	%f61, %f58, %f60;
	ld.const.f32 	%f62, [p+148];
	sub.f32 	%f63, %f59, %f62;
	mul.f32 	%f64, %f61, %f61;
	mad.f32 	%f65, %f63, %f63, %f64;
	mov.f32 	%f66, 0f02081cea;    	// 1e-037
	max.f32 	%f67, %f65, %f66;
	rsqrt.approx.f32 	%f68, %f67;
	ld.const.f32 	%f69, [p+160];
	mul.f32 	%f70, %f61, %f60;
	mad.f32 	%f71, %f62, %f63, %f70;
	mul.f32 	%f72, %f68, %f69;
	mul.f32 	%f73, %f71, %f72;
	ld.const.f32 	%f74, [p+8];
	mov.f32 	%f75, 0f00000000;    	// 0
	setp.gt.f32 	%p5, %f73, %f75;
	@!%p5 bra 	$Lt_4_10498;
	.loc	16	97	0
	mul.f32 	%f76, %f73, %f74;
	.loc	16	98	0
	mov.f32 	%f77, %f74;
	bra.uni 	$Lt_4_10242;
$Lt_4_10498:
	.loc	16	101	0
	mov.f32 	%f78, 0f3f800000;    	// 1
	add.f32 	%f79, %f73, %f78;
	mul.f32 	%f77, %f74, %f79;
	mov.f32 	%f76, 0f00000000;    	// 0
$Lt_4_10242:
	.loc	16	106	0
	mov.f32 	%f80, 0f3f800000;    	// 1
	div.approx.f32 	%f81, %f80, %f68;
	mul.f32 	%f82, %f69, %f81;
	.loc	16	108	0
	mov.f32 	%f83, 0f3f800000;    	// 1
	sub.f32 	%f84, %f83, %f76;
	mul.f32 	%f76, %f82, %f84;
	.loc	16	109	0
	mov.f32 	%f85, 0f3f800000;    	// 1
	sub.f32 	%f86, %f85, %f77;
	mul.f32 	%f77, %f82, %f86;
	.loc	16	127	0
	ld.const.f32 	%f87, [p+16];
	mul.f32 	%f88, %f87, %f76;
	cvt.sat.f32.f32 	%f89, %f88;
	ld.const.f32 	%f90, [p+156];
	sub.f32 	%f91, %f57, %f90;
	div.approx.f32 	%f92, %f91, %f90;
	mov.f32 	%f93, 0f00000000;    	// 0
	setp.lt.f32 	%p6, %f92, %f93;
	ld.const.f32 	%f94, [p+20];
	mul.f32 	%f95, %f94, %f92;
	ld.const.f32 	%f96, [p+24];
	mul.f32 	%f97, %f96, %f92;
	selp.f32 	%f98, %f95, %f97, %p6;
	cvt.sat.f32.f32 	%f99, %f98;
	add.f32 	%f100, %f89, %f99;
	mul.f32 	%f101, %f99, %f89;
	sub.f32 	%f102, %f100, %f101;
	.loc	16	131	0
	ld.const.f32 	%f103, [p+28];
	mov.f32 	%f104, 0f3f800000;   	// 1
	mov.f32 	%f105, 0f02081cea;   	// 1e-037
	max.f32 	%f106, %f102, %f105;
	div.approx.f32 	%f107, %f104, %f106;
	cvt.f64.f32 	%fd1, %f107;
	mov.f64 	%fd2, 0dbff0000000000000;	// -1
	add.f64 	%fd3, %fd1, %fd2;
	cvt.rn.f32.f64 	%f108, %fd3;
	mul.f32 	%f109, %f103, %f108;
	mad.f32 	%f110, %f109, %f91, %f57;
	.loc	16	133	0
	ld.const.f32 	%f111, [p+32];
	mul.f32 	%f102, %f111, %f102;
	.loc	16	134	0
	ld.const.f32 	%f112, [p+36];
	sub.f32 	%f102, %f102, %f112;
	.loc	16	135	0
	cvt.sat.f32.f32 	%f102, %f102;
	.loc	16	139	0
	ld.const.f32 	%f113, [p+12];
	mul.f32 	%f114, %f113, %f77;
	ld.const.f32 	%f115, [p+0];
	mul.f32 	%f116, %f115, %f77;
	cvt.sat.f32.f32 	%f117, %f114;
	cvt.sat.f32.f32 	%f118, %f116;
	mov.f32 	%f119, 0f02081cea;   	// 1e-037
	max.f32 	%f120, %f118, %f119;
	mov.f32 	%f121, 0f3f800000;   	// 1
	div.approx.f32 	%f122, %f121, %f120;
	mov.f32 	%f123, 0fbf800000;   	// -1
	add.f32 	%f124, %f122, %f123;
	ld.const.f32 	%f125, [p+4];
	mad.f32 	%f126, %f124, %f125, %f124;
	mad.f32 	%f127, %f126, %f61, %f58;
	mad.f32 	%f128, %f63, %f126, %f59;
	mul.f32 	%f129, %f117, %f127;
	mul.f32 	%f130, %f117, %f128;
	mul.f32 	%f131, %f129, %f110;
	mul.f32 	%f132, %f130, %f110;
	ld.const.f32 	%f133, [p+112];
	ld.const.f32 	%f134, [p+104];
	mul.f32 	%f135, %f134, %f131;
	ld.const.f32 	%f136, [p+100];
	mad.f32 	%f137, %f132, %f136, %f135;
	ld.const.f32 	%f138, [p+108];
	mad.f32 	%f139, %f110, %f138, %f137;
	add.f32 	%f140, %f133, %f139;
	.loc	16	50	0
	cvt.sat.f32.f32 	%f140, %f140;
	mov.f32 	%f141, 0f3b4d2e1c;   	// 0.0031308
	setp.lt.f32 	%p7, %f140, %f141;
	@!%p7 bra 	$Lt_4_11010;
	.loc	16	52	0
	mov.f32 	%f142, 0f414eb852;   	// 12.92
	mul.f32 	%f140, %f140, %f142;
	bra.uni 	$Lt_4_10754;
$Lt_4_11010:
	.loc	16	54	0
	mov.f32 	%f143, 0fbd6147ae;   	// -0.055
	lg2.approx.f32 	%f144, %f140;
	mov.f32 	%f145, 0f3ed55476;   	// 0.41666
	mul.f32 	%f146, %f144, %f145;
	ex2.approx.f32 	%f147, %f146;
	mov.f32 	%f148, 0f3f870a3d;   	// 1.055
	mad.f32 	%f140, %f147, %f148, %f143;
$Lt_4_10754:
	.loc	16	140	0
	ld.const.f32 	%f149, [p+116];
	ld.const.f32 	%f150, [p+92];
	mul.f32 	%f151, %f150, %f131;
	ld.const.f32 	%f152, [p+88];
	mad.f32 	%f153, %f132, %f152, %f151;
	ld.const.f32 	%f154, [p+96];
	mad.f32 	%f155, %f110, %f154, %f153;
	add.f32 	%f156, %f149, %f155;
	.loc	16	50	0
	cvt.sat.f32.f32 	%f156, %f156;
	mov.f32 	%f157, 0f3b4d2e1c;   	// 0.0031308
	setp.lt.f32 	%p8, %f156, %f157;
	@!%p8 bra 	$Lt_4_11522;
	.loc	16	52	0
	mov.f32 	%f158, 0f414eb852;   	// 12.92
	mul.f32 	%f156, %f156, %f158;
	bra.uni 	$Lt_4_11266;
$Lt_4_11522:
	.loc	16	54	0
	mov.f32 	%f159, 0fbd6147ae;   	// -0.055
	lg2.approx.f32 	%f160, %f156;
	mov.f32 	%f161, 0f3ed55476;   	// 0.41666
	mul.f32 	%f162, %f160, %f161;
	ex2.approx.f32 	%f163, %f162;
	mov.f32 	%f164, 0f3f870a3d;   	// 1.055
	mad.f32 	%f156, %f163, %f164, %f159;
$Lt_4_11266:
	.loc	16	141	0
	ld.const.f32 	%f165, [p+120];
	ld.const.f32 	%f166, [p+80];
	mul.f32 	%f167, %f166, %f131;
	ld.const.f32 	%f168, [p+76];
	mad.f32 	%f169, %f132, %f168, %f167;
	ld.const.f32 	%f170, [p+84];
	mad.f32 	%f171, %f110, %f170, %f169;
	add.f32 	%f172, %f165, %f171;
	.loc	16	50	0
	cvt.sat.f32.f32 	%f172, %f172;
	mov.f32 	%f173, 0f3b4d2e1c;   	// 0.0031308
	setp.lt.f32 	%p9, %f172, %f173;
	@!%p9 bra 	$Lt_4_12034;
	.loc	16	52	0
	mov.f32 	%f174, 0f414eb852;   	// 12.92
	mul.f32 	%f172, %f172, %f174;
	bra.uni 	$Lt_4_11778;
$Lt_4_12034:
	.loc	16	54	0
	mov.f32 	%f175, 0fbd6147ae;   	// -0.055
	lg2.approx.f32 	%f176, %f172;
	mov.f32 	%f177, 0f3ed55476;   	// 0.41666
	mul.f32 	%f178, %f176, %f177;
	ex2.approx.f32 	%f179, %f178;
	mov.f32 	%f180, 0f3f870a3d;   	// 1.055
	mad.f32 	%f172, %f179, %f180, %f175;
$Lt_4_11778:
	.loc	16	142	0
	mul.f32 	%f181, %f102, %f4;
	.loc	16	382	0
	mov.f32 	%f182, %f181;
	ld.param.s8 	%r23, [__cudaparm_cuda_kernel_composite_remap];
	mov.u32 	%r24, 0;
	setp.eq.s32 	%p10, %r23, %r24;
	@%p10 bra 	$Lt_4_12290;
	.loc	16	383	0
	ld.const.f32 	%f183, [p+144];
	ld.const.f32 	%f184, [p+140];
	mad.f32 	%f185, %f181, %f184, %f183;
	cvt.sat.f32.f32 	%f182, %f185;
$Lt_4_12290:
	.loc	16	384	0
	ld.param.s32 	%r25, [__cudaparm_cuda_kernel_composite_dstPitch];
	mul24.lo.s32 	%r26, %r10, %r25;
	cvt.u64.s32 	%rd6, %r26;
	ld.param.u64 	%rd7, [__cudaparm_cuda_kernel_composite_dstFrame];
	add.u64 	%rd8, %rd6, %rd7;
	add.u64 	%rd9, %rd1, %rd8;
	st.global.v4.f32 	[%rd9+0], {%f140,%f156,%f172,%f182};
$LBB27_cuda_kernel_composite:
	.loc	16	386	0
	exit;
$LDWend_cuda_kernel_composite:
	} // cuda_kernel_composite

	.entry cuda_kernel_showAlpha (
		.param .u64 __cudaparm_cuda_kernel_showAlpha_srcFrame,
		.param .u64 __cudaparm_cuda_kernel_showAlpha_dstFrame,
		.param .s32 __cudaparm_cuda_kernel_showAlpha_width,
		.param .s32 __cudaparm_cuda_kernel_showAlpha_height,
		.param .s32 __cudaparm_cuda_kernel_showAlpha_srcPitch,
		.param .s32 __cudaparm_cuda_kernel_showAlpha_dstPitch)
	{
	.reg .u32 %r<26>;
	.reg .u64 %rd<11>;
	.reg .f32 %f<4>;
	.reg .pred %p<3>;
	.loc	16	389	0
$LBB1_cuda_kernel_showAlpha:
	.loc	16	392	0
	cvt.s32.u16 	%r1, %ctaid.x;
	cvt.s32.u16 	%r2, %ntid.x;
	mul24.lo.s32 	%r3, %r1, %r2;
	cvt.s32.u16 	%r4, %ctaid.y;
	cvt.s32.u16 	%r5, %ntid.y;
	mul24.lo.s32 	%r6, %r4, %r5;
	cvt.u32.u16 	%r7, %tid.x;
	add.u32 	%r8, %r3, %r7;
	cvt.u32.u16 	%r9, %tid.y;
	add.u32 	%r10, %r6, %r9;
	ld.param.s32 	%r11, [__cudaparm_cuda_kernel_showAlpha_width];
	set.le.u32.s32 	%r12, %r11, %r8;
	neg.s32 	%r13, %r12;
	ld.param.s32 	%r14, [__cudaparm_cuda_kernel_showAlpha_height];
	set.le.u32.s32 	%r15, %r14, %r10;
	neg.s32 	%r16, %r15;
	or.b32 	%r17, %r13, %r16;
	mov.u32 	%r18, 0;
	setp.eq.s32 	%p1, %r17, %r18;
	@%p1 bra 	$Lt_5_1282;
	bra.uni 	$LBB4_cuda_kernel_showAlpha;
$Lt_5_1282:
	.loc	16	400	0
	mov.s32 	%r19, 16;
	mul24.lo.s32 	%r20, %r8, %r19;
	cvt.u64.s32 	%rd1, %r20;
	ld.param.s32 	%r21, [__cudaparm_cuda_kernel_showAlpha_srcPitch];
	mul24.lo.s32 	%r22, %r10, %r21;
	cvt.u64.s32 	%rd2, %r22;
	ld.param.u64 	%rd3, [__cudaparm_cuda_kernel_showAlpha_srcFrame];
	add.u64 	%rd4, %rd2, %rd3;
	add.u64 	%rd5, %rd1, %rd4;
	ld.global.f32 	%f1, [%rd5+12];
	.loc	16	405	0
	ld.param.s32 	%r23, [__cudaparm_cuda_kernel_showAlpha_dstPitch];
	mul24.lo.s32 	%r24, %r10, %r23;
	cvt.u64.s32 	%rd6, %r24;
	ld.param.u64 	%rd7, [__cudaparm_cuda_kernel_showAlpha_dstFrame];
	add.u64 	%rd8, %rd6, %rd7;
	add.u64 	%rd9, %rd1, %rd8;
	mov.f32 	%f2, 0f3f800000;     	// 1
	st.global.v4.f32 	[%rd9+0], {%f1,%f1,%f1,%f2};
$LBB4_cuda_kernel_showAlpha:
	.loc	16	406	0
	exit;
$LDWend_cuda_kernel_showAlpha:
	} // cuda_kernel_showAlpha

	.entry cuda_kernel_showColor (
		.param .u64 __cudaparm_cuda_kernel_showColor_srcFrame,
		.param .u64 __cudaparm_cuda_kernel_showColor_dstFrame,
		.param .s32 __cudaparm_cuda_kernel_showColor_width,
		.param .s32 __cudaparm_cuda_kernel_showColor_height,
		.param .s32 __cudaparm_cuda_kernel_showColor_srcPitch,
		.param .s32 __cudaparm_cuda_kernel_showColor_dstPitch)
	{
	.reg .u32 %r<26>;
	.reg .u64 %rd<11>;
	.reg .f32 %f<7>;
	.reg .pred %p<3>;
	.loc	16	409	0
$LBB1_cuda_kernel_showColor:
	.loc	16	412	0
	cvt.s32.u16 	%r1, %ctaid.x;
	cvt.s32.u16 	%r2, %ntid.x;
	mul24.lo.s32 	%r3, %r1, %r2;
	cvt.s32.u16 	%r4, %ctaid.y;
	cvt.s32.u16 	%r5, %ntid.y;
	mul24.lo.s32 	%r6, %r4, %r5;
	cvt.u32.u16 	%r7, %tid.x;
	add.u32 	%r8, %r3, %r7;
	cvt.u32.u16 	%r9, %tid.y;
	add.u32 	%r10, %r6, %r9;
	ld.param.s32 	%r11, [__cudaparm_cuda_kernel_showColor_width];
	set.le.u32.s32 	%r12, %r11, %r8;
	neg.s32 	%r13, %r12;
	ld.param.s32 	%r14, [__cudaparm_cuda_kernel_showColor_height];
	set.le.u32.s32 	%r15, %r14, %r10;
	neg.s32 	%r16, %r15;
	or.b32 	%r17, %r13, %r16;
	mov.u32 	%r18, 0;
	setp.eq.s32 	%p1, %r17, %r18;
	@%p1 bra 	$Lt_6_1282;
	bra.uni 	$LBB4_cuda_kernel_showColor;
$Lt_6_1282:
	.loc	16	420	0
	mov.s32 	%r19, 16;
	mul24.lo.s32 	%r20, %r8, %r19;
	ld.param.s32 	%r21, [__cudaparm_cuda_kernel_showColor_srcPitch];
	mul24.lo.s32 	%r22, %r10, %r21;
	cvt.u64.s32 	%rd1, %r20;
	cvt.u64.s32 	%rd2, %r22;
	ld.param.u64 	%rd3, [__cudaparm_cuda_kernel_showColor_srcFrame];
	add.u64 	%rd4, %rd2, %rd3;
	add.u64 	%rd5, %rd1, %rd4;
	ld.global.v4.f32 	{%f1,%f2,%f3,%f4}, [%rd5+0];
	.loc	16	421	0
	mul.f32 	%f1, %f1, %f4;
	.loc	16	422	0
	mul.f32 	%f2, %f2, %f4;
	.loc	16	423	0
	mul.f32 	%f3, %f3, %f4;
	.loc	16	425	0
	ld.param.s32 	%r23, [__cudaparm_cuda_kernel_showColor_dstPitch];
	mul24.lo.s32 	%r24, %r10, %r23;
	cvt.u64.s32 	%rd6, %r24;
	ld.param.u64 	%rd7, [__cudaparm_cuda_kernel_showColor_dstFrame];
	add.u64 	%rd8, %rd6, %rd7;
	add.u64 	%rd9, %rd1, %rd8;
	mov.f32 	%f5, 0f3f800000;     	// 1
	st.global.v4.f32 	[%rd9+0], {%f1,%f2,%f3,%f5};
$LBB4_cuda_kernel_showColor:
	.loc	16	426	0
	exit;
$LDWend_cuda_kernel_showColor:
	} // cuda_kernel_showColor

